cuda数组

http://blog.csdn.net/xfortius/article/details/8998534

刚刚看了一些cuda的东西，打算写个程序，结果遇到一堆问题。首先就是host和device上的数组传递问题，被搞得有点晕，看了一些资料，在这里总结如下。

1：问题是怎么来的

在device上要用一维数组、二维数组、三维数组，对于一维数组用了cudaMalloc和cudaMemcpy进行内存分配和赋值，但是对于二维和三维的分配本想这样转换成一维进行。但是这样感觉赋值又不方便，刚刚看了一个例子，这样做的：

[cpp] view plain copy

1. arr[n][n];  
2.   
3. cudaMalloc((void**) &dst, sizeof(float) * n * n);   
4.   
5. cudaMemcpy2D(dst, sizeof(float) * n, arr, sizeof(float) * n,  
6.         sizeof(float) * n, n, cudaMemcpyHostToDevice);  

       这里先在c上面定义一个二维数组arr[n][n]，二维数组是连续的，相当于在CPU上申请了一片连续的空间。后面要在GPU申请空间了，并且要把CPU的值复制到GPU中去。那么现在GPU上用cudaMalloc申请空间，注意这里申请的空间也是连续的。这样可以用把值从CPU复制到GPU了，这里通过cudaMemcpy2D来进行赋值。在这个函数中最重要的是参数pitch这个概念，这里相当于图像中的补齐概念。但是还是有点不同，在图像中补齐是必须的，但是在这里分配的数据中补齐与否是看自己决定，可以补齐也可以不补齐。

      在上面这个例子中，明显是没有补齐的。但是因为在GPU上分配的空间是连续的，所以这里也不需要补齐而直接用连续空间。但是这样做的坏处是什么呢？在GPU中对以256的整数倍开头的访问比较快速，所以一般都是补齐成256的整数倍。如果不进行补齐，访问的时候就可能慢些了。

2：对于二维数组赋值与访问

其实先是不清楚上面的概念的，直接按照网上一般的通用写法如下：

[cpp] view plain copy

1. cutilSafeCall(cudaMallocPitch((void**)&GPU_InputData->m_prData1,&pitch,width*sizeof(float),height));  
2. cutilSafeCall(cudaMemcpy2D(GPU_InputData->m_prData1,pitch,CPU_InputData->m_prData1,width,width,height,cudaMemcpyHostToDevice));  

其中第一句是分配空间，第二句是复制参数。那么下面在分析一下具体过程。

对于第一句，是分配二维数组。数组的长宽为width*sizeof(float)与height，这里为了快速范围所以进行了补齐。补齐过后行的大小为pitch，这里非常类似图像里面的概念。

对于第二句，是复制参数。把CPU的数据复制到GPU中，因为在GPU中有补齐行为，所以在GPU中的空间有可能是不完全填满的，也即GPU中空间不是连续的。所以在数据访问的时候要注意。

3：对于三维数组赋值与访问

[cpp] view plain copy

1. cudaExtent extent=make_cudaExtent(Col,Row,8); //cuda上数组的大小  
2. cudaMalloc3D(GPU_InputData->InitDeformPr,extent); //分配空间  
3. cudaMemcpy3DParms HostToDev={0}; //host与device间相互复制的参数，先赋为0  
4. //host上面的参数,将普通指针封装成cudaPitchedPtr类型指针  
5. HostToDev.srcPtr=make_cudaPitchedPtr((void*)CPU_InputData.InitDeformPr,Col*sizeof(float),Col,Row);    
6. HostToDev.dstPtr=GPU_InputData->InitDeformPr;  
7. HostToDev.extent=extent;  
8. HostToDev.kind=cudaMemcpyHostToDevice;  
9. cutilSafeCall(cudaMemcpy3D(&HostToDev)); //cudaMemcpy3D就是在host和device上进行三维数据传输，具体数据和方向由其参数cudaMemcpy3DParas决定  

这里和二维的其实是一样的，只不过三维的参数复杂一点，所以引入了一些辅助参数，比如关于大小的extent，关于三维复制参数的cudaMemcpy3DParms来进行设置。具体的不再细说。

4：关于cuda数组

在网上搜索的时候，看到了也有用cuda数组的，也就是cudaArray。看了一些资料后才知道，其实cudaArray本身就是为了cuda中的纹理存储器设计的。其就是为了把数据和纹理存储器相联系用于加速访问速度。其不论分配还是访问过程都和前面的类似，稍微不同的就是一些具体的参数设置。但是内存分配和赋值、访问的思想是一样的。

5：常用数组形式对比

常用的数组形式，对于我来说就是三种具体表现为：p[i][j][k]、cudaPitchedPtr、cudaArray。其中p[i][j][k]这种模式的是一般CPU中用的，其本身是连续分配，在cuda中也可以用，但是因为没有对齐，所以访问效率可能比较底下。cudaPitchedPtr是cuda中常用的吧，其进行了对齐。而cudaArray是cuda中为了利用纹理内存而引入的。

好了暂时遇到的问题就着一些，后面要加快进程了。。加油。。