是否可以为内核(内部或外部)分配共享内存并在从内核调用的其他设备函数中使用它? 对我来说特别有趣的是,如果/如何将它用作返回的参数/数组。
在设备函数中使用共享内存作为输入参数似乎没有问题(至少我没有遇到任何问题、错误或意外结果。
当我将它用作返回参数时,我遇到了几个问题:
我可以运行从调试配置构建的程序。
但我无法调试它 -> 当我使用共享内存时它在设备功能中崩溃
我也得到错误
cuda-memcheck
- >“无效__global__
阅读”
因为地址越界,它从共享地址空间读取
那么是否可以使用共享内存将数组从设备函数返回给内核?
编辑:
我写了一个非常简单的例子来排除我做的其他错误。
#define CUDA_CHECK_RETURN(value) { \
cudaError_t _m_cudaStat = (value); \
if (_m_cudaStat != cudaSuccess) { \
printf( "Error %s at line %d in file %s\n", \
cudaGetErrorString(_m_cudaStat), __LINE__, __FILE__); \
exit(-1); \
} }
__device__ void Function( const int *aInput, volatile int *aOutput )
{
for( int i = 0; i < 10; i++ )
aOutput[i] = aInput[i] * aInput[i];
}
__global__ void Kernel( int *aInOut )
{
__shared__ int aShared[10];
for(int i=0; i<10; i++)
aShared[i] = i+1;
Function( aShared, aInOut );
}
int main( int argc, char** argv )
{
int *hArray = NULL;
int *dArray = NULL;
hArray = ( int* )malloc( 10*sizeof(int) );
CUDA_CHECK_RETURN( cudaMalloc( (void**)&dArray, 10*sizeof(int) ) );
for( int i = 0; i < 10; i++ )
hArray[i] = i+1;
CUDA_CHECK_RETURN( cudaMemcpy( dArray, hArray, 10*sizeof(int), cudaMemcpyHostToDevice ) );
cudaMemcpy( dArray, hArray, 10*sizeof(int), cudaMemcpyHostToDevice );
Kernel<<<1,1>>>( dArray );
CUDA_CHECK_RETURN( cudaMemcpy( hArray, dArray, 10*sizeof(int), cudaMemcpyDeviceToHost ) );
cudaMemcpy( hArray, dArray, 10*sizeof(int), cudaMemcpyDeviceToHost );
free( hArray );
CUDA_CHECK_RETURN( cudaFree( dArray ) );
cudaFree( dArray );
return 0;
}
我用一个线程块执行内核,每个块一个线程。构建程序并运行它是没有问题的。我得到了预期的结果。 但是如果程序用
cuda-memcheck
测试它终止内核并出现以下日志:
Error unspecified launch failure at line 49 in file ../CuTest.cu
========= Invalid __global__ read of size 4
========= at 0x00000078 in /home/strautz/Develop/Software/CuTest/Debug/../CuTest.cu:14:Function(int const *, int volatile *)
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x01000000 is out of bounds
========= Device Frame:/home/strautz/Develop/Software/CuTest/Debug/../CuTest.cu:25:Kernel(int*) (Kernel(int*) : 0xd0)
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/lib/libcuda.so (cuLaunchKernel + 0x34b) [0x55d0b]
========= Host Frame:/usr/lib/libcudart.so.5.0 [0x8f6a]
=========
========= Program hit error 4 on CUDA API call to cudaMemcpy
========= Saved host backtrace up to driver entry point at error
========= Host Frame:/usr/lib/libcuda.so [0x24e129]
========= Host Frame:/usr/lib/libcudart.so.5.0 (cudaMemcpy + 0x2bc) [0x3772c]
========= Host Frame:[0x5400000]
=========
========= ERROR SUMMARY: 2 errors
共享内存是否必须对齐,我是否必须做其他事情或者可以忽略它——不这么认为吗?
正如这里所描述的,这只是一个驱动程序问题。在我更新到当前版本后,一切正常。
查看CUDA 5.0安装文件/usr/local/cuda-5.0/samples/6_Advanced/reduction/doc/reduction.ppt
sdata
是设备函数warpReduce()
的局部变量。它存储共享内存的地址。共享内存可以由设备函数中的地址读取/写入。然后从外部的共享内存中读取最终的归约结果warpReduce()
template <unsigned int blockSize>
__device__ void warpReduce(volatile int *sdata, unsigned int tid) {
if (blockSize >= 64) sdata[tid] += sdata[tid + 32];
if (blockSize >= 32) sdata[tid] += sdata[tid + 16];
if (blockSize >= 16) sdata[tid] += sdata[tid + 8];
if (blockSize >= 8) sdata[tid] += sdata[tid + 4];
if (blockSize >= 4) sdata[tid] += sdata[tid + 2];
if (blockSize >= 2) sdata[tid] += sdata[tid + 1];
}
template <unsigned int blockSize>
__global__ void reduce6(int *g_idata, int *g_odata, unsigned int n) {
extern __shared__ int sdata[];
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockSize*2) + tid;
unsigned int gridSize = blockSize*2*gridDim.x;
sdata[tid] = 0;
while (i < n) { sdata[tid] += g_idata[i] + g_idata[i+blockSize]; i += gridSize; }
__syncthreads();
if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
if (blockSize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); }
if (tid < 32) warpReduce(sdata, tid);
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}