GPU GEMS 3からの並列プレフィックスアルゴリズムで使用されるCONFLICT_FREE_OFFSETマクロ

StackOverflow https://stackoverflow.com//questions/9689185

  •  13-12-2019
  •  | 
  •  

質問

まず、ここにアルゴリズムへのリンクがあります：

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
.

任意の助けが高く評価されています。ありがとう

役に立ちましたか？

解決

そのコードと記事を共書したと書いています、そして私はあなたがスキャンアルゴリズムについて学習するためだけに記事を使うことを要求し、そのコードを使わないでください。 CUDAが新品だったときに書かれており、私はCUDAにとって新しいものでした。 CUDAで現代のスキャンの実装を使用する場合は、銀行の紛争回避を必要としません。

簡単な方法でスキャンしたい場合は、 thrust::inclusive_scan または grefodicetagcode 。

本当にスキャンを実装したい場合は、このような最近の記事を参照してください。 [1] 。またはより速いコードを持つ本物のOpusのために、それはもう少し研究を必要とするでしょう、これはこの1つ [2] 。または Sean Baxterのチュートリアルを読みます（ただし、後者はスキャンアルゴリズムのセミラル作業の引用を含みません）。

[1] Shubhabrata Sengupta、Mark Harris、Michael Garland、John D. Owens。 「多くのコアGPU用の効率的な並列スキャンアルゴリズム」 Jakub Kurzak、David A. Bader、Jack Dongarra、編集者、マルチコアおよびアクセラレータ、Chapman＆Hall / CRCの計算科学、第19章、413-442ページ。 http://www.idav.ucdavis.edu/publications/print_pub?pub_id.= 1041

[2]メリル、D.およびGrimshaw、A。ストリームアーキテクチャの並列スキャン。 Virginia大学コンピュータサイエンス学科技術報告書CS2009-14。 2009年12月。

ライセンス： CC-BY-SA と 帰属
所属していません StackOverflow
scroll top