0

PTXマニュアル（バージョン2.3）（http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/ptx_isa_2.3.pdf）6.4.2の状態：

配列要素には、明示的に計算されたバイトアドレスを使用するか、角括弧表記を使用して配列にインデックスを付けることでアクセスできます。角括弧内の式は、定数整数、レジスタ変数、または単純な「定数オフセット付きレジスタ」式のいずれかです。オフセットは、レジスタ変数に加算または減算される定数式です。より複雑なインデックスが必要な場合は、使用前にアドレス計算として記述する必要があります。

 ld.global.u32  s, a[0]; 
 ld.global.u32  s, a[N-1]; 
 mov.u32  s, a[1];  // move address of a[1] into s

これを試してみると、バージョンポインタとバイトオフセットのみが機能します[a+0]。

このコードはロードに失敗します：

.reg    .f32 f<1>;
.global .f32 a[10];
ld.global.f32 f0,a[0];

これは正常にロードされますが：

.reg    .f32 f<1>;
.global .f32 a[10];
ld.global.f32 f0,[a+0];

バイトオフセットバージョンの問題は、それが実際にはバイトオフセットであるということです。したがって、型の基本的なサイズを考慮に入れる必要があります。つまり、2番目の要素は[a+4]です。あなたa[1]のためにこれを解決することになっているのに対して。

何が問題になっているのか？

編集

そして、ここにはさらに深刻な問題があります。上記のテキストは、レジスタ変数を使用して配列にインデックスを付けることができると述べています。

ld.global.f32 f0,a[u0];

ここで、u0はおそらく.reg.u32または他の互換性のある整数です。

ただし、ポインタとバイトオフセットの方法では、これは不可能です。次のようなことをすることは違法です。

mul.u32 u1,u0,4;
ld.global.f32 f0,[a+u1]; // here a reg variable is not allowed.

現在、これは厳しい制限です。ただし、loadステートメントの前に別のアドレス計算を行うことができます。しかし、これは物事を複雑にします。

4

1 に答える 1

1

これは、引用した PTX ドキュメントには適合しないようですが、配列内のアイテムのサイズに対応する乗数を追加できます。たとえば、10 番目の 32 ビット ワードを取得するには、次のようにします。

ld.const.u32 my_u32, [my_ptr + 10 * 4];
于 2012-11-09T15:06:59.067 に答える