共享内存中的PTX可变长度缓冲区

Question

我正在尝试在PTX中实现一个全局reduction内核，它使用共享内存在线程块中进行reduction (就像所有CUDA C示例一样)。在CUDA C中，on可以在共享内存中定义可变长度数组

extern __shared__ float sdata[];

我怎样才能在PTX中获得等同的东西？

看起来不合适的是固定长度的数组，比如

.shared .f32 sdata[ LENGTH ];

因为我希望内核对于不同的输入数组长度是可重用的。

我能做的就是定义一个变量

.shared .f32 sdata;

并将其用作数组的基地址。希望它是在共享内存开始时分配的。然后我可以访问数组元素，如下所示

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

此外，这看起来有点滑稽，因为sdata被定义为float。但它实际上是一个浮点数的地址。从这个意义上说，上面这一行确实是正确的。然而，我不确定这是否可以保证是正确的，比如说，只要偏移量不大于内核启动时指定的共享内存大小。

PTX手册没有解释共享内存中的可变长度缓冲区。

有人知道如何在PTX中实现可变长度缓冲区吗？

Answer 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];
}

Answer 2

CUDA C on中的

可以在共享内存中定义可变长度数组

外部__shared__ float sdata[]

这不是usual sense of that term中的可变长度数组--它只是用于访问动态有限的共享内存量的语法，共享内存量是在内核启动时设置的。

事实是，CUDA编译器引入了.extern定义，这是一个不幸的实现细节- nVIDIA不幸地将其暴露为CUDA语法的一部分。