В качестве школьного проекта мы работаем над параллельным трассировщиком лучей с 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 может быть чем угодно
