1

作为一个学校项目,我们 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 可以是任何东西

4

1 回答 1

0

我可以完全离开这里,但在我看来,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 和 3prefix_sum[state_id]数组都为 0,并且您的本地 WG 大小 >= 4,则会出现数据竞争。

此外,是否有一些非常好的理由必须为输入和输出使用相同的缓冲区?在我看来,如果你有单独的输入/输出缓冲区,它会变得不那么复杂。

EDIT2:我刚刚注意到你说“光线总是在缓冲区中移动到比原来更小的索引”(对不起,我错过了)。这很好,但还不够——它们总是移动到比同一个本地 WG 中的任何其他光线索引更小的索引吗?如果是的话,很好,但还有我提到的其他数据竞赛。

于 2018-03-22T16:32:11.967 回答