CUDA在执行过程中组合线程独立(??)变量

| 伙计们,如果标题令人困惑,我深表歉意。我虽然费时又辛苦,却无法想出正确的方式在一行中表达问题。因此,这里有更多详细信息。我正在做一个基本的图像减法,其中第二幅图像已被修改,我需要找到对该图像进行了多少更改的比率。为此,我使用了以下代码。两张图片均为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 counders独立于线程,因为每次我尝试使用它们时,似乎它们只会增加一次。 当我刚开始使用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];
    }
}
第一个节执行您的串行代码所执行的操作-计算差和小于错误的元素的线程局部总数。请注意,我已经编写了此版本,以便每个线程都被设计为处理输入数据的多个条目。这样做是为了帮助抵消随后并行缩减的计算成本,并且其想法是,与输入数据集条目相比,您将使用更少的块和线程。 第二个节是归约本身,它是在共享内存中完成的。它实际上是一种“树状”操作,其中将单个线程块中的线程局部小计集的大小首先求和为32个小计,然后将这些小计合并,直到有该块的最终小计为止,然后存储的是该块的总数。最后,您会看到一小部分总数的清单,每个启动的块一个,可以复制回主机,并在此处计算最终结果。 请注意,我在浏览器中对此代码进行了编码,但尚未编译,可能会出现错误,但是它应该使您了解您尝试执行的“高级”版本如何工作。     
分母非常简单,因为它只是大小。 分子比较麻烦,因为给定线程的值取决于所有先前的值。您将不得不顺序执行该操作。 您正在寻找的东西可能是atomicAdd。不过,它非常慢。 我认为您会发现此问题相关。您的num基本上是全局数据。 CUDA阵列到阵列总和 或者,您可以将错误检查的结果转储到数组中。然后可以对结果计数进行并行化。这有点棘手,但我认为类似的事情会扩大:http://tekpool.wordpress.com/2006/09/25/bit-count-parallel-counting-mit-hakmem/     

要回复问题请先登录注册