使用OpenCL优化GPU内核代码

Question

目前，我的GPU在内核执行时间方面比我的CPU慢。我以为可能是因为我测试的样本较小，CPU由于启动开销较小而最终更快完成。然而，当我用近10倍于样本大小的数据测试内核时，CPU仍然完成得更快，GPU落后近400ms。

2.39MB文件运行时间 CPU：43.511ms GPU：65.219ms

32.9MB文件运行时间 CPU：289.541ms GPU：605.400ms

我尝试使用本地内存，虽然我100%确定我用错了，遇到了两个问题。内核在任何设置本地工作大小的情况下都会在1000-3000ms之间完成，或者我会遇到状态码-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;
}

我尝试设置本地变量时使用的参考链接： 如何在OpenCL中使用本地内存？ 用于查找kernelWorkGroupSize的链接（这就是为什么我在kernelArg中设置了1024）： 1GB VRAM中有200万个浮点数的CL_OUT_OF_RESOURCES？ 我看到其他人也有类似的问题，其中GPU比CPU更慢，但对于他们中的许多人来说，他们正在使用clEnqueueKernel而不是clEnqueueNDRangeKernel。

如果您需要有关此内核的更多信息，请参阅我的以前的问题： FIFO实现在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;
}

- VedhaR

---

我非常确定你真的不想在内部的 for 循环中使用 if() 语句。一个聪明的编译器也许能够将 if 提升出循环，但 GPU 驱动程序可能没有足够的时间和智能来有效地执行此操作。 - EOF

你正在解决/实现哪个问题/算法？ - mfa

@mfa 我正在尝试实现一个低通FIR滤波器。滤波器效果很好，现在只是需要减少内核执行时间的问题。 - VedhaR

@VedhaR：你到底为什么需要一个switch语句？那可能更糟糕。让我们简单地看一下条件语句，好吗？if (globalId+i > 63)显然等同于if (globalId+i >= 64)，而这显然等同于if (i >= 64 - globalId)，这可以通过改变for循环的初始化语句轻松消除：for (int i=64-globalId; i<65; i++)。 - EOF

@EOF 如果我在CPU上运行它，它仍然会崩溃。调试器告诉我：“在65elementmult.exe的0x003002E9处发生未处理的异常：0xC0000005：访问位置0x0BF6BFFC时违规。”这条消息会针对其余的消息重复出现，但地址位置不同。如果我在GPU上运行它，一切都会冻结，显示驱动程序也会失败。我将添加导致此问题的代码。您需要哪些其他信息？在循环中，我只是用i替换了索引，因为我试图调试它为什么会崩溃。 - VedhaR

Answer 1

运行以下内核以处理2400万元素数组

__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;
}

对于一个25个计算单位的设备池，处理速度在200毫秒以内完成，但对于一个8核CPU则需要超过500毫秒。

可能是您拥有高端CPU和低端GPU，或者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;
}

每个cpu（8个计算单元）的计算时间不超过150毫秒，每个gpu（25个计算单元）的计算时间不超过80毫秒。每项工作只需65次操作。使用__constant和__read_only和__write_only参数说明符以及一些整数工作减少，可以很容易地加速这个低操作数。

对于数组和输出，使用float4而不是float类型应该会使您的cpu和gpu速度提高80％，因为它们都是SIMD类型和向量计算单元。

此内核的瓶颈包括：

• 每个线程仅有65个乘法和65个求和。
• 但数据仍然通过pci-express接口传输，速度慢。
• 每个浮点运算有1个条件检查（i < max_i），需要展开循环。
• 所有东西都是标量，尽管您的cpu和gpu是基于向量的。

通常情况下：

• 第一次运行内核会触发opencl的即时编译器优化，速度较慢。至少运行5-10次以获得精确的时间。
• __constant空间仅为10-100kB，但比__global更快，非常适合amd的hd5000系列。
• 内核开销为100微秒，而65个缓存操作小于该时间，并被内核开销时间（甚至更糟的是，由于pci-e延迟）所掩盖。
• 工作项太少会使占用率降低，速度慢。

另外：

• 4核Xeon @ 3 GHz比32核gpu @ 600 MHz快得多（1/4 vliw5），因为具有分支预测、总缓存带宽、指令延迟和无pcie延迟。
• HD5000系列amd卡已经过时，与gimped相同。
• HD5450具有166 GB/s的常量内存带宽
• 它还只有83 GB/s的LDS（本地内存）带宽

• 除非您计划升级计算机（针对数组），否则让它在__global驱动程序优化上工作，而不是LDS。也许，从LDS中取奇数元素，从__global中取偶数元素，可以获得83 + 83 = 166 GB/s的带宽。您可以尝试。也许两两结对比交替更好，以避免银行冲突。

• 将系数用作__constant（166 GB/s），将数组用作__global，应该为您提供166 + 83 = 249 GB/s的组合带宽。

• 每个系数元素仅在每个线程中使用一次，因此我不建议使用私有寄存器（499 GB/s）

- huseyin tugrul buyukisik

---

优化后的代码速度提高了3倍，但不确定是否有正确的输出。 - huseyin tugrul buyukisik

另外，HD5450是一种矢量架构，而您的内核是标量类型，因此CPU和GPU都未得到充分利用。您应该将其更改为矢量版本。我会尝试同时进行。但是矢量类型使其非常困难，新的GPU技术今天是标量。 - huseyin tugrul buyukisik

它的常量内存比本地内存更快，因此系数应该使用常量内存而不是本地内存。也许通过将系数共享到两个内存中同时使用可以得到最佳结果，但当存在PCI-E瓶颈时这并不值得。 - huseyin tugrul buyukisik

此外，在我们将if语句移除的情况下，将系数从全局变量改为常量也会增加执行时间。我猜它们是在特定情况下使用的？或者可能需要在内核代码中进行适应（我的意思是你不能仅仅从__global更改为__constant，内核代码本身还需要更改一些其他变量）。 - VedhaR

这是使用 CPU 还是 GPU？ - huseyin tugrul buyukisik

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;
}

附注：可能存在一些错误，例如全局到本地索引的重新计算等（我现在要去睡觉了 :)）尽管如此，上述实现应该让你朝着如何开始使用本地内存的正确方向前进。

- doqtor

---

谢谢你的回答！我可以说，去掉if语句使内核绑定时间缩短了150毫秒。然而，添加本地内存几乎使其跳到了900毫秒（是原来的两倍）。但是，使用你提供的最后一种实现方式，我想现在我可以制作狂欢音乐了哈哈，它以最奇怪的方式改变了歌曲。 - VedhaR

但是我明白了，与其使用全局内存引用系数，我可以将这些值带入本地并以此方式使用（应该更快且系数不会改变）。然而，在这种情况下，localId的值是多少？ - VedhaR

在你的例子中，'localWorkSize [0] = 64;'，我也使用了相同的值。为了将数据从__global复制到__local缓冲区，前64个工作项会分别复制前64个值（每个工作项复制一个值，因为__local表示该缓冲区是共享/可见的），然后第一个工作项复制最后一个值。 - doqtor

好的，但是如果我有多个__local变量声明，并且每个变量都有不同的大小，get_local_id(x)会返回按它们初始化顺序声明的变量的值吗？在这种情况下，x是被初始化的第x个元素。抱歉这个问题有点奇怪，如果不清楚，请告诉我。 - VedhaR

这篇博客文章应该能为您澄清相关问题。 - doqtor

顺便说一下，我刚刚注意到了一些奇怪的事情。我决定将最终结果输出到文本文件中，我注意到前十几个输出值变化正常，但之后就一直停留在0.99985，直到最后十几行才再次变化。我很确定我以前也遇到过这个问题。 - VedhaR