CUDA 在执行期间结合线程独立（？？）变量答案

【问题标题】：CUDA combining thread independent(??) variables during executionCUDA 在执行期间结合线程独立（？？）变量
【发布时间】：2012-03-12 04:26:34
【问题描述】：

各位，如果标题令人困惑，我深表歉意。我虽然漫长而艰难，但无法想出正确的方法来用一行来表达这个问题。所以这里有更多细节。我正在做一个基本的图像减法，其中第二个图像已被修改，我需要找到对图像进行了多少更改的比率。为此，我使用了以下代码。两张图片都是 128x1024。

for(int i = 0; i < 128; i++)
{
    for(int j = 0; j < 1024; j++)
    {
        den++;
        diff[i * 1024 + j] = orig[i * 1024 + j] - modified[i * 1024 + j];
        if(diff[i * 1024 + j] < error)
        {
            num++;
        }
    }
}
ratio = num/den;

上述代码在 CPU 上运行良好，但我想尝试在 CUDA 上执行此操作。为此，我可以设置 CUDA 对图像进行基本减法（下面的代码），但我不知道如何执行条件 if 语句来得出我的比率。

__global__ void calcRatio(float *orig, float *modified, int size, float *result)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if(index < size)
        result[index] = orig[index] - modified[index];
}

所以，到目前为止，它可以工作，但我无法弄清楚如何并行化每个线程中的 num 和 den 计数器来计算所有线程执行结束时的比率。对我来说，感觉 num 和 den 计数器是独立于线程的，因为每次我尝试使用它们时，它们似乎只增加一次。

任何帮助都将不胜感激，因为我刚刚开始使用 CUDA，而且我在网上看到的每个示例似乎都不适用于我需要做的事情。

编辑：修正了我幼稚的代码。忘记在代码中键入主要条件之一。这是漫长的一天。

for(int i = 0; i < 128; i++)
{
    for(int j = 0; j < 1024; j++)
    {
        if(modified[i * 1024 + j] < 400.0)  //400.0 threshold value to ignore noise
        {
            den++;  
            diff[i * 1024 + j] = orig[i * 1024 + j] - modified[i * 1024 + j];
            if(diff[i * 1024 + j] < error)
            {
                num++;
            }
        }
    }
}
ratio = num/den;

【问题讨论】：

标签： c++ visual-studio-2008 cuda

【解决方案1】：

您需要用于跨所有线程执行全局求和的操作称为“并行归约”。虽然您可以使用原子操作来执行此操作，但我不推荐它。 CUDA SDK里面有reduction kernel和一篇非常好的论文讨论这个技术，值得一读。

如果我正在编写代码来做你想做的事，它可能看起来像这样：

template <int blocksize>
__global__ void calcRatio(float *orig, float *modified, int size, float *result, 
                            int *count, const float error)
{
    __shared__ volatile float buff[blocksize];

    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    int count = 0;
    for(int i=index; i<n; i+=stride) {
        val = orig[index] - modified[index];
        count += (val < error);
        result[index] = val;
    }

    buff[threadIdx.x] = count;
    __syncthreads();


    // Parallel reduction in shared memory using 1 warp
    if (threadId.x < warpSize) {

        for(int i=threadIdx.x + warpSize; i<blocksize; i+= warpSize) {
            buff[threadIdx.x] += buff[i];

        if (threadIdx.x < 16) buff[threadIdx.x] +=buff[threadIdx.x + 16];
        if (threadIdx.x < 8)  buff[threadIdx.x] +=buff[threadIdx.x + 8];
        if (threadIdx.x < 4)  buff[threadIdx.x] +=buff[threadIdx.x + 4];
        if (threadIdx.x < 2)  buff[threadIdx.x] +=buff[threadIdx.x + 2];
        if (threadIdx.x == 0) count[blockIdx.x] = buff[0] + buff[1];
    }
}

第一节执行您的串行代码所做的事情 - 计算差异和 线程本地 小于错误的元素总数。注意我编写了这个版本，以便每个线程都被设计为处理多个输入数据条目。这样做是为了帮助抵消随后并行减少的计算成本，其想法是您将使用比输入数据集条目更少的块和线程。

第二节是归约本身，在共享内存中完成。它实际上是一种“树状”操作，其中首先将单个线程块内的线程局部小计集的大小汇总为 32 个小计，然后将小计组合起来，直到有块的最终小计，并且然后存储的是 block 的总数。您最终会得到一小部分计数小计，每个您启动的块都有一个，可以复制回主机并在那里计算您需要的最终结果。

请注意，我在浏览器中对此进行了编码并没有对其进行编译，可能会出现错误，但它应该让您了解您正在尝试执行的“高级”版本如何工作。

【讨论】：

感谢 talonmies。但我刚刚意识到我在键入代码时犯了一个巨大的错误。经过漫长的一天，我从记忆中输入了这个，所以忘记了主要部分。编辑了我上面的帖子。
这对代码几乎没有影响。您添加的条件进入我发布的内核第一节的循环中。
噢噢噢。这大概是我的想法……但表达得更清楚了。你知道这个操作的名称。这在未来可能会被证明是有用的。

【解决方案2】：

分母很简单，因为它只是大小。

分子更麻烦，因为它对给定线程的值取决于所有先前的值。您将不得不连续执行该操作。

您正在寻找的东西可能是 atomicAdd。不过速度很慢。

我想你会发现这个问题很相关。您的 num 基本上是全局数据。 CUDA array-to-array sum

或者，您可以将错误检查的结果转储到一个数组中。然后可以并行计算结果。这会有点棘手，但我认为这样的事情会扩大：http://tekpool.wordpress.com/2006/09/25/bit-count-parallel-counting-mit-hakmem/

【讨论】：