__grid_constant__ 参数限定符有什么作用？

Question

几周前，NVIDIA 的 Stephen Jones 进行了一场名为“CUDA：新功能及超越”的 GTC 演讲，其中他介绍了 CUDA v11.7 中即将推出的功能：名为

__grid_constant__

的内核参数限定符/装饰器。我不明白这是什么意思的解释。

具体来说，

__grid_constant__  int x

与

int x

有何不同？它们不是都只是由线程从常量内存中读取吗？

Answer 1

NVidia 更新编程指南总是需要一点时间。

现在已在 CUDA C++ 编程指南

中完成此操作

7.2.4。

__grid_constant__

计算架构的

__grid_constant__

注释大于或 等于 7.0 注释 const 限定的

__global__

函数参数 非引用类型：

• 具有网格的使用寿命，

• 对于网格来说是私有的，即，主机线程和来自其他网格（包括子网格）的线程无法访问该对象，

• 每个网格有一个不同的对象，即网格中的所有线程都看到相同的地址，

• 是只读的，即修改grid_constant对象或其任何子对象是未定义的行为，包括可变成员。

这意味着编译器不会阻止您更改数据，但这样做会导致错误的行为。

要求：

• 用 grid_constant 注解的内核参数必须具有 const 限定的非引用类型。

• 所有函数声明必须与任何 _grid_constant 参数匹配。

• 函数模板特化必须与任何 grid_constant 参数的主模板声明相匹配。

• 函数模板实例化指令必须与任何 grid_constant 参数的主模板声明相匹配。

根据 @paleonix 评论，这允许优化，因为编译器知道它可以遵循比 C++ 更严格的语义。

如果采用

__global__

函数参数的地址，则 编译器通常会在线程中复制内核参数 本地内存并使用副本地址，部分支持 C++ 语义，允许每个线程修改自己的本地副本 函数参数。使用注释

__global__

函数参数

__grid_constant__

确保编译器不会在线程本地内存中创建内核参数的副本，而是使用 参数本身的通用地址。避免本地副本可能 从而提高性能。

__device__ void unknown_function(S const& s);
__global__ void kernel(const __grid_constant__ S s) {    
    s.x += threadIdx.x; // Undefined Behavior: tried to modify read-only memory 
    // Compiler will _not_ create a per-thread thread local copy of "s": 
    unknown_function(s); }

请注意，如果

s.x

被标记为可变，C++ 将允许您更改它。 C++ 还允许您抛弃

s

的常量性。这个注解假设你不会做这样的事情并进行相应的优化。