3

Cuda-SDKでnbodyコードを読んだとき、コードのいくつかの行を調べたところ、GPUGems3「FastN-BodySimulationwithCUDA」の論文とは少し異なることがわかりました。

私の質問は次のとおりです。まず、次のコードで記述されているように、blockIdx.xがグローバルからメモリをロードしてメモリを共有することに関与しているのはなぜですか？

for (int tile = blockIdx.y; tile < numTiles + blockIdx.y; tile++)
{
    sharedPos[threadIdx.x+blockDim.x*threadIdx.y] =
        multithreadBodies ?
        positions[WRAP(blockIdx.x + q * tile + threadIdx.y, gridDim.x) * p + threadIdx.x] : //this line
        positions[WRAP(blockIdx.x + tile,                   gridDim.x) * p + threadIdx.x];  //this line

    __syncthreads();

    // This is the "tile_calculation" function from the GPUG3 article.
    acc = gravitation(bodyPos, acc);

    __syncthreads();
}

紙によるとこんな感じじゃないですか。なんでだろう

    sharedPos[threadIdx.x+blockDim.x*threadIdx.y] =
        multithreadBodies ?
        positions[WRAP(q * tile + threadIdx.y, gridDim.x) * p + threadIdx.x] :
        positions[WRAP(tile,                   gridDim.x) * p + threadIdx.x];

第二に、ボディごとの複数のスレッドで、なぜthreadIdx.xがまだ関与しているのですか？合計がthreadIdx.yのみによるため、修正値であるか、まったく関与していないと想定されていませんか？

if (multithreadBodies)
{
    SX_SUM(threadIdx.x, threadIdx.y).x = acc.x; //this line
    SX_SUM(threadIdx.x, threadIdx.y).y = acc.y; //this line
    SX_SUM(threadIdx.x, threadIdx.y).z = acc.z; //this line

    __syncthreads();

    // Save the result in global memory for the integration step
    if (threadIdx.y == 0)
    {
        for (int i = 1; i < blockDim.y; i++)
        {
            acc.x += SX_SUM(threadIdx.x,i).x; //this line
            acc.y += SX_SUM(threadIdx.x,i).y; //this line
            acc.z += SX_SUM(threadIdx.x,i).z; //this line
        }
    }
}

誰かが私にこれを説明できますか？それはより速いコードのためのある種の最適化ですか？

4

2 に答える 2

5

私はこのコードと論文の著者です。番号付きの回答は、番号付きの質問に対応しています。

  1. WRAPこれはマイクロ最適化であるため、マクロへの blockIdx.x オフセットはこのペーパーでは言及されていません。もう価値があるかどうかさえわかりません。目的は、異なる SM が異なる DRAM メモリ バンクにアクセスし、すべてが同時に同じバンクにアクセスするのではなく、これらのロード中のメモリ スループットを最大化することでした。オフセットがないと、blockIdx.x同時に実行されているすべてのスレッド ブロックが同じアドレスに同時にアクセスします。全体的なアルゴリズムは帯域幅の制限ではなく計算であるため、これは間違いなく小さな最適化です。悲しいことに、それはコードをより混乱させます。

  2. あなたが言うように、合計は acrossthreadIdx.yですが、各スレッドは個別の合計を実行する必要があります (各スレッドは個別のボディの重力を計算します)。したがってthreadIdx.x、(概念的には 2D の) 共有メモリ配列の右側の列にインデックスを付けるために を使用する必要があります。

彼の（実際には正しくない）答えでSystmDの質問に答えるにgridDim.yは、（デフォルト/共通）1Dブロックの場合は1つだけです。

于 2012-09-12T02:04:58.393 に答える
0

1) 配列 SharedPos は、各ブロックのスレッドの同期 (__syncthreads() を使用) の前に、各ブロック (つまり、各タイル) の共有メモリにロードされます。blockIdx.x は、アルゴリズムによるタイルのインデックスです。

各スレッド (index threadIdx.x threadIdx.y) は共有配列 SharedPos の一部をロードします。blockIdx.x は、タイルのインデックスを参照します (マルチスレッドなし)。

2) acc はボディ インデックス blockIdx.x * blockDim.x + threadIdx.x の float3 です (integrateBodies 関数の冒頭を参照)

q>4 (128 body,p =16,q=8 gridx=8) を使用したこの合計中に、 multithreadBodies=true でいくつかの問題が見つかりました。(GTX 680 を使用)。blockDim.y 全体で一部の合計が実行されませんでした ...

それを避けるためにコードを変更しました。動作しますが、その理由はわかりません...

if (multithreadBodies)
{
    SX_SUM(threadIdx.x, threadIdx.y).x = acc.x;
    SX_SUM(threadIdx.x, threadIdx.y).y = acc.y;
    SX_SUM(threadIdx.x, threadIdx.y).z = acc.z;

    __syncthreads();


        for (int i = 0; i < blockDim.y; i++) 
        {
            acc.x += SX_SUM(threadIdx.x,i).x;
            acc.y += SX_SUM(threadIdx.x,i).y;
            acc.z += SX_SUM(threadIdx.x,i).z;
        }

}

別の質問: 最初のループで:

for (int tile = blockIdx.y; tile < numTiles + blockIdx.y; tile++) 
{
}

grid.y=1 であるため、blockIdx.y が使用される理由がわかりません。

3) コードを高速化するために、非同期の H2D および D2D メモリ コピーを使用します (私のコードでは重力カーネルのみを使用します)。

于 2012-07-10T09:19:57.110 に答える