简单描述一下自己要做的事情：（1）CPU三维的vector--->（2）CPU三维数组--->（3）转换到GPU中的三维数组--->（4）转换到CPU中的三维数组，而其中问题主要出在第3、4步。
主要是没有理解一个问题，那就是“cuda的各种拷贝一定要是内存连续的”。而自己在申请三维数组的时候用的是new或者malloc，这种在申请一维数组的时候是连续的，但是在申请多维数组就会出现不连续，因此在这里犯了致命错误。
http://hpcbbs.it168.com/thread-7366-1-1.html这个帖子给了很好的建议，“vector<vector<float> > 并不是二维数组吧，它只是实现了二维数组的操作（比如[][]).内存是不连续的。要用cudaMemcpy还是得定义 float 2darray[N][M] 或者 直接 float *2darray = new float(M*N);”。反正就是这样，纸上得来终觉浅，自己多亲身力为一下。
1 #include "example1.cuh" 2 #include "Struct.h" 3 /************************************************************************/ 4 /* 转换成设备可以识别的 */ 5 /************************************************************************/ 6 void InitCPUData(DataMatrix &datamatrix,std::vector<std::vector<std::vector<float > > > vec3D1, 7 std::vector<std::vector<std::vector<float > > > vec3D2,int width,int height,int depth) 8 { 9 int i,j,k; 10 for (i=0;i<depth;i++) 11 { 12 for (j=0;j<height;j++) 13 { 14 for (k=0;k<width;k++) 15 { 16 datamatrix.Mat3D1[i][j][k]=vec3D1[i][j][k]; 17 datamatrix.Mat3D2[i][j][k]=vec3D2[i][j][k]; 18 } 19 } 20 } 21 } 22 23 /************************************************************************/ 24 /* 分配并且赋值 */ 25 /************************************************************************/ 26 __host__ void AllocDataAndVal(DataStruct &datastruct,DataMatrix datamatrix,int width,int height,int depth) 27 { 28 //分配内存 29 cudaExtent extent=make_cudaExtent(sizeof(float)*6,7,8); 30 cutilSafeCall(cudaMalloc3D(&(datastruct.Vec3D1),extent)); 31 cutilSafeCall(cudaMalloc3D(&(datastruct.Vec3D2),extent)); 32 //赋值 33 cudaMemcpy3DParms Parms3D1={0}; 34 cudaMemcpy3DParms Parms3D2={0}; 35 Parms3D1.dstPtr=datastruct.Vec3D1; 36 Parms3D2.dstPtr=datastruct.Vec3D2; 37 Parms3D1.srcPtr=make_cudaPitchedPtr((void*)datamatrix.Mat3D1,width*sizeof(float),width,height); 38 Parms3D2.srcPtr=make_cudaPitchedPtr((void*)datamatrix.Mat3D2,width*sizeof(float),width,height); 39 Parms3D1.extent=extent; 40 Parms3D2.extent=extent; 41 Parms3D1.kind=cudaMemcpyHostToDevice; 42 Parms3D2.kind=cudaMemcpyHostToDevice; 43 cudaMemcpy3D(&Parms3D1); 44 cudaMemcpy3D(&Parms3D2); 45 } 46 47 48 /************************************************************************/ 49 /* 核函数 */ 50 /************************************************************************/ 51 __global__ void kernel(DataStruct datastruct,int width,int height,int depth) //实现类中两个数组的相加，保持到第一个数组中 52 { 53 char* devPtr1=(char*)datastruct.Vec3D1.ptr; //起始地址 54 char* devPtr2=(char*)datastruct.Vec3D2.ptr; 55 int pitch=datastruct.Vec3D1.pitch; //pitch，相当于宽度 56 int SlicePitch=pitch*height; 57 //用线程 58 int xid=threadIdx.x; 59 int yid=threadIdx.y; 60 int zid=threadIdx.z; 61 if (xid<width&&yid<height&&zid<depth) 62 { 63 ((float*)((char*)(devPtr1+zid*SlicePitch)+yid*pitch))[zid]=((float*)((char*)(devPtr1+zid*SlicePitch)+yid*pitch))[zid]+ 64 ((float*)((char*)(devPtr2+zid*SlicePitch)+yid*pitch))[zid]; 65 } 66 } 67 68 /************************************************************************/ 69 /* 返回到主机上 */ 70 /************************************************************************/ 71 __host__ void GPU2CPU(DataStruct &datastruct,DataMatrix datamatrix, int width,int height,int depth) 72 { 73 cudaExtent extent=make_cudaExtent(sizeof(float)*6,7,8); 74 cudaMemcpy3DParms Parms3D1={0}; 75 cudaMemcpy3DParms Parms3D2={0}; 76 Parms3D1.srcPtr=datastruct.Vec3D1; 77 Parms3D2.srcPtr=datastruct.Vec3D2; 78 Parms3D1.dstPtr=make_cudaPitchedPtr((void*)datamatrix.Mat3D1,width*sizeof(float),width,height); 79 Parms3D2.dstPtr=make_cudaPitchedPtr((void*)datamatrix.Mat3D2,width*sizeof(float),width,height); 80 Parms3D1.extent=extent; 81 Parms3D2.extent=extent; 82 Parms3D1.kind=cudaMemcpyDeviceToHost; 83 Parms3D2.kind=cudaMemcpyDeviceToHost; 84 cudaMemcpy3D(&Parms3D1); 85 cudaMemcpy3D(&Parms3D2); 86 87 }
主函数：
1 // 说明：在cu中host和device的虽然写在一起,但是是分开编译的，这个在一起只是形式上的。如果函数前面有__global__由主机调用设备执行，__device__设备调用设备执行，__host__主机调用主机执行。其分别对应三种形式为核函数、核函数中的函数、一般函数。 2 3 #include <iostream> 4 #include <vector> 5 #include <algorithm> 6 #include "example1.cuh" 7 #include "Struct.h" 8 9 int main() 10 { 11 int i,j,k; 12 int width=6; 13 int height=7; 14 int depth=8; 15 std::vector<std::vector<std::vector<float > > > vec3D1(width); //建立6*7*8的三维数组，范文depth-height-width 16 std::vector<std::vector<std::vector<float > > > vec3D2(width); 17 18 vec3D1.resize(depth); 19 vec3D2.resize(depth); 20 for (i=0;i<depth;i++) 21 { 22 vec3D1[i].resize(height); 23 vec3D2[i].resize(height); 24 for (j=0;j<height;j++) 25 { 26 vec3D1[i][j].resize(width); 27 vec3D2[i][j].resize(width); 28 for (k=0;k<width;k++) 29 { 30 vec3D1[i][j][k]=i+j+k; 31 vec3D2[i][j][k]=i*j*k; 32 } 33 } 34 } 35 36 ////////////////////////////////////////////////////////////////////////// 37 //将数据转换成设备可以接受的形式，为赋值做准备，这个是在主机上进行 38 DataMatrix datamatrix; 39 InitCPUData(datamatrix,vec3D1,vec3D2,width,height,depth); 40 41 ////////////////////////////////////////////////////////////////////////// 42 //给设备分配内存并且赋值，这个是在设备上进行 43 DataStruct datastruct; 44 AllocDataAndVal(datastruct,datamatrix,width,height,depth); 45 46 ////////////////////////////////////////////////////////////////////////// 47 //调用核函数 48 dim3 dimBlock(8,7,6); 49 kernel<<<1,dimBlock>>>(datastruct,width,height,depth); 50 51 ////////////////////////////////////////////////////////////////////////// 52 //返回到主机，并显示出来 53 GPU2CPU(datastruct,datamatrix,width,height,depth); 54 for (i=0;i<depth;i++) 55 { 56 for (j=0;j<height;j++) 57 { 58 for (k=0;k<width;k++) 59 { 60 printf("%f ",datamatrix.Mat3D1[i][j][k]); 61 } 62 printf(" "); 63 } 64 printf(" "); 65 printf(" "); 66 } 67 68 69 //释放空间 70 cudaFree(&(datastruct.Vec3D1)); 71 cudaFree(&(datastruct.Vec3D2)); 72 73 }