OpenCL内核优化答案

【问题标题】：Optimization of OpenCL kernelOpenCL内核优化
【发布时间】：2011-05-04 16:28:18
【问题描述】：

我在包含大量无符号字符的内核中工作，我使用 clCreateBuffer 创建内存对象。比我通过 clEnqueueWriteBuffer 复制一大块无符号字符到这个内存对象。然后我在循环中调用从这个内存对象读取的内核，执行一些逻辑并将新数据写入同一个地方（在这个循环中我不调用 clEnqueueWriteBuffer 或 clEnqueueReadBuffer）。这是内核代码：

__kernel void test(__global unsigned char *in, unsigned int offset) {
    int grId = get_group_id(0);
    unsigned char msg[1024];
    offset *= grId;

    // Copy from global to private memory
    size_t i;
    for (i = 0; i < 1024; i++)
        msg[i] = in[ offset + i ];

    // Make some computation here, not complicated logic    

    // Copy from private to global memory
    for (i = 0; i < 1024; i++)
        in[ offset + i ] = msg[i];
}

当循环完成时（循环运行 cca 1000 次）然后我通过 clEnqueueReadBuffer 从内存对象读取结果。

这段代码可以优化吗？

【问题讨论】：

您可以使用 clCreateBuffer 从数组中创建一个缓冲区，然后使用 clEnqueueMapBuffer 将其映射到您的原始内存中吗？

标签： c optimization opencl

【解决方案1】：

一些建议：

在内核的开头做一个in += get_group_id(0) * offset。
一次读取 4 个字符（处理 uchar4 或 uint）。
如果可能，也一次处理 4 个字符。
在每个线程中使用 1K 私有数组，工作组大小和占用率将受到严重限制，运行更多线程处理更少字符可能更有效。
似乎每组中的所有线程都将处理完全相同的数据；这可能不是您的想法。

【讨论】：

【解决方案2】：

您可以尝试使用矢量版本（uchar8 而不是 uchar），但编译器可能会以这种方式对其进行优化。最重要的是始终分析您的代码并进行实验。

编辑

现在似乎甚至支持 uchar16： http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/vectorDataTypes.html

【讨论】：

【解决方案3】：

首先想到的是展开循环可以帮助您跳过条件评估。您可以使用this pragma 使其更容易。

在 Nvidia 芯片上使用共享内存也有很大帮助（如果您当前的本地内存默认不使用共享内存）

【讨论】：

【解决方案4】：

对于优化，您需要说明您进行了何种计算。通过将计算分组到工作组并让它们在本地内存上工作，可以获得最大的性能优势。您需要非常注意私有内存（最小）和本地内存（小）的大小。

您的内核多久调用一次？所有内核都使用相同的数据吗？可以想象一个本地内存缓冲区，工作组中的所有线程都将部分数据读入本地内存，然后共享数据。您需要注意同步。

我建议查看 SDK 供应商的示例。我只知道 nVidia SDK。那里的示例非常复杂，但阅读起来非常有趣。

对 float4 等矢量类型的更改应适用于 ATI 板。据说 nVidia 最适合使用标量和内部编译器优化。这是稍后使用分析器进行微调的东西。您可以通过内存优化获得巨大的性能。

【讨论】：