在OpenCL中執行掃描

Question

我一直試圖讓一個簡單的掃描工作很長一段時間。對於小問題，輸出是正確的，但是對於大輸出，我有時只能得到正確的結果。我檢查了Apple's OpenCL example，我基本上做同樣的事情（除了銀行衝突，我忽略了atm）。因此，這裏的第一個階段的代碼：

__kernel void 
scan_init(__global int * input, 
      __global int * sums) 
{ 
    int gid = get_global_id(0); 
    int lid = get_local_id(0); 
    int chunk_size = get_local_size(0)*2; 

    int chunk = gid/chunk_size; 
    int offset = chunk*chunk_size; 

    reduction(input, offset); 

    // store sums 
    if(lid==0) 
    { 
    sums[chunk] = input[(chunk+1)*chunk_size-1]; 
    } 

    downsweep(input, offset); 
} 

，還原功能本身：

void reduction(__global int * input, 
     int offset) 
{ 
int stride = 1; 
int grp_size = get_local_size(0); 
int lid = get_local_id(0); 

for(int d = grp_size; d > 0; d>>=1) 
{ 
    barrier(CLK_GLOBAL_MEM_FENCE); 

    if(lid < d) 
    { 
    int ai = stride*(2*lid+1)-1+offset; 
    int bi = stride*(2*lid+2)-1+offset; 
    input[bi] += input[ai]; 
    } 

    stride *= 2; 
    } 
} 

在第二階段，泛音資金用於構建總和爲每個元素：

void downsweep(__global int * input, 
     const unsigned int offset) 
{ 
    int grp_size = get_local_size(0); 
    int lid = get_local_id(0); 
    int stride = grp_size*2; 

    for(int d = 1; d <= grp_size; d *=2) 
    { 
    barrier(CLK_GLOBAL_MEM_FENCE); 

    stride >>=1; 

    if(lid+1 < d) 
    { 
     int src = 2*(lid + 1)*stride-1+offset; 
     int dest = src + stride; 
     input[dest]+=input[src]; 
    } 
    } 
} 

輸入被填充爲本地工作量的倍數。每個工作組都可以掃描兩倍大小的塊。我將每個塊的總和保存在總和數組中，我用它來檢查結果。以下是1的陣列的輸出輸入尺寸4000：

Chunk size: 1024 
Chunks: 4 
Scan global size: 4096 
Local work size: 512 
Sum size: 4 
0:1024 1:1120 2:2904 3:928 

然而，預期的結果將是

0:1024 1:1024 2:1024 3:928 

如果我再次運行該代碼，我得到：

0:1056 1:5376 2:1024 3:928 
0:1024 1:1088 2:1280 3:992 
0:5944 1:11156 2:3662 3:1900 
0:7872 1:1056 2:2111 3:1248 

調用內核如下：

clEnqueueNDRangeKernel(cl_ctx->queue, scan_init, 1, NULL, &scan_global_size, &local_work_size, 0, NULL, NULL); 

如果全球規模是4096和當地大小爲512。如果我限制了本地工作組大小爲64，輸出如下：

0:128 1:128 2:128 3:288 4:128 5:128 6:192 7:192 
8:192 9:254 10:128 11:256 12:128 13:360 14:128 15:128 
16:128 17:128 18:128 19:288 20:128 21:128 22:128 23:128 
24:192 25:128 26:128 27:192 28:128 29:128 30:128 31:32 

如果我改變輸入尺寸爲512和任何塊大小，一切都很好！最後，當使用輸入大小513和256的組大小（也就是說，我有兩個塊，每個塊都有512個元素，第二個塊只有第一個元素設置爲1）時，第一個元素的結果階段是：

0:1 1:2 2:1 3:6 4:1 5:2 6:1 7:14 
8:1 9:2 10:1 11:6 12:1 13:2 14:1 15:28 
16:1 17:2 18:1 19:6 20:1 21:2 22:1 23:14 
24:1 25:2 26:1 27:6 28:1 29:2 30:1 31:56 
32:1 33:2 34:1 35:6 36:1 37:2 38:1 39:14 
40:1 41:2 42:1 43:6 44:1 45:2 46:1 47:28 
48:1 49:2 50:1 51:6 52:1 53:2 54:1 55:14 
56:1 57:2 58:1 59:6 60:1 61:2 62:1 63:148 

，它應該是：

0:1 1:2 2:1 3:4 4:1 5:2 6:1 7:8 
8:1 9:2 10:1 11:4 12:1 13:2 14:1 15:16 
16:1 17:2 18:1 19:4 20:1 21:2 22:1 23:8 
24:1 25:2 26:1 27:4 28:1 29:2 30:1 31:32 
32:1 33:2 34:1 35:4 36:1 37:2 38:1 39:8 
40:1 41:2 42:1 43:4 44:1 45:2 46:1 47:16 
48:1 49:2 50:1 51:4 52:1 53:2 54:1 55:8 
56:1 57:2 58:1 59:4 60:1 61:2 62:1 63:64 

我的猜測是，它是由不同的線程同時訪問相同數據的問題，但是，這不應該是這樣，因爲每個工作組正在處理不同的輸入數據塊。任何關於此事的幫助將不勝感激！

Answer 1

我懷疑問題與barrier（）不是一個工作組間同步。每個工作組都有自己的障礙，而且您對工作組本身的排序沒有任何保證。當您將輸入集大小更改爲512時，您可能會讓您的所有工作組在同一個多處理器上運行，並因此被偶然同步。

你的塊變量是get_group_id（0）/ 2，這意味着你有兩個完整的工作組分配給同一個塊。你可能想要另一種方式。如果它們碰巧以鎖步方式運行，它們會簡單地覆蓋對方的工作，因爲它們的加載存儲庫依賴關係會匹配。否則，它們可能會或可能不會干擾，總是在多次求和值的方向上進行干擾。

在這個問題上的一個提示是在你自己的問題中：「每個工作組都可以掃描兩倍大小的塊。」這應該意味着數組大小的一半是足夠的。

下降掃描（）中的循環也具有奇特性。第一次迭代什麼也不做;蓋+ 1> = 1，並且d從1開始。這可能是一個無足輕重的多餘迭代，但是在規劃中它是一個偏差。