3

私の質問は、CUDA の配列の動的に変化する要素のセットへの結合されたグローバル書き込みに関するものです。次のカーネルを検討してください。

__global__ void
kernel (int n, int *odata, int *idata, int *hash)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    odata[hash[i]] = idata[i];
}

ここnで、配列の最初の要素には、 の最初の要素から更新されるhashのインデックスが含まれます。明らかに、これはひどい合体の欠如につながります。私のコードの場合、あるカーネル呼び出しでのハッシュは、別のカーネル呼び出しでのハッシュとはまったく無関係です (そして、他のカーネルは別の方法でデータを更新します)。そのため、この特定のケンレルを最適化するためにデータを単純に並べ替えるという選択肢はありません。odatanidata

この状況のパフォーマンスを改善できる CUDA の機能はありますか? テクスチャ メモリについてはよく耳にしますが、読んだ内容をこの問題の解決策に変換することはできませんでした。

4

2 に答える 2

3

テクスチャリングは読み取り専用のメカニズムであるため、GMEMへの分散書き込みのパフォーマンスを直接向上させることはできません。代わりに、このように「ハッシュ」した場合：

odata[i] = idata[hash[i]];

（おそらくあなたのアルゴリズムは変換できますか？）

次に、テクスチャメカニズムを検討することにはいくつかの利点があるかもしれません。（あなたの例は本質的に1Dのようです）。

また、共有メモリ/L1分割がキャッシュに向けて最適化されていることを確認することもできます。ただし、これは散在する書き込みにはあまり役立ちません。

于 2012-10-17T16:16:23.510 に答える
0

ハッシュ結果の範囲を制限できますか? たとえば、スレッドの最初の 1K 反復は、0 から 8K の範囲odataのみにアクセスすることがわかっている場合があります。

その場合、共有メモリを使用できます。共有メモリのブロックを割り当てて、一時的に共有メモリに高速分散書き込みを行うことができます。次に、結合されたトランザクションで共有メモリ ブロックをグローバル メモリに書き戻します。

于 2012-10-18T16:28:58.180 に答える