0

我正在编写一个 CUDA 内核，我必须在这个设备上执行：

name: GeForce GTX 480
CUDA capability: 2.0
Total global mem:  1610285056
Total constant Mem:  65536
Shared mem per mp:  49152
Registers per mp:  32768
Threads in warp:  32
Max threads per block:  1024
Max thread dimensions:  (1024, 1024, 64)
Max grid dimensions:  (65535, 65535, 65535)

内核的最小形式是： 

_global__ void CUDAvegas( ... )
{
devParameters p;
extern __shared__ double shared[];
int width = Ndim * Nbins;
int ltid = p.lId;
while(ltid < 2* Ndim){
shared[ltid+2*width] = ltid;
ltid += p.lOffset; //offset inside a block
}
__syncthreads();
din2Vec<double> lxp(Ndim, Nbins);

__syncthreads();
for(int i=0; i< Ndim; i++){
  for(int j=0; j< Nbins; j++){
    lxp.v[i][j] = shared[i*Nbins+j];
    }
}
}// end kernel

其中 Ndim=2, Nbins=128, devParameters 是一个类，其方法 p.lId 用于计算本地线程的 id（在块内），而 din2Cec 是一个类，用于使用新命令（在它的析构函数中，我实现了对应的 delete[])。nvcc 输出为：

nvcc -arch=sm_20   --ptxas-options=-v   file.cu -o file.x
ptxas info    : Compiling entry function '_Z9CUDAvegas4LockidiiPdS0_S0_P7sumAccuP17curandStateXORWOWS0_i' for 'sm_20'
ptxas info    : Function properties for  _Z9CUDAvegas4LockidiiPdS0_S0_P7sumAccuP17curandStateXORWOWS0_i
               0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 22 registers, 116 bytes cmem[0], 51200 bytes cmem[2]

线程数与多处理器限制兼容：最大共享内存、每个线程的最大寄存器和 MP 和每个 MP 的扭曲。如果我启动 64 个线程 X 30 个块（每块共享内存为 4128），没关系，但如果使用超过 30 个块，我会收到错误：

cudaCheckError() failed at file.cu:508 : unspecified launch failure
========= Invalid __global__ read of size 8
=========     at 0x000015d0 in CUDAvegas
=========     by thread (0,0,0) in block (1,0,0)
=========     Address 0x200ffb428 is out of bounds

我认为这是分配单线程内存的问题，但我不明白每个 MP 和总块的限制是什么......有人可以帮助我或提醒正确的话题吗？

PS：我知道提供的内核什么都不做，但这只是为了了解我的极限问题。

4

1 回答 1

1

我认为您收到的错误是解释性的。指出存在大小为 8 的数据类型的全局越界读取。负责越界读取的是块 (1,0,0) 中的线程 (0,0,0) . 我怀疑负责任的指令lxp.v[i][j] = shared[i*Nbins+j];在最后一个嵌套for循环中。可能您分配的全局内存量与您启动的块数无关，因此当您启动太多块时，您会收到这样的错误。

于 2013-02-24T21:03:27.513 回答