CUDA:在减少翘曲和易失性关键字中
CUDA: In warp reduction and volatile keyword
阅读了以下
问题及其答案链接
我脑子里还剩下一个问题。从我在 C/C++ 的背景;我知道使用volatile
有它的缺点。答案中还指出,在 CUDA 的情况下,如果不使用关键字volatile
优化可以将共享数组替换为寄存器以保留数据。
我想知道在计算(总和)减少时可能遇到的性能问题是什么。
例如__device__ void sum(volatile int *s_data, int tid)
{
if (tid < 16)
{
s_data[tid] += s_data[tid + 16];
s_data[tid] += s_data[tid + 8];
s_data[tid] += s_data[tid + 4];
s_data[tid] += s_data[tid + 2];
s_data[tid] += s_data[tid + 1];
}
}
我正在使用减少翘曲。由于所有带有 in warp 的线程都是同步的,因此我认为没有必要使用syncthreads()
结构。
我想知道删除关键字volatile
弄乱我的总和(由于 cuda 优化)吗?我可以在没有volatile
关键字的情况下使用这样的缩减吗?
由于我多次使用此缩减功能,volatile
关键字会导致性能下降吗?
从该代码中删除易失性关键字可能会破坏费米和开普勒 GPU 上的该代码。这些 GPU 缺少直接在共享内存上运行的指令。相反,编译器必须向寄存器发出加载/存储对,从寄存器发出加载/存储对。
在此上下文中,volatile 关键字的作用是使编译器遵守加载-操作-存储周期,而不是执行将s_data[tid]
值保留在寄存器中的优化。保持寄存器中的总和累积会破坏使该扭曲级别共享内存求和正常工作所需的隐式内存同步。
相关文章:
- 易失性sig_atomic_t的内存安全性
- C++易失性:保证 32 位访问?
- 避免易失性和非易失性成员函数的代码重复
- 当 2 个线程共享同一物理内核时,具有错误共享的易失性增量在发布中的运行速度比在调试中慢
- 如何访问常量易失性 std::array?
- 为什么在 C++20 中弃用易失性?
- 根据 MSVC,具有易失性成员的结构不再是 POD
- 是否允许编译器优化掉局部易失性变量
- 访问共享内存而不使用易失性、std::atomic、信号量、互斥锁和自旋锁
- 如何避免对无锁程序使用易失性?
- C++:易失性实例中的易失性成员函数 - 将数组分配给指针是无效的转换?
- g++ 6.3,avx 内联函数上的 Kahan 求和用易失性关键字进行序列化
- 是什么让这种易失性打破了结构的指针算法?
- 如果不需要易失性,为什么 std::atomic 方法会提供易失性重载
- *(易失性无符号整数 *) 的含义 0x00 = 0x00;
- 使用易失性 c 字符串和 std::cout
- 我可以使用互斥锁或关键字(静态)代替C++中的易失性吗?
- 易失性关键字和 RAII 习语 (C++)
- 易失性关键字在C++中成员函数声明中的位置
- C/C++ 更罕见的关键字 - 寄存器、易失性、外部、显式