OpenCL 优化答案

【问题标题】：OpenCL OptimizationOpenCL 优化
【发布时间】：2017-04-10 11:43:38
【问题描述】：

我是 OpenCL 的新手。我写了一个 OpenCL 内核来计算灰度。如何优化该代码，可能吗？为什么计算时间浮动这么多？有时我不会加速其他人。我做错了什么？

kernel code:
kernel void grayscale(__global unsigned char *input)
{
    size_t i = get_global_id(0);    

    float grayscaleValue = (input[i*3] * 0.299F) + (input[i*3+1] * 0.587F) + (input[i*3+2] * 0.114F);

    input[i*3] = grayscaleValue;
    input[i*3+1] = grayscaleValue;
    input[i*3+2] = grayscaleValue;   
}


cpu code:
void GrayScaleCPU(struct PPMFile *ppmStruct)
{    
    for (int i = 0; i < ppmStruct->imageSize; i+=3)
    {
        float greyscaleValue = (ppmStruct->data[i] * 0.299F) + (ppmStruct->data[i+1] * 0.587F) + (ppmStruct->data[i+2] * 0.114F);
        ppmStruct->out[i] = greyscaleValue;
        ppmStruct->out[i+1] = greyscaleValue;
        ppmStruct->out[i+2] = greyscaleValue;
    }
}

int main(void)
{
    struct timespec tS1, tS2;

    tS1.tv_sec = 0;
    tS1.tv_nsec = 0;

    tS2.tv_sec = 0;
    tS2.tv_nsec = 0;

    ...

    clock_settime(CLOCK_REALTIME, &tS1);
    GrayScaleCPU(ppmf);
    clock_gettime(CLOCK_REALTIME, &tS1);

    printf ("Timming took %.12lu seconds to run.\n", tS1.tv_nsec);

    ...

    clock_settime(CLOCK_REALTIME, &tS2);
    GrayScaleOpenCL(ppmf2);
    clock_gettime(CLOCK_REALTIME, &tS2);

    printf ("Timming took %.12lu seconds to run.\n", tS2.tv_nsec);

    float time2 = tS2.tv_nsec;
    float time1 = tS1.tv_nsec;
    float speedup = time2/time1;

    printf ("Speed UP OpenCL/CPU %.20f.\n", speedup);
    return 0;
}

【问题讨论】：

您在测量什么以及如何测量？
我正在测量两个函数中的时间。在 opencl 的情况下，我正在测量与准备午餐内核有关的所有事情。使用结构时间规范。这是正确的？ @Jovasa
所以您正在衡量创建命令队列、上下文和构建程序？如果是这样，你做错了，你也应该测量十次左右的执行，放弃第一次，因为它通常比较慢，然后取平均值。
是的，我测量一切，不应该那样做吗？我应该只测量内核“clEnqueueNDRangeKernel”的执行吗？关于设备之间的内存传输我也应该丢弃吗？
内存事务可以包含在内，但在现代 GPU 上，它们与计算并行发生（第二次迭代计算可能与第一次迭代数据被读回同时发生）。但是是的，你应该只测量内核执行时间和数据传输检查这个stackoverflow.com/questions/7980090/… 并且还有一些其他问题也应该对你有所帮助

标签： opencl

【解决方案1】：

尝试将全局内存缓冲到线程内存中：

unsigned char l_input0 = input[i*3];
unsigned char l_input1 = input[i*3 + 1];
unsigned char l_input2 = input[i*3 + 2];
//compute grayscale using l_input0,1,2
input[i*3] = grayscale;
input[i*3 + 1] = grayscale;
input[i*3 + 2] = grayscale;

此外，如果您在调用内核时数据没有正确间隔，您最终可能会在每个 unsigned char 上执行，而不是像 for 循环示例中那样每第三个 unsigned char 执行。

然后，您可以进一步使用本地内存和工作组，并以块的形式进行计算，但这更具挑战性，因为本地工作大小非常特定于设备并且需要是全局工作大小的倍数。我发现 16、32 和 64 的本地工作大小适用于大多数设备。

最后，您对 OpenCL 进行基准测试，确保您测量的是内核性能，而不是内核排队时间。最简单的方法是启动一个计时器，将内核排入队列，在队列上调用 clainish，然后停止计时器。大多数 OpenCL 设备都内置了时序和分析功能，由队列处理。

【讨论】：

为什么我不应该测量排队时间？它属于进程，不是吗？ @奥斯汀
当然，你可以测量它，但它是 A) 微不足道的，B) 仅仅因为内核入队并不意味着它已经完成，所以要小心你是如何收集你的计时数据的。这就是警告。