OpenCL

Question

我正在尝试编写一个 OpenCL 代码来对多维数组执行逐元素操作。

我知道 OpenCL 缓冲区是扁平的，这使得索引有点棘手。我在处理二维数组时成功了，但是对于 3+ 维数组，我要么索引错误，要么结果错误。

更令人惊讶的是，我使用与 2D 情况下相同的索引 principle/formula。

二维案例:

__kernel void test1(__global int* a, __global int* b, __global int* c, const int height) {
    int i = get_global_id(0);
    int j = get_global_id(1);
    c[i + height * j] = a[i + height * j] + b[i + height * j];
}

正确。

3D案例:

__kernel void test1(__global int* a, __global int* b, __global int* c, const int dim1, const int dim2) {
    int i = get_global_id(0);
    int j = get_global_id(1);
    int k = get_global_id(2);

    int idx = i + dim1 * j + dim1 * dim2 * k;
    c[idx] = a[idx] + b[idx];
}

错误的结果（通常是一个充满非常接近 0 的值的输出缓冲区）。

4D案例:

__kernel void test1(__global int* a, __global int* b, __global int* c, const int dim1, const int dim2, const int dim3) {
    int i = get_global_id(0);
    int j = get_global_id(1);
    int k = get_global_id(2);
    int l = get_global_id(3);

    int idx = i + dim1 * j + dim1 * dim2 * k + l * dim1 * dim2 * dim3;
    c[idx] = a[idx] + b[idx];
}

这是索引错误：enqueue_knl_test1 pyopencl._cl.LogicError: clEnqueueNDRangeKernel failed: INVALID_WORK_DIMENSION

Answer 1

在 4D 情况下，您只是错误地使用了 API。 OpenCL 不支持无限数量的全局/局部维度。最多 3 个。

在 2D 情况下，您的索引似乎有误。假设 行优先 数组。应该是 i + j * width 而不是 i + j * height.

在 3D 情况下，内核内部的索引似乎没问题，假设内存布局以行为主，dim1 等于 cols（宽度），dim2 等于 rows（高度）。但无论如何，你的问题缺乏上下文：

输入缓冲区分配和初始化。
内核调用代码（参数、工作组和全局大小）。
结果合集。同步。
您可能正在访问超出分配的缓冲区大小。应该检查一下。

错误地执行这些步骤很容易导致意想不到的结果。即使你的内核代码没问题。

如果您想调试索引问题，最简单的方法就是编写一个简单的内核来输出计算出的索引。

__kernel void test1(__global int* c, const int dim1, const int dim2) {
    int i = get_global_id(0);
    int j = get_global_id(1);
    int k = get_global_id(2);

    int idx = i + dim1 * j + dim1 * dim2 * k;
    c[idx] = idx;
}

然后您应该会得到一个线性增加值的结果。我将从一个工作组开始，然后继续使用多个工作组。

此外，如果在数组之间执行简单的逐元素操作，那么使用一维索引会简单得多。您可以简单地使用一维工作组和等于元素数量的全局大小（四舍五入以适应工作组暗淡）：

__kernel void test1(__global int* a, __global int* b, __global int* c, const int total) {
    // no need for complex indexing for elementwise operations
    int idx = get_global_id(0);
    if (idx < total)
    {
       c[idx] = a[idx] + b[idx];
    }
}

您可能会将 local_work_size 设置为硬件允许的最大大小（例如 Nvidia 为 512，AMD 为 256），并将 global_work_size 设置为四舍五入为 [= 的倍数的元素总数14=]。参见 clEnqueueNDRangeKernel。

2D 和 3D dims 通常用于访问 2D / 3D 中相邻元素的操作 space。比如图像卷积。

OpenCL - 4D 数组的逐元素操作

OpenCL - Element-wise operations on 4D array

arrays

parallel-processing

gpu

pyopencl