OpenCL：3D 数组处理 - 全局大小限制答案

【问题标题】：OpenCL: 3D array processing - Globale size limitOpenCL：3D 数组处理 - 全局大小限制
【发布时间】：2018-01-30 04:38:56
【问题描述】：

我正在处理尺寸为 xdim=49、ydim=1024 和 zdim=64 的 3D 数组。我的 DEVICE_MAX_WORK_ITEM_SIZES 只有 512/512/512。如果我声明我的

size_t global_work_size = {xdim, ydim, zdim}; 并启动 3D 内核，

由于我的 ydim > 512，我得到了错误的结果。如果我的所有尺寸都低于 512，我得到了预期的结果。请让我知道是否有替代方案？

【问题讨论】：

标签： opencl

【解决方案1】：

假设您提供的维度是数据的大小，您可以通过让每个 GPU 线程计算更多数据来减小全局工作大小。我的意思是，你的情况下的每个线程都会做一个计算，如果你改变你的内核来做 y 维度的 2 个计算，那么你可以将你正在触发的线程数减半。 global_work_size 决定了您在每个方向上执行的线程数。举个例子吧：

假设您有一个要进行一些计算的数组，并且您拥有的数组大小为 2048。如果您按以下方式编写内核，您将需要 2048 作为 global_work_size：

__kernel void calc (__global int *A, __global int *B)
{
  int i = get_global_id(0);
  B[i] = A[i] * 5;
}

在这种情况下，全局工作大小将是：

size_t global_work_size = {2048, 1, 1};

但是，如果您将内核更改为以下内核，您也可以降低全局工作大小：()

__kernel void new_calc (__global int *A, __global int *B)
{
  int i = get_global_id(0);
  for (int ind = 0; ind < 8; ind++)
    B[i*8 + ind] = A[i*8 + ind] * 5;
}

那么这样，你可以使用全局大小为：

size_t global_work_size = {256, 1, 1};

同样使用第二个内核，您的每个线程将执行更多工作，从而提高利用率。

【讨论】：

【解决方案2】：

CL_DEVICE_MAX_WORK_ITEM_SIZES 只限制工作组的大小，而不是全局工作项的大小（是的，这是一个可怕的常量名称）。您受到 CL_DEVICE_MAX_WORK_GROUP_SIZE 的严格限制，这是一个工作组中允许的项目总数（由于乘法，您通常会比 CL_DEVICE_MAX_WORK_ITEM_SIZES 早得多。

因此，继续发布您的全局工作大小为 49、1024、64。它应该可以工作。如果不是，您使用的是get_local_id 而不是get_global_id 或有其他错误。我们定期推出 4096 x 4096 全局工作大小的 2D 内核。

另见Questions about global and local work size

如果您不使用共享本地内存，则无需担心本地工作组的大小。实际上，您可以传递 NULL 而不是指向 local_work_size 大小数组的指针，并让运行时选择一些东西（如果您的全局维度很容易被小数整除，这会有所帮助）。

【讨论】：