提前退出线程是否会破坏块中 CUDA 线程之间的同步？ [复制]答案

【问题标题】：Does early exiting a thread disrupt synchronization among CUDA threads in a block? [duplicate]提前退出线程是否会破坏块中 CUDA 线程之间的同步？ [复制]
【发布时间】：2012-07-24 03:31:20
【问题描述】：

我正在使用 CUDA 实现某种图像处理算法，我对整体线程同步问题有一些疑问。

手头的问题可以这样解释：

我们有一个大小为 W*H 的图像。对于图像的每个像素，我需要运行 9 个相同数据的并行进程，每个进程都会给出一个值数组作为结果（整个算法的数组长度相同，比如说 N，大约是 20 或 30 ）。对于每个像素，这 9 个过程在完成计算后会将其结果累积到最终数组（每个像素的单个数组）中。

为了实现并行化，我设计了以下结构：我生成尺寸为 (10,10,9) 的块，这意味着每个线程块将处理 10*10 大小的子图像，并且每个线程将为单个像素处理 9 个相同进程中的 1 个。在这种情况下，网格尺寸将为 (W/10,H/10,1)。对于线程块，我将分配一个长度为 100*N 的共享内存数组，每个线程将根据其当前像素的坐标写入适当的共享内存位置。所以，我需要在这里与 atomicAdd 和 __synchthreads() 同步。

这里的问题是，如果一个像素的值为零，那么我们根本不需要处理它，所以我想为这样的像素退出，否则我会做不必要的工作，因为很大一部分图像由零（背景）组成。所以，我想写如下内容：

//X and Y are the coordinates of the current pixel in the input image.
//threadIdx.z gives the index of the process among the 9 for the current pixel. 

int X=blockIdx.x * blockDim.x + threadIdx.x;
int Y=blockIdx.y * blockDim.y + threadIdx.y;
int numOfProcessForTheCurrPixel=threadIdx.z;
int linearIndexOfPixelInBlock=threadIdx.y * blockDim.x + threadIdx.x;

unsigned short pixelValue=tex2D(image,X,Y);
//Here, threads processing zero-pixels will exit immediately.
if(pixelValue==0)
 return;

float resultArray[22];
//Fill the result array according to our algorithm, mostly irrelevant stuff.
ProcessPixel(resultArray,X,Y,numOfProcessForTheCurrPixel);

for(int i=0;i<22;i++)
    atomicAdd(&__sharedMemoryArray[22*linearIndexOfPixelInBlock + i],resultArray[i]);

 __syncthreads(); 
 //Then copy from the shared to the global memory and etc.

在这种情况下让我担心的是编程指南所说的内容：

__syncthreads() 允许在条件代码中使用，但前提是条件在整个线程块中的计算结果相同，否则代码执行可能会挂起或产生意外的副作用。

所以在我的例子中，如果一个 10*10 线程块中的一些像素是零并且一些或不是，那么属于零像素的线程将在开始时立即退出，其他线程将继续它们的处理.在这种情况下同步怎么样，它会继续正常工作还是会像编程指南所说的那样产生未定义的行为？我想过让零像素线程处理垃圾数据以保持它们忙碌，但是如果我们有完全由零组成的块（并且我们经常有它们），这将不必要地增加处理时间。这种情况下怎么办？

【问题讨论】：

您的代码是导致死锁的秘诀。请参阅链接的主题以获得全面的答案。
我明白了，所以经线中的所有线程实际上都必须命中屏障指令。那么在我的情况下可以做什么，我怎样才能从零像素块中退出而不做不必要的工作，同时避免死锁和同步问题？
还有一个问题，在链接线程中，它说命中屏障指令的线程数是由扭曲大小而不是活动线程数增加的。因此，调度程序可能会从当前线程块开始新的扭曲过程，直到它也遇到障碍。因此，在这种情况下，系统如何检查块中的所有线程是否都击中障碍，是否将到达计数与某个值进行比较，例如“（ceil（（块中的线程数）/（warp_size）） +1)*warp_size" 还是什么？我不明白这如何导致死锁？
请记住，GPU 上的多个线程不会彼此独立执行。 warp 中的所有线程同时执行相同的指令。在 if 语句中，如果一个线程使用 if 子句，而所有其他线程都使用 else 子句，则一个线程将在其他线程空闲时执行 if 子句，然后 else 子句线程将在一个线程空闲时执行。在 if 语句结束时，线程重新同步执行相同的指令。
我知道我现在应该如何对内核进行编码，但我仍然对“因此，如果 warp 中的任何线程执行 bar 指令，就好像 warp 中的所有线程已执行 bar 指令。”指南报价的一部分。让我们假设一个线程处理了 if 子句的 else 部分，而其他线程采用了 if 方式，并且我们在 else 部分有一个障碍。因此，根据引用的句子，假设经线中的所有线程都遇到了障碍，并且通过经线大小增加了到达计数，因此所有线程都被视为被阻塞。那么，这怎么会导致死锁呢？

标签： parallel-processing cuda synchronization

【解决方案1】：

为避免产生死锁，所有线程都需要无条件地触发 _synchthreads()。在您的示例中，您可以通过将 return 替换为一个 if 语句来执行此操作，该语句跳过函数的大部分并直接针对 _syncthreads() 用于零像素情况。

unsigned short pixelValue=tex2D(image,X,Y);
//If there's nothing to compute, jump over all the computation stuff
if(pixelValue!=0)
{

    float resultArray[22];
    //Fill the result array according to our algorithm, mostly irrelevant stuff.
    ProcessPixel(resultArray,X,Y,numOfProcessForTheCurrPixel);

    for(int i=0;i<22;i++)
        atomicAdd(&__sharedMemoryArray[22*linearIndexOfPixelInBlock + i],resultArray[i]);

}

__syncthreads(); 

if (pixelValue != 0)
{
    //Then copy from the shared to the global memory and etc. 
}

【讨论】：