在CUDA中，如何在有条件执行的代码(例如，__shfl_sync或<cg>.shfl )中获得这个翘曲的线程掩码？
我正在尝试更新一些旧的CUDA代码(前CUDA 9.0)，而且我在更新warp (例如，__shfl)的使用时遇到了一些困难。
基本上，内核的相关部分可能如下所示：
int f = d[threadIdx.x];
int warpLeader = <something in [0,32)>;
// Point being, some threads in the warp get removed by i < stop
for(int i = k; i < stop; i+=skip)
{
// Point being, potentially more threads don't see the shuffle below.
if(mem[threadIdx.x + i/2] == foo)
{
// Pre CUDA 9.0.
f = __shfl(f, warpLeader);
}
}也许这并不是最好的例子(真正的代码太复杂，无法发布)，但是用旧的本质很容易完成的两件事是：
- 洗牌/广播任何线程在这个时候碰巧在这里。
- 仍然可以使用翘曲相对线程索引。
我不知道如何做上述的CUDA 9.0后。
这个问题在这里几乎/部分地得到了回答：如何在CUDA中的条件while语句中同步翘曲内的线程？，但我认为该帖子有一些未解决的问题。
我不相信__shfl_sync(__activemask(), ...)会起作用。这个问题和许多其他网站都提到了这一点。
链接的问题说要使用coalesced_group，但我的理解是，这种类型的cooperative_group重新排列线程，所以如果您在脑海中有一个warpLeader (on [0，32))，我不确定是否有办法“确定”它在coalesced_group中的新级别。
(此外，基于链接问题中简短的注释会话，似乎不清楚coalesced_group是否只是__activemask()的一个很好的包装器.)
可以像链接问题中所描述的那样，迭代地使用__ballot_sync构建一个掩码，但是对于类似于上面的代码，这可能会变得非常乏味。这是我们取得CUDA > 9.0的唯一途径吗？
回答 1
Stack Overflow用户
发布于 2022-03-22 19:09:15
我不相信
__shfl_sync(__activemask(), ...)会起作用。这个问题和许多其他网站都提到了这一点。
链接的问题没有显示任何这样的用法。此外，规范博客特别指出，使用符合以下要求：
洗牌/广播任何线程在这个时候碰巧在这里。
博客指出，这是不正确的用法：
//
// Incorrect use of __activemask()
//
if (threadIdx.x < NUM_ELEMENTS) {
unsigned mask = __activemask();
val = input[threadIdx.x];
for (int offset = 16; offset > 0; offset /= 2)
val += __shfl_down_sync(mask, val, offset);(在概念上类似于链接问题中给出的用法。)
但是对于该博客中定义的“机会主义”用法，它实际上给出了清单9中的使用示例，类似于您声明“无法工作”的用法。当然，它确实遵循了您给出的定义：
洗牌/广播任何线程在这个时候碰巧在这里。
如果您的算法意图正是这样，它应该可以正常工作。然而，在许多情况下，这并不是对算法意图的正确描述。在这种情况下，博客建议一个逐步的过程来找到一个正确的掩码：
- 不要仅仅使用FULL_MASK (即32个线程的0xffffff)作为掩码值。如果经纱中的所有线程都不能按照程序逻辑到达原语，那么使用FULL_MASK可能会导致程序挂起。
- 不要仅仅使用__activemask()作为掩码值。__activemask()告诉您在调用函数时哪些线程碰巧是收敛的，这可能与您在集体操作中想要的不同。
- 分析程序逻辑，了解成员资格要求。根据您的程序逻辑计算前面的掩码。
- 如果您的程序进行了机会主义的扭曲同步编程，请使用“检测”函数(如__activemask()和__match_all_sync() )来找到正确的掩码。
- 使用__syncwarp()将具有内部依赖关系的操作分离。不要假设锁步执行.
注意，步骤1和步骤2与其他注释并不矛盾。如果您确实知道您打算参与整个翘曲(通常在“机会主义”设置中不为人所知)，那么使用硬编码的完全掩码是非常好的。
如果您确实打算使用您给出的机会主义定义，那么使用__activemask()来提供掩码没有什么问题，事实上，博客给出了一个使用示例，步骤4也证实了这种情况下的使用。
https://stackoverflow.com/questions/71577356
复制相似问题
腾讯云开发者
