在opencl中为GPU优化内核代码

Question

截至目前,在内核执行时,我的GPU比我的CPU慢.我想也许是因为我正在测试一个小样本,因为较小的启动开销,CPU最终完成得更快.但是,当我使用几乎是样本大小10倍的数据测试内核时,CPU仍然完成得更快,GPU几乎落后400毫秒.

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

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

我尝试使用本地内存,虽然我100%肯定我使用它错了,并遇到了两个问题.内核在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;
}

尝试设置局部变量时使用的参考链接: 如何在OpenCL中使用本地内存？

用于查找kernelWorkGroupSize的链接(这就是为什么我在kernelArg中设置了1024个): CL_OUT_OF_RESOURCES为2百万个浮点数和1GB VRAM？

我见过其他人有类似的问题,其中GPU比CPU慢,但对于其中许多人来说,他们使用的是clEnqueueKernel而不是clEnqueueNDRangeKernel.

如果您需要有关此内核的更多信息,请继续我之前的问题: 在内核OpenCL中实现FIFO实现的最佳方法

为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接口卡在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参数说明符可以非常轻松地加速这种少量操作,并减少一些整数工作.

对于cpu和gpu,使用float4而不是float类型可以将cpu和gpu的速度提高%80,因为它们是SIMD类型和向量计算单元.

这个内核的瓶颈是:

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

通常:

• 第一次运行内核会及时触发opencl编译器优化,缓慢.至少运行5-10次以获得准确的时间.
• __恒定空间仅为10 - 100 kB但比__global快,对amd的hd5000系列有益.
• 内核开销为100微秒,而65个缓存操作小于此值,并且被内核开销时间所影响(更糟糕的是,由pci-e延迟).
• 工作项目太少会使占用率降低,变慢.

也:

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

• 其中还有83 GB/s的L1和L2缓存带宽,所以只需让它在__global驱动程序优化而不是LDS上运行,除非你打算升级你的计算机.(对于Array ofcourse)也许,来自LDS的奇数元素,甚至来自__global的元素也可能有83 + 83 = 166 GB/s带宽.你可以试试.在银行冲突方面,可能两个比两个更好.

• 使用系数作为__constant(166 GB/s)和数组作为__global,可以为您提供166 + 83 = 249 GB/s的组合带宽.

• 每个系数元素每个线程只使用一次,所以我不打算使用私有寄存器(499 GB/s)