2

私は CUDA コードを書いており、GForce 9500 GT グラフィック カードを使用しています。

20000000 整数要素の配列を処理しようとしていますが、使用しているスレッド番号は 256 です

ワープ サイズは 32 です。計算能力は 1.1 です。

これはハードウェアですhttp://www.geforce.com/hardware/desktop-gpus/geforce-9500-gt/specifications

今、ブロック番号 = 20000000/256 = 78125 ?

この音は正しくありません。ブロック数の計算方法を教えてください。どんな助けでも大歓迎です。

私のCUDAカーネル関数は次のとおりです。アイデアは、各ブロックがその合計を計算し、次に各ブロックの合計を加算することによって最終的な合計が計算されるというものです。

__global__ static void calculateSum(int * num, int * result, int DATA_SIZE)
{
    extern __shared__ int shared[];
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;

    shared[tid] = 0;
    for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
        shared[tid] += num[i];
    }

    __syncthreads();
    int offset = THREAD_NUM / 2;
    while (offset > 0) {
        if (tid < offset) {
            shared[tid] += shared[tid + offset];
        }
        offset >>= 1;
        __syncthreads();
    }

    if (tid == 0) {
        result[bid] = shared[0];

    }
}

そして、私はこの関数を

calculateSum <<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>> (gpuarray, result, size);

THREAD_NUM = 256 で、GPU 配列のサイズは 20000000 です。

ここでは、ブロック番号を 16 として使用していますが、それが正しいかどうかはわかりません。最大の並列処理が達成されるようにするにはどうすればよいですか?

これが私の CUDA Occupancy Calculator の出力です。ブロック番号が 8 の場合、占有率は 100% になります。つまり、ブロック番号が 8 でスレッド番号が 256 の場合に最大の効率が得られるということです。あれは正しいですか？

CUDA占有計算 ありがとう

4

4 に答える 4

3

各スレッドが1つの要素を処理し、各ブロックに256スレッドがある場合、20000000スレッドを実行する必要があり、結果として正確に78125ブロックになります。これは完全に有効な番号です。

ただし、少し問題があります。CC1.1デバイスは手元にありませんが、CC1.3では次のようになります。

Maximum sizes of each dimension of a grid:     65535 x 65535 x 1

したがって、データのさまざまな部分に対してカーネルを数回実行するか、2Dグリッドを作成して、スレッドの2Dアドレスを配列要素の1Dアドレスに簡単に変換する必要があります。

于 2012-05-02T12:17:26.383 に答える
2

あなたの場合、スレッドの総数（20000000）は、ブロックあたりのスレッド数（256）で均等に除算されるため、その数（78125）を使用できます。数値が均等に除算されない場合、通常の整数除算はそれを切り捨て、最終的に必要なスレッド数よりも少なくなります。したがって、その場合は、除算の結果を次のような関数で切り上げる必要があります。

int DivUp(int a, int b) {
  return ((a % b) != 0) ? (a / b + 1) : (a / b);
}

この関数は要素よりも多くのスレッドを提供する可能性があるため、最後のいくつかのスレッドで計算を中止するために、カーネルにテストを追加する必要もあります。

int i(blockIdx.x * blockDim.x + threadIdx.x);
if (i >= n_items) {
  return;
}

ただし、追加の障害があります。ハードウェアは、グリッドの各次元で最大65535ブロックに制限され、2つの次元（xとy）に制限されます。したがって、DivUp（）を使用した後、それよりも高いカウントになる場合は、2つのオプションがあります。ワークロードを分割してカーネルを複数回実行するか、2次元を使用することができます。

2次元を使用するには、それぞれがハードウェア制限よりも低く、乗算すると、必要な実際のブロック数になる2つの数値を選択します。次に、カーネルの上部にコードを追加して、2つの次元（xとy）を1つのインデックスに結合します。

于 2012-05-02T12:35:44.017 に答える
2

投稿したカーネル コードは、起動するブロックの数に関係なく、任意の入力データ サイズを処理できます。選択は単にパフォーマンスに依存する必要があります。

経験則として、この種のカーネルでは、1 つのマルチプロセッサで同時に実行されるブロックの数に、カード上のマルチプロセッサの数を掛けた数が必要です。最初の数値は CUDA ツールキットに同梱されている CUDA 占有スプレッドシートを使用して取得できますが、上限はマルチプロセッサあたり 8 ブロックになり、2 番目の数値はお使いのデバイスでは 4 になります。これは、可能な限り最大の並列処理を実現するために必要なブロックが 32 個を超えないことを意味しますが、正確に答えるには、現在持っていないコンパイラにアクセスする必要があります。

また、ベンチマークを使用して、4、8、12、16、20、24、28、または 32 ブロック (カード上のマルチプロセッサの数であるため、4 の倍数) のいずれかを使用して、最適なブロック数を実験的に決定することもできます。

于 2012-05-02T17:50:35.143 に答える
1

カーネルでグリッドの x 次元のみを使用しています。したがって、cc 1.1 を使用すると 65535 ブロックに制限されます。

20000000/256 = 78125 は正しいです!

したがって、間違いなく1ブロック以上必要です。

カーネル：

//get unique block index
const unsigned int blockId = blockIdx.x //1D
    + blockIdx.y * gridDim.x //2D

//terminate unnecessary blocks
if(blockId >= 78124)
    return;

//... rest of kernel

最も単純なアプローチは、2 つの y ブロックを使用し、カーネルでブロック ID をチェックすることです。

dim3 gridDim = dim3(65535, 2);

これにより、52945 ブロック以上が役に立たなくなります。オーバーヘッドが何であるかはわかりませんが、最初に x 次元、次に y および z 次元を埋めると、特に z 次元に達すると、非常に多くの未使用ブロックが作成される可能性があります!

(Nvidia は、ここでのケースのように、カーネル内の一意のブロックの使用に最適なグリッドの使用を取得するユーティリティ関数を明確に提供する必要があります)

この簡単な例では、x と y を使用してルートを計算してみませんか。

grid(280, 280) = 78400 blocks //only 275 blocks overhead, less is not possible

これは、コンピューティング機能 3.0 の大きな利点の 1 つです。各ブロックの 32 ビット範囲により、多くの場合、作業が楽になります。なぜ 65535 に限定されたのか、私には理解できませんでした。

しかし、私は依然として下位互換性を好みます。

@talonmies のバリエーションもテストします。

于 2012-05-03T00:37:44.873 に答える