OpenCL 2.0 - CPU 和 GPU 上的工作组操作

Question

我正在测试以下代码，以便在 OpenCL 2.0 中使用 work_groups 内置函数执行并行数组元素添加。 (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”（这只是每个工作组中第一个位置的元素）

我错过了什么吗？

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

1. 在“if”之前添加“barrier(CLK_LOCAL_MEM_FENCE)”

2. 更改局部大小和/或数组大小/全局大小。

注意事项：

• 我使用 clEnqueueWriteBuffer 传递输入数组，然后使用 clEnqueueReadBuffer 读取总和

• CPU：i5 6200u

• GPU：英特尔核芯显卡 520

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

Answer 1

您得到不同的结果是因为您使用了错误的work_group_reduce_add。

OpenCL 2.0 spec 说：

这个内置函数必须在一个 执行内核的工作组。

当您致电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

规范中的第二件事：

注意：浮点运算的顺序不保证 work_group_reduce_、work_group_scan_inclusive_ 和 work_group_scan_exclusive_ 操作的内置函数 half、float 和 double 数据类型。这些浮点数的顺序 对于给定的工作组，操作也是不确定的。

因此，我计算的包容性扫描后的结果可能不一定相同，这就是您观察到的 GPU 返回的内容（GPU 返回 0、4、8、12，这恰好是每个缓冲区的最后一个值)。

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