cudaMallocPitch( void** devPtr，size_t* pitch，size_t widthInBytes，size_t height )中参数pitch到底什么意思呢？？

[beyond_jhf]

不知道pitch到底指的是什么，用什么好处，没理解透，请大虾们多指教，先谢过了。。

[无心人_过过小日子]

c语言申请2维内存时，一般是连续存放的。a[y][x]存放在第y*widthofx*sizeof(元素)+x*sizeof(元素)个字节。
但在cuda的global memory访问中，从256字节对齐的地址(addr=0, 256, 512, ...)开始的连续访问是最有效率的。
这样，为了提高内存访问的效率，有了cudaMallocPitch函数。
cudaMallocPitch函数分配的内存中，数组的每一行的第一个元素的开始地址都保证是对齐的。因为每行有多少个
数据是不确定的，widthofx*sizeof(元素)不一定是256的倍数。故此，为保证数组的每一行的第一个元素的开始地址
对齐，cudaMallocPitch在分配内存时，每行会多分配一些字节，以保证widthofx*sizeof(元素)+多分配的字节是
256的倍数(对齐)。这样，上面的y*widthofx*sizeof(元素)+x*sizeof(元素)来计算a[y][x]的地址就不正确了。
而应该是y*[widthofx*sizeof(元素)+多分配的字节]+x*sizeof(元素)。
而函数中返回的pitch的值就是widthofx*sizeof(元素)+多分配的字节。

补充：widthInBytes作为输入参数，应该是widthofx*sizeof(元素)。

[beyond_jhf1]

还不是很明白，你看看这个程序的测试结果：
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
int main()
{
float* devPtr;
size_t pitch;
int width=1000;
int height=1000;
cudaMallocPitch((void**)&devPtr, &pitch,width * sizeof(float), height);
printf("\npitch=%d\n",pitch);
return 0;
}
为什么返回的pitch是4032，比1000大那么多，不是很浪费内存空间吗？？而且4032也不是256的倍数啊。。

[无心人_过过小日子]

1)要看清pitch的单位是字节。sizeof(float)=4。
2)LZ的例子中对齐在32字节上。(上面说256字节是个例子，关键是对齐)。

[无心人_过过小日子]

[Quote=引用 3 楼 l7331014 的回复:]
2)LZ的例子中对齐在32字节上。(上面说256字节是个例子，关键是对齐)。
[/Quote]

抱歉，计算错误，对齐在64字节上。

[beyond_jhf1]

[Quote=引用 4 楼 l7331014 的回复:]
引用 3 楼 l7331014 的回复:
2)LZ的例子中对齐在32字节上。(上面说256字节是个例子，关键是对齐)。


抱歉，计算错误，对齐在64字节上。
[/Quote]

64字节是怎么计算出来的呢？？还要麻烦你了

[beyond_jhf1]

对了，为了对齐多分配的那些内存还可以用于其他不？？？

[无心人_过过小日子]

[Quote=引用 6 楼 beyond_jhf1 的回复:]
对了，为了对齐多分配的那些内存还可以用于其他不？？？
[/Quote]

理论上应该是可以的。不过，意义不大。(以上面的例子来说每行多了32字节而已)。
global memory该不会紧张到这个地步。

[无心人_过过小日子]

[Quote=引用 5 楼 beyond_jhf1 的回复:]
64字节是怎么计算出来的呢？？还要麻烦你了
[/Quote]

4032=63*64。

[beyond_jhf]

[Quote=引用 8 楼 l7331014 的回复:]
引用 5 楼 beyond_jhf1 的回复:
64字节是怎么计算出来的呢？？还要麻烦你了


4032=63*64。
[/Quote]

还是不明白为什么写成4032=63*64的形式，请兄台明示了。。

[无心人_过过小日子]

[Quote=引用 9 楼 beyond_jhf 的回复:]
还是不明白为什么写成4032=63*64的形式，请兄台明示了。。
[/Quote]

看编程手册的第5章性能指南的5.1.2.1全局存储器这节，其中提到了存取器事务(一次访问宽度，如不对齐的话
估计访问宽度就浪费了)大小为32字节(1.2能力及以上),64字节，128字节。

LZ例子中，pitch最小为实际存储需要的1000*4=4000字节。
如是32字节对齐的，则pitch应该为4000(32的倍数)
现在是4032=63*64，则是64字节对齐了。至于128，4032不能被128整除。

[beyond_jhf]

谢谢了

[bengkui123456]

请教怎样利用pitch访问二维数组

[keyboymaxwell119]

謝謝，我剛好也這裡不懂的地方，有幫助到我的解惑。

[hyyingyi1003613]

感谢，正好看到这一部分，解惑了

[Iriszps]

真的太谢谢了，我也解惑了不少

[ajiao05240625]

[Quote=引用 10 楼 l7331014 的回复:]
引用 9 楼 beyond_jhf 的回复:
还是不明白为什么写成4032=63*64的形式，请兄台明示了。。


看编程手册的第5章性能指南的5.1.2.1全局存储器这节，其中提到了存取器事务(一次访问宽度，如不对齐的话
估计访问宽度就浪费了)大小为32字节(1.2能力及以上),64字节，128字节。

LZ例子中，pitch最小为实际存储需要的1000*4=4000字节。
如是……
[/Quote]

那我可不可以自己设定合并访问的宽度呢？
还有，如果既是32的倍数，也是64的倍数，按那个宽度访问呢？