雙重減少opencl教程

Question

我是一個OpenCl的新手。

我需要在一維雙精度數組上運算一個約化（和運算符）。

我一直在網上游蕩，但我發現的例子很混亂。 任何人都可以發佈一個容易閱讀（也可能是高效的）教程實施？

附加信息： - 我可以訪問一個GPU設備; - 我使用C爲內核代碼

Answer 1

您提到您的問題涉及60k雙打，這將不適合您的設備的本地內存。我把一個內核放在一起，將你的矢量減少到10-30個左右的值，你可以將它與你的主機程序相加。我在我的機器上遇到雙打問題，但是如果您啓用雙打併將「浮動」更改爲「雙倍」，則此內核應該可以正常工作。我將調試我遇到的雙重問題，併發布更新。

PARAMS：

• 全球浮動* inVector - 花車源總結
• 全球浮動* outVector - 花車的列表，每個工作組
• const int的inVectorSize - 總inVector持有的花車數量
• 本地浮動* resultScratch - 每個工作組使用的本地內存。您需要爲組中的每個工作項目分配一個浮點數。預期的大小= sizeof（cl_float）* get_local_size（0）。例如，如果您爲每個組使用64個工作項，則這將是64個浮點數= 256個字節。切換到雙打將使其512字節。由openCL規範定義的最小LDS大小爲16kb。 See this question有關傳遞本地（NULL）參數的更多信息。

用法：

1. 分配內存爲輸入和輸出緩衝器。
2. 爲設備上的每個計算單元創建一個工作組。
3. 確定最佳工作組大小，並使用它來計算'resultScratch'的大小。
4. 調用內核，讀出向量返回主機
5. 循環遍歷outVector的副本並添加以獲得最終總和。

潛在的優化：

1. 像往常一樣，你要調用了大量數據的內核。太少的數據並不值得傳輸和設置時間。

2. 使inVectorSize（和矢量）爲（工作組大小）*（工作組數）的最高倍數。只用這些數據調用內核。內核均勻地分割數據。在等待回調期間計算主機上任何剩餘數據的總和（或者，爲CPU設備構建相同的內核並僅傳遞剩餘數據）。在上面的步驟＃5中添加outVector時，從這個總和開始。這種優化應該保持工作組在整個計算過程中均勻飽和。

   __kernel void floatSum(__global float* inVector, __global float* outVector, const int inVectorSize, __local float* resultScratch){ 
       int gid = get_global_id(0); 
       int wid = get_local_id(0); 
       int wsize = get_local_size(0); 
       int grid = get_group_id(0); 
       int grcount = get_num_groups(0); 
   
       int i; 
       int workAmount = inVectorSize/grcount; 
       int startOffest = workAmount * grid + wid; 
       int maxOffest = workAmount * (grid + 1); 
       if(maxOffset > inVectorSize){ 
        maxOffset = inVectorSize; 
       } 
       resultScratch[wid] = 0.0; 
       for(i=startOffest;i<maxOffest;i+=wsize){ 
         resultScratch[wid] += inVector[i]; 
       } 
       barrier(CLK_LOCAL_MEM_FENCE); 
   
       if(gid == 0){ 
         for(i=1;i<wsize;i++){ 
           resultScratch[0] += resultScratch[i]; 
         } 
         outVector[grid] = resultScratch[0]; 
       } 
   

   }

此外，啓用雙打：

#ifdef cl_khr_fp64 
#pragma OPENCL EXTENSION cl_khr_fp64 : enable 
#else 
#ifdef cl_amd_fp64 
#pragma OPENCL EXTENSION cl_amd_fp64 : enable 
#endif 
#endif 

更新：AMD APP KernelAnalyzer得到了一個更新（V12），它的顯示，這個內核的雙精度版本實際上，ALU綁定在5870和6970卡上。