OpenCL/CUDA：兩階段約簡算法

可以通過調用__reduce（）來減少大數組。多次。OpenCL/CUDA：兩階段約簡算法

下面的代碼但只使用兩個階段，並記錄在案here：

但是我無法理解這一兩個階段減少的算法。有人可以給出一個更簡單的解釋嗎？

__kernel 
void reduce(__global float* buffer, 
     __local float* scratch, 
     __const int length, 
     __global float* result) { 

    int global_index = get_global_id(0); 
    float accumulator = INFINITY; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
     float element = buffer[global_index]; 
     accumulator = (accumulator < element) ? accumulator : element; 
     global_index += get_global_size(0); 
    } 

    // Perform parallel reduction 
    int local_index = get_local_id(0); 
    scratch[local_index] = accumulator; 
    barrier(CLK_LOCAL_MEM_FENCE); 
    for(int offset = get_local_size(0)/2; offset > 0; offset = offset/2) { 
     if (local_index < offset) { 
      float other = scratch[local_index + offset]; 
      float mine = scratch[local_index]; 
      scratch[local_index] = (mine < other) ? mine : other; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
    } 
    if (local_index == 0) { 
     result[get_group_id(0)] = scratch[0]; 
    } 
}

它也可以使用CUDA很好地實現。

來源

2012-06-11 gpuguy

此並行簡化代碼是規範CUDA簡化的簡單OpenCL端口（缺少OpenCL無法完成的幾項優化）。它由NVIDIA的[Mark Harris]（http://stackoverflow.com/users/749748/harrism）編寫。您可以在CUDA SDK [縮小示例]（http://developer.nvidia.com/cuda-cc-sdk-code-samples#reduction）中找到一個非常有教育意義的白皮書。閱讀完白皮書後，編輯您的問題以解釋您不明白的內容，並且有人可能會進一步提供幫助。 – talonmies

感謝您指向正確的方向。那麼我很難理解兩階段和多階段縮減內核之間的區別，因爲[在這裏]（http://developer.amd.com/Membership/Print.aspx?ArticleID=221&web=http://developer.amd .com /文檔/文章） – gpuguy

您創建N線程。第一個線程查看位置0，N，2 * N處的值......第二個線程查看值1，N + 1，2 * N + 1，...這是第一個循環。它將length值減少爲N個值。

然後每個線程將其最小值保存在共享/本地內存中。然後你有一個同步指令（barrier(CLK_LOCAL_MEM_FENCE)）。然後你在共享/本地內存中的標準減少。完成後，本地ID爲0的線程將其結果保存在輸出數組中。

總而言之，您的值從length減少到N/get_local_size(0)的值。此代碼執行完畢後，您需要執行最後一次傳遞。但是，這可以完成大部分工作，例如，您可能具有長度〜10^8，N = 2^16，get_local_size（0）= 256 = 2^8，並且此代碼將10^8個元素減少爲256個元素。

哪部分你不明白？

來源

2012-06-11 07:28:41 user434507

嗨，我也卡在減少過程中。那麼我試圖在OpenCL中實現排序算法。問題是，我可以舒適地將減少應用在當地記憶中，這對許多人來說並不是什麼汗水。我無法做的是將存在於本地數組中的已排序數組減少/合併/合併到一個大的完全排序數組中。 Offcourse最後我可以在全局內存上運行排序算法，但我不認爲這是要做的。請幫助我...... – Yash

我也有理解這個問題。我的理解是這個內核將問題簡化爲'get_num_groups（0）'元素。這將是'get_global_size（0）/ get_local_size（0）'並且由內核排隊時設置的NDRange定義。所以主機代碼對於這個內核是必不可少的，但遺憾的是從未在原始的AMD頁面上提供。我認爲匹配'CL_DEVICE_MAX_COMPUTE_UNITS * CL_DEVICE_MAX_WORK_GROUP_SIZE'的NDRange可以完成順序部分的最大並行性。糾正我，如果我錯了。 – Armin

OpenCL/CUDA：兩階段約簡算法

回答

相關問題