opencl 中的本地和全局工作大小
local and global work sizes in open cl
我正在尝试学习 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
？
这是否意味着每个 thread/kernell 都是在类似 64 通道 simd 的波前执行的？ （如果是这样，我可以称之为双重并行化，但也许我混淆了一些东西）
有人 clarify/help 能理解这一点吗？[还可以添加一些关键提示，在调度此类内核时需要注意什么？]
工作组大小为 64 意味着整个工作负载将被分成 64 个组。wavefront 比工作组大小更底层，并且 hardware/driver 受控——即开发人员可以' t 改变波前尺寸。
工作组可以在其内部共享本地内存，但不能与外部组共享。本地内存通常比设备的全局内存快一个数量级。每当您需要多次读取一个值时，将该值缓存在本地内存中是值得的。
我打算列出一个长篇大论的例子，说明本地内存和工作组何时有助于提高性能，但它偏离了你的问题范围。如果你愿意，我仍然可以 post。
编辑：
如果没有示例，我觉得这些信息不是很有用。假设您需要对 1920x1080 图像进行 运行 模糊滤镜。假设每个像素（一个 5x5 框）和 32 位像素（类型：float）周围有 2 个像素半径。
选项 1 让每个工作项加载一个正方形像素区域，处理目标输出像素，并将结果写回全局内存。这将生成正确的结果，但效率低下。源图像中的每个像素将被各种工作项读取多达 25 次，并且您的全局写入也会分散。我将列出的其他选项不描述全局写入，只描述读取性能提升。
请注意，在幕后，opencl 实现仍在将您的工作划分为工作组，即使您没有在代码中使用它们。
选项 2a 利用本地内存和工作组绕过全局内存瓶颈。与现代设备上的全局内存一样快，本地内存仍然快一个数量级。本地内存通常 运行s 等于或接近设备的时钟速度。如果您使用大小为 64 的工作组处理同一图像，则工作项可以协作处理输出图像的 8x8 区域。工作项可以一次将 12x12 像素区域加载到本地内存中，处理一个像素所需的 25 次读取都是从本地内存中读取的，从而大大加快了性能。每个像素的平均全局读取数为 1.497（对比 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 性能限制。
我正在尝试学习 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
？
这是否意味着每个 thread/kernell 都是在类似 64 通道 simd 的波前执行的？ （如果是这样，我可以称之为双重并行化，但也许我混淆了一些东西）
有人 clarify/help 能理解这一点吗？[还可以添加一些关键提示，在调度此类内核时需要注意什么？]
工作组大小为 64 意味着整个工作负载将被分成 64 个组。wavefront 比工作组大小更底层，并且 hardware/driver 受控——即开发人员可以' t 改变波前尺寸。
工作组可以在其内部共享本地内存，但不能与外部组共享。本地内存通常比设备的全局内存快一个数量级。每当您需要多次读取一个值时，将该值缓存在本地内存中是值得的。
我打算列出一个长篇大论的例子，说明本地内存和工作组何时有助于提高性能，但它偏离了你的问题范围。如果你愿意，我仍然可以 post。
编辑： 如果没有示例，我觉得这些信息不是很有用。假设您需要对 1920x1080 图像进行 运行 模糊滤镜。假设每个像素（一个 5x5 框）和 32 位像素（类型：float）周围有 2 个像素半径。
选项 1 让每个工作项加载一个正方形像素区域，处理目标输出像素，并将结果写回全局内存。这将生成正确的结果，但效率低下。源图像中的每个像素将被各种工作项读取多达 25 次，并且您的全局写入也会分散。我将列出的其他选项不描述全局写入，只描述读取性能提升。 请注意，在幕后，opencl 实现仍在将您的工作划分为工作组，即使您没有在代码中使用它们。
选项 2a 利用本地内存和工作组绕过全局内存瓶颈。与现代设备上的全局内存一样快，本地内存仍然快一个数量级。本地内存通常 运行s 等于或接近设备的时钟速度。如果您使用大小为 64 的工作组处理同一图像，则工作项可以协作处理输出图像的 8x8 区域。工作项可以一次将 12x12 像素区域加载到本地内存中，处理一个像素所需的 25 次读取都是从本地内存中读取的，从而大大加快了性能。每个像素的平均全局读取数为 1.497（对比 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 性能限制。