作为一个学校项目，我们正在使用 OpenCL 开发并行光线追踪器。 这是我们第一个使用 OpenCL 的项目，因此我们可能对它有一些不理解。
我们正在尝试实现并行缓冲区压缩，以删除完成的光线或未与任何东西碰撞的光线，以便下一次迭代需要处理的数据更少。 基本上，我们有一个缓冲区，其中包含渲染、跟踪它们、获取碰撞数据、压缩缓冲区所需的尽可能多的s_ray_states，以便只有光线与里面的对象碰撞，然后对它们进行着色。
因此，我们有一个缓冲区uint *prefix_sum其中包含必须在缓冲区s_ray_state *ray_states中移动到每个s_ray_state的索引，以减少发送到着色内核的光线数量，以及跟踪/着色内核的下一次迭代。
可悲的是，下面的ray_sort内核似乎无法正常工作，我们验证了输入prefix_sum数据，这是 100% 正确的，对于ray_states缓冲区也是如此，但我们在输出中得到了不需要的数据。
我们正在启动一个工作组(全局工作大小 = 本地工作大小)，光线总是在缓冲区中移动到比其原始索引更小的索引。我们已经设置了障碍，并使用s_ray_state *tmp缓冲区来防止并行执行写入彼此的数据，但它似乎不起作用，即使消除了障碍，我们也会得到相同的结果。
我们俩已经为此工作了4天，并且已经向其他学生寻求帮助，但似乎没有人能够弄清楚出了什么问题。 我们可能没有足够的了解障碍/mem围栏来确保这实际上可以工作。
我们已经尝试过让单个工作组中的单个工作项对整个数组进行排序，这很有效，甚至提供了更好的性能。
下面的代码应该工作吗？根据我们对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()，否则返回 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
X可以是任何东西
我可能完全不在这里，但在我看来，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：我刚刚注意到你说"光线总是在缓冲区中移动到比原始索引更小的索引"(对不起，我错过了)。这很好，但还不够 - 它们是否总是移动到比同一本地工作组中任何其他射线的索引更小的索引？如果是，很好，但我提到的其他数据竞赛仍然存在。