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

      在完成了二中的histogramBinsBuf的统计后，需要对其进行进一步的处理以得到基数排序的最终偏移信息。这个过程是本例子中最复杂的一部分。其中根据分组的情况包含3个内核调用或者5个内核调用。
      这里需要说明下，内核本身执行时需要对数据进行预处理才能放到GPU上运算。这个例子的使用上面，采用的规则是：

         数据的个数必须是64*256=16384的倍数。

      我们先以单倍，即16384个数据为例来分析该例子。

     先看第一个调用的内核函数scanArrayKerneldim2思想解析。其并行结构的划分可以简单用下图表示：



      也就是说利用二维内核来计算，全局的内核分配为根据numElem/256划分x轴的block个数。当然因为我们分析的是最基本的单倍情况，这里只有一个组。而y轴的维度坐标情况为从0~255。那么这个内核用来计算什么呢？首先看下数据通过内核读取时的排序情况：

      在x为0时，y从下到上的索引为0~255;x为1时，y从下到上的索引为256~511……

      在内核scanArrayKerneldim2处理完成后，生成两个结果内存对象：

1.      sumBufferin对象用于存储每行的统计和；即对入桶时索引相同的二中的结果以64个为一组进行加和。因为以单倍方式，所以sumBufferin得到256个加和结果；

2.      scanedHistogramBinsBuf以纵向y轴方向排序，x=0时，结果为0，x=1时记录前面的组的结果即x=0时的统计结果；同理x=2时记录(x=0)+(x=1)的统计结果。

     还是有点乱，没关系，我们上图说话：

 


       这个就是sumBufferin的加和计算模型了。是不是清晰些了？那么还有一个scanedHistogramBinsBuf究竟是怎么回事呢？下面是一个超级简图，因为画起来比较困难，这里简化了很多，B
,0<=n<=63是分块后的索引，B[0]表示histogram中0号块的统计结果,其他依次类推。



      也就是说除了索引为0，每个都是前面块的结果的累加和。

      内核函数为：

__kernel void ScanArraysdim2(__global uint *output,
__global uint *input,
__local uint* block, //64*1
const uint block_size,//64
const uint stride, //64
__global uint* sumBuffer)
{

int tidx = get_local_id(0);
int tidy = get_local_id(1);
int gidx = get_global_id(0);
int gidy = get_global_id(1);
int bidx = get_group_id(0);
int bidy = get_group_id(1);

int lIndex = tidy * block_size + tidx;
//因为实际是以256个为一组的，以y为偏移量，以x为基准量进行总的position的计算
int gpos = (gidx << RADIX) + gidy;
//以行为单位标记组ID
int groupIndex = bidy * (get_global_size(0)/block_size) + bidx;

/* Cache the computational window in shared memory */
//取出每行的数据
block[tidx] = input[gpos];
barrier(CLK_LOCAL_MEM_FENCE);

uint cache = block[0];

/* build the sum in place up the tree */
//就是缩减树算法
for(int dis = 1; dis < block_size; dis *=2)
{
if(tidx>=dis)
{
cache = block[tidx-dis]+block[tidx];
}
//等待第一次缩减完成
barrier(CLK_LOCAL_MEM_FENCE);

block[tidx] = cache;
//等待本次统计结果完成
barrier(CLK_LOCAL_MEM_FENCE);
}

/* store the value in sum buffer before making it to 0 */
//统计每行的加和结果
sumBuffer[groupIndex] = block[block_size-1];

/*write the results back to global memory */

if(tidx == 0)
{

output[gpos] = 0;
}
else
{

output[gpos] = block[tidx-1];
}

}
      这里其中有一个缩减树，可能有些难理解，但是可以通过查资料找到。对核心代码已经做了简略的注释，可以根据图示进行理解。