erly

Ну, я ж специально на атомных функциях делал. Вообще, производительность получилась очень хорошая, как ни странно. Надо придумать, почему так получилось.. Как вариант, такое обоснование: 1. Отфильтрованных точек существенно меньше, чем число точек в исходном массиве (примерно 50к против 900к), поэтому атомная функция вызывалась не очень часто. 2. Компоновка выходного массива выполняется параллельно и с хорошей слитностью адресов памяти. Возможно, как раз второй пункт имеет решающее значение для производительности. Потому что в непосредственно фильтрации ничего ресурсоемкого нет.

__global__ void TubeFilter(float *in, float *out, float *limits, uint *out_size, uint in_size) { __shared__ uint save_count; // count of successfully filtered items in current block if (threadIdx.x == 0) { save_count = 0; // initialize by some one thread } __syncthreads(); uint read_idx = blockIdx.y * blockDim.x * blockDim.y * gridDim.x + blockIdx.x * blockDim.x * blockDim.y + threadIdx.x; if (read_idx >= in_size) { return; } uint r3 = read_idx * 3; // each thread works with all 3 dimensions of a point float x = in[r3]; float y = in[r3 + 1]; float z = in[r3 + 2]; bool save_flag = false; // true if the filter conditions are complied uint save_idx = 0; // the thread will write to this index of out array if (sqrt(x*x + y*y) <= limits[0] && z >= limits[1] && z <= limits[2]) { save_flag = true; save_idx = atomicAdd(&save_count, 1); } __syncthreads(); __shared__ uint first; // first write position of current block if (threadIdx.x == 0) { // increase the common size of out array in global memory first = atomicAdd(out_size, save_count); } if (save_flag) { uint w3 = (first + save_idx) * 3; out[w3] = x; out[w3 + 1] = y; out[w3 + 2] = z; } } Вроде получилось ядро фильтрации. На двух atomicAdd, один из них для значения в shared памяти, второй - для global памяти. Проревьюйте, пожалуйста, кому интересно. Здесь реализован фильтр координат, попадающих в цилиндрическую область заданного радиуса и высоты. in - входной массив 3D координат - float x, y, z координаты. out - выходной массив 3D координат. В него пишутся только точки, удовлетворяющие фильтру. limits - параметры фильтрации: [0] - радиус, [1] - минимальное значение по оси Z, [2] - максимальное значение по Z. out_size - текущий размер выходного массива. Аллоцирован в global памяти. Изначально установлен в host-функции в 0, затем увеличивается в каждом блоке на количество отфильтрованных точек. in_size - размер входного массива. Не меняется.

Спасибо, на него тоже посмотрю. Thrust по любому на базовых CUDA примитивах построен, чудес наверное не стоит от него ждать.

Ну да, свертка.. редукция - это ж одно и то же. Я про них и говорю. Для суммирования они подходят - итеративно общий результат так получать удобно. А в моей фильтрации по идее итерировать не надо: можно за один проход все вычислить, результат зависит только от одного элемента. Зависимости между тредами возникают только из-за адреса, в который результат надо записать. Например: задача отфильтровать значения по порогу 100. Для входного массива [15, 131, 618] результат будет [131, 618]. То есть вторая нить (которой досталось отфильтровать значение 131) должна как-то понять, что писать в выходной массив надо по индексу 0, т.к. первая нить (фильтрующая 15) в выходной массив свое значение не пропускает. Смотрю сейчас на atomic functions. Думаю, что на них надо делать, но не понимаю, как они с абсолютным проседанием производительности борятся. Ну либо все еще не улавливаю всех возможностей редукции..

Почитал про свертку еще. Так и не понял, как ее использовать в описанной выше задаче. Свертка во всех примерах используется для операций типа подсчета суммы всех элементов массива, т.е. когда размер выходного массива заранее известен - фиксированный, или кратен размеру входного массива, или строго зависит от него. Если не сложно, опиши немного подробнее как применить свертку для фильтрации данных.

Такая презентация на эту тему встречалась? http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf

В составе PCL есть модуль компрессии. http://pointclouds.org/documentation/tutorials/compression.php Попробую его. Но с форматом видео контейнера вопрос все еще полностью открыт. Насколько я понимаю, тут тоже сжатие с потерями реализовано - теряется точность координат и количество сохраненных точек. В параметрах указывается сколько точек сохранять на единицу объема (на кубический сантиметр или миллиметр).

Спасибо, большое. Видел такое раньше, но не понял до конца значит. Пошел изучать..

Коллеги, кто имеет опыт в CUDA, подскажите, пожалуйста. Возможно ли в принципе в кернел-фукнции выполнить фильтрацию данных? Так, чтобы размер данных на выходе функции был переменным. Пример задачи: На входе массив значений float, в результирующий массив надо записать только те значения, которые превышают заданный порог. Как такая задача решается? Не могу сообразить, как заставить треды выполнить такую синхронизированную запись в выходной массив. Тут получается, что индекс в выходном массиве априорно не известен, зависит от реультатов работы остальных нитей.

Друзья, как сейчас обстоит дело с форматами записи/стриминга Point Cloud video? Попробовал сделать сохранение потока напрямую как есть, в тривиальном бинарном формате: количество точек в облаке (фиксированная длина), массив данных для каждого кадра (переменная длина, зависимая от первого поля). Оно конечно работает в каком-то смыле - записать и воспроизвести последовательность облаков можно. Но размер файлов получается огромный, т.к. запись идет без сжатия. И естественно, с позиционированием в файле проблема. Может есть какие-то стандарты или даже готовые библиотеки для этого?

Сделал обе эти функции на CUDA, отлично работают. Готовых в librealsense и pcl не нашел.

Поставил наконец-то драйвера. В итоге драйвер заработал после отключения Secure Boot . Теперь следующий шаг - подцепить CUDA к обработке потока от камеры. После гугления решил отдать предпочтение CUDA вместо OpenCL. Мне надо сделать преобразование системы координат для облака точек от камеры и его фильтрацию по координатам. Изучаю сорцы librealsense и pcl по работе с CUDA... должно же в них где-то уже быть такое реализовано.

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

Думаю, нет. У меня же хостовая ОС была - Windows. А сейчас Ubuntu - вторая хостовая ось. Пытаюсь пока хотя бы в безвиртуальном режиме драйвера видюхи подцепить..

Начал биться с OpenCL ). В моем ноуте стоит видяха Nvidia MX 150, думал ее использовать. Работал в Ubuntu в виртуалке, хост - винда. Соответственно, к видеокарте нормального доступа нет. Поставил Ubuntu второй хостовой системой, пытаюсь драйвера прикрутить, чтобы opencl нашел видеокарту. Пока безуспешно, продолжаю.