AMD OpenCL例子阅读笔记系列之Radix_Sort(四)_radix sort gpu hlsl-优快云博客

本文深入探讨了一维扫描内核scanArrayKerneldim1的实现细节，该内核用于处理sumBufferin数据，通过缩减树算法完成偏移量整理。作为GPU并行计算的一部分，此步骤对前文所述二维扫描后的结果进行进一步处理。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

在三中我们又得到了两个中间结果sumBufferin以及scanedHistogramBinsBuf两个结果。那么接下来如何呢？我们继续循着程序执行路径往下看，因为numGroups=64,所以runFixOffsetKernel中间的if部分判断的内容将不会执行。也就是我们说的单倍形式下只需要3个内核。

那么下一个内核通过查看我们发现是scanArrayKerneldim1，弄完了2维现在开始弄一维的结果了。还是老样子先看下总体的结构划分，全局和块分配上都是RADICES，也就是256。因为单倍和多倍情况不太相同，我们理一下单倍情况下的内核参数为：summaryBUfferout、sumBufferin、RADICES * sizeof(cl_uint)的本地内存空间、 groupS = RADICES（内核块大小）。显然sumBufferin是我们上一讲也就是三中得到的组的加和结果。我们先贴出来内核函数：

__kernel void ScanArraysdim1(__global uint *output,
                         __global uint *input,
                         __local uint* block,
                         const uint block_size
                         )
  {
    int tid = get_local_id(0);
      int gid = get_global_id(0);
       int bid = get_group_id(0);
      
         /* Cache the computational window in shared memory */
       block[tid]     = input[gid];
 
       uint cache = block[0];
 
    /* build the sum in place up the tree */
       for(int stride = 1; stride < block_size; stride *=2)
       {
             
              if(tid>=stride)
              {
                     cache = block[tid-stride]+block[tid];
              }
              barrier(CLK_LOCAL_MEM_FENCE);
 
              block[tid] = cache;
              barrier(CLK_LOCAL_MEM_FENCE);
             
       }
    /*write the results back to global memory */
       if(tid == 0)
       {    
              output[gid]     = 0;
       }
       else
       {
              output[gid]     = block[tid-1];
       }
  }

通过查看代码，我们可以看到很熟悉的缩减树算法。这也是我为什么直接贴代码的原因。那么最终的输出结果也就是说对sumBufferin的结果进行偏移量整理，这个过程和scanedHistogramBinsBuf雷同，只不过一个是2维，一个是1维的。所以这里关于这个内核函数不做过多讲解。下一讲将对调整偏移量的最后一个内核函数进行讲解同时对该函数中的3个内核函数进行下整理。