parallel-processing - 提早退出线程是否会破坏 block 中CUDA线程之间的同步？

This question already has answers here:

Can I use __syncthreads() after having dropped threads?

(2个答案)

4年前关闭。

我正在使用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线程块中的某些像素为零，而某些像素为零，则属于零像素的线程将在开始时立即退出，而其他线程将继续进行处理。在这种情况下，同步又如何呢？它会仍然正常工作，还是会产生 undefined 的行为，如《编程指南》所述？我曾想过让零像素线程处理垃圾数据以使它们忙碌，但是如果我们拥有完全由零组成的块(并且我们经常拥有它们)，这无疑会增加处理时间。在这种情况下可以做什么？

最佳答案

为了避免产生死锁，所有线程都需要无条件地命中_synchthreads()。您可以在示例中通过用if语句替换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. 
}

关于parallel-processing - 提早退出线程是否会破坏 block 中CUDA线程之间的同步？，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/11654171/

上一篇：ruby-on-rails - 使用accepts_nested_attributes_for和多态时，无法批量分配 protected 属性

下一篇：visual-studio-2010 - 如何在VS2010中将.scss文件与CSS intellisense关联

相关文章：

c++ - 我可以将 std::transform 与并行执行策略一起使用吗？

templates - 从 .cpp 文件调用模板化 CUDA 内核

java - 如何与已调用监听器的父线程通信？

多核中的R caret nnet软件包

multithreading - 设置线程在特定内核上工作的优势？

math - CUDA tgamma函数: is documentation wrong or tgamma isn't a gamma?

c++ - OpenCV 是在没有 CUDA 视频解码支持的情况下构建的

C# 精确显示标签 X 秒并隐藏它 Y 秒

ios - 通过 Wifi 网络连接 iOS 一组 iOS 设备

delphi - 如何从Delphi xe2中的OTL终止parallel.foreach