私は次のように簡単なCUDAカーネルを書きました：
__global__ void cudaDoSomethingInSharedMemory(float* globalArray, pitch){
__shared__ float sharedInputArray[1088];
__shared__ float sharedOutputArray[1088];
int tid = threadIdx.x //Use 1D block
int rowIdx = blockIdx.x //Use 1D grid
int rowOffset = pitch/sizeof(float);//Offset in elements (not in bytes)
//Copy data from global memory to shared memory (checked)
while(tid < 1088){
sharedInputArray[tid] = *(((float*) globalArray) + rowIdx*rowOffset + tid);
tid += blockDim.x;
__syncthreads();
}
__syncthreads();
//Do something (already simplified and the problem still exists)
tid = threadIdx.x;
while(tid < 1088){
if(tid%2==1){
if(tid == 1087){
sharedOutputArray[tid/2 + 544] = 321;
}
else{
sharedOutputArray[tid/2 + 544] = 321;
}
}
tid += blockDim.x;
__syncthreads();
}
tid = threadIdx.x;
while(tid < 1088){
if(tid%2==0){
if(tid==0){
sharedOutputArray[tid/2] = 123;
}
else{
sharedOutputArray[tid/2] = 123;
}
}
tid += blockDim.x;
__syncthreads();
}
__syncthreads();
//Copy data from shared memory back to global memory (and add read-back for test)
float temp = -456;
tid = threadIdx.x;
while(tid < 1088){
*(((float*) globalArray) + rowIdx*rowOffset + tid) = sharedOutputArray[tid];
temp = *(((float*) globalArray) + rowIdx*rowOffset + tid);//(1*) Errors are found.
__syncthreads();
tid += blockDim.x;
}
__syncthreads();
}
コードは、「sharedOutputArray」を「interlaced」から「clustered」に変更することです。「123 321 123 321 ... 123 321」は、「123 123 123 .. 123 321 321 321 ... 321」に変更され、クラスター化されたものを出力します。結果はグローバルメモリ配列「globalArray」になります。「globalArray」は「cudaMallocPitch（）」によって割り当てられます
このカーネルは、2D配列を処理するために使用されます。考え方は単純です。1行に1ブロック（つまり、1Dグリッドとブロック数は行数に等しい）、各行にNスレッドです。行番号は1920、列番号は1088です。したがって、1920ブロックがあります。
問題は、N（1ブロック内のスレッド数）が64、128、または256の場合、すべてが正常に機能する（少なくとも機能しているように見える）ことです。ただし、Nが512の場合（CUDA計算機能2.0でGTX570を使用しており、1ブロックの各次元の最大サイズは1024です）、エラーが発生しました。
エラーは次のとおりです。位置256から287までのグローバルメモリ内の行の要素（それぞれが4バイトの浮動小数点数）（インデックスは0から始まり、エラーストリップの長さは32要素、128ビット）は123ではなく0です。 。「123123123...0 0 0 0 0 ...0123123...」のように見えます。上記の行（1 *）を確認したところ、これらの要素は「sharedOutputArray」で123であり、要素（たとえば、tid == 270）が（1 *）で読み込まれると、「temp」は0を示しました。「tid」を確認しようとしました==255"および"tid== 288 "で、要素は123（corrent）でした。このタイプのエラーは、ほぼすべての1920行で発生しました。
スレッドを「同期」（おそらくすでに過剰同期）しようとしましたが、機能しませんでした。私を混乱させているのは、64、128、または256スレッドが正常に機能したのに、512が機能しなかった理由です。512スレッドを使用すると、パフォーマンスが最適化されない可能性があることを知っています。どこでミスをしたのかを知りたいだけです。
前もって感謝します。