并行缩减无法正常工作

Posted 2023-03-29

【中文标题】并行缩减无法正常工作

【英文标题】：Parallel Reduction does not work correctly

【发布时间】：2021-12-22 22:40:23

【问题描述】：

我在 OpenCL 上编写了以下并行内核缩减。我只想对 BlockSum 数组中的所有值求和。在使用work_group_reduce_add(BlockSum[GetIndex]); 时，它可以正常工作，但使用我从https://www.fz-juelich.de/SharedDocs/Downloads/IAS/JSC/EN/slides/opencl/opencl-05-reduction.pdf?__blob=publicationFile （幻灯片11）读取的优化代码无法正常工作。这里的错误似乎是什么？ global_work_size 设置为 16,16 以及 local_work_size（意味着每个工作组总共有 256 个线程）。在 work_group_reduce_add 的情况下，我得到 255，这是正确的，但使用优化的代码，我得到 0

__kernel void Reduction()

        unsigned char GetThreadX = get_local_id(0); //it takes values from 0..15
        unsigned char GetThreadY = get_local_id(1); //it takes values from 0..15
        unsigned char GetGroup   = get_local_size(0); //16
        unsigned short  BlockSum[256];      
        int SumOfAll= 0;            
        
        unsigned short GetIndex = GetThreadX + (GetGroup * GetThreadY); // takes values 0..255, group=16        
        
        BlockSum[GetIndex] = 1;             
        barrier(CLK_LOCAL_MEM_FENCE);       
        
        SumOfAll= work_group_reduce_add(BlockSum[GetIndex]); //works great  
        
        // BUT CODE BELOW DOES NOT SUM CORRECTLY
        /*
        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];       
        barrier(CLK_LOCAL_MEM_FENCE);
        */
        printf("SumOfAll=%d\n",SumOfAll);

【问题讨论】：

unsigned char GetThreadX = get_local_id(0); 为什么是char 而不是int？

@AndreasHadjigeorgiou Cos get_local_id(0) 从 0..15 获取值

嗯，我不知道你能做到，有趣！！！

@AndreasHadjigeorgiou 为什么白白浪费宝贵的 GPU 私有内存？

是的，这是一个好点！寄存器确实很珍贵。

【参考方案1】：

好的，问题已解决。 BlockSum[256]; 没有声明为 __local 而是作为私有内存（没有__local 地址空间限定符），这意味着每个线程（或核心）都有自己的这些数据的副本，但优化的缩减代码正在寻找对于线程之间的共享本地内存数据，汇总值。此外，变量 int SumOfAll; 也应声明为 __local 并进行初始化或 private 在我的情况下之前没有任何初始化。你选择。

所以工作内核现在看起来像这样。

我希望这种类型的错误能帮助像我这样不谨慎的人。

__kernel void Reduction()

        unsigned char GetThreadX = get_local_id(0); //it takes values from 0..15
        unsigned char GetThreadY = get_local_id(1); //it takes values from 0..15
        unsigned char GetGroup   = get_local_size(0); //16

        //*********************************************************
        //below was the offending code and the root of the problem 
        //**********************************************************
        __local unsigned short  BlockSum[256];      
        int SumOfAll;           
        //**********************************************************
        
        unsigned short GetIndex = GetThreadX + (GetGroup * GetThreadY); // takes values 0..255, group=16        
        
        BlockSum[GetIndex] = 1;             
        barrier(CLK_LOCAL_MEM_FENCE);       
        
        //SumOfAll = work_group_reduce_add(BlockSum[GetIndex]); 
        
        // OPTIMIZED CODE BELOW NOW SUM UP CORRECTLY
        
        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];       
        barrier(CLK_LOCAL_MEM_FENCE);
        
        printf("SumOfAll=%d\n",SumOfAll);