PTX數組作爲不工作

Question

的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]; 

與字節偏移版本的問題是，它確實是一個字節偏移。因此，必須考慮該類型的基本尺寸，即第二個元素是[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. 

現在，這是一個嚴重的限制。但是，可以在加載語句之前執行另一個地址計算。但是這使事情變得複雜。

Answer 1

這似乎不符合您所引用的PTX文檔，但您可以添加與數組中項目大小相對應的乘數。例如，獲得第10個32位字：

ld.const.u32 my_u32, [my_ptr + 10 * 4];