OpenCL 奇怪的行为答案

【问题标题】：OpenCL strange behaviorOpenCL 奇怪的行为
【发布时间】：2015-06-03 01:58:51
【问题描述】：

早安，

我想我已经尝试了一切来找出问题所在，但我做不到。我有以下主机代码：

cl_mem cl_distances = clCreateBuffer(context, CL_MEM_READ_WRITE, 2 * sizeof(cl_uint), NULL, NULL);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_distances);

cl_event event;
clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_workers, &local_workers, 0, NULL, &event);

clWaitForEvents(1, &event);

对于设备：

__kernel void walk(__global uint *distance_results)
{
    uint global_size = get_global_size(0);
    uint local_size = get_local_size(0);

    uint global_id = get_global_id(0);
    uint group_id = get_group_id(0);
    uint local_id = get_local_id(0);

    for (uint step = 0; step < 500; step++) {
        if (local_id == 0) {
            distance_results[group_id] = 0;
        }

        barrier(CLK_LOCAL_MEM_FENCE);

        for (uint n = global_id; n < 1000; n += global_size) {
            if (local_id == 0) {
                atomic_add(&distance_results[group_id], 1);
            }
        }

        barrier(CLK_GLOBAL_MEM_FENCE);

        if (global_id == 0) {
            for (uint i = 0; i < (global_size / local_size); i++) {
                printf("step: %d; group: %d; data: %d\n", step, i, distance_results[i]);
            }
        }

        barrier(CLK_GLOBAL_MEM_FENCE);
    }
}

因此，在每个“步骤”中，我只需将每个组的距离 [group_id] 加一个 1 1000 次。然后我只是从 global_id == 1 的线程中读取结果。在每一步我都应该有以下文本：

步骤：59；组：0；数据：500
步骤：59；组：1；数据：500

但实际上有很多字符串有错误的数据：

步骤：4；组：0；数据：500
第4步;组：1；数据：210
步骤：5；组：0；数据：500
步骤：5；组：1；数据：214

如果我将 global_workers 设置为 1，将 local_workers 设置为 1，那么一切正常。但是，如果我将 global_workers 设置为 2 并将 local_workers 设置为 1，那么我就会出现这种奇怪的行为。

您知道为什么会发生这种情况吗？

【问题讨论】：

标签： opencl gpgpu

【解决方案1】：

这里发生了一些事情，但我认为核心问题来自对 OpenCL 的一个非常普遍的误解。这个电话：

barrier(CLK_GLOBAL_MEM_FENCE);

这不是一个全球性的障碍。它是具有全局内存栅栏的局部屏障。换句话说，它仍然只在单个工作组中的工作项之间进行同步，而不是在其他工作组中的工作项之间进行同步。

打印结果的代码中的循环只有工作组 0 的正确值，因为它只在工作组 0 中运行。如果你真的希望这段代码工作，打印结果的循环必须位于单独的 NDRange 中，并在 NDRange 之间进行适当的同步。

内存栅栏只控制将哪些类型的内存写入提交到内存。在这种情况下，您需要对两者都设置全局栅栏，因为您正在尝试屏蔽全局内存写入，而不是本地内存写入。

【讨论】：