OpenCL 2.0 - work_group 在 CPU 和 GPU 上的操作

Question

我正在测试以下代码，以便使用 work_groups 内置函数在 OpenCL 2.0 中执行并行数组元素添加。（在本例中为 inclusive_add 和 reduce_add）

kernel void(global const float *input,
            global float *sum,
            local float *scratch)
{
    uint local_id = get_local_id(0);
    scratch[local_id] = work_group_inclusive_add(input[get_global_id(0)]);

    if (local_id == get_local_size(0)-1)
        sum[get_group_id(0)] = work_group_reduce_add(scratch[local_id]);
}

如果我用一个从 0 到 15 的浮点数数组测试它，步长为 1，global_size = 16 和 local_size = 4 我期望结果是“6.0 22.0 38.0 54.0 " 如果我选择我的 CPU 作为设备，这很好用。

但是一旦我选择 GPU 和运行相同的代码，我就会得到“0.0 4.0 8.0 12.0”（这只是每个工作组第一个位置的元素）

我是不是漏掉了什么？

我尝试做但没有影响的事情：

在"if"
改变局部大小and/or数组大小/全局大小。

备注：

我正在使用 clEnqueueWriteBuffer 传递输入数组，然后使用 clEnqueueReadBuffer
CPU: i5 6200u
GPU：英特尔核芯显卡 520

（是的，它们支持 OpenCL 2.0，我可以像在运行时构建程序时一样，通过 ioc64 传递 -cl-std=CL2.0 成功构建内核）

Answer 1

你得到不同的结果是因为你使用 work_group_reduce_add 错误的方式。

OpenCL 2.0 spec 说：

This built-in function must be encountered by all work-items in a work-group executing the kernel.

调用 work_group_reduce_add 时情况并非如此。您需要从那里完全删除 if 语句。通过添加只允许一个工作项访问它的 if 语句，您正在计算一个值的总和。然后还给你。

在work_group_scan_inclusive_add之后的数字应该如下：

w1: 0,1,2,3 -> 0,1,3,6
w2: 4,5,6,7 -> 4,9,15,22
w3: 8,9,10,11 -> 8,17,27,38
w4: 12,13,14,15 -> 12,25,39,54

work_group_reduce_add之后：

w1: 10
w2: 50
w3: 90
w4: 130

规范中的第二件事：

NOTE: The order of floating-point operations is not guaranteed for the work_group_reduce_, work_group_scan_inclusive_ and work_group_scan_exclusive_ built-in functions that operate on half, float and double data types. The order of these floating-point operations is also non-deterministic for a given workgroup.

所以我计算的包容性扫描后的结果可能不一定相同，这就是您观察到的 GPU 返回的结果（GPU 返回 0,4,8,12，这恰好是每个缓冲区）。

总结：删除 work_group_reduce_add 之前的 if 语句应该可以解决问题。

OpenCL 2.0 - work_group 在 CPU 和 GPU 上的操作

OpenCL 2.0 - work_group operations on CPU and GPU

c

gpu

intel

opencl