OpenCL在数组中找到最大值答案

【问题标题】：OpenCL find max in arrayOpenCL在数组中找到最大值
【发布时间】：2016-04-08 04:40:41
【问题描述】：

我试图通过归约运算符在一维数组中找到最大值。我曾参考过该方法：OpenCL™ Optimization Case Study: Simple Reductions

以下是我的代码：

    __kernel void Normallize(__global float* input, __global float* output,__global float* cmax, int rows, int cols){

    int g_idx = get_global_id(0);

    for(int i=0 ; i< get_global_size(0) ; i++) cmax[i] = 0;

    barrier(CLK_GLOBAL_MEM_FENCE);

    for(int offset =  get_global_size(0)/2 ; offset >0 ; offset--){
        if(g_idx < offset){
            float pre = input[g_idx];
            float next = input[g_idx + offset];
            cmax[g_idx] = (pre > next) ? pre:next;
        }
        barrier(CLK_GLOBAL_MEM_FENCE);
    }

   output[g_idx] = cmax[0];
}

经过一番研究，我仍然无法找出代码中的问题。

【问题讨论】：

标签： opencl

【解决方案1】：

您是说这个吗（amd gpu 的 %60 VALU 利用率）？：

__kernel void maxping(__global __read_only float * a, __global __write_only float *b){
                        int threadId=get_global_id(0);
                        int localThreadId=get_local_id(0);
                        int localSize=get_local_size(0);
                        __local float fastMem[256];
                        fastMem[localThreadId]=a[threadId];
                        barrier(CLK_GLOBAL_MEM_FENCE|CLK_LOCAL_MEM_FENCE);

                        for(int i=localSize/2;i>=1;i/=2)
                        {
                            if(localThreadId<i)
                            {
                                if(fastMem[localThreadId]<fastMem[localThreadId+i])
                                    fastMem[localThreadId]=fastMem[localThreadId+i];
                            }
                            barrier(CLK_LOCAL_MEM_FENCE);
                        }
                        if(localThreadId==0)
                            b[threadId]=fastMem[localThreadId];
}

每个组（256 个线程）在本地内存中减少，并将每个组中的第一个值设置为其组的最大值。这个例子有从 0 到 4095 的 4096 个元素。

对于上层内核，VALU 的用法类似于：

x: idle thread
o: thread in process,  m: thread in memory operation

**  :      m m m m m m m m m m m m m m m m
i=0 :      o o o o o o o o x x x x x x x x
i=1 :      o o o o x x x x x x x x x x x x
i=2 :      o o x x x x x x x x x x x x x x
i=3 :      o x x x x x x x x x x x x x x x
**  :      m m m m m m m m m m m m m m m m

但我计算的步数更多，每行跨越 250 个单位。

__kernel void maxpong(__global __write_only  float * a, __global __read_only float *b){
                        int threadId=get_global_id(0);
                        int localSize=get_local_size(0);
                        int maxGroups=4096/localSize;
                        if(threadId==0)
                        {
                            float maxv=FLT_MIN;
                            for(int i=0;i<maxGroups;i++)
                            {
                                if(maxv<b[i*localSize])
                                    maxv=b[i*localSize];
                            }
                            a[0]=maxv;

                        }
}

只有第一个线程（cpu 中最好）执行简单的 max(0,1,2,...,M) 并将 a 的第一个元素设置为 max(a)。

第一个内核完成 255/256 的总计算。但它使每个计算单元的一半核心保持不变。因此，您可以在另一半核心中排序另一件事。这可能是另一个阵列被 max()'ed 或相同阵列的 min()'ed 甚至相同阵列的相同最大值，但在其中一半工作，而其他核心在另一半工作。

具有不同初始内核的 max(a) 的 %73 VALU 利用率：

            __kernel void maxping(__global __read_only float * a, __global __write_only float *b){
                int threadId=get_global_id(0);
                int localThreadId=get_local_id(0);
                int localSize=get_local_size(0);
                __local float fastMem[256];
                __local float fastMem2[256];
                fastMem[localThreadId]=a[threadId];
                fastMem2[localThreadId]=a[threadId+2048];

                barrier(CLK_GLOBAL_MEM_FENCE|CLK_LOCAL_MEM_FENCE);

                for(int i=localSize/2;i>=1;i/=2)
                {
                    if(localThreadId<i)
                    {
                        // sorting first part
                        if(fastMem[localThreadId]<fastMem[localThreadId+i])
                            fastMem[localThreadId]=fastMem[localThreadId+i];
                    }
                    else if(localThreadId>localSize-i)
                    {
                        // sorting second part
                        if(fastMem2[localThreadId]<fastMem2[localThreadId-i])
                            fastMem2[localThreadId]=fastMem2[localThreadId-i];
                    }
                    else
                    {
                         // idle thread. Free compute slot. 
                         // can squeeze some geometry computing
                         // or up-sweep scan of another reduction type
                    }
                    barrier(CLK_LOCAL_MEM_FENCE);
                }
                if(localThreadId==0)
                    b[threadId]=(fastMem[localThreadId]>fastMem2[255]?fastMem[localThreadId]:fastMem2[255]);
            }

这对 4096 个元素的数组使用 2048 个线程。将第 0、256、512、.. 元素设置为各自的组最大值，然后您可以轻松检查主机端哪个更大。

仍有未使用的核心。

对于上层内核，VALU 的用法类似于：

x: idle thread
o: thread in process, m: thread doing memory operation

**  :      m m m m m m m m m m m m m m m m
i=0 :      o o o o o o o o o o o o o o o o
i=1 :      o o o o x x x x x x x x o o o o
i=2 :      o o x x x x x x x x x x x x o o
i=3 :      o x x x x x x x x x x x x x x o
**  :      m m m m m m m m m m m m m m m m

但是 i 步数为 log2(256) 次，因此有更多的“i”步数，并且 AMD 硬件有 64 个内核，即使在一个步骤中有 64 个线程时也可以完全服务。当我们将此循环的所有线程使用情况相加时，它不会给出 %73，但是当其他“扭曲”（其中 40 个）流向同一计算单元时，会填充更多的孔，因此更多的向量算术逻辑单元变得更经常使用。即使是本地内存分配部分也很重要，因为所有核心的内存操作单元都保持忙碌（全局到本地，本地到全局），而其他扭曲保持比较单元忙碌。

编辑：如果您不需要全局大小的 256 的倍数，那么您可以在本地内存操作之后添加全局 id 检查，这样它就不会执行未定义的行为。也许你可以用额外的 FLT_MIN 值填充数组。

【讨论】：