终于到最后出结果的时刻了,有点小激动呢哼。通过5的分析,实际上再看最后一个内核函数就会很轻松了。最后一个内核函数是permute函数,以64*256个元素为组进行处理。每个work-item处理256个数据。在单倍情况下就只有64*256个数据处理,这时候全局和局部都是64,一维情况。
看了下,似乎没有什么难以理解的地方,主要用到了局部内存和内核同步等知识。直接上内核函数。
__kernel
void 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] = s