在GPU中優化opencl中的內核代碼

Question

截至目前，就內核執行時間而言，我的GPU比我的CPU慢。我想也許因爲我正在測試一個小樣本，由於較小的啓動開銷，CPU最終以更快的速度完成。但是，當我測試內核的數據幾乎是樣本大小的10倍時，CPU仍然以更快的速度完成，而GPU幾乎落後了400毫秒。

運行與2.39MB文件 CPU：43.511ms GPU：65.219ms

運行與32.9MB文件 CPU：289.541ms GPU：605.400ms

我使用本地內存試過，但我我確信我錯誤地使用了它，並且遇到了兩個問題。內核在1000-3000ms之間的任何地方完成（取決於我爲localWorkSize設置的大小），還是運行狀態代碼-5，即CL_OUT_OF_RESOURCES。

這是一位SO成員幫助我的內核。

__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output) { 

int globalId = get_global_id(0); 
float sum=0.0f; 
for (int i=0; i< 65; i++) 
{ 
    float tmp=0; 
    if (globalId+i > 63) 
    { 
     tmp=Array[i+globalId-64]*coefficients[64-i];  

    } 

    sum += tmp; 

} 
Output[globalId]=sum; 
} 

這是我嘗試使用本地內存。第一位將是主機代碼的一個片段，以下部分是內核。

//Set the size of localMem 
status |= clSetKernelArg(
    kernel, 
    2, 
    1024, //I had num_items*(float) but it gave me a -5. Num items is the amount of elements in my array (around 1.2 million elements) 
    null); 
printf("Kernel Arg output status: %i \n", status); 

//set a localWorkSize 
localWorkSize[0] = 64; 

//execute the kernel with localWorkSize included 
status = clEnqueueNDRangeKernel(
    cmdQueue, 
    kernel, 
    1, 
    NULL, 
    globalWorkSize, 
    localWorkSize, 
    0, 
    NULL, 
    &someEvent); 


//Here is what I did to the kernel*************************************** 
__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output, __local float *localMem) { 

int globalId = get_global_id(0); 
int localId = get_local_id(0); 

localMem[localId] = globalId[globalId]; 

float sum=0.0f; 
for (int i=0; i< 65; i++) 
{ 
    float tmp=0; 
    if (globalId+i > 63) 
    { 
     tmp=localMem[i+localId-64]*coefficients[64-i]; 

    } 

    sum += tmp; 

} 
Output[globalId]=sum; 
} 

參考鏈接我想設置局部變量時使用：用於查找kernelWorkGroupSize How do I use local memory in OpenCL?

鏈接（這就是爲什麼我在kernelArg 1024集）： CL_OUT_OF_RESOURCES for 2 millions floats with 1GB VRAM?

我已經看到其他人在GPU比CPU慢的情況下也存在類似的問題，但是對於其中的許多人來說，他們使用clEnqueueKernel而不是clEnqueueNDRangeKernel。

繼承人我剛纔的問題，如果你需要對這個內核的詳細信息： Best approach to FIFO implementation in a kernel OpenCL

發現GPU的藏漢一些優化技巧。 https://developer.amd.com/wordpress/media/2012/10/Optimizations-ImageConvolution1.pdf

已編輯的代碼;錯誤仍然存​​在

__kernel void lowpass2(__global float *Array, __global float *coefficients, __global float *Output) { 

int globalId = get_global_id(0); 
float sum=0.0f; 
float tmp=0.0f; 
for (int i=64-globalId; i< 65; i++) 
{ 

tmp = 0.0f; 
tmp=Array[i]*coefficients[i];  
sum += tmp; 

} 
Output[globalId]=sum; 
} 

Answer 1

運行下面的內核爲24000000個元件陣列

__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output) { 

int globalId = get_global_id(0); 
float sum=0.0f; 
for (int i=0; i< 65; i++) 
{ 
    float tmp=0; 
    if (globalId+i > 63) 
    { 
     tmp=Array[i+globalId-64]*coefficients[64-i];  

    } 

    sum += tmp; 

} 
Output[globalId]=sum; 
} 

下200毫秒對於25計算單元設備池完成，但500毫秒爲一個8芯的CPU。

要麼你有一個高端的CPU和一個低端的CPU，或者GPU驅動已經被縮小，或者GPU的PCI-E接口卡在PCI-E 1.1 @ 4X帶寬，所以主機和設備之間的陣列副本是有限的。

在另一方面，這個優化版本：

__kernel void lowpass(__global __read_only float *Array,__constant float *coefficients, __global __write_only float *Output) { 

     int globalId = get_global_id(0); 
     float sum=0.0f; 
     int min_i= max(64,globalId)-64; 
     int max_i= min_i+65; 
     for (int i=min_i; i< max_i; i++) 
     { 
      sum +=Array[i]*coefficients[globalId-i];  
     } 
     Output[globalId]=sum; 
} 

具有下150毫秒爲CPU（8計算單元）和80毫秒爲GPU（25計算單元）計算下次。每件作品只有65次。使用__constant和__read_only和__write_only參數說明符以及一些減少的整數工作，可以非常輕鬆地加速這種少量的操作。

對於陣列和輸出，使用float4而不是浮點類型應該爲cpu和gpu提高速度80％，因爲它們是SIMD類型和矢量計算單位。這個內核的

瓶頸是：

• 只有65乘法和每線65個求和。
• 但是數據仍然通過pci-express接口傳輸，速度很慢。
• 另外1個條件檢查（i < max_i）每次浮動操作都很高，需要循環展開。
• 儘管您的cpu和gpu都是基於矢量的，但一切都是標量。

一般：

• 首次運行的內核只是在OpenCL中，慢的時候編譯器優化觸發。精確計時至少運行5-10次。
• __constant空間僅爲10 - 100 kB，但其速度超過__global，適用於amd的hd5000系列。
• 內核開銷爲100微秒，而65次緩存操作小於此時間並受到內核開銷時間（甚至更糟糕的是，通過pci-e延遲）的影響。
• 工作項目太少會使職業比例變小，變慢。

另外：

• 4核至強@ 3 GHz的比GPU的16 *（VLIW5的1/4）2（計算單元）= 32芯快得多@因爲分支的600兆赫預測，總緩存帶寬，指令延遲和無計時延遲。
• HD5000系列amd卡是遺留的，與gimped相同。
• HD5450具有166 GB/s的恆定存儲器帶寬
• 其中也有僅83 GB/s的LDS（本地存儲器）帶寬

• 其中也有83 GB/s的L1和L2高速緩存帶寬，使得就讓它除非你計劃升級你的計算機，否則你可以使用__全局驅動程序優化，而不是LDS。（對於Array of course而言）也許，來自LDS的奇怪元素甚至來自__global的元素可能具有83 + 83 = 166 GB/s的帶寬。你可以試試。在銀行衝突方面，或許兩兩比較好。

• 使用係數作爲__constant（166 GB/s）和Array作爲__global應該會給你166 + 83 = 249 GB/s的組合帶寬。

• 每個係數元素用於一次每個線程，所以我不建議使用專用寄存器（499 GB /秒）

Answer 2

引入本地內存讓我們先來舉if聲明圈外前：

__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output) 
{ 
int globalId = get_global_id(0); 
float sum=0.0f; 
int start = 0; 
if(globalId < 64) 
    start = 64-globalId; 
for (int i=start; i< 65; i++) 
    sum += Array[i+globalId-64] * coefficients[64-i];  
Output[globalId]=sum; 
} 

然後引進本地存儲可以這樣來實現：

__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output) 
{ 
    int globalId = get_global_id(0); 
    int local_id = get_local_id(0); 

    __local float local_coefficients[65]; 
    __local float local_array[2*65]; 

    local_coefficient[local_id] = coefficients[local_id]; 
    if(local_id == 0) 
     local_coefficient[64] = coefficients[64]; 
    for (int i=0; i< 2*65; i+=get_local_size(0)) 
    { 
     if(i+local_id < 2*65) 
      local_array[i+local_id] = Array[i+global_id]; 
    } 
    barrier(CLK_LOCAL_MEM_FENCE); 

    float sum=0.0f; 
    int start = 0; 
    if(globalId < 64) 
     start = 64-globalId; 
    for (int i=start; i< 65; i++) 
     sum += local_array[i+local_id] * local_coefficient[64-i];  
    Output[globalId]=sum; 
} 

PS可能會出現一些錯誤，如全局到本地索引重新計算等（我現在即將入睡:)）儘管如此，上面的實現應該使您如何開始使用本地內存正確的方向。