OpenCL: вывод массива переменной длины - PullRequest
Новая
       4

OpenCL: вывод массива переменной длины

2 голосов
/ 20 декабря 2011

Мы работаем над заданием для курса GPGPU. Мы выбрали алгоритм, внедрили его в процессор и теперь конвертируем его в OpenCL.

Алгоритм, который мы выбрали, загружает модель в виде набора треугольников и растеризует их до вокселей. Воксели определяются как VBO точечных данных. Затем мы используем геометрический шейдер для преобразования этих точек в воксели в виде треугольников.

Таким образом, нашей программе OpenCL необходимо взять список треугольников и вывести список переменных точек.

А вывод массива переменной длины кажется проблемой.

Решение, которое мы нашли, состоит в том, чтобы атомарно увеличивать счетчик и использовать этот счетчик в качестве индекса как выходного массива, так и конечного размера массива. За исключением ... оба наших графических процессора не поддерживают расширение для атомарных операций.

Это то, что мы имеем до сих пор: 

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable

#define POS1      i0 * 3 + 0
#define POS2      i0 * 3 + 1
#define POS3      i0 * 3 + 2

void WritePosition( __global float* OutBuffer, uint inIndex, __global float* inPosition )
{
    OutBuffer[ inIndex * 3 ] = inPosition[0];
    OutBuffer[ inIndex * 3 + 1] = inPosition[1];
    OutBuffer[ inIndex * 3 + 2] = inPosition[2];
}

__kernel void Voxelize( 
    __global float* outPointcloudBuffer, 
    __global float* inTriangleBuffer, 
    __global uint* inoutIndex
)
{
    size_t i0 = get_global_id(0);
    size_t i1 = get_local_id(0);

    WritePosition( outPointcloudBuffer, inIndex[0], &inTriangleBuffer[ i0 ] );

    //atomic_inc(inoutIndex[0]);
    inoutIndex[0] = max(inoutIndex[0], i0);
}

И результат этого очень странный. Мы тестируем очень маленькую модель (12 треугольников, 36 позиций, 108 поплавков), и в результате мы получаем 31, 63 или 95. Всегда кратно 16 минус 1.

Как мы можем получить длину нашего выходного массива переменной длины?

Заранее спасибо.

1 Ответ

5 голосов
/ 20 декабря 2011

Я думаю, это обычно решается следующим образом:

  • Первый проход: вычисление необходимого размера массива на графическом процессоре с использованием примитива scan (сумма параллельного префикса). Выше ссылка содержит пример реализации от Apple.
  • Выделите необходимые ресурсы на стороне хоста, используя результат алгоритма сканирования. Обратите внимание, что результат алгоритма сканирования часто можно использовать в качестве подсказки индекса для результатов отдельных рабочих элементов.
  • Второй проход (необязательно): сжатие массива до тех элементов, которые необходимо учитывать при третьем проходе.
  • Третий проход: перезапустите алгоритм, передавая индексы назначения и выделенный массив.

Возможно, вы захотите взглянуть на марширующие кубы NVIDIA OpenCL реализация , где реализованы все три вышеупомянутых прохода.

Бест, Кристоф

...