在 CUDA 中实现减少的问题

Question

我正在尝试在 cuda 中实现 reduce，我在其中找到数组中的最大元素。我已经使用这个内核来找到最小值，它可以工作，但是当我尝试找到最大值时它不起作用。我已经反复遍历算法，找不到错误。任何帮助将不胜感激。 （另外，当我在那里取消注释打印语句时，我得到不同的输出，这也很令人头疼....）

__global__
void findMaxAndMin(const float* const d_logLuminance, float* reduceCopy, int length, float* min_logLum, float* max_logLum){
    int idx = threadIdx.x + blockDim.x*blockIdx.x;
    if(idx >= length){
        return;
    }
    reduceCopy[idx] = d_logLuminance[idx];
    __syncthreads();

    //do a reduction with max

    for(int offset = 1;offset < length;offset = offset*2){
        if(idx % (offset*2) == 0){
            int compIdx = idx + offset;
            if(compIdx < length){
                float newVal = a_max(reduceCopy[idx], reduceCopy[compIdx]);
                if(idx == 0){
                    //printf("val %f \n", newVal);
                }
                __syncthreads();
                reduceCopy[idx] = newVal;
                __syncthreads();
            }
        }
        __syncthreads();
    }
    __syncthreads();
    if(idx == 0){
        *max_logLum = reduceCopy[0];
    }

}

Answer 1

代码有几个问题。如果它在最低限度内起作用，那么你很幸运。

• 我假设你启动了多个区块（你使用blockIdx.x）。您在另一个块中重复使用一个块的结果 - reduceCopy[compIdx] 可能由另一个块设置。您不能依赖这一点：您无法预测块的执行顺序或同步它们。 __syncthreads() 是一个仅在单个区块内有效的屏障！

• if(idx >= length) return 很危险，因为它会导致并非所有线程都会到达以下__syncthreads。

• 您在发散条件 if(compIdx < length) 中有 __syncthreads()。

• a_max 未定义。请记住始终包含最小的工作示例。我可以猜到这个函数应该做什么，但也许其中还潜伏着另一个错误？

在我看来，您在理论上对并行减少有很好的理解，但是由于 CUDA 特定的行为，实现失败了。

我建议您阅读一些关于如何专门在 CUDA 上进行并行缩减的示例。