Проблема барьера сжатия параллельного буфера OpenCL

Question

Как школьный проект, мы работаем над параллельным raytracer с 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

Xможет быть что угодно

Answer 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] не имеет дубликатов в каждом идентификаторе локального размера, иначе все равно будет гонка данных. Так, например, если для обоих state_id-s 1 и 3 prefix_sum[state_id] массив имеет 0, и ваш локальный размер рабочей группы>= 4, будет гонка данных.

Кроме того, есть ли действительно веская причина, по которой вы должны использовать один и тот же буфер для ввода и вывода? мне кажется, было бы намного проще, если бы у вас были отдельные входные / выходные буферы.

EDIT2: я только что заметил, что вы сказали, что "лучи всегда перемещаются в буфере с меньшим индексом, чем их оригинал" (извините, я пропустил это). Это хорошо, но недостаточно - всегда ли они перемещаются к меньшему индексу, чем индекс любого другого луча в той же локальной рабочей группе? если да, хорошо, но есть еще другая гонка данных, которую я упомянул.