如何添加额外的工作项以使全局工作大小成为本地工作大小的倍数

Question

我正在编写一个 OpenCL 程序，但是我的全局工作量不是本地工作量的倍数。在 OpenCL 中，全局工作大小必须可以被本地工作大小整除，所以我读到的解决方案是添加一些额外的工作项，这些工作项不会对全局工作大小进行四舍五入，并使其可以被所选的本地工作大小整除。

例如，假设本地工作大小为 4，全局工作大小为 62（您有 62 个元素需要内核对其进行操作）

这里的想法是再添加 2 个空闲的工作项，以使全局工作大小为 64。因此，由于 64 可以被 4 整除，所以一切都很好。

关于如何准确地实现这样的空闲工作项有什么想法吗？如果我只是将全局工作大小增加到 64，我的内核会额外执行两次，这会改变程序正在执行的计算结果，最终会产生错误的结果。

Answer 1

这是一种标准方法，round-up 全局工作量是本地工作量的倍数。在这种情况下，我们必须在内核中添加绑定检查，以确保只有那些工作项执行落在有效数据范围内的计算。可以通过将实际数据大小指定为内核参数并将其与工作项的全局索引进行比较来完成。示例内核如下所示：

__kernel void example_kernel(__global int* input, __global int* output, int dataSize)
{
    int index = get_global_id(0);

    if (index < dataSize)
    {
        /*
        rest of the kernel...
        */
    }
}

Answer 2

从 OpenCL 2.0 开始，不再需要全局工作大小是本地工作大小的倍数。

最好保留本地工作大小 NULL 除非有真正的性能优势。

您可以向下舍入 gws 并在边缘工作迭代器中进行额外处理

gws = (old_gws/lws) * lws;
leftover = old_gws - gws;

在内核中

if(get_global_id(0) == (get_global_size(0)-1))
    // do computation for rest of the work-items (leftover)

如何添加额外的工作项以使全局工作大小成为本地工作大小的倍数

How to add extra work items in order to make global work size a multiple of the local work size

opencl