3

AMD Radeon HD 7700 GPU を使用しています。次のカーネルを使用して、ウェーブフロント サイズが 64 であることを確認します。

__kernel
void kernel__test_warpsize(
        __global T* dataSet,
        uint size
        )
{   
    size_t idx = get_global_id(0);

    T value = dataSet[idx];
    if (idx<size-1)
        dataSet[idx+1] = value;
}

メイン プログラムでは、128 要素の配列を渡します。初期値は dataSet[i]=i です。カーネルの後、次の値を期待します: dataSet[0]=0 dataSet[1]=0 dataSet[2]=1 ... dataSet[63]=62 dataSet[64]=63 dataSet[65]=63 dataSet [66]=65 ... データセット[127]=126

しかし、dataSet[65] が 63 ではなく 64 であることがわかりました。これは私の予想とは異なります。

私の理解では、最初のウェーブフロント (64 スレッド) は dataSet[64] を 63 に変更する必要があります。したがって、2 番目のウェーブフロントが実行されると、スレッド #64 は 63 を取得し、それを dataSet[65] に書き込む必要があります。しかし、dataSet[65] はまだ 64 のままです。なぜですか?

4

2 に答える 2

2

未定義の動作を呼び出しています。ワークグループ内の別のスレッドが書き込んでいるメモリにアクセスしたい場合は、バリアを使用する必要があります。

さらに、GPU が 2 つのウェーブフロントを同時に実行しているとします。次に、dataSet[65] には実際に正しい値が含まれており、最初のウェーブフロントはまだ完了していません。

また、すべてのアイテムを 0 として出力することも、仕様に従って有効な結果です。それは、すべてを完全に連続して実行することもできるからです。だからこそ障壁が必要なのです。

あなたのコメントに基づいて、私はこの部分を編集しました：

インストールhttp://developer.amd.com/tools-and-sdks/heterogeneous-computing/codexl/ 読む: http://developer.amd.com/download/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide.pdf

一定量のスレッド内で分岐を最適化することは、最適化のほんの一部にすぎません。AMD HW がワークグループ内でウェーブフロントをスケジュールする方法と、(ワークグループ内で) ウェーブフロントの実行をインターリーブすることによってメモリ レイテンシを隠す方法について読む必要があります。分岐は、ワークグループ全体の実行にも影響を与えます。それを実行するための有効な時間は、基本的に最も長く実行されている単一のウェーブフロントを実行する時間と同じです (グループ内のすべてが終了するまでローカルメモリなどを解放できないため、別のスケジュールを立てることができません)。ワークグループ)。ただし、これはローカル メモリやレジスタの使用状況などにも依存します。実際に何が起こるかを確認するには、CodeXL を取得して GPU プロファイリングを実行します。これにより、デバイスで何が起こっているかが正確に表示されます。

そして、これは現世代のハードウェアにのみ適用されます。そのため、この概念は OpenCL 仕様自体にはありません。これらのプロパティは大きく変化し、ハードウェアに大きく依存します。

しかし、AMD のウェーブフロント サイズを本当に知りたい場合、答えはほぼ 64 です ( OpenCL プログラミング ガイドについては、http://devgurus.amd.com/thread/159153 を参照してください) 。現在のラインナップ全体を構成するすべてのGCNデバイスで64です。一部の古いデバイスには 16 または 32 があるかもしれませんが、現在はすべてが 64 です (nvidia の場合は一般的に 32 です)。

于 2013-11-09T10:38:59.363 に答える