2

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

extern __shared__ float sdata[];

如何在 PTX 中获得等价物？

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

.shared .f32 sdata[ LENGTH ];

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

我能做的是定义一个变量

.shared .f32 sdata;

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

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

这看起来也有点好笑，因为sdata它被定义为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 回答
-1

在 CUDA C 中，可以在共享内存中定义一个可变长度数组

extern __shared__ float sdata[];

这不是该术语通常意义上的可变长度数组- 它只是用于访问动态限制数量的共享内存的语法，它是在内核启动期间设置的。

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

于 2020-02-05T18:52:05.827 回答