Jump to content
Compvision.ru

erly

Пользователи
  • Content count

    33
  • Joined

  • Last visited

Everything posted by erly

  1. Ну, я ж специально на атомных функциях делал. Вообще, производительность получилась очень хорошая, как ни странно. Надо придумать, почему так получилось.. Как вариант, такое обоснование: 1. Отфильтрованных точек существенно меньше, чем число точек в исходном массиве (примерно 50к против 900к), поэтому атомная функция вызывалась не очень часто. 2. Компоновка выходного массива выполняется параллельно и с хорошей слитностью адресов памяти. Возможно, как раз второй пункт имеет решающее значение для производительности. Потому что в непосредственно фильтрации ничего ресурсоемкого нет.
  2. Коллеги, кто имеет опыт в CUDA, подскажите, пожалуйста. Возможно ли в принципе в кернел-фукнции выполнить фильтрацию данных? Так, чтобы размер данных на выходе функции был переменным. Пример задачи: На входе массив значений float, в результирующий массив надо записать только те значения, которые превышают заданный порог. Как такая задача решается? Не могу сообразить, как заставить треды выполнить такую синхронизированную запись в выходной массив. Тут получается, что индекс в выходном массиве априорно не известен, зависит от реультатов работы остальных нитей.
  3. __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 - размер входного массива. Не меняется.
  4. Спасибо, на него тоже посмотрю. Thrust по любому на базовых CUDA примитивах построен, чудес наверное не стоит от него ждать.
  5. Ну да, свертка.. редукция - это ж одно и то же. Я про них и говорю. Для суммирования они подходят - итеративно общий результат так получать удобно. А в моей фильтрации по идее итерировать не надо: можно за один проход все вычислить, результат зависит только от одного элемента. Зависимости между тредами возникают только из-за адреса, в который результат надо записать. Например: задача отфильтровать значения по порогу 100. Для входного массива [15, 131, 618] результат будет [131, 618]. То есть вторая нить (которой досталось отфильтровать значение 131) должна как-то понять, что писать в выходной массив надо по индексу 0, т.к. первая нить (фильтрующая 15) в выходной массив свое значение не пропускает. Смотрю сейчас на atomic functions. Думаю, что на них надо делать, но не понимаю, как они с абсолютным проседанием производительности борятся. Ну либо все еще не улавливаю всех возможностей редукции..
  6. Почитал про свертку еще. Так и не понял, как ее использовать в описанной выше задаче. Свертка во всех примерах используется для операций типа подсчета суммы всех элементов массива, т.е. когда размер выходного массива заранее известен - фиксированный, или кратен размеру входного массива, или строго зависит от него. Если не сложно, опиши немного подробнее как применить свертку для фильтрации данных.
  7. Такая презентация на эту тему встречалась? http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf
  8. В составе PCL есть модуль компрессии. http://pointclouds.org/documentation/tutorials/compression.php Попробую его. Но с форматом видео контейнера вопрос все еще полностью открыт. Насколько я понимаю, тут тоже сжатие с потерями реализовано - теряется точность координат и количество сохраненных точек. В параметрах указывается сколько точек сохранять на единицу объема (на кубический сантиметр или миллиметр).
  9. Друзья, как сейчас обстоит дело с форматами записи/стриминга Point Cloud video? Попробовал сделать сохранение потока напрямую как есть, в тривиальном бинарном формате: количество точек в облаке (фиксированная длина), массив данных для каждого кадра (переменная длина, зависимая от первого поля). Оно конечно работает в каком-то смыле - записать и воспроизвести последовательность облаков можно. Но размер файлов получается огромный, т.к. запись идет без сжатия. И естественно, с позиционированием в файле проблема. Может есть какие-то стандарты или даже готовые библиотеки для этого?
  10. Спасибо, большое. Видел такое раньше, но не понял до конца значит. Пошел изучать..
  11. Сделал обе эти функции на CUDA, отлично работают. Готовых в librealsense и pcl не нашел.
  12. Всем привет. Коллеги, есть у кого-нибудь опыт работы с 3D камерой Intel Realsense 435i? Пытаюсь освоить, но идет тяжело. Примеры откомпилились, запускаются. Работаю под Ubuntu 18.04. Ближайшая задача: научиться задавать положение камеры в реальном мире, фильтровать получившееся облако точек по заданному объему в пространстве. Дальнейшая задача: с помощью библиотеки PCL распознавать положение людей, перемещающихся перед камерой. Глобальная задача: совместная обработка данных от нескольких таких камер, размещенных вокруг целевого объема пространства. В этом объеме перемещаются люди, все так же требуется вычислять положение их тел/рук/ног. Может кто-либо помочь мне с этими вопросами? PS. Пакет Nuitrak для скелетного распознавания, поддерживающий эту камеру, я видел. Он классный. Но распознает только людей, стоящих на земле. В моем случае он полностью бессилен - люди могут принимать произвольные положения, например висеть вниз головой.
  13. Поставил наконец-то драйвера. В итоге драйвер заработал после отключения Secure Boot . Теперь следующий шаг - подцепить CUDA к обработке потока от камеры. После гугления решил отдать предпочтение CUDA вместо OpenCL. Мне надо сделать преобразование системы координат для облака точек от камеры и его фильтрацию по координатам. Изучаю сорцы librealsense и pcl по работе с CUDA... должно же в них где-то уже быть такое реализовано.
  14. Драйверы пытаюсь ставить проприетарные. Пока это сделать не удалось, гуглю как обходить возникшие трудности. Вечером продолжу.
  15. Думаю, нет. У меня же хостовая ОС была - Windows. А сейчас Ubuntu - вторая хостовая ось. Пытаюсь пока хотя бы в безвиртуальном режиме драйвера видюхи подцепить..
  16. Начал биться с OpenCL ). В моем ноуте стоит видяха Nvidia MX 150, думал ее использовать. Работал в Ubuntu в виртуалке, хост - винда. Соответственно, к видеокарте нормального доступа нет. Поставил Ubuntu второй хостовой системой, пытаюсь драйвера прикрутить, чтобы opencl нашел видеокарту. Пока безуспешно, продолжаю.
  17. Ок, добавился туда, почитываю. Спасибо!
  18. Друзья, а может кто знает русскоязычный ресурс по компьютерному зрению, на котором активность зашкаливает и жизнь бурлит? )) Куда все ушли с этого форума?
  19. Насчет позиционирования я имел в виду, что координаты камеры и углы ее поворотов известны. Надо только их куда-то ввести в API. Куда? По API документация не особо обширная, автогенеренная по коду. Поэтому так сразу не видно, можно ли где-то в свойствах девайса задать его положение. Автоопределять положение камеры буду на поздних этапах, сейчас бы хоть как-то сделать.. ) По поводу точности вычисления глубины.. да, уже ощутил, что дрожание матрицы глубины на удалении метра 4-5 уже очень существенные. Когда покупал камеру, была мысль брать прямо с нее облако точек и по нему скелет определять. Сейчас вижу, что это вряд ли получится. Надо более хитро выкручиваться. Кроме того, обнаружилась большая зависимость дрожания от освещения. Как ни странно, при хорошей яркой сцене глубина определяется фигово, наверное потому что засветки от собственного лазерного проектора не видно. И наоборот, при темной схеме качество улучшается. Но при темной схеме будет проблема с динамичным перемещением объекта съемки - экспозиция требуется слишком большая.
  20. Обычно как бывает... Сначала не понимаешь ничего в теме, ищешь кто б объяснил. А когда сам въедешь, то уже лень другим с нуля объяснять ). Надеюсь, что так все и произошло и все стали профи.
  21. Поиск стекла на фотографии

    Фото выглядит так, будто искать надо прямоугольник, а не стекло.
  22. Попробовал применять трансформацию снаружи, т.е. после получения облака точек... тяжелая операция, проц сильно тормозит на ней. Так что применение поворотов внутри камеры было бы очень кстати.
  23. Добрый день. Расскажите в двух словах, пожалуйста. 1. Есть ли способ преобразования видео от камер, подключенных по GigE, в формат H.264/265 налету? Для записи/стриминга получившегося потока каким-либо медиа-сервером. 2. Бывает ли такое преобразование потоков для линейных камер? 3. Бывает ли оно же для гиперспектральных камер?
  24. Спасибо за подсказку. И тем не менее... Может быть знаете, как задать положение камеры? В принципе, вращать координаты можно уже после получения облака точек. А нельзя ли раньше, чтобы не тратить процессор на это? Чтобы в выходном потоке камеры сразу честные 3D координаты присутствовали.
×