本篇教程中，我們學習一下如何用opencl有效實作數組求和，也就是通常所說的reduction問題。
在程式中，我們設定workgroup size為256，kernel的輸入、輸出緩沖參數都用uint4的格式，這樣我們原始求和的數組大小為256*4的倍數，資料類型為uint。我們設定每個workgroup處理處理512個uint4，即2048個uint
為了簡便期間，我們輸出數組長度定為4096，即需要2個workgruop來處理。
kernel代碼如下：
__kernel void reduce(__global uint4* input, __global uint4* output, __local uint4* sdata)
{
// 把資料裝入lds
unsigned int tid = get_local_id(0);
unsigned int bid = get_group_id(0);
unsigned int gid = get_global_id(0);
unsigned int localSize = get_local_size(0);
unsigned int stride = gid * 2;
sdata[tid] = input[stride] + input[stride + 1];
barrier(CLK_LOCAL_MEM_FENCE);
// 在lds中進行reduction操作，得到數組求和的結果
for(unsigned int s = localSize >> 1; s > 0; s >>= 1)
{
if(tid < s)
{
sdata[tid] += sdata[tid + s];
}
barrier(CLK_LOCAL_MEM_FENCE);
}
// 把一個workgroup計算的結果輸出到輸出緩沖，是一個uint4，還需要在host端再進行一次reduction過程
if(tid == 0) output[bid] = sdata[0];
}
在程式中，global和local的NDRange，我們都用一維的形式。下面以圖的方式看下kernel代碼是如何執行的：
對第一個workgroup中的第一個thread的來說，它首先進行一次reduction操作，把兩個uint4相加，放到lds(shared memory)中，然後再在lds中進行reduction操作，此時要從global memory中取資料，可以看出連續的thread通路連續的global memory，這時可以利用合并讀寫。
申請的shared memory大小為groupsize*sizeof(uint4)，相加後uint4放入32bank的lds中，放置的方式應該是如下圖所示，因為放入的是uint4，是以會放入連續的4個bank中(每個bank都是dword寬)，可見隻能同時有8個thread通路lds，是以會有一定程式的bank conflit。從App profiler session，我們可以看到：
接下來，kernel會通過一個for循環疊代執行reduction操作，求得一個workgroup中的uint4的和。
疊代的第一次s=128,這時會執行如下圖的兩兩相加，workgroup中同時執行的thread為128,thread local id大于等于128的線程都不會做什麼事情，在每個循環的末尾，有一個barrier來同步所有thread，以便所有thread都完成這次循環後再進入下一次循環。
第二次疊代的時候，隻剩下前面128個uint4，workgroup中同時執行的thread為64。最後，當s=1時候，完成疊代reduction操作，然後把thread0（第一個thread)的結果輸出。
在host段，我們還要做一次相加操作，把不同workgroup得到的uint4，拆分成uint，并相加求得最終的結果。
//在cpu reduction各個workgroup的結果以及uint4分量 reduction
output = 0;
for(int i = 0; i < numBlocks * VECTOR_SIZE; ++i)
output += outMapPtr[i];
printf("gpu reduction result:%d\n", output);
if(refOutput==output) printf("passed\n");
程式執行後結果如下：
完整的代碼請參考：
工程檔案gclTutorial11
代碼下載下傳：
稍後提供