OpenCL：縮小示例，並保留內存對象/將cuda代碼轉換爲openCL

Question

我已經經歷了幾個例子，將一組元素減少爲一個元素，但沒有成功。有人在NVIDIA論壇上發佈了這個消息。我已經從浮點變量更改爲整數。

__kernel void sum(__global const short *A,__global unsigned long *C,uint size, __local unsigned long *L) { 
      unsigned long sum=0; 
      for(int i=get_local_id(0);i<size;i+=get_local_size(0)) 
        sum+=A[i]; 
      L[get_local_id(0)]=sum; 

      for(uint c=get_local_size(0)/2;c>0;c/=2) 
      { 
        barrier(CLK_LOCAL_MEM_FENCE); 
        if(c>get_local_id(0)) 
          L[get_local_id(0)]+=L[get_local_id(0)+c]; 

      } 
      if(get_local_id(0)==0) 
        C[0]=L[0]; 
      barrier(CLK_LOCAL_MEM_FENCE); 
} 

這種看法合適？第三個參數「大小」是否應該是當地工作規模或全球工作規模？

設置我的論點是這樣，

clSetKernelArg(ocReduce, 0, sizeof(cl_mem), (void*) &DevA); 
clSetKernelArg(ocReduce, 1, sizeof(cl_mem), (void*) &DevC); 
clSetKernelArg(ocReduce, 2, sizeof(uint), (void*) &size); 
clSetKernelArg(ocReduce, 3, LocalWorkSize * sizeof(unsigned long), NULL); 

第一個參數是輸入，我想從之前推出的內核的輸出保留。

clRetainMemObject(DevA); 
clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocKernel, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL); 
//the device memory object DevA now has the data to be reduced 

clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocReduce, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL); 
clEnqueueReadBuffer(hCmdQueue[Plat-1][Dev-1],DevRE, CL_TRUE, 0, sizeof(unsigned long)*512,(void*) RE , 0, NULL, NULL); 

今天我打算嘗試將以下cuda縮減示例轉換爲openCL。

__global__ voidreduce1(int*g_idata, int*g_odata){ 
extern __shared__ intsdata[]; 

unsigned int tid = threadIdx.x; 
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; 
sdata[tid] = g_idata[i] + g_idata[i+blockDim.x]; 
__syncthreads(); 


for(unsigned int s=blockDim.x/2; s>0; s>>=1) { 
if (tid < s) { 
sdata[tid] += sdata[tid + s]; 
} 
__syncthreads(); 
} 

// write result for this block to global mem 
if(tid == 0) g_odata[blockIdx.x] = sdata[0]; 
} 

還有一個更優化的，（完全展開+多個每個線程元素）。

http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf

這可能使用的OpenCL？

灰熊給了我這個建議有一天，

」 ......使用的工作n個要素上，並將它們減少了像N/16（或任何其他數字）減少內核然後你。迭代調用該內核，直到你下降到一個元素，這是你的結果「

我想試試這個，但我不完全知道從哪裏開始，我想先得到一些東西上班。

Answer 1

只要您只有一個工作組正在處理縮減工作，您提供的第一個縮減代碼應該工作（所以get_global_size(0) == get_local_size(0)）。在這種情況下，內核的size參數將是A中的元素數（與全局或本地工作無關）。雖然這是一個可行的解決方案，但在進行約簡時讓大部分空閒空間看起來非常浪費，這正是我提出反覆調用縮減內核的原因。這將成爲可能只用輕微修改的代碼：

__kernel void sum(__global const short *A, __global unsigned long *C, uint size, __local unsigned long *L) { 
     unsigned long sum=0; 
     for(int i=get_global_id(0); i < size; i += get_global_size(0)) 
       sum += A[i]; 
     L[get_local_id(0)]=sum; 

     for(uint c=get_local_size(0)/2;c>0;c/=2) 
     { 
       barrier(CLK_LOCAL_MEM_FENCE); 
       if(c>get_local_id(0)) 
         L[get_local_id(0)]+=L[get_local_id(0)+c]; 

     } 
     if(get_local_id(0)==0) 
       C[get_group_id(0)]=L[0]; 
     barrier(CLK_LOCAL_MEM_FENCE); 
} 

與調用此一GlobalWorkSize較小然後size（例如4）將由的4*LocalWorkSize因素之一，它可以迭代減少A輸入（由使用輸出緩衝區作爲下一個調用sum的輸入，並使用不同的輸出緩衝區。其實這並不完全正確，因爲第二次（以及所有後續）迭代需要A爲global const unsigned long*類型，所以您實際上需要內核，但您明白了。關於cuda縮減示例：爲什麼要麻煩轉換它，它的工作原理與我上面發佈的opencl版本完全相同，除了每次迭代只減少一個硬編碼大小（2*LocalWorkSize insted爲size/GlobalWorkSize*LocalWorkSize）。

我個人使用幾乎相同的方式進行的減少，雖然我已經分成兩個部分內核，只使用本地存儲的最後一次迭代使用路徑：

__kernel void reduction_step(__global const unsigned long* A, __global unsigned long * C, uint size) { 
     unsigned long sum=0; 
     for(int i=start; i < size; i += stride) 
       sum += A[i]; 
     C[get_global_id(0)]= sum; 
} 

爲全面的最後一步使用了工作組內部減少的版本。當然，你需要reduction step的第二個版本，取global const short*，這個代碼是你的代碼的一個未經測試的適應（我不能發佈我自己的版本，可惜）。這種方法的優點是內核執行大部分工作的複雜程度要小得多，並且由於分支不同而導致數量減少。這使得它比另一個變體快一點。然而，對於最新的編譯器版本和最新的硬件，我都沒有結果，所以這個觀點可能會也可能不會是正確的（儘管我懷疑這可能是由於發散分支數量的減少）。

現在你鏈接到的文件：當然可以使用opencl中該文件中建議的優化，除了使用opencl不支持的模板外，所以blocksize必須被硬編碼。當然，opencl的版本已經爲每個內核做了多次添加，如果按照我上面提到的方法進行操作，將不會從通過本地內存展開縮減中受益，因爲這隻在最後一步完成，不應該花費整個計算時間的很大一部分足夠大。此外，我發現在展開的實現中缺乏同步有點麻煩。這隻適用於該部分中的所有線程屬於同一個warp。然而，在當前的nvidia卡（未來的nvidia卡，amd卡和cpus）上執行任何硬件時，這並不是必須的（儘管我認爲它應該適用於當前的amd卡和當前的cpu實現，但我不一定會依靠它）），所以我會遠離這一點，除非我需要速度的絕對最後一點的減少（然後仍然提供一個通用的版本，並切換到，如果我不承認硬件或類似的東西）。

Answer 2

縮減內核看起來對我來說是正確的。在縮小中，大小應該是輸入數組A的數字元素。該代碼在sum中累積每個線程的部分總和，然後執行本地內存（共享內存）減少並將結果存儲到C。您將在每個本地工作組中獲得一份C的部分款項。要麼再次用一個工作組調用內核以獲得最終答案，要麼在主機上累積部分結果。