提前退出线程是否会破坏CUDA线程块之间的同步？

Question

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

问题可以这样解释：

我们有一个大小为W * H的图像。 对于图像的每个像素，我需要运行9个相同数据的并行处理，并且每个进程都会产生一个值数组作为结果（整个算法中数组的长度相同，假设为N，大约为20或30）。对于每个像素，这9个进程将在完成计算后将它们的结果累积到最终数组中（每个像素只有一个单独的数组）。

为了并行化这个过程，我设计了以下结构：我生成尺寸为（10,10,9）的块，这意味着每个线程块将处理一个10×10大小的子图像，每个线程将处理一个单独像素的9个相同进程之一。 在这种情况下，网格尺寸将是（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的线程块中的一些像素为零而另一些不为零，则属于零像素的线程将在开始时立即退出，而其他线程将继续处理。在这种情况下，同步是否仍然正常工作，还是会像编程指南所说的那样生成未定义的行为？我考虑让零像素线程处理垃圾数据以保持它们忙碌，但如果我们有完全由零组成的块（我们经常遇到这种情况），这将不必要地增加处理时间。在这种情况下应该怎么办？

- Ufuk Can Bicici

---

你的代码是死锁的配方。请查看链接的帖子获取详细答案。 - talonmies

请记住，GPU 上的多个线程不是独立执行的。一个 warp 中的所有线程同时执行相同的指令。在 if 语句中，如果一个线程执行 if 子句，而其他所有线程执行 else 子句，则一个线程将执行 if 子句，而其他线程将空闲，然后 else 子句线程将执行，而一个线程将空闲。在 if 语句结束时，线程再次同步执行相同的指令。 - dthorpe

我现在知道如何编写内核代码，但我仍然困惑于指南引用中的“因此，如果warp中的任何线程执行bar指令，则就好像warp中的所有线程都执行了bar指令。”部分。假设一个线程处理了if子句的else部分，而其他线程则采用了if方式，并且我们在else部分有一个屏障。根据引用的句子，假定warp中的所有线程都命中了屏障并增加了到达计数器的warp大小，因此所有线程都被视为已阻塞。那么，这会导致死锁吗？ - Ufuk Can Bicici

（继续上一条评论）我问了系统如何检测到屏障已经完成。如果有一个到达计数器，那么这个计数器必须等于块中的线程数，以显示所有线程都完成了屏障。如果按每个线程递增该计数器，则可能会创建死锁，因为在if语句的情况下，该数字永远不会达到总线程计数。但是由于计数器使用warp大小进行更新，因此它就像所有线程都已经到达了屏障一样。 - Ufuk Can Bicici

由于您的代码基于像素值而不是线程索引，因此无法保证每个warp中至少有一个线程会触发__syncthread，因此这段代码确实会导致死锁，正如talonmie所说。但是，如果您可以保证这一点（例如使用线程索引），那么您可以提前退出。请参见同一问题的我的新答案。 - chappjc

Answer 1

为避免创建死锁，所有线程都需要无条件地触发 _synchthreads()。您可以通过将返回语句替换为 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. 
}

- dthorpe