线程同步在“易失性__shared__`标志”上的必要性
我的问题是在阅读B.5.记忆栅栏函数的最后一个例子时出现的。我理解这个例子中的标志，检查最后的块处理和。
在我的想象中，如果将标志声明为volatile __shared__，而不使用__syncthreads()，线程迟早会看到标志true，更有可能不同时看到带有true的标志。
因此，我开始认为，如果我可以不使用__syncthreads()**?**来做同样的事情，那么在我们有一个易失性共享内存标志的情况下，我们有什么替代或解决办法吗？
(我在下面的部分给出了我的虚拟设计，它们都不起作用。)
尝试1：
__device__ unsigned int count = 0;
volatile __shared__ bool isLastBlockDone; /// as volatile
__global__ void sum(const float* array, unsigned int N,
volatile float* result)
{
float partialSum = calculatePartialSum(array, N);
if (threadIdx.x == 0) {
result[blockIdx.x] = partialSum;
__threadfence();
unsigned int value = atomicInc(&count, gridDim.x);
isLastBlockDone = (value == (gridDim.x - 1));
}
while (!isLastBlockDone) __threadfence_block(); /// No __syncthreads();
if (isLastBlockDone) {
float totalSum = calculateTotalSum(result);
if (threadIdx.x == 0) {
result[0] = totalSum;
count = 0;
}
}
}其动机是使用繁忙的等待循环来查找标志isLastBlockDone为真的时刻。但是，在启动配置中，此代码不能处理超过一个块。我坚信螺纹发散发生在这里。那么，__syncthreads();是唯一必要的方法吗？
尝试2：
__global__ void sum(const float* array, unsigned int N,
volatile float* result)
{
volatile __shared__ unsigned int count; /// as __shared__
volatile __shared__ bool isLastThreadDone; /// as volatile
count = 0;
isLastThreadDone = false;
__syncthreads(); /// Init
float partialSum = calculatePartialSum(array, N);
result[blockIdx.x] = partialSum;
__threadfence();
unsigned int value = atomicInc(&count, blockDim.x);
isLastThreadDone = (value == (blockDim.x - 1));
while (!isLastThreadDone) __threadfence_block(); /// No __syncthreads();
if (isLastThreadDone) {
float totalSum = calculateTotalSum(result);
if (threadIdx.x == 0) {
result[0] = totalSum;
count = 0;
}
}
}其动机是将原始示例从检查最后一个块迁移到，检查每个块中的最后一个线程。我之所以产生这种动机，是因为我更喜欢volatile __shared__ bool，而不是只在一个街区内保存簿记，所以我让count和isLastThreadDone成为块本地__shared__。
然而，问题仍然是，程序仍然有一个僵局。再一次，我想螺纹发散发生在这里。那么，__syncthreads();是唯一能让它工作的方法吗？
此外，我还举了上面的例子来解决线程同步问题，并让程序在没有死锁的情况下被终止。总结的正确性不是我问题的主要焦点。
回答 1
Stack Overflow用户
发布于 2021-12-29 16:25:47
在某些情况下，应该可以用共享内存自旋等待替换__syncthreads()语句，但是我怀疑它的效用。首先，我没有数据表明它更高效/更快/更好，其次，我认为它不需要在代码中包含__syncthreads()，它只是将语句的需求推到了其他地方(第三，如下所示，它触发了racecheck工具)。
您的try1案例没有意义，因为这里的一般策略是线程块耗尽的减少(请参阅threadFenceReduction示例代码)。线程块排水机制使用atomics来指示哪个线程块是最后到达"end“的线程块。只有一个线程块将获得此信号；其他所有线程块只需退出。因此，由于只有一个线程块将接收isLastBlockDone的布尔isLastBlockDone条件，所有其他线程块(通常只会退出)都会在后续的while语句中永久等待。你有可预见的僵局。这与线程发散无关。所有线程(在所有线程块中)在while语句中的行为都是相同的，但是只满足一个线程块退出while循环所需的条件。
您的try2示例更简单。由于您已经将所有的使用都转换为__shared__内存，因此不可能使用这种方法进行线程间的通信，因此它完全不适合协调多个线程块的行为。
要想得到任何类似于您的请求的内容，我们需要结合这两种方法的元素：
- 保留用于线程块之间通信的原子机制。
- 用等效的自旋等待替换
__syncthreads()。
基本的(1D)线程块耗尽方法如下所示：
__device__ int count = 0;
....
__shared__ bool i_am_last_block;
perform_threadblock_level_reduction();
if (!threadIdx.x){
int value = atomicAdd(&count, 1);
i_am_last_block = (value == (gridDim.x - 1));}
__syncthreads(); // **
if (i_am_last_block)
perform_final_stage_reduction();如果出于某种原因，应该可以替换(或至少移动)上面用__syncthreads()描述的// **语句，方法是使用旋转等待：
__device__ int count = 0;
....
__shared__ volatile bool ready;
__shared__ volatile bool i_am_last_block;
if (!threadIdx.x) ready = false;
__syncthreads();
perform_threadblock_level_reduction();
if (!threadIdx.x){
int value = atomicAdd(&count, 1);
i_am_last_block = (value == (gridDim.x - 1));
ready = true;}
while (!ready);
if (i_am_last_block)
perform_final_stage_reduction();注意，__syncthreads()语句已经从线程代码的末尾“移动”到开头。但是，在__syncthreads()语句中从ready的初始化到线程0的使用之间的任何if使用都将满足需求；严格来说，可能没有必要显式地添加另一个__syncthreads()语句。换句话说，我所熟悉的任何合理的perform_threadblock_level_reduction() (为线程块产生一个单独的部分和)都必然涉及到__syncthreads()的使用。因此，有了这个附带条件，我所描述的“额外”__syncthreads()就没有必要了。
我认为所有这些都不太可能在代码性能方面产生有用的结果。下面是一个完整的例子：
$ cat t1942.cu
#include <iostream>
const int nTPB = 256; // must be power-of-2
template <typename T>
__device__ void perform_threadblock_level_reduction(T *in, T *out, size_t n, size_t oidx, size_t start, size_t stride){
__shared__ T sdata[nTPB];
sdata[threadIdx.x] = 0;
for (size_t i = start; i < n; i += stride) sdata[threadIdx.x] += in[i];
for (int i = nTPB>>1; i > 0; i >>= 1){
__syncthreads();
if (threadIdx.x < i) sdata[threadIdx.x] += sdata[threadIdx.x+i];}
if (!threadIdx.x) {out[oidx] = sdata[0]; __threadfence();}
}
__device__ int count = 0;
template <typename T>
__global__ void r(T *in, T *out, size_t n){
__shared__ volatile bool ready;
__shared__ volatile bool i_am_last_block;
if (!threadIdx.x) ready = false;
perform_threadblock_level_reduction(in, out, n, blockIdx.x, blockIdx.x*blockDim.x+threadIdx.x, gridDim.x*blockDim.x);
if (!threadIdx.x){
int value = atomicAdd(&count, 1);
i_am_last_block = (value == (gridDim.x - 1));
ready = true;}
while (!ready);
if (i_am_last_block)
perform_threadblock_level_reduction(out, out, gridDim.x, 0, threadIdx.x, blockDim.x);
}
typedef int mt;
const size_t s = 1048576;
const int nBLK = 64;
int main(){
mt *in, *out;
cudaMallocManaged(&in, s*sizeof(mt));
cudaMallocManaged(&out, nBLK*sizeof(mt));
for (size_t i = 0; i < s; i++) in[i] = 1;
r<<<nBLK, nTPB>>>(in, out, s);
cudaDeviceSynchronize();
std::cout << out[0] << std::endl;
}
$ nvcc -o t1942 t1942.cu -lineinfo
$ compute-sanitizer ./t1942
========= COMPUTE-SANITIZER
1048576
========= ERROR SUMMARY: 0 errors
$注意，在上述情况下，cuda-memcheck或compute-sanitizer工具racecheck将报告错误。这是因为线程间通信使用共享变量，没有插入的__syncthreads()语句(基本上是您的请求)。然而，无论如何，while循环加上volatile的使用会使代码工作。
这种用旋转等待代替__syncthreads()的方法还有一个额外的挑战或缺点。当翘曲完全到达__syncthreads()时，在__syncthreads()等待时，它们就是失速。这一点很重要，因为这意味着warp调度程序无法为调度/发布选择那些翘曲。当翘曲连续到达一个__syncthreads()时，调度程序的注意力将“集中”在其余的翘曲上。这为这些偏差提供了有益的前进动力，有效地确保了__syncthreads()最终会得到满足，如果可能的话。
旋转等待方法没有这样的优势。翘曲“等待”在旋转-等待仍然是可调度的。虽然在实践中似乎不太可能发生，但至少从理论上讲，经纱调度程序可以“总是”选择处于旋转阶段的翘曲--等待，而不是仍然需要向前推进的翘曲。(至少，我们可以说排除这种可能性并不是一件小事。)这将造成僵局。因此，与旋转等待相比，使用__syncthreads()有明显的优势，即使旋转等待似乎是可行的。
https://stackoverflow.com/questions/70502529
复制相似问题
腾讯云开发者
