首页 > 代码库 > cuda数组的拷贝

cuda数组的拷贝

原文链接

简单描述一下自己要做的事情:(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-width16     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("\n");63         }64             printf("\n");65             printf("\n");66     }67 68 69     //释放空间70     cudaFree(&(datastruct.Vec3D1));71     cudaFree(&(datastruct.Vec3D2));72 73 }

 

cuda数组的拷贝