0

基本的な削減によってベクトル累積を実行できる単純な CUDA カーネルがあります。複数のブロックに分割することで、より大きなデータを処理できるようにスケールアップしています。ただし、カーネルが使用する適切な量の共有メモリを割り当てるという私の仮定は、不正なメモリ アクセスで失敗しています。この制限を増やすと消えますが、その理由を知りたいです。これが私が話しているコードです：

コアカーネル:

    __global__ static
    void vec_add(int *buffer,
               int numElem,    //  The actual number of elements
               int numIntermediates)   //  The next power of two of numElem
    {
        extern __shared__ unsigned int interim[];

        int index = blockDim.x * blockIdx.x + threadIdx.x;

        //  Copy global intermediate values into shared memory.
        interim[threadIdx.x] =
          (index < numElem) ? buffer[index] : 0;

        __syncthreads();

        //  numIntermediates2 *must* be a power of two!
        for (unsigned int s = numIntermediates / 2; s > 0; s >>= 1) {
            if (threadIdx.x < s) {
                interim[threadIdx.x] += interim[threadIdx.x + s];
            }
            __syncthreads();
        }

        if (threadIdx.x == 0) {
            buffer[blockIdx.x] = interim[0];
        }
    }

そして、これは呼び出し元です：

void accumulate (int* buffer, int numElem)
{
    unsigned int numReductionThreads =
      nextPowerOfTwo(numElem); // A routine to return the next higher power of 2.

    const unsigned int maxThreadsPerBlock = 1024;  // deviceProp.maxThreadsPerBlock

    unsigned int numThreadsPerBlock, numReductionBlocks, reductionBlockSharedDataSize;

    while (numReductionThreads > 1) {

        numThreadsPerBlock = numReductionThreads < maxThreadsPerBlock ?           
            numReductionThreads : maxThreadsPerBlock;

        numReductionBlocks = (numReductionThreads + numThreadsPerBlock - 1) / numThreadsPerBlock;

        reductionBlockSharedDataSize = numThreadsPerBlock * sizeof(unsigned int);

        vec_add <<< numReductionBlocks, numThreadsPerBlock, reductionBlockSharedDataSize >>>
            (buffer, numElem, numReductionThreads);

        numReductionThreads = nextPowerOfTwo(numReductionBlocks);
    }

}

次の構成の GPU で 1152 要素のサンプル セットを使用してこのコードを試しました: タイプ: Quadro 600 MaxThreadsPerBlock: 1024 MaxSharedMemory: 48KB

出力：

Loop 1: numElem = 1152, numReductionThreads = 2048, numReductionBlocks = 2, numThreadsPerBlock = 1024, reductionBlockSharedDataSize = 4096
Loop 2: numElem = 1152, numReductionThreads = 2, numReductionBlocks = 1, numThreadsPerBlock = 2, reductionBlockSharedDataSize = 8
CUDA Error 77: an illegal memory access was encountered

「暫定的な」共有メモリが不正なメモリ アクセスを引き起こしているのではないかと疑い、次の行で共有メモリを勝手に 2 倍に増やしました。

reductionBlockSharedDataSize = 2 * numThreadsPerBlock * sizeof(unsigned int);

そして、カーネルが正常に動作し始めました!

私が理解していないのは、問題を解決するためにこの追加の共有メモリを提供しなければならなかった理由です (一時的に)。

このマジック ナンバーを確認するためのさらなる実験として、6912 ポイントのはるかに大きなデータセットを使用してコードを実行しました。今回は、2X や 4X でも役に立ちませんでした。

Loop 1: numElem = 6912, numReductionThreads = 8192, numReductionBlocks = 8, numThreadsPerBlock = 1024, reductionBlockSharedDataSize = 16384

Loop 2: numElem = 6912, numReductionThreads = 8, numReductionBlocks = 1, numThreadsPerBlock = 8, reductionBlockSharedDataSize = 128
CUDA Error 77: an illegal memory access was encountered

しかし、共有メモリのサイズを 8 倍に増やすと、問題は再び解消されました。

もちろん、48KB の共有メモリの制限をすぐに使い果たしてしまうため、ますます大きなデータ セットに対してこの倍率を任意に選択することはできません。だから私は自分の問題を解決する正当な方法を知りたい.

4

1 に答える 1

3

インデックス外アクセスを指摘してくれた @havogt に感謝します。問題は、vec_add メソッドの numIntermediates として間違った引数を使用していたことです。意図は、カーネルがスレッド数とまったく同じ数のデータ ポイントで動作することでした。これは常に 1024 である必要があります。numThreadsPerBlock を引数として使用して修正しました。

vec_add <<< numReductionBlocks, numThreadsPerBlock, reductionBlockSharedDataSize >>>
        (buffer, numElem, numThreadsPerBlock);
于 2016-11-13T22:57:49.837 に答える