c - 为 GPU 优化 opencl 中的内核代码

Question

截至目前，就内核执行时间而言，我的 GPU 比 CPU 慢。我想也许因为我是用一个小样本进行测试，所以 CPU 由于启动开销较小而最终完成得更快。然而，当我用几乎是样本大小 10 倍的数据测试内核时，CPU 的完成速度仍然更快，而 GPU 落后了将近 400 毫秒。

2.39MB 文件的运行时中央处理器:43.511 毫秒GPU:65.219 毫秒

32.9MB 文件的运行时中央处理器:289.541 毫秒GPU:605.400 毫秒

我尝试使用本地内存，但我 100% 确定我用错了，但遇到了两个问题。内核在 1000-3000 毫秒之间的任何时间完成(取决于我为 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;
}

我尝试设置局部变量时使用的引用链接: How do I use local memory in OpenCL?

用于查找 kernelWorkGroupSize 的链接(这就是我在 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

为 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 驱动程序被 gimped 或 gpu 的 pci-e 接口(interface)停留在 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 类型的 Array 和 Output 应该将您的 cpu 和 gpu 的速度提高 %80，因为它们是 SIMD 类型和 vector 计算单元。

这个内核的瓶颈是:

• 每个线程只有 65 次乘法和 65 次求和。
• 但数据仍然通过 pci-express 接口(interface)传输，速度很慢。
• 还有 1 个条件检查(i < max_i)每个浮点操作很高，需要循环展开。
• 尽管您的 cpu 和 gpu 是基于 vector 的，但一切都是标量。

通常:

• 第一次运行内核会触发 opencl 的即时编译器优化，速度很慢。至少运行 5-10 次以确保准确的时间。
• __constant space 只有 10 - 100 kB，但它比 __global 更快，适合 amd 的 hd5000 系列。
• 内核开销为 100 微秒，而 65 次缓存操作少于此，并且被内核开销时间(甚至更糟糕的是 pci-e 延迟)所掩盖。
• 工作项太少导致占用率低、慢。

还有:

• 4 核 Xeon @ 3 GHz 比 16(vliw5 的 1/4)*2(计算单元)=32 个 gpu 核 @600 MHz 快得多，因为分支预测、总缓存带宽、指令延迟和无- pcie 延迟。
• HD5000 系列 amd 卡是传统的，与 gimped 相同。
• HD5450 具有 166 GB/s 的恒定内存带宽
• 也只有 83 GB/s LDS(本地内存)带宽

• 它还有 83 GB/s 的 L1 和 L2 缓存带宽，所以让它在 __global 驱动程序优化而不是 LDS 上工作，除非你计划升级你的计算机。(当然是阵列)也许，来自 LDS 的奇怪元素，甚至来自 __global 的元素也可以有 83+83 = 166 GB/s 的带宽。你可以试试。就银行冲突而言，两个两个比交替更好。

• 将系数用作 __constant (166 GB/s) 并将 Array 用作 __global 应该可以得到 166 + 83 = 249 GB/s 的组合带宽。

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