为 GPU 优化 opencl 中的内核代码

Question

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

2.39MB 文件的运行时间 CPU：43.511 毫秒 GPU：65.219 毫秒

32.9MB 文件的运行时间 CPU：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;
}

参考link我在尝试设置局部变量时使用的： How do I use local memory in OpenCL?

Link 用于查找 kernelWorkGroupSize（这就是我在 kernelArg 中设置 1024 的原因）： CL_OUT_OF_RESOURCES for 2 millions floats with 1GB VRAM?

我看到其他人也有类似的问题，其中 GPU 比 CPU 慢，但对于他们中的许多人来说，他们使用的是 clEnqueueKernel 而不是 clEnqueueNDRangeKernel。

如果您需要有关此内核的更多信息，请参阅我之前的问题：

还发现了一些针对 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

在介绍本地内存之前，让我们先将 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;
}

P.S。那里可能存在一些错误，例如全局到本地索引的重新计算等。（我现在要睡觉了:)）尽管如此，上面的实现应该让您了解如何开始使用本地内存的正确方向。

Answer 2

运行为 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 是基于向量的，但一切都是标量。

一般：

运行ning内核第一次触发opencl的just in time编译器优化，速度慢。运行至少 5-10 次才能准确计时。
__constant space 只有 10 - 100 kB，但比 __global 快，适合 amd 的 hd5000 系列。
内核开销为 100 微秒，而 65 次缓存操作少于此值，并且被内核开销时间（甚至更糟的是 pci-e 延迟）所掩盖。
工作项太少占用率低，速度慢

另外：

由于分支预测、总缓存带宽、指令延迟和无 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) 和数组作为 __global 应该给你 166 + 83 = 249 GB/s 组合带宽。
每个系数元素每个线程只使用一次所以我不建议使用私有寄存器 (499 GB/s)

为 GPU 优化 opencl 中的内核代码

Optimizing kernel code in opencl for a GPU

c

opencl