Cuda _sync 函数，如何处理未知的线程掩码？
Cuda _sync functions, how to handle unknown thread mask?
这道题是关于适应从锁步到独立程序计数器的语义变化。本质上，我可以将 int __all(int predicate); 之类的调用更改为 for volta.
例如，int __all_sync(unsigned mask, int predicate);
具有语义：
Evaluate predicate for all non-exited threads in mask and return non-zero if and only if predicate evaluates to non-zero for all of them.
文档假设调用者知道哪些线程处于活动状态，因此可以准确地填充掩码。
a mask must be passed that specifies the threads participating in the call
我不知道哪些线程处于活动状态。这是一个内联到用户代码中不同位置的函数。这使得以下之一具有吸引力：
__all_sync(UINT32_MAX, predicate);
__all_sync(__activemask(), predicate);
第一个类似于在 https://forums.developer.nvidia.com/t/what-does-mask-mean-in-warp-shuffle-functions-shfl-sync/67697 被宣布为非法的案例，从那里引用：
For example, this is illegal (will result in undefined behavior for warp 0):
if (threadIdx.x > 3) __shfl_down_sync(0xFFFFFFFF, v, offset, 8);
第二个选择，这次引用自
The __activemask() operation has no such reconvergence behavior. It simply reports the threads that are currently converged. If some threads are diverged, for whatever reason, they will not be reported in the return value.
操作语义似乎是：
- 有 N 个线程的 warp
- 编译时控制流启用了 M (M <= N) 个线程
- D（M的D个子集）线程收敛，作为运行时属性
- __activemask returns哪些线程恰好收敛了
建议同步线程然后使用 activemask，
__syncwarp();
__all_sync(__activemask(), predicate);
一个 nvidia 博客 post 说这也是未定义的，https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/
Calling the new __syncwarp() primitive at line 10 before __ballot(), as illustrated in Listing 11, does not fix the problem either. This is again implicit warp-synchronous programming. It assumes that threads in the same warp that are once synchronized will stay synchronized until the next thread-divergent branch. Although it is often true, it is not guaranteed in the CUDA programming model.
这标志着我的想法结束了。同一篇博客最后给出了一些关于选择掩码值的指导：
- Don’t just use FULL_MASK (i.e. 0xffffffff for 32 threads) as the mask value. If not all threads in the warp can reach the primitive according to the program logic, then using FULL_MASK may cause the program to hang.
- Don’t just use __activemask() as the mask value. __activemask() tells you what threads happen to be convergent when the function is called, which can be different from what you want to be in the collective operation.
- Do analyze the program logic and understand the membership requirements. Compute the mask ahead based on your program logic.
但是，我无法计算掩码应该是什么。这取决于调用站点的控制流，包含 __all_sync 的代码被内联到其中，我不知道。我不想更改每个函数以采用无符号掩码参数。
如何在没有全局转换的情况下检索语义正确的行为？
TL;DR：总而言之，正确的编程方法很可能是做您声明不想做的事情。
更长：
此 blog 特别建议了一种处理未知线程掩码的机会主义方法：在所需操作之前使用 __activemask() 并将其用于所需操作。也就是说（从博客中逐字摘录）：
int mask = __match_any_sync(__activemask(), (unsigned long long)ptr);
那应该是完全合法的。
你可能会问“博客末尾提到的第2条呢？”我想如果你仔细阅读并考虑到我刚刚摘录的以前的用法，它建议“不要只使用 __activemask()”如果你想要不同的东西。从那里的全文来看，这种阅读似乎很明显。这并没有废除先前构造的合法性。
您可能会问“一路上偶然或强制的分歧怎么办？” （即在处理从 elsewhwere 调用的函数期间）
我想你只有两个选择：
在函数入口处获取__activemask()的值。稍后调用所需的同步操作时使用它。这是您对调用环境的意图的最佳猜测。 CUDA 不保证这将是 正确的 ，但是如果您在同步函数调用。
明确调用环境的意图 - 向您的函数添加掩码参数并在所有地方重写代码（您已声明您不想这样做）。
如果您允许在进入您的函数之前出现扭曲发散的可能性，这会掩盖调用环境的意图，则无法从您的函数中推断出调用环境的意图。需要明确的是，具有 Volta 执行模型的 CUDA 允许随时发生 warp 发散的可能性。因此，正确的做法是重写代码，使调用点的意图明确，而不是试图从被调用函数中推断出来。
这道题是关于适应从锁步到独立程序计数器的语义变化。本质上，我可以将 int __all(int predicate); 之类的调用更改为 for volta.
例如，int __all_sync(unsigned mask, int predicate);
具有语义：
Evaluate predicate for all non-exited threads in mask and return non-zero if and only if predicate evaluates to non-zero for all of them.
文档假设调用者知道哪些线程处于活动状态，因此可以准确地填充掩码。
a mask must be passed that specifies the threads participating in the call
我不知道哪些线程处于活动状态。这是一个内联到用户代码中不同位置的函数。这使得以下之一具有吸引力：
__all_sync(UINT32_MAX, predicate);
__all_sync(__activemask(), predicate);
第一个类似于在 https://forums.developer.nvidia.com/t/what-does-mask-mean-in-warp-shuffle-functions-shfl-sync/67697 被宣布为非法的案例，从那里引用：
For example, this is illegal (will result in undefined behavior for warp 0):
if (threadIdx.x > 3) __shfl_down_sync(0xFFFFFFFF, v, offset, 8);
第二个选择，这次引用自
The __activemask() operation has no such reconvergence behavior. It simply reports the threads that are currently converged. If some threads are diverged, for whatever reason, they will not be reported in the return value.
操作语义似乎是：
- 有 N 个线程的 warp
- 编译时控制流启用了 M (M <= N) 个线程
- D（M的D个子集）线程收敛，作为运行时属性
- __activemask returns哪些线程恰好收敛了
建议同步线程然后使用 activemask，
__syncwarp();
__all_sync(__activemask(), predicate);
一个 nvidia 博客 post 说这也是未定义的，https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/
Calling the new __syncwarp() primitive at line 10 before __ballot(), as illustrated in Listing 11, does not fix the problem either. This is again implicit warp-synchronous programming. It assumes that threads in the same warp that are once synchronized will stay synchronized until the next thread-divergent branch. Although it is often true, it is not guaranteed in the CUDA programming model.
这标志着我的想法结束了。同一篇博客最后给出了一些关于选择掩码值的指导：
- Don’t just use FULL_MASK (i.e. 0xffffffff for 32 threads) as the mask value. If not all threads in the warp can reach the primitive according to the program logic, then using FULL_MASK may cause the program to hang.
- Don’t just use __activemask() as the mask value. __activemask() tells you what threads happen to be convergent when the function is called, which can be different from what you want to be in the collective operation.
- Do analyze the program logic and understand the membership requirements. Compute the mask ahead based on your program logic.
但是，我无法计算掩码应该是什么。这取决于调用站点的控制流，包含 __all_sync 的代码被内联到其中，我不知道。我不想更改每个函数以采用无符号掩码参数。
如何在没有全局转换的情况下检索语义正确的行为？
TL;DR：总而言之，正确的编程方法很可能是做您声明不想做的事情。
更长：
此 blog 特别建议了一种处理未知线程掩码的机会主义方法：在所需操作之前使用 __activemask() 并将其用于所需操作。也就是说（从博客中逐字摘录）：
int mask = __match_any_sync(__activemask(), (unsigned long long)ptr);
那应该是完全合法的。
你可能会问“博客末尾提到的第2条呢？”我想如果你仔细阅读并考虑到我刚刚摘录的以前的用法，它建议“不要只使用 __activemask()”如果你想要不同的东西。从那里的全文来看，这种阅读似乎很明显。这并没有废除先前构造的合法性。
您可能会问“一路上偶然或强制的分歧怎么办？” （即在处理从 elsewhwere 调用的函数期间）
我想你只有两个选择：
在函数入口处获取
__activemask()的值。稍后调用所需的同步操作时使用它。这是您对调用环境的意图的最佳猜测。 CUDA 不保证这将是 正确的 ，但是如果您在同步函数调用。明确调用环境的意图 - 向您的函数添加掩码参数并在所有地方重写代码（您已声明您不想这样做）。
如果您允许在进入您的函数之前出现扭曲发散的可能性，这会掩盖调用环境的意图，则无法从您的函数中推断出调用环境的意图。需要明确的是，具有 Volta 执行模型的 CUDA 允许随时发生 warp 发散的可能性。因此，正确的做法是重写代码，使调用点的意图明确，而不是试图从被调用函数中推断出来。