问题标签 [gpu-warp]

我想在 cuda 中实现关键部分。我阅读了很多关于这个主题的问题和答案，答案通常涉及 atomicCAS 和 atomicExch。

但是，这在 warp 级别不起作用，因为 warp 中的所有线程在 atomicCAS 之后都获取相同的锁，从而导致死锁。

我认为有一种方法可以通过使用 warp __ballot 或 __any 指令来真正锁定 cuda。

但是，经过多次尝试，我没有得到令人满意的（阅读工作）解决方案。

这里有人对此有很好的答案吗？

ps：我知道扭曲发散很糟糕，所以不要告诉我改变我的算法。

我对过滤器内核有一个奇怪的性能反转，无论有没有分支。带分支的内核比不带分支的内核运行速度快约 1.5 倍。

基本上我需要对一堆辐射射线进行排序，然后应用交互内核。由于附带的数据很多，我不能多次使用诸如推力::sort_by_key() 之类的东西。

算法思路：

1. 为所有可能的交互类型（五个）运行一个循环
2. 在每个循环中，warp 线程都会为其交互类型投票
3. 循环完成后，每个 warp 线程都知道另一个具有相同交互类型的线程
4. 线程选举他们的领导者（每个交互类型）
5. 领导者使用 atomicAdd 更新交互偏移表
6. 每个线程将其数据写入相应的偏移量

我使用了这篇 Nvidia 帖子中描述的技术https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/

我的第一个内核在循环内包含一个分支并运行约 5 毫秒：

我的第二个内核使用两个元素的查找表，不使用分支并运行约 8 毫秒：

共同部分：

这个怎么可能？有什么方法可以实现这样的过滤器内核，使其运行速度比分支更快吗？

UPD：完整的源代码可在https://bitbucket.org/radiosity/engine/src的filter_kernel cuda_equation/radiance_cuda.cu 中找到

在 CUDA 9 中，nVIDIA 似乎有了这种“合作组”的新概念；出于某种原因，我并不完全清楚，__ballot()现在（= CUDA 9）不推荐使用__ballot_sync(). 那是别名还是语义发生了变化？

...对于现在已__sync()添加到其名称中的其他内置函数的类似问题。

在最近的 nVIDIA GPU uarchitectures 中，单个流式多处理器似乎被分解为 4 个子单元；它们中的每一个都有 8 个“正方形”的水平或垂直“条”，对应于不同的功能单元：整数运算、32 位触发器、64 位触发器和加载/存储。一个单一的扭曲调度器似乎与每个这样的“季度 SM”相关联。

现在，在 CUDA 编程模型中，每个 warp 的线程（= 32 个线程）被指令锁定在一起。但是，当实际执行工作时，并且在这种情况下，比如说，warp 中只有后半部分或后四分之一的线程处于活动状态 - 这些 sub-warp 是否可以安排到 2 或 3 个季度 SM，而另一季度-SM 在做其他工作吗？

使用 CUDA 的 shfl.idx 指令，我们执行本质上是一个经线内聚集：每个通道提供一个基准和一个原点通道，并获取原点通道的基准。

逆运算，scatter呢？我的意思是，不是分散到记忆中，而是分散到车道上。也就是说，每条车道都提供了一个基准和一个目的地车道，并且对于恰好有另一条车道瞄准它们的车道 - 它们最终具有目标车道的值；其他通道以未定义/任意值结束。

我很确定 PTX 没有这样的东西。它可能以某种方式存在于SASS中吗？如果没有，是否有比通过通道索引分散到共享内存和从共享内存加载更好的方法来实现这一点？

我目前正在开展一个项目，在该项目中我正在展开减少的最后一个扭曲。我已经完成了上面的代码；但是，一些修改是通过猜测完成的，我想解释一下原因。我写的代码只有函数kernel4

这是一个缩减算法，其余代码已经提供。

代码：

我的第一个问题是：在声明时

我想验证我对 volatile 的理解。Volatile 防止编译器错误地优化我的代码，并承诺加载/存储是通过缓存完成的，而不仅仅是寄存器（如果有错误请纠正我）。对于归约，如果部分归约和仍存储在寄存器中，为什么会出现问题？

我的第二个问题是：在进行实际的翘曲减少时

在没有 (n > 64) 和 (n > 32) 条件的情况下，归约和将产生不正确的结果。我得到的结果是：

经过 5 次试验，GPU 缩减始终产生 0.0204 的误差。我很谨慎地认为这是一个浮点运算错误。

老实说，我的老师的助手建议进行此更改以添加 (n > 64) 和 (n > 32) 条件，但没有解释为什么它会修复代码。

由于我的试验中的 n 超过 64，为什么这个条件会改变结果。我很难追溯问题，因为我不能像在 CPU 中那样使用打印功能。

我需要做一个看起来像这样的经纱改组：

在这张图片上，线程数被限制为8使其可读。如果我阅读了 Nvidia SDK 和 ptx 手册，shuffle 指令应该可以完成这项工作，特别是shfl.idx.b32 d[|p], a, b, c;ptx 指令。

从我读到的手册：

因此，提供适当的band值c，我应该能够通过编写这样的函数来做到这一点（灵感来自 CUDA SDK__shufl原始实现）。

}

如果可能，srcLane和的常数是c多少？我无法确定它们（我使用的是 CUDA 8.0）。

最好的，

蒂莫咖啡

我最近阅读了GK110 白皮书，其中声称每个 SM 有 4 个 warp 调度器，每个调度器都有双指令调度单元。在每个周期，每个 warp 调度程序都会选择一个合格的 warp 来为其执行指令。

我的问题是在 GK110 中，每个 SM 包含 192 个 CUDA 核心（SP），但 SM 每个周期只能调度 4 个 warp，即 4 x 32 = 128 个核心将被使用（假设所有线程只需要单精度单元），那么其他 64 个核心会做什么呢？

假设我有一个 local_size=8*8*8 的 OpenGL 计算着色器。调用如何映射到 nVidia GPU 扭曲？相同的调用gl_LocalInvocationID.x会在同一个扭曲中吗？还是你？还是z？我不是指所有调用，我只是指一般聚合。

我之所以这样问是因为有一段时间的优化，并不是所有的调用都有工作要做，所以我希望它们处于同一个扭曲中。

在阅读了 CUDA 开发人员博客上的这篇文章后，我很难理解什么时候可以安全\正确使用__activemask().__ballot_sync()

在Active Mask Query部分，作者写道：

这是不正确的，因为它会导致部分和而不是总和。

之后，在Opportunistic Warp-level Programming部分中，他们使用该函数__activemask()是因为：

如果您想在库函数中使用 warp 级编程，但您无法更改函数接口，这可能会很困难。