提问
当将thread_group
类型的对象传递给设备函数时,是否有通过引用传递与通过值传递的偏好?
- 其中一个是“正确的”吗
- 每种方法有什么不同
- 每种方法在什么时候应该是首选
示例
编程指南和developer blog中的类似示例似乎以不同的方式处理此问题。
编程指南
__device__
int sum(const thread_block& g, int *x, int n) {
// ...
g.sync()
return total;
}
开发者博客
__device__
int sum(thread_block block, int *x, int n) {
...
block.sync();
...
return total;
}
附加信息
编程指南也有关于构造隐式组的说法:
虽然可以在代码中的任何位置创建隐式组,但这样做很危险。为隐式组创建句柄是一项集体操作-组中的所有线程都必须参与。如果组是在并非所有线程都到达的条件分支中创建的,则可能导致死锁或数据损坏。出于这个原因,建议您预先为隐式组创建一个句柄(尽可能早,在发生任何分支之前),并在整个内核中使用该句柄。由于同样的原因,组句柄必须在声明时初始化(没有默认的构造函数),并且不鼓励复制构造它们。
这会让我相信通过引用传递它们是更好的,但我承认有足够多的细节隐藏在各种合作组中,我可能错过了一些细微差别。按值传递是否会被认为是“复制构造”,因此不鼓励?
我没有注意到任何性能或结果的差异使用任何一个,但我可能只是没有测试正确的边缘情况;或者“未定义的行为”可能只是以一种不会引起问题的方式进行。
1条答案
按热度按时间j8ag8udp1#
首先是一些观察:
1.您引用的博客文章来自2017年预览该功能时,文档是最新的。仅在此基础上,您就应该倾向于使用const通过引用传递习惯用法,因为它的源代码比较新。
1.正如你自己所证明的那样,因为CUDA使用了C对象模型的高度精简的实现,而编译器喜欢内联函数扩展以提高性能,所以你不太可能在真实的世界中发现编译器会为这两种情况生成不同的代码。
因此,我认为const通过引用传递版本是您应该使用的,这既来自 C 正确性POV,也因为当前文档建议您应该使用。可能会有一些极端情况,在某个地方,某个时候,有人在按值传递版本中被复制构造烧毁,但我怀疑你必须非常努力才能发生这种情况。Caveat emptor和所有这些...