新手求教，关于cudaMallocPitch的问题

[redsmoke007]

假如要使一个使2D矩阵内每一个元素都加1。如果我用cudaMallocPitch函数（不用cudaMalloc）来分配线性存储区。
我如何在核函数中通过threadIdx变量和blockIdx变量来索引矩阵元素。
我看cuda Programming Guide中的那个例子似乎是在一个线程中就遍历了矩阵。如果我仅仅想一个线程索引一个矩阵元素呢？

新手求指教！

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void myKernel(float *ptr,size_t pitch,int width,int height)
{
	int x=threadIdx.x;
	int y=blockIdx.x;
	int offset=x+y*blockDim.x;
	if (offset<width*height)
	{
		//如何通过offset来索引线性存储器ptr，从而使矩阵内的每一个元素加1。
	}
}

int main(void)
{
	float myMatrix[2000][1500];
	int width=2000;
	int height=1500;
	for (int i=0;i<width;i++)
	{
		for(int j=0;j<height;j++)
		{
			myMatrix[i][j]=i+j;
		}
	}
	float *dev_mat;
	size_t pitch;
	cudaMallocPitch(&dev_mat,&pitch,width*sizeof(float),height);
	cudaMemcpy(dev_mat,myMatrix,width*height*sizeof(float),cudaMemcpyHostToDevice);
	myKernel<<<(2000*1500+256-1)/256,256>>>(dev_mat,pitch,width,height);
	cudaMemcpy(myMatrix,dev_mat,width*height*sizeof(float),cudaMemcpyDeviceToHost);
	cudaFree(dev_mat);
	return 0;

}

[wszlb1987]

同问。。。。

[JH_kanwairen]

假如要使一个使2D矩阵内每一个元素都加1。如果我用cudaMallocPitch函数（不用cudaMalloc）来分配线性存储区。
我如何在核函数中通过threadIdx变量和blockIdx变量来索引矩阵元素。

对于这个问题 我的理解是用PItch方式分配的只是为了方便对齐，并不影响threadIdx和blockIdx变量对矩阵元素的索引。

[cai_niao_yi_zhi]

cudaMallocPitch(void**,int*,widthInBytes,height):这个函数是在线性内存中分配二维数组，因此在使用时，还是用一维的方式使用。注意，width的单位是字节，而height单位是数据类型，而第二个参数的含义下面一段话说明了。 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(元素)；这样的话，复制内容时也要作相应的修改。