CUDA在执行期间结合线程独立(?? ?)变量

CUDA combining thread independent(??) variables during execution

本文关键字：变量 独立 线程 执行期 结合 CUDA      更新时间：2023-10-16

如果标题让人困惑，我很抱歉。我想了很久，想不出一个合适的方法来用一句话来表达这个问题。这里有更多的细节。我正在做一个基本的图像减法，其中第二张图像被修改了，我需要找到对图像进行了多少更改的比率。为此，我使用了以下代码。两个图像都是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;

需要用于跨所有线程执行全局求和的操作称为"并行缩减"。虽然您可以使用原子操作来做到这一点，但我不推荐这样做。在CUDA SDK中有一个简化内核和一篇非常好的论文讨论了该技术，值得一读。

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

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];
    }
}

第一节做你的串行代码所做的——计算一个差值和一个线程局部元素的小于error的总数。注意，在编写这个版本时，每个线程都设计为处理输入数据的多个条目。这样做是为了帮助抵消随之而来的并行减少的计算成本，其思想是，您将使用比输入数据集条目更少的块和线程。

第二节是缩减本身，在共享内存中完成。它实际上是一种"树状"操作，其中单个线程块中的线程局部小计集合的大小首先被求和为32个小计，然后将小计组合直到块的最终小计，然后存储为块的总数。您将在count中得到一个小的汇总列表，每个启动的块一个，可以将其复制回主机，并在那里计算您需要的最终结果。

请注意，我在浏览器中编写了这个代码，并没有编译它，可能会有错误，但它应该给你一个关于你正在尝试做的"高级"版本如何工作的想法。

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

分子更麻烦，因为对于给定线程，它的值取决于之前的所有值。你必须连续地做那个操作。

你要找的东西可能是atomicAdd。但是它非常慢。

我想你会发现这个问题很相关。num基本上是全局数据。CUDA数组到数组sum

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

相关文章：