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

    终于到最后出结果的时刻了，有点小激动呢哼。通过5的分析，实际上再看最后一个内核函数就会很轻松了。最后一个内核函数是permute函数，以64*256个元素为组进行处理。每个work-item处理256个数据。在单倍情况下就只有64*256个数据处理，这时候全局和局部都是64，一维情况。

        看了下，似乎没有什么难以理解的地方，主要用到了局部内存和内核同步等知识。直接上内核函数。

__kernelvoid permute(__global const uint* unsortedData,             __global const uint* scanedBuckets,             uint shiftCount,             __local ushort* sharedBuckets,             __global uint* sortedData){    size_t groupId = get_group_id(0);    size_t localId = get_local_id(0);    size_t globalId = get_global_id(0);    size_t groupSize = get_local_size(0);            /* Copy prescaned thread histograms to corresponding thread shared block */    for(int i = 0; i < RADICES; ++i)    {        uint bucketPos = groupId * RADICES * groupSize + localId * RADICES + i;        sharedBuckets[localId * RADICES + i] = scanedBuckets[bucketPos];    }    barrier(CLK_LOCAL_MEM_FENCE);        /* Premute elements to appropriate location */    for(int i = 0; i < RADICES; ++i)    {        uint value = unsortedData[globalId * RADICES + i];        value = (value >> shiftCount) & 0xFFU;        uint index = sharedBuckets[localId * RADICES + value];        sortedData[index] = unsortedData[globalId * RADICES + i];        sharedBuckets[localId * RADICES + value] = index + 1;    barrier(CLK_LOCAL_MEM_FENCE);    }}

      前一个循环把所有的桶内容移到局部内存，第二个循环把数据放到合适的位置，每个work-item处理RADICES=256个数据，因每组是64个work-item所以每组就处理256*64个数据，就是我们前面说的单倍数据量了。有人说那还有两个附加的内核没讲吧？对是的。但是我感觉有些东西留着自己研究会好些。都讲了反而不美。留几个问题思考下：

1.在内核函数中多次使用了局部内存，这样做有什么好处？

2.多内核函数的调用问题，这个例子里面也算是典型了，每个内核或者几个内核组合得到中间结果，然后交由后续内核继续进行求解。

3.留的两个内核函数，及多倍情况下的数据处理。

4.这显然只是一个例子，如果需要实际应用，也就是说拿来用，怎么用？

5.在RadixSort中两次用到了缩减树的算法，这个算法能不能进行优化一下？怎么优化？

       好了，就这些吧，RadixSort就到这里啦。