1

まず、アルゴリズムへのリンクは次のとおりです。

GPU Gems 3、第39章：CUDAを使用した並列プレフィックス合計（スキャン）。

バンクの競合を回避するために、NUM_BANKS（つまり、計算可能性2.xのデバイスの場合は32）要素ごとに共有メモリ配列にパディングが追加されます。これは（図39-5のように）によって行われます。

int ai = offset*(2*thid+1)-1
int bi = offset*(2*thid+2)-1
ai += ai/NUM_BANKS
bi += ai/NUM_BANKS
temp[bi] += temp[ai]

ai/NUM_BANKSがマクロとどのように同等であるかわかりません。

   #define NUM_BANKS 16  
   #define LOG_NUM_BANKS 4  
   #define CONFLICT_FREE_OFFSET(n) \  
          ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))

等しいではありませんか

n >> LOG_NUM_BANKS

どんな助けでも大歓迎です。ありがとう

4

1 に答える 1

11

私はそのコードを書き、記事を共同執筆しました。この記事はスキャンアルゴリズムの学習にのみ使用し、コードは使用しないでください。CUDAが新しく、私がCUDAを初めて使用したときに書かれました。CUDAで最新のスキャンの実装を使用する場合、銀行の競合を回避する必要はありません。

簡単な方法でスキャンを実行する場合は、thrust::inclusive_scanまたはを使用しますthrust::exclusive_scan。

本当にスキャンを実装したい場合は、この記事[1]などの最新の記事を参照してください。または、コードが高速であるが、もう少し研究が必要な実際の作品の場合、これは[2]です。または、 Sean Baxterのチュートリアルを読んでください（ただし、後者にはスキャンアルゴリズムに関する独創的な研究の引用は含まれていません）。

[1] Shubhabrata Sengupta、Mark Harris、Michael Garland、およびJohnD.Owens。「メニーコアGPU用の効率的な並列スキャンアルゴリズム」。Jakub Kurzak、David A. Bader、およびJack Dongarraの編集者、マルチコアおよびアクセラレータを使用した科学コンピューティング、Chapman＆Hall / CRC Computational Science、第19章、413〜442ページ。テイラーアンドフランシス、2011年1月。

[2] Merrill、D.およびGrimshaw、A.ストリームアーキテクチャの並列スキャン。テクニカルレポートCS2009-14、バージニア大学コンピュータサイエンス学部。2009年12月。

于 2012-03-14T01:25:34.483 に答える