AMD OpenCL例子阅读笔记系列之Radix_Sort(二)

        因为Radix Sort本身比较大，这里分块对系统做阐述。在上一个博文中，已经向大家介绍了该例子的主机部分。这里向大家重点讲述下内核的并行思想。至于内存对象及调用顺序等需要大家结合AMD的例子来看，都弄到博文会很长。

今天来讲一下例子中的比较简单的内核函数应用histogram内核应用。

       Histogram内核用于分组统计随机生成的数据。统计的规则为：

1.      将所有数据分成以256为一份的组；

2.      以一维的方式对数据进行运算；

3.      对于全局的线程数设定为与需要处理的数据个数相同；

4.      对于每个块的线程数设置成256;

5.      每个块统计对应的分配数据中按偏移量的数据，统计的方式为按索引统计。

      具体方式可以简化用下图表示：

     这里将所有数据分成256个数为一组，然后根据基数排序的特点。并且程序是以LSB的方式进行，拿第一次循环为例，首先将最低8位取出，取出后将索引位置处的值加1。

内核函数如下：

__kernel
void histogram(__global const uint* unsortedData,
               __global uint* buckets,
               uint shiftCount,
               __local uint* sharedArray)
{
    size_t localId = get_local_id(0);
    size_t globalId = get_global_id(0);
    size_t groupId = get_group_id(0);
    size_t groupSize = get_local_size(0);
   
    uint numGroups = get_global_size(0) / get_local_size(0);
  
    /* Initialize shared array to zero */
   
    sharedArray[localId] = 0;
 
    barrier(CLK_LOCAL_MEM_FENCE);
   
    /* Calculate thread-histograms */
   uint value = unsortedData[globalId];
   value = (value >> shiftCount) & 0xFFU;
   atomic_inc(sharedArray+value);
   
   
    barrier(CLK_LOCAL_MEM_FENCE);
   
    /* Copy calculated histogram bin to global memory */
   
    uint bucketPos = groupId  * groupSize + localId ;
   //uint bucketPos = localId * numGroups + groupId ;
    buckets[bucketPos] = sharedArray[localId];
   
}

    这个内核比较简单，最终的结果保存在histogramBinsBuf的cl_mem对象中。里面用到了内核中同步块内内存以及原子操作的知识。