OpenCL内核的优化

Question

我在内核中使用大量无符号字符工作，我用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内存对象的结果。

可以优化这段代码吗？

Answer 1

你可以尝试矢量版本（uchar8而不是uchar），但是编译器可以用这种方式优化它。 最重要的是您的代码始终在进行实验。

编辑

好像连uchar16现在支持： http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/vectorDataTypes.html

Answer 2

东西，首先穿越心灵，是您展开循环可以帮助你跳过状态评估。您可以使用this pragma使其更容易。

使用NVIDIA的芯片共享存储器还可以极大地帮助（如果您的当前本地MEM不使用共享内存默认情况下）

Answer 3

对于你需要解释你做什么样的计算优化。性能的最大好处可以通过将计算分组到工作组中并让它们在本地内存上工作来获得。您需要特别注意私有内存（最小）和本地内存（小）的大小。

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

我建议看看SDK供应商的样品。我只知道nVidia SDK。那里的样本非常复杂，但非常有趣。

像float4shall这样的向量类型的改变适用于ATI板。据说nVidia在标量和内部编译器优化方面效果最佳。这是稍后用分析器进行微调的情况。你可以通过内存优化获得大量的性能。

Answer 4

几点建议：

• 在内核的开始做单in += get_group_id(0) * offset。
• 一次读取4个字符（在uchar4或uint上工作）。
• 如果可能，一次也处理4个字符。
• 在每个线程中都有一个1K专用阵列，工作组大小和占用率将受到严重限制，运行更多处理少量字符的线程效率可能更高。
• 似乎每个组中的所有线程都会处理完全相同的数据;它可能不是你想到的。