1

次のカーネルは音圧フィールドを計算します。各スレッドは、pressureベクトルの独自のプライベート インスタンスを計算し、グローバル メモリに合計する必要があります。pressureベクトルを計算するコードが正しいと確信していますが、これで期待どおりの結果が得られるようにするのにまだ問題があります。

int gid       = get_global_id(0);
int lid       = get_local_id(0);
int nGroups   = get_num_groups(0);
int groupSize = get_local_size(0);
int groupID   = get_group_id(0);

/* Each workitem gets private storage for the pressure field.
 * The private instances are then summed into local storage at the end.*/
private float2    pressure[HYD_DIM_TOTAL];
local   float2    pressure_local[HYD_DIM_TOTAL];

/* Code which computes value of 'pressure' */

//wait for all workgroups to finish accessing any memory
barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE);

/// sum all results in a workgroup into local buffer:
for(i=0; i<groupSize; i++){

    //each thread sums its own private instance into the local buffer
    if (i == lid){
        for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){
            pressure_local[iHyd] += pressure[iHyd];
        }
    }
    //make sure all threads in workgroup get updated values of the local buffer
    barrier(CLK_LOCAL_MEM_FENCE);
}

/// copy all the results into global storage
//1st thread in each workgroup writes the group's local buffer to global memory
if(lid == 0){
    for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){
        pressure_global[groupID +nGroups*iHyd] = pressure_local[iHyd];
    }
}

barrier(CLK_GLOBAL_MEM_FENCE);

/// sum the various instances in global memory into a single one
// 1st thread sums global instances
if(gid == 0){

    for(iGroup=1; iGroup<nGroups; iGroup++){

        //we only need to sum the results from the 1st group onward
        for(iHyd=0; iHyd<HYD_DIM_TOTAL; iHyd++){

            pressure_global[iHyd] += pressure_global[iGroup*HYD_DIM_TOTAL +iHyd];
            barrier(CLK_GLOBAL_MEM_FENCE);
        }
    }
}

データの次元に関するいくつかの注意事項: スレッドの総数は 100 から 2000 の間で変化しますが、場合によってはこの範囲外になることがあります。
groupSizeハードウェアに依存しますが、現在、1(cpu) から 32(gpu) の間の値を使用しています。
HYD_DIM_TOTALコンパイル時に認識され、4 から 32 の間で変化します (通常は 2 の累乗になりますが、必ずしもそうではありません)。

このリダクション コードに明らかに間違っている点はありますか?

PS: AMD APP SDK 2.8 を搭載した i7 3930k と NVIDIA GTX580 でこれを実行しました。

4

1 に答える 1

4

ここで 2 つの問題に気付きました。1 つは大きく、もう 1 つは小さいです。

  • このコードは、バリアの機能を誤解していることを示唆しています。バリアは、複数のワークグループ間で同期されることはありません。ワークグループ内でのみ同期します。CLK_GLOBAL_MEM_FENCE はグローバル同期のように見えますが、実際にはそうではありません。そのフラグは、現在の作業項目のグローバル メモリへのアクセスをすべてフェンスするだけです。したがって、未処理の書き込みは、このフラグを使用したバリアの後でグローバルに監視できます。ただし、バリアの同期動作は変更されません。これは、ワークグループの範囲内にあるだけです。別の NDRange またはタスクを起動する以外に、OpenCL にはグローバル同期はありません。
  • 最初の for ループにより、複数の作業項目が互いの計算を上書きします。iHyd を使用した pressure_local のインデックス作成は、同じ iHyd を使用する各作業項目によって行われます。これにより、未定義の結果が生成されます。

お役に立てれば。

于 2013-02-19T18:46:40.693 に答える