在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可能会出现一些错误，如全局到本地索引重新计算等（我现在即将入睡:)）尽管如此，上面的实现应该使您如何开始使用本地内存正确的方向。