The OpenNET Project / Index page

[ новости /+++ | форум | теги | ]

форумы  помощь  поиск  регистрация  вход/выход  слежка  RSS
"v4l2 cuda ускорение копирования кадра в память GPU"
Вариант для распечатки  
Пред. тема | След. тема 
Форум Программирование под UNIX (C/C++)
Изначальное сообщение [ Отслеживать ]

"v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 19-Июн-16, 16:13 
Всем привет,

Пишу код на Ubuntu 14, Tegra 3, камера с разрешением 4224x3156, yuv422, 26Мб. Забираю кадр каждые 70мс, далее memcpy в pinned memory host, cudaMemcpy из pinned host в divice memory и далее обработка в GPU.

Memcpy занимает 200мс. Это больше чем частота кадров. Как можно ускорить загрузку буфера кадра в память GPU?

Можно ли использовать v4l2 IO_METHOD_USERPTR в качестве pinned memory, что бы не тратить время на memcpy.

С уважением Виктор.

Ответить | Правка | Cообщить модератору

Оглавление

Сообщения по теме [Сортировка по времени | RSS]


1. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от fail on 20-Июн-16, 18:31 
> Всем привет,

...
> Можно ли использовать v4l2 IO_METHOD_USERPTR в качестве pinned memory, что бы не
> тратить время на memcpy.
> С уважением Виктор.

шаг 1.

https://linuxtv.org/downloads/v4l-dvb-apis/
&
https://linuxtv.org/downloads/v4l-dvb-apis/io.html

шаг 2.

смотрим в /usr/include/*
что-то вроде video*.h на предмет V4L2_MEMORY_USERPTR

Ответить | Правка | ^ к родителю #0 | Наверх | Cообщить модератору

2. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 20-Июн-16, 22:20 
Спасибо за ответ.
Делал по примеру отсюда: https://linuxtv.org/downloads/v4l-dvb-apis/capture-example.html

Метод MMAP:

В process_image всего одна строка memcpy( pinned, p, size ) //200ms

Инициализирую так: cudaMallocHost( pinned, size )

Метод USERPTR:
Пробовал вместо: buffers[n_buffers].start = malloc(buffer_size)
делать: cudaMallocHost( buffers[n_buffers].start, buffer_size )

Прога зависает здесь:
case IO_METHOD_USERPTR:
                for (i = 0; i < n_buffers; ++i) {
                        struct v4l2_buffer buf;

                        CLEAR(buf);
                        buf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
                        buf.memory = V4L2_MEMORY_USERPTR;
                        buf.index = i;
                        buf.m.userptr = (unsigned long)buffers[i].start;
                        buf.length = buffers[i].length;

     ------------------->>>>>>if (-1 == xioctl(fd, VIDIOC_QBUF, &buf))
                                errno_exit("VIDIOC_QBUF");
                }
                type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
                if (-1 == xioctl(fd, VIDIOC_STREAMON, &type))
                        errno_exit("VIDIOC_STREAMON");
                break;
        }

Можете пояснить?

Ответить | Правка | ^ к родителю #1 | Наверх | Cообщить модератору

3. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от fail on 20-Июн-16, 23:05 
> Спасибо за ответ.
> Делал по примеру отсюда: https://linuxtv.org/downloads/v4l-dvb-apis/capture-example.html

...

> Можете пояснить?

- давно несколько в другом направлении копаюсь
- с cuda вообще не пересекался - есть возможность работы с памятью выделенной через стандартный malloc  или надо строго через cudaMallocHost( pinned, size ) работать?
- как вариант(?), проверка на ошибки пропущена, кусок кода оттуда же

                if (-1 == xioctl(fd, VIDIOC_DQBUF, &buf)) {
                        switch (errno) {
                        case EAGAIN:
                                return 0;

                        case EIO:
                                /* Could ignore EIO, see spec. */

                                /* fall through */

                        default:
                                errno_exit("VIDIOC_DQBUF");
                        }
                }


Ответить | Правка | ^ к родителю #2 | Наверх | Cообщить модератору

4. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 20-Июн-16, 23:18 
Разумеется буферы выделены корректно, проверяю.
До Вашего куска не доходит.
Зависает в start_capturing( в режиме IO_METHOD_USERPTR)

Да, надо в cuda иначе все совсем медленно.

Ответить | Правка | ^ к родителю #3 | Наверх | Cообщить модератору

5. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 20-Июн-16, 23:21 
Сорри поторопился.
Сразу проверить код ошибки в start_capturing не пробовал.
Попробую отпишусь.

Ответить | Правка | ^ к родителю #4 | Наверх | Cообщить модератору

6. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от fail on 21-Июн-16, 11:57 
> Сорри поторопился.
> Сразу проверить код ошибки в start_capturing не пробовал.
> Попробую отпишусь.

пока идет локальная пересборка,
ради интереса глянул в сети инфу по cuda,
http://ecee.colorado.edu/~siewerts/extra/code/example_code_a.../

немного статистики
14 страница в CUDA_Getting_Started_Guide_For_Linux.pdf

GTX 670, memory pinned:
host2device: 3 G/s
device2host: 2 G/s

+

имхается, надо глянуть примеры:
.../cuda_work/samples/0_Simple/simpleStreams/
.../cuda_work/samples/0_Simple/simpleZeroCopy/


P.S.:

а ус-во видео захвата на какой шине сидит:
USB(2,3), PCI(32,64), etc ?

Ответить | Правка | ^ к родителю #5 | Наверх | Cообщить модератору

7. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 21-Июн-16, 22:11 
Прибор Jetson TK1, немного солгал про Tegra 3, наверное Tegra 4.
В вики пишут, что 64бит шина. Причем RAM у GPU и CPU общая, отсюда еще больше вопросов к memcpy. 200мс/26Мб это 130Мб/сек. SSD быстрее пишет :)
На выхах буду анализировать ссылки и ответы присланные вами.
Как вариант, запланировал эксперимент про мультипоточность, т.к. в момент memcpy занято одно ядро из 4-х.
Отпишусь по результатам.
Ответить | Правка | ^ к родителю #6 | Наверх | Cообщить модератору

8. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от fail_ on 22-Июн-16, 21:12 
> Прибор Jetson TK1, немного солгал про Tegra 3, наверное Tegra 4.
> В вики пишут, что 64бит шина. Причем RAM у GPU и CPU
> общая, отсюда еще больше вопросов к memcpy. 200мс/26Мб это 130Мб/сек. SSD
> быстрее пишет :)

...

думается, сырой захват идет через USB 2
+
как вариант, видео "пожатое" идет через тотже USB..
имхается, поэтому такая задержка...


Ответить | Правка | ^ к родителю #7 | Наверх | Cообщить модератору

9. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 22-Июн-16, 21:35 
Камера через MIPI CS2 4lane. Тут все нормуль. Данные в yuv422.


Ответить | Правка | ^ к родителю #8 | Наверх | Cообщить модератору

10. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от fail_ on 23-Июн-16, 20:32 
> Камера через MIPI CS2 4lane. Тут все нормуль. Данные в yuv422.

bandwidth ? (dma ?)
yuv422 - это raw data

Ответить | Правка | ^ к родителю #9 | Наверх | Cообщить модератору

11. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 23-Июн-16, 22:42 
Реально не знаю.
По вики много.


Ответить | Правка | ^ к родителю #10 | Наверх | Cообщить модератору

12. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 27-Июн-16, 00:11 
Пробовал memcpy из простого буфера в простой:

uchar * pIn{ new uchar[ size ]};
uchar * pOut{ new uchar[ size ]};
memcpy( pOut, pIn, size );
delete [] pIn;
delete [] pOut;

Тоже около 200мс.

Пока получается, что на Tegra TK1 4224х3156х2 байт копируются с такой скоростью.
Но это учитывая, то, что там драйвер камеры и подсистема V4L2 делает.
Буду разбираться.

Сделал вариант с OpenMP, реально на 4 поделилось время:

static void memcpy_openmp( uchar * pOut, const uchar * pIn, const int size ) {

    //omp_set_dynamic( 0 );
    omp_set_num_threads( 4 );
    //int blockSize{ 1666368 };
    int blockSize{ size / 4 };
    int th_id, nthreads;
    #pragma omp parallel for
        for( int i = 0; i < size / blockSize; i++ ){
            int offset{ i * blockSize };
            memcpy( pOut + offset, pIn + offset, blockSize );
            th_id = omp_get_thread_num();
            nthreads = omp_get_num_threads();
            qDebug() << th_id << nthreads;
        }

}

Где-то 70-80мс.

Вопрос открыт.

Ответить | Правка | ^ к родителю #11 | Наверх | Cообщить модератору

13. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 27-Июн-16, 01:46 
По V4L2:

https://devtalk.nvidia.com/default/topic/894783/jetson-tk1/v...-/

IO_METHOD_USERPTR не поддерживается. Только MMAP.

Ответить | Правка | ^ к родителю #12 | Наверх | Cообщить модератору

14. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от fail_ on 27-Июн-16, 07:54 
> По V4L2:

https://devtalk.nvidia.com/default/topic/894783/jetson-tk1/v...-/
> IO_METHOD_USERPTR не поддерживается. Только MMAP.

остаются: v4l2_memory_mmap & v4l2_memory_dmabuf
может что-то из этого зафурычит..


Ответить | Правка | ^ к родителю #13 | Наверх | Cообщить модератору

15. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 27-Июн-16, 22:21 
Спасибо. Все ваши советы обязательно проверю. Не было времени пока.


Ответить | Правка | ^ к родителю #14 | Наверх | Cообщить модератору

16. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 03-Июл-16, 23:10 
v4l2_memory_dmabuf не поддерживается.

Только MMAP.

Пока сделал параллельно. Укладываюсь в 70мс. Но позже надо будет в mp4 жать, надеюсь CUDA на себя все возьмет.

Ответить | Правка | ^ к родителю #15 | Наверх | Cообщить модератору

17. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 04-Июл-16, 23:28 
Сделал тест:

- memcpy from userspace pointer to userspace pointer is about 33ms per 26mb
- memcpy from userspace pointer to pinned pointer is about 33ms per 26mb
- memcpy from mmap v4l pointer to pinned is about 200ms per 26mb

Mmap медленный. Интересно почему..

Тут у человека такая же проблема: https://devtalk.nvidia.com/default/topic/948258/jetson-tk1/p.../

Ответить | Правка | ^ к родителю #16 | Наверх | Cообщить модератору

18. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от vsw (ok) on 05-Июл-16, 23:53 
Решилось с помощью V4L2_MEMORY_USERPTR.

В найденном варианте v4l2 нашел почти похожий пример(http://www.friendlyarm.net/forum/topic/1006), как в стандартном примере, но пользовательский буфер выделяется не malloc, а memalign из malloc.h. Иначе Invalid argument error 22.

Настраиваемся на IO_METHOD_USERPTR и в init_userp:

        unsigned int page_size;

        page_size = getpagesize ();
        buffer_size = (buffer_size + page_size - 1) & ~(page_size - 1);

..................
..................

  for (n_buffers = 0; n_buffers < 4; ++n_buffers) {
                buffers[n_buffers].length = buffer_size;
                buffers[n_buffers].start = memalign (/* boundary */page_size,buffer_size);

                if (!buffers[n_buffers].start) {
          fprintf (stderr, "Out of memory\n");
                exit (EXIT_FAILURE);
    }
        }


И далее работаем с буфером buf.m.userptr размера buf.length.

Что в отличии от mmap буфера, позволяет быстро копировать данные кадра в pinned memory и далее в cuda.

Думаю что вопрос закрыт.

Всем спасибо, особенно пользователю fail :-)

С уважением Виктор.

Ответить | Правка | ^ к родителю #17 | Наверх | Cообщить модератору

19. "v4l2 cuda ускорение копирования кадра в память GPU"  +/
Сообщение от fail on 08-Июл-16, 14:43 
> Решилось с помощью V4L2_MEMORY_USERPTR.
> В найденном варианте v4l2 нашел почти похожий пример(http://www.friendlyarm.net/forum/topic/1006),
> как в стандартном примере, но пользовательский буфер выделяется не malloc, а
> memalign из malloc.h. Иначе Invalid argument error 22.

мда,
насчет memory alignment - совсем из головы вылетело - хотя:

- специализированные {C,G}PU
- прочая эмбедовка

к этому фактору бывают очень чуствительны..

...

Ответить | Правка | ^ к родителю #18 | Наверх | Cообщить модератору

Архив | Удалить

Рекомендовать для помещения в FAQ | Индекс форумов | Темы | Пред. тема | След. тема




Партнёры:
PostgresPro
Inferno Solutions
Hosting by Hoster.ru
Хостинг:

Закладки на сайте
Проследить за страницей
Created 1996-2024 by Maxim Chirkov
Добавить, Поддержать, Вебмастеру