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 は、dim Ndim*Nbins の Vector を新しいコマンドで作成するためのクラスです (そのデストラクタで、対応する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]

スレッドの数は、MultiProcessors の制限と互換性があります: 最大共有メモリ、スレッドごとの最大レジスタ、および 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 に答える