【代码学习】关于数组和核函数输入参数的问题

有人在论坛提交了一个问题：

楼主编写了一个核函数A和输入数据缓冲区p1，p1为全局内存，采用如下方式定义： cufftComplex * p1; 并用cudaMalloc函数为缓冲区分配了一片显存空间。 cudaMalloc((void**)&p1, sizeof(cufftComplex)*L1); 在核函数A中，对p1做了一个简单的赋值。 __global__ void A(cufftComplex * input_buf,cufftComplex * Out_Buf) { const unsigned int id=(blockIdx.x*blockDim.x+threadIdx.x); Out_Buf[id].x=input_buf[id].x; __syncthreads(); } 运算结果正确！ 后来楼主又想：每次调用A函数的时候，都要输入一次输入参数p1，而且是从host拷贝到device。而p1是设备端的内存，按说GPU线程是认识的，不用作为输入参数，少一个输入参数没准可以提高运行速度。于是改成如下方式： __global__ void A(cufftComplex * Out_Buf) { const unsigned int id=(blockIdx.x*blockDim.x+threadIdx.x); Out_Buf[id].x=p1[id].x; __syncthreads(); } 编译报错，说不认识p1。 再将p1的定义改为： __device__ cufftComplex * p1; 编译通过了，可是计算结果为全0，不对了。

我们工程师给与的回复：

回复：首先需要说明的是，直接传递参数，只是一个指针（大小为4B或者8B，根据你的32-bit或者64-bit的当前针对的架构不同）。此指针参数将被放置到constant cache中的，因此不会像你想象的那样提高速度的。 （实际上，给kernel传递一个指针，然后用此指针访存是常见做法的，楼主你想多了...您那样做不能提高速度。。。） 其次，您的形式为__device__ cufftComplex * p1;不能指针使用（因为它没有被正确的初始化）。 您的确可以保留您的__device__的p1指针不动的，但是您需要在host上分配一个同样的host_p1指针（用cudaMalloc()), 然后再用cudaMemcpy将此host上指针的值赋值给您写的那个__device__上的p1，它才能正确的使用。 反之，试图直接在host上使用您的这个p1进行赋值（例如cudaMalloc)， 或者不赋值直接使用，均是不对的。 最后，您的kernel中的__syncthreads()没有存在的必要，可以去掉。 以及，计算结果全0很多时候代表您的kernel挂掉了，后续的复制直接没有进行。您可以通过同步的cudaMemcpy或者cudaDeviceSynchronize()的返回值来判定这一点。

提问者回复：

按照版主的方法，终于将device端数组用起来了，并比较了核函数输入指针参数和直接使用device端数组的运行效率： 1：结论：使用核函数输入指针参数（该参数其实为host端可见的,cudamalloc的指针）比在核函数内直接使用设备端数组还快百分之几，所以，以后还是老老实实用指针参数吧。。。 2：带device前缀的，设备端数组应该用cudaMemcpyToSymbol来赋值（注意必须用cudaMemcpyToSymbol，用cudaMemcpy的话还是会崩溃，运算结果全0），具体代码如下： __device__ cufftComplex * p1; cufftComplex * p1_host; cudaMalloc((void**)&p1_host, sizeof(cufftComplex)*L1); cudaMemcpyToSymbol(p1,&p1_host;,sizeof(cufftComplex *),0,cudaMemcpyHostToDevice); 之后，可以在核函数A里面直接使用p1了，不过，运行了上千万次加法的对比测试结果， 方法一：__global__ void A() { const unsigned int id=(blockIdx.x*blockDim.x+threadIdx.x); p1[id].x=p1[id].x+10; } A<<<B,T,0>>>(); 比 方法二：__global__ void A(cufftComplex * Out_Buf) { const unsigned int id=(blockIdx.x*blockDim.x+threadIdx.x); Out_Buf[id].x=Out_Buf[id].x+10; } A<<<B,T,0>>>(p1_host); 还要慢上个百分之几，具体数值： 方法一：直接用device端数组： 3637ms 3564ms 3719ms 3688ms 3647ms 3677ms 3519ms 3599ms 方法二：核函数输入指针参数 3374ms 3504ms 3420ms 3565ms

本案例中，用cudaMemcpy显然是可以的，但需要获得一下该symbol的address. 实际上，cudaMemcpyToSymbol()是cudaGetSymbolAdress() + cudaMemcpy()的合体。