Opencl Kernel - 每个工作项都会覆盖全局内存？

Question

我正在尝试编写一个内核来获取字符串的字符频率。

首先，这是我现在的内核代码：

_kernel void readParallel(__global char * indata, __global int * outdata)
{
        int startId = get_global_id(0) * 8;
        int maxId = startId + 7;

        for (int i = startId; i < maxId; i++)
        {
            ++outdata[indata[i]];
        }
 }

变量inData保存全局内存中的字符串，outdata是全局内存中256个int值的数组。每个工作项从字符串中读取 8 个符号，并应在数组中增加相应的 ASCII 码。代码编译并执行，但outdata 包含的总出现次数少于inData 中的字符数。我认为问题在于工作项覆盖了全局内存。如果你能给我一些解决这个问题的建议，那就太好了。

顺便说一句。我是 OpenCL 的新手 ;-)，是的，我在其他问题中寻找解决方案。

Answer 1

您正在体验您对全局内存的使用不是原子（C++-oriented description of what those are 或another description by the Intel TBB folks）的影响。按时间顺序发生的事情是：

某些工作组“线程”将outData[123] 加载到某个寄存器r1 中

...很多工作，阅读和写作，发生，包括 outData[123]...

同一工作组“线程”递增r1

...很多工作，阅读和写作，发生，包括 outData[123]...

同一工作组“线程”将r1 写入outData[123]

因此，写入outData[123] 的值在读取和写入之间的时间段内“丢弃”更新（我忽略了并行写入相互破坏的可能性，而不是其中一个胜出）。

你需要做的是：

• 使用原子操作 - 对代码的修改量最少，但效率很低，因为它在很大程度上序列化了您的工作，或者
• 使用 work-item-specific、warp-specific 和/或 work-group-specific 部分结果，这需要更少/更便宜的同步，并在完成大量工作后最终将它们组合起来他们。

---

在不相关的注释中，正如@huseyintugrulbuyukisik 正确指出的那样，您的代码使用签名的char 值来索引数组。要解决此问题，请执行以下操作之一：

• 将那些 char 重新解释为数组索引的 unsigned chars（并在读取数组时重新解释。
• 将 char 值向上转换为更大的整数类型，并加上 128 以获得 outArray 的偏移量。
• 将您的内核定义为仅支持 ASCII 字符（不超过 127 个），在这种情况下您可以忽略此问题（尽管如果输入无效，这将是一个潜在的崩溃。
• 如果您只关心printable characters 的频率（但也可以在输入中包含非打印字符），您可以在计算字符之前执行运行时检查。