提前退出线程会干扰块中的CUDA线程之间的同步吗？ [英] Does early exiting a thread disrupt synchronization among CUDA threads in a block?

问题描述

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

I am implementing a certain image processing algorithm with CUDA and I have some questions about the thread synchronization issue overall.

现在的问题可以这样解释：

The problem at hand can be explained like that:

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

We have an image with the size of W*H. For each pixel of the image I need to run 9 identical-data parallel processes and each process gives an array of values as the result (the arrays are of the same length for the whole algorithm, lets say N, which is around 20 or 30). For each pixel, these 9 processes will accumulate their results in a final array (a single array for each pixel) after they finish their calculations.

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

In order to parallelise this, I designed the following structure: I generate blocks with the dimensions of (10,10,9), this means each thread block will process a 10*10 sized sub image and each thread will process 1 of the 9 identical processes for a single pixel. The grid dimension will be (W/10,H/10,1) in this case. For a thread block, I will allocate a shared memory array with the length of 100*N and each thread will write to the appropriate shared memory location according to the coordinates of its current pixel. So, I need a synchronization with an atomicAdd and __synchthreads() here.

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

The problem here is, if a pixel has a value of zero, then we don't need to process it at all, so I want to exit for such pixels, otherwise I will do unnecessary work since a large portion of the image consists of zeroes (background). So, I thought of writing something like the following:

//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. 

在这种情况下是编程指南所说的：

What worries me in this situation is what the Programming Guide is saying:

在条件代码中允许使用__ syncthreads（），但只有条件整个线程块，否则代码执行可能会挂起或产生意想不到的副作用。

__syncthreads() is allowed in conditional code but only if the conditional evaluates identically across the entire thread block, otherwise the code execution is likely to hang or produce unintended side effects.

在10×10线程块中的像素为零并且一些或者不是，则属于零像素的线程将在开始时立即退出，并且其他线程将继续它们的处理。在这种情况下，同步是什么，它将仍然正常工作或将生成未定义的行为，如编程指南所说？我想使零像素线程处理垃圾数据，以保持他们忙，但这将不可避免地增加处理时间，如果我们有块由零完全（我们有它们经常）。在这种情况下可以做什么？

So in my case, if some of the pixels in a 10*10 thread block are zero and some or not, then the threads belonging to the zero pixels will exit immediately at the beginning and the other threads will continue to their processing. What about the synchronization in this case, will it work still properly or will it generate undefined behavior like the Programming Guide says? I thought of making zero pixel threads processing garbage data to keep them busy but this will unnecesarilly increase the processing time if we have blocks which consist of zeroes entirely (and we have them very often). What can be done in this case?

推荐答案

为了避免创建死锁，所有线程需要无条件地击中_synchthreads你可以在你的例子中通过用一个if语句替换返回值，该函数跳过大部分函数，​​直接用于零像素情况的_syncthreads（）。

To avoid creating a deadlock, all threads need to hit the _synchthreads() unconditionally. You can do that in your example by replacing the return with an if statement that jumps over the bulk of the function and heads straight for the _syncthreads() for the zero pixel case.

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. 
}

这篇关于提前退出线程会干扰块中的CUDA线程之间的同步吗？的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！