隣接する配列要素に共有メモリを使用しますか？

Question

私はCUDAで画像を処理したいと思います。各ピクセルの新しい値は、1つの行の2つの隣接ピクセルに基づいて計算されます。各値は2回だけ使用されるため、ピクセル値には__shared__のメモリを使用するのは意味がありますか？問題の構造に合わないので、タイルもそれを行うには間違った方法ではありませんか？私のアプローチは、各ピクセル上でスレッドを実行し、スレッドごとに隣接ピクセル値を読み込むことです。

Answer 1

現在サポートされているすべてのCUDAアーキテクチャにはキャッシュがあります。 コンピューティング機能3.5以降では、これらは特に読み取り専用データに対して効率的です（読み書きデータはL2にのみキャッシュされるため、L1キャッシュは読み取り専用データに限定されます）。入力データへのポインタをconst __restrict__とマークすると、コンパイラはほとんどvia the L1 texture cacheをロードします。 __ldg() builtinを明示的に使用して強制することもできます。

共有メモリを介して隣りのピクセルからのデータの再利用を明示的に管理することは可能ですが、これはキャッシュに依存するだけの利点ではないことがわかります。

もちろん、共有メモリを使用するかどうかにかかわらず、x方向のブロックサイズを最大化し、最適なアクセスローカリティのためにblockSize.yを1にしたいとします。

Answer 2

統合メモリアクセスを利用して共有メモリを使用して結合します。イメージが行単位で格納されるようにするだけです。各ブロックは線形アレイのチャンクを処理する。データの再利用（最初と最後のピクセルを除くすべてのピクセルは3回の処理に関与する）のため、処理されるすべてのピクセルの値を共有メモリにコピーするのがカーネルの初期段階で有益です。