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

       首先我们先整理下我们目前通过上面几讲的内容得到的结果：

1.      sumBufferin:组加和结果

2.      summaryBUfferout:组加和结果进行偏移量整理的结果

3.      scanedHistogramBinBuf：以组为单位进行偏移量计算结果

       那么这些结果与最后我们需要的偏移量有什么关系，我们先来看下最后一个内核函数：

FixOffsetkernel。同样先看下整体的结构，整体的分配情况和scanArrayKerneldim2相同在前面我们已经画过结构，这里不再重复。在这个内核前面有这样一段英文注释：

 /*run fixoffset kernel: for each row of the 2nd dimention ,add thesummary of the previous

row computed bythe ScanArraysdim1 kenel ,(the 1st row needn't be changed!), now the correct  position has been computed out*/

       翻译过来大概就是这样：运行fixoffset内核：对于每一个2维中的行，加上在ScanArraysdim1内核中得到的行就是我们最终需要的偏移量结果（通过查看代码实际上是列值相加更加符合）。云里雾里，为什么加上之后就是我们需要的了？凭啥？是吧。呵呵，不急，我们先回想下我们得到的结果，然后看看内核代码先：

  __kernel void FixOffset(__global uint* input,__global uint* output)

   {

         int gidx = get_global_id(0);

         int gidy = get_global_id(1);

         int gpos = gidy + (gidx << RADIX );

         output[gpos] += input[gidy];

  }         

       这内核代码贼拉简单了，果然就是以列为单位，每一列的对应元素加上summaryBUffout的对应元素就完了。干嘛用的？

       我们先从最初的状态来看下，首先我们把数据分成了以256个数据为一组。既然分组，那么如果需要对这256个数据以组为单位进行基数排序，就应该需要知道这一组中每个索引的对应偏移量，其中索引为0的偏移量无疑是0，索引为1的偏移量应该是多少？这应该是所有问题的关键所在。

      事实上，这里有全局计数和局部计数两个概念（这是我想的，只是为了描述方便）。全局计数就是索引为0~256的全局桶的统计结果（以偏移量存储，对应summaryBUffout）。局部计数则需要以已经完成了全局桶的多少个排序为偏移。举一个简单的例子，估计大家就明白了。对于第二组block[1]中的block[1][1]也就是第二个256个数据的索引为[1]的偏移量来说，它需要加和的数据首先是summary[1],那么summary[1]里面存储的是什么呢？是全局总共有多少个索引为0的数据。因为是在block[1]所以，已经完成的可以看做block[0],也就是scanedHistogramBinsBuf（没有经过最终计算的结果,即相对前block块的偏移量加和结果）中的block[0][1]的结果。（这里为了描述，自己定义了一些变量来说明，只是为了帮助理解）。

      所以FixOffset与RadixSort的思路也就逐渐清晰了，这里做一下总结：

      首先对所有数据进行预处理，预处理后数据个数必须为256*64=16384的倍数。然后将数据以256个为一组进行分组，分组后通过几个内核函数，计算出每一组中相对全局的偏移量结果。得到每一组的偏移量结果后，根据偏移量结果轮询每组数据放到排序后对应的位置。循环进行32/8=4此上述操作直到排序完成。