OpenCL手动实现并行归约优化代码计算结果错误问题求助

问题排查结果 #

你代码存在两个核心错误导致手动实现的归约逻辑异常：

• 第一，BlockSum数组未声明为工作组共享内存
  当前你写的unsigned short BlockSum[256];属于线程私有局部变量，每个工作项都单独持有一份256长度的数组，仅自身对应GetIndex位置被赋值为1，其余位置都是未初始化的随机值，无法实现跨线程的数据交换完成归约。而内置的work_group_reduce_add是工作组级原语，会自动收集同组所有工作项的输入值做求和，因此返回结果正确。
  修复方案：给BlockSum添加__local修饰符，让整个工作组共用同一份内存数组。
• 第二，仅0号线程对SumOfAll做了赋值，其余线程的SumOfAll仍为初始值0
  你的逻辑里只有GetIndex == 0的工作项会把最终求和结果写入SumOfAll，其余255个工作项的SumOfAll还是初始值0，如果直接让所有线程执行printf，会看到大量0输出，很容易误判为归约失败。修复时要么仅让0号线程打印结果，要么在赋值完成后调用work_group_broadcast把结果同步给组内所有线程。

修正后的完整代码 #

__kernel void Reduction()
{
        unsigned char GetThreadX = get_local_id(0); // 取值范围0~15
        unsigned char GetThreadY = get_local_id(1); // 取值范围0~15
        unsigned char GetGroup   = get_local_size(0); // 16
        __local unsigned short  BlockSum[256];      
        int SumOfAll= 0;            
        
        unsigned short GetIndex = GetThreadX + (GetGroup * GetThreadY); // 取值范围0~255      
        
        BlockSum[GetIndex] = 1;             
        barrier(CLK_LOCAL_MEM_FENCE);       
        
        // 手动归约逻辑
        for(unsigned short stride=128; stride>1; stride >>= 1) {
            if(GetIndex < stride)
                BlockSum[GetIndex] += BlockSum[GetIndex + stride];          
            barrier(CLK_LOCAL_MEM_FENCE);           
        }               
        if(GetIndex==0) {
            SumOfAll = BlockSum[0] + BlockSum[1];
            // 可选：把结果同步给所有线程
            // SumOfAll = work_group_broadcast(SumOfAll, 0);
            printf("SumOfAll=%d\n",SumOfAll);
        }
}

内容的提问来源于stack exchange，提问作者Maverick。