opencl 中的本地和全局工作大小

Question

我正在尝试学习 open cl，但有一个 混乱的根源我不明白 现在，它与这些行有关

size_t global_item_size = LIST_SIZE; // Process the entire lists 
size_t local_item_size = 64; // Divide work items into groups of 64 
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, 
        &global_item_size, &local_item_size, 0, NULL, NULL); 

我知道内核在 LIST_SIZE 线程数上被调用（我的意思是我在执行中得到了 LIST_SIZE 内核，希望是并行的）[对吗？] 但是什么意思

size_t local_item_size = 64; // Divide work items into groups of 64 

？

这是否意味着每个线程/内核都在 64 通道类似 simd 的波前执行？ （如果是这样，我可以称之为双重并行化，但我可能会混淆一些东西）

有人可以澄清/帮助理解这一点吗？，[也可能添加一些重要的提示，在安排扭动此类内核时要注意什么？]

Answer 1

工作组大小为 64 意味着整个工作负载将被分成 64 个组。波前比工作组大小更底层，并且是由硬件/驱动程序控制的——即开发人员无法更改波前大小。

工作组可以在其内部共享本地内存，但不能与外部组共享。本地内存通常比设备的全局内存快一个数量级。每当您需要多次读取一个值时，将该值缓存在本地内存中是值得的。

我将列出一个冗长的示例，说明本地内存和工作组何时有助于提高性能，但这与您的问题范围有点偏离。如果你愿意，我仍然可以发布。

---

编辑： 如果没有示例，我觉得这些信息不是很有用。假设您需要在 1920x1080 图像上运行模糊滤镜。假设每个像素周围有 2 个像素半径（一个 5x5 框）和 32 位像素（类型：float）。

选项 1 让每个工作项加载一个正方形像素区域，处理目标输出像素，并将结果写回全局内存。这将产生正确的结果，但效率低下。源图像中的每个像素将被各种工作项读取多达 25 次，并且您的全局写入也将分散。我将列出的其他选项不描述全局写入，仅描述读取性能提升。 请注意，在幕后，opencl 实现仍在将您的工作划分为工作组，即使您没有在代码中使用它们。

选项 2a 利用本地内存和工作组绕过全局内存瓶颈。与现代设备上的全局内存一样快，本地内存仍然快一个数量级。本地内存通常以或接近设备的时钟速度运行。如果使用大小为 64 的工作组处理同一图像，则工作项可以协同处理输出图像的 8x8 区域。工作项可以一次将一个 12x12 的像素区域加载到本地内存中，并且从本地内存中读取处理一个像素所需的 25 次读取，从而大大提高了性能。每个像素的平均全局读取次数为 1.497（vs 25！）。 240x135 的工作组可以通过这种方式处理整个 1920x1080 的图像。

选项 2b 使用 64 的工作组大小，您的组可以处理更大区域的图像。上面描述的图 2a 使用工作组来处理 8x8 输出区域，但它仅使用 12*12*sizeof(float)（=576 字节）的本地内存。通过增加工作组的输出区域，可以减少双读或四读的边界像素。 opencl 规范说设备需要有 32kb 的可用本地内存。一些数学计算确定每个工作组可以安全地处理 90x90 方形像素区域 - floor(sqrt(32768/sizeof(float)) = 90。现在每个像素的平均全局读取数为 1.041。只需要 22x12 个工作组。

当您能够实现几乎 24:1 的全局读取数量减少时，即使是低端 GPU 也可能突然受到 ALU 性能的限制。