0

自分でCUDAを使用してベクトル和の削減を実装しようとしましたが、修正できるエラーが発生しましたが、実際の問題が何であるかがわかりませんでした。

以下のカーネルを実装しました。これは、NVIDIAのサンプルで使用されているものとほとんど同じです。

__global__ 
void reduce0(int *input, int *output)
{
    extern __shared__ int s_data[];

    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    s_data[tid] = input[i];
    __syncthreads();

    for( int s=1; s < blockDim.x; s *= 2) {
        if((tid % 2*s) == 0) {
            s_data[tid] += s_data[tid + s];
        }

        __syncthreads();
    }

    if(tid == 0) {
        output[blockIdx.x] = s_data[0];
    }
}

さらに、ホスト側で以下のように共有メモリスペースを計算しました

int sharedMemSize = numberOfValues * sizeof(int);

使用されているスレッドのブロックが複数ある場合、コードは正常に実行されます。1つのブロックのみを使用すると、上記の範囲外のインデックスエラーで終了します。ホストコードを次の行を見つけた例の1つと比較して、エラーを探します。

int smemSize = (threads <= 32) ? 2 * threads * sizeof(T) : threads * sizeof(T);

ブロック/グリッドの設定を少し試してみると、次の結果が得られました。

  1. ブロック、任意の数のスレッド=>コードがクラッシュする
  2. > 2ブロック、任意の数のスレッド=>コードは正常に実行されます
  3. 1ブロック、任意の数のスレッド、共有メモリサイズ2 *＃threads=>コードは正常に実行されます

これについては数時間考えていますが、使用するスレッドまたはブロックの数が少なすぎると、範囲外のエラーが発生する理由がわかりません。

更新：要求に応じてカーネルを呼び出すホストコード

int numberOfValues = 1024 ;
int numberOfThreadsPerBlock = 32;
int numberOfBlocks = numberOfValues / numberOfThreadsPerBlock;

int memSize = sizeof(int) * numberOfValues;

int *values = (int *) malloc(memSize);
int *result = (int *) malloc(memSize);

int *values_device, *result_device;
cudaMalloc((void **) &values_device, memSize);
cudaMalloc((void **) &result_device, memSize);

for(int i=0; i < numberOfValues ; i++) {
    values[i] = i+1;
}

cudaMemcpy(values_device, values, memSize, cudaMemcpyHostToDevice);

dim3 dimGrid(numberOfBlocks,1);
dim3 dimBlock(numberOfThreadsPerBlock,1);
int sharedMemSize = numberOfThreadsPerBlock * sizeof(int);

reduce0 <<< dimGrid, dimBlock, sharedMemSize >>>(values_device, result_device);

if (cudaSuccess != cudaGetLastError())
        printf( "Error!\n" );

cudaMemcpy(result, result_device, memSize, cudaMemcpyDeviceToHost);
4

1 に答える 1

1

問題は、モジュロと乗算の優先順位である可能性があります。 tid % 2*s等しいが(tid % s)*2、あなたが望むtid % (s*2)

int smemSize = (threads <= 32) ? 2 * threads * sizeof(T) : threads * sizeof(T)少数のスレッドに使用する必要がある理由は、範囲外のインデックス作成が原因です。これが発生する1つの例は、29スレッドを起動する場合です。いつtid=28とs=2分岐が発生し28 % (2*2) == 0、インデックスを作成しs_data[28+2]ますが、共有メモリは29スレッドにしか割り当てられていません。

于 2013-01-04T10:35:05.057 に答える