1

次のチュートリアルから N 体の問題をシミュレートする OpenCL コードを研究しています。

http://www.browndeertechnology.com/docs/BDT_OpenCL_Tutorial_NBody-rev3.html

私の主な問題はカーネルコードに依存しています:

   for(int jb=0; jb < nb; jb++) { /* Foreach block ... */

19          pblock[ti] = pos_old[jb*nt+ti]; /* Cache ONE particle position */
20          barrier(CLK_LOCAL_MEM_FENCE); /* Wait for others in the work-group */

21          for(int j=0; j<nt; j++) { /* For ALL cached particle positions ... */
22             float4 p2 = pblock[j]; /* Read a cached particle position */
23             float4 d = p2 - p;
24             float invr = rsqrt(d.x*d.x + d.y*d.y + d.z*d.z + eps);
25             float f = p2.w*invr*invr*invr;
26             a += f*d; /* Accumulate acceleration */
27          }

28          barrier(CLK_LOCAL_MEM_FENCE); /* Wait for others in work-group */
29       }

実行時に正確に何が起こるかわかりません。カーネル コードは n 回実行されます。ここで、n は作業項目の数 (スレッドの数でもあります) ですが、コードの上記の部分では、ローカル メモリを使用します。ワークグループごとに (nb 個のワークグループがあるようです)

では、実行時に、最初の「バリア」まで、 pblock 配列を pos_old のグローバル値でローカルに埋めますか?

バリア ? の前は jb=0 であるため、常に最初のバリアまで、別のワーク グループの場合、pblock 配列には他のワーク グループの配列と同じ値が含まれます。

これらの配列をすべてのワークグループで共有する方法のようですが、これは私には完全には明らかではありません。

どんな助けでも大歓迎です。

4

1 に答える 1

0

カーネルコード全体を投稿していただけますか? params とプライベート変数について仮定する必要があります。

グループには nt 個の作業項目があり、ti は現在の作業項目を表します。ループが実行されると、グループ内の各アイテムは単一の要素のみをコピーします。通常、このコピーはグローバル データ ソースからのものです。最初のバリアは、他のアイテムがコピーを作成するまでワークアイテムを強制的に待機させます。これが必要なのは、グループ内のすべての作業項目が他のすべての作業項目からコピーされたデータを読み取る必要があるためです。ti は作業項目ごとに異なる必要があるため、値は同じであってはなりません。(ただし、jb*nt は最初のループではゼロに等しくなります)

カーネルコード全体は次のとおりです。

__kernel
void 
nbody_sim(
    __global float4* pos ,
    __global float4* vel,
    int numBodies,
    float deltaTime,
    float epsSqr,
    __local float4* localPos,
    __global float4* newPosition,
    __global float4* newVelocity)
{
    unsigned int tid = get_local_id(0);
    unsigned int gid = get_global_id(0);
    unsigned int localSize = get_local_size(0);

    // Number of tiles we need to iterate
    unsigned int numTiles = numBodies / localSize;

    // position of this work-item
    float4 myPos = pos[gid];
    float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f);

    for(int i = 0; i < numTiles; ++i)
    {
        // load one tile into local memory
        int idx = i * localSize + tid;
        localPos[tid] = pos[idx];

        // Synchronize to make sure data is available for processing
        barrier(CLK_LOCAL_MEM_FENCE);

        // calculate acceleration effect due to each body
        // a[i->j] = m[j] * r[i->j] / (r^2 + epsSqr)^(3/2)
        for(int j = 0; j < localSize; ++j)
        {
            // Calculate acceleartion caused by particle j on particle i
            float4 r = localPos[j] - myPos;
            float distSqr = r.x * r.x  +  r.y * r.y  +  r.z * r.z;
            float invDist = 1.0f / sqrt(distSqr + epsSqr);
            float invDistCube = invDist * invDist * invDist;
            float s = localPos[j].w * invDistCube;

            // accumulate effect of all particles
            acc += s * r;
        }

        // Synchronize so that next tile can be loaded
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    float4 oldVel = vel[gid];

    // updated position and velocity
    float4 newPos = myPos + oldVel * deltaTime + acc * 0.5f * deltaTime * deltaTime;
    newPos.w = myPos.w;
    float4 newVel = oldVel + acc * deltaTime;

    // write to global memory
    newPosition[gid] = newPos;
    newVelocity[gid] = newVel;
}

各ワークグループに「localSize」ワークアイテムを持つ「numTiles」ワークグループがあります。

「gid」はグローバル インデックスで、「tid」はローカル インデックスです。

ループ「for(int i = 0; i < numTiles; ++i)」と「i=0」の最初の反復から始めましょう。

たとえば、次のようになります。

numTiles = 4、localSize = 25、および numBodies = 100 = ワークアイテムの数。

次に、実行時に、gid = 80 の場合、tid = 5、idx = 5 となり、最初の代入は localPos[5] = pos[5] になります。

ここで、gid = 5、次に tid = 5 および idx = 5 を使用します。localPos[5] = pos[5] と同じ代入があります。

したがって、私が理解していることから、最初の反復と最初の「バリア」の後、各ワークアイテムには同じローカル配列「localPos」、つまり最初のグローバルブロックのサブ配列「pos[0: 24]」。

これは何が起こるかの良い説明ですか?

于 2012-09-07T22:44:04.767 に答える