2

CUDA を使用して、潜在的に大規模な 3D データ セットの計算を行っています。最初に短いコード スニペットを表示するのが最善だと思います。

void launch_kernel(/*arguments . . . */){
    int bx = xend-xstart, by = yend-ystart, bz = zend-zstart;

    dim3 blocks(/*dimensions*/);
    dim3 threads(/*dimensions*/);
    kernel<<blocks, threads>>();
}

セルの 3D セットがあり、それぞれを計算するためにカーネルを起動する必要があります。問題は、入力サイズが GPU、特にスレッドの能力を超える可能性があることです。したがって、次のようにコードします。

void launch_kernel(/*arguments . . . */){
       int bx = xend-xstart, by = yend-ystart, bz = zend-zstart;

       dim3 blocks(bx,by,1);
       dim3 threads(bz);
       kernel<<blocks, threads>>();
   }

...うまくいきません。寸法が 1000x1000x1000 の場合はどうでしょうか。- ブロックごとに 1000 スレッドを起動できません。あるいは、寸法が 5x5x1000 の場合はどうでしょうか? - 現在、ブロックをほとんど起動していませんが、カーネルはハードウェアの 5x5x512 b/c で起動する必要があり、各スレッドは 2 つの計算を実行します。また、カーネル内の次元を識別できるようにするために、z の一部をブロックに、一部をスレッドに配置して、すべての次元をマッシュアップすることもできません。現在：

__global__ void kernel(/*arguments*/){
    int x = xstart + blockIdx.x;
    int y = ystart + blockIdx.y;
    int z = zstart + threadIdx.x;
    if(x < xend && y < yend && z < zend){
        //calculate
    }
}

これらの変数を把握するための確実で効率的な方法が必要です。

ブロック x の次元、ブロック y の次元、スレッド x (および y? と z?)、blockIdx と threadIdx を介してカーネルに入った後の x、y、z、および入力がハードウェアを超える場合は、 「ステップ」カーネル計算内の for ループで各次元を取得します。

ご不明な点がございましたら、お尋ねください。これは難しい質問であり、私を悩ませてきました (特に、起動するブロック/スレッドの量がパフォーマンスの主要な要素であるため)。このコードは、さまざまなデータ セットの決定を自動化する必要がありますが、それを効率的に行う方法がわかりません。前もって感謝します。

4

2 に答える 2

3

あなたはここで物事を非常に複雑にしすぎていると思います。基本的な問題は、1000 x 1000x1000の計算ドメインでカーネルを実行する必要があることのようです。したがって、1000000000スレッドが必要です。これは、すべてのCUDA互換ハードウェアの機能の範囲内です。したがって、少なくとも計算を実行するために必要なスレッド数を備えた標準の2D CUDA実行グリッドを使用し（実行方法がわからない場合はコメントを残し、回答に追加します）、カーネル呼び出し内で使用します。次のような小さなセットアップ機能：

__device__ dim3 thread3d(const int dimx, const int dimxy)
{
    // The dimensions of the logical computational domain are (dimx,dimy,dimz)
    // and dimxy = dimx * dimy
    int tidx = threadIdx.x + blockIdx.x * blockDim.x;
    int tidy = threadIdx.y + blockIdx.y * blockDim.y;
    int tidxy = tidx + gridDim.x * tidy;

    dim3 id3d;
    id3d.z = tidxy / dimxy;
    id3d.y = tidxy / (id3d.z * dimxy);
    id3d.x = tidxy - (id3d.z * dimxy - id3d.y * dimx);

    return id3d;
}

[免責事項：ブラウザで記述され、コンパイルされたり、実行されたり、テストされたりすることはありません。自己責任で使用してください]。

この関数は、CUDA 2D実行グリッドから3Dドメイン（dimx、dimy、dimz）の「論理」スレッド座標を返します。カーネルの最初で次のように呼び出します。

__global__ void kernel(arglist, const int dimx, const int dimxy)
{
    dim3 tid = thread3d(dimx, dimxy);

    // tid.{xyx} now contain unique 3D coordinates on the (dimx,dimy,dimz) domain
    .....
}

そのグリッドを設定することには整数の計算オーバーヘッドがたくさんあることに注意してください。そのため、なぜ本当に3Dグリッドが必要なのかを考えたいと思うかもしれません。実際には必要ない回数に驚かれることでしょう。その設定されたオーバーヘッドの多くは回避できます。

于 2012-06-08T06:24:38.487 に答える
1

最初cudaGetDeviceProperties()に GPU の計算能力を見つけるために使用します。これにより、GPU で許可されているブロックあたりのスレッド数を正確に知ることができます (CUDA 対応デバイスで実行できるようにプログラムを一般化する必要がある場合)。

次に、その数値を使用して、入力の次元をテストする大きなネストされたifステートメントを作成します。すべての次元が十分に小さければ、(bx,by,bz) スレッドのブロックを 1 つ持つことができます (ありそうにありません)。それがうまくいかない場合は、1 つのブロックに収まる最大の次元 (または 2 つの次元) を見つけ、それに応じて分割します。(MAX_NUMBER_THREADS_PER_BLOCK,1,1)それがうまくいかない場合は、最小の次元を分割して、その一部が 1 つの(bx/MAX_NUMBER)THREADS_PER_BLOCK,by,bz)ブロックに収まるようにする必要がbx<by<bzありbx>MAX_NUMBER_THREADS_PER_BLOCKます。

ケースごとに異なるカーネルが必要になります。これは少し面倒ですが、結局のところ、実行可能な仕事です。

于 2012-06-08T00:37:19.623 に答える