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对象中。里面用到了内核中同步块内内存以及原子操作的知识。