OpenCL 并行缓冲区压缩障碍问题

Posted 2023-03-17

【中文标题】OpenCL 并行缓冲区压缩障碍问题

【英文标题】：OpenCL parallel buffer compaction barrier issue

【发布时间】：2018-08-31 05:07:42

【问题描述】：

作为一个学校项目，我们 4 正在使用 OpenCL 开发并行光线追踪器。 这是我们第一个使用 OpenCL 的项目，所以我们可能对它有一些不理解。

我们正在尝试实现并行缓冲区压缩以移除完成的光线或未与任何物体发生碰撞的光线，以便下一次迭代处理更少的数据。 基本上，我们有一个足够多的s_ray_states 缓冲区，用于渲染、跟踪它们、获取碰撞数据、压缩缓冲区，以便只有光线与其中的对象碰撞，然后对它们进行着色。

所以我们有一个缓冲区uint *prefix_sum，其中包含每个s_ray_state必须移动到缓冲区s_ray_state *ray_states中的索引，以减少发送到着色内核的光线数量，以及下一个跟踪/阴影内核的迭代。

遗憾的是，下面的 ray_sort 内核似乎无法正常工作，我们验证了输入 prefix_sum 数据，这是 100% 正确的，对于 ray_states 缓冲区也是如此，但我们在其中收到了不需要的数据输出。

我们正在启动一个工作组（全局工作大小 = 局部工作大小），光线总是在缓冲区中移动到比其原始索引更小的索引。我们设置了障碍，并使用s_ray_state *tmp 缓冲区来防止并行执行写入彼此的数据，但它似乎不起作用，即使移除障碍我们也会得到相同的结果。

我们俩都已经搞了 4 天了，已经向其他同学求助了，但似乎没有人能够弄清楚哪里出了问题。 我们可能对障碍/内存栅栏的了解不够，无法确保这实际上可以工作。

我们已经尝试让单个工作组中的单个工作项对整个数组进行排序，这很有效，甚至可以提供更好的性能。

下面的代码应该可以工作吗？以我们对 OpenCL 的理解，它应该可以工作，我们做了很多研究，但从未真正得到任何明确的答案..

kernel void ray_sort(
    global read_only uint *prefix_sum,
    global read_write struct s_ray_state *ray_states,
    global read_only uint *ray_states_size,
    local read_write struct s_ray_state *tmp
)

    int l_size = get_local_size(0);
    int l_id = get_local_id(0);
    int group_id = -1;
    int group_nb = *ray_states_size / l_size;
    int state_id;

    while (++group_id < group_nb)
    
        state_id = group_id * l_size + l_id;
        tmp[l_id] = ray_states[state_id];
        barrier(CLK_LOCAL_MEM_FENCE);
        if (did_hit(tmp[l_id]))
            ray_states[prefix_sum[state_id]] = tmp[l_id];
        barrier(CLK_GLOBAL_MEM_FENCE);
    

ray_states 长度为ray_states_size

prefix_sum 包含每个 ray_states 元素必须移动到的索引

tmp 是大小为 local_work_size 的本地缓冲区

local_work_size = global_work_size

did_hit() 如果射线击中一个物体返回 1，否则返回 0

我们希望将 ray_states 元素移动到包含在 prefix_sum 中的索引

示例：每个ray_states[id] 都被移动到prefix_sum[id] 索引中 ray_states

prefix_sum: 0 | 0 | 1 | 1 | 2 | 3 | 3 | 3 | 4

did_hit(ray_states[id]): 0 | 1 | 0 | 1 | 1 | 0 | 0 | 1 | 0

did_hit(output[id]): 1 | 1 | 1 | 1 | X | X | X | X | X

Xs 可以是任何东西

【问题讨论】：

您好 elXor，我觉得您需要编辑问题以详细说明“我们在输出中收到了不需要的数据。” -- 即“我们期待（特定的东西）并得到（特定的东西）” -- 这样人们就可以专注于这个问题。

@LeonBambrick 我觉得我们得到的输出对于这个问题并不重要，我们只是期望 ray_states 中的元素被重新组织到 prefix_sum 中包含的索引。我们只是得到了一些没有按照预期方式重新组织的东西（元素没有移动到正确的索引）。

【参考方案1】：

我可以完全离开这里，但在我看来，did_hit(ray_states[state_id]) 您正在读取同一块全局内存，您将其放入本地内存缓冲区 tmp 仅上面 2 行。这不会有问题，除非您将该缓冲区用于输入和输出。

在我看来，硬件上实际发生的情况是这样的：

    tmp[l_id] = ray_states[state_id];
    tmp[l_id] = ray_states[state_id];
    tmp[l_id] = ray_states[state_id];
    tmp[l_id] = ray_states[state_id];
    tmp[l_id] = ray_states[state_id];

       ... local-work-size times

    barrier(CLK_LOCAL_MEM_FENCE);

    if (did_hit(ray_states[state_id]))
        ray_states[prefix_sum[state_id]] = tmp[l_id];
    if (did_hit(ray_states[state_id]))
        ray_states[prefix_sum[state_id]] = tmp[l_id];
    if (did_hit(ray_states[state_id]))
        ray_states[prefix_sum[state_id]] = tmp[l_id];
    if (did_hit(ray_states[state_id]))
        ray_states[prefix_sum[state_id]] = tmp[l_id];

      ... again local-work-size times

考虑到 WItem 并行执行顺序是不确定的（硬件可以选择它想要的任何顺序），这将导致随机结果。你可以试试这个吗：

    if (did_hit(tmp[l_id]))
        ray_states[prefix_sum[state_id]] = tmp[l_id];

顺便说一句，如果ray_states_size 只是一个简单的整数，您可以直接传递它，通过将参数设为“uint ray_states_size”。没有必要在那里玩转缓冲区。

EDIT1：我的建议只有在 prefix_sum[state_id] 在每个本地工作大小的 id 中都没有任何重复项时才有效，否则仍然会有数据竞争。所以例如如果state_id-s 1 和 3 的 prefix_sum[state_id] 数组都为 0，并且您的本地 WG 大小 >= 4，则会出现数据竞争。

另外，是否有一些非常好的理由必须使用相同的缓冲区进行输入和输出？在我看来，如果你有单独的输入/输出缓冲区，它会简单得多。

EDIT2：我刚刚注意到您说“光线总是在缓冲区中移动到比原始索引更小的索引”（对不起，我错过了）。这很好，但还不够——它们是否总是移动到比同一本地 WG 中的 任何其他光线的索引更小的索引？如果是，那很好，但我还提到了另一个数据竞赛。

【讨论】：

我解决了你提到的 did_hit() 的问题（在实施和我的问题中），遗憾的是，仍然没有改变任何东西..

我不知道你可以将值传递给内核，即使在某些示例中我看到它们是通过指针传递的，我会调查一下，谢谢。 prefix_sum 确实有重复，但did_hit(state_id) 只能为其中一个返回 1，永远不会更多，永远不会更少。我不能使用 2 个不同的缓冲区，因为这会在我们学校的 GPU 上占用太多内存。我尝试了 2 个缓冲区，但它也不起作用..

光线总是被移动到比其原始索引更小的索引，或者根本不移动。它们可能会被移动到与同一本地工作组中的其他光线相同的索引，但 tmp 本地不应该关心这个吗？

tmp 本地将为您提供有效的可读副本，因此之后写入ray_states 是安全的。但问题是最终写了什么。如果本地 WG 的任意两个不同 WItems 的 did_hit(tmp[l_id]) 条件评估为真，那么这些 WItems 的 prefix_sum[state_id] 可以相同吗？如果是，那么您有一个问题：最终存储的是这两个 WItem 中最后一个写入的内容，但您不知道最后执行的是哪一个。