2

スレッド ブロック内のリダクションに共有メモリを使用するグローバル リダクション カーネルを PTX に実装しようとしています (そこにあるすべての CUDA C の例のように)。CUDA C では、共有メモリに可変長配列を定義する可能性があります。

extern __shared__ float sdata[];

PTX で同等のものを取得するにはどうすればよいですか?

適切と思われないのは、次のような固定長配列です

.shared .f32 sdata[ LENGTH ];

カーネルをさまざまな入力配列の長さで再利用できるようにしたいので。

私ができることは、1つの変数を定義することです

.shared .f32 sdata;

配列のベースアドレスとして使用します。共有メモリの先頭に割り当てられることを期待して。次に、次のような配列要素にアクセスできます

ld.shared.f32 %r4,[sdata + <offset>]

sdataまた、は として定義されているため、これは少しおかしく見えますfloat。しかし、実際には float のアドレスです。この意味で、上記の行は確かに正しいです。ただし、オフセットがカーネルの起動時に指定された共有メモリ サイズを超えない限り、これが正しいことが保証されているかどうかはわかりません。

PTX のマニュアルでは、共有メモリ内の可変長バッファーについて説明していません。

PTX に可変長バッファを実装する方法を知っている人はいますか?

4

2 に答える 2

1

これは機能します。externただし、リンケージ変数が導入されるため、完全なソリューションではありません。

.version 2.3
.target sm_20
.extern .shared .align 4 .b8 sdata[];
.entry func (.param .s32 param0,...)
{
 //
 // Base addresses
 mov.u64 w2,sdata;  // shared memory
 ld.shared.s32 i9,[w2+0];
}
于 2012-10-29T15:33:33.543 に答える