Знакомство с OpenGL Interoperability

Всем доброго дня,

Надеюсь, при прочтении этого блока в своём ридере, моя картинка вас не напугала. Но сегодня, я хочу описать применение взаимодействия технологии CUDA с OpenGL на примере моего небольшого pet-примера, первую версию которого я описывал в статье ранее. Тех, кому интересен раздел, известный под английским названием CUDA and OpenGL interoperability, прошу кликать сюда.

Все статьи в Гугле рассказывают о теории следующим образом: по классике, перед рендерингом кадра, программа формирует логику и инициализацию сцены в CPU, и потом уже GPU занимается её рендерингом. А теперь, представьте, что вы ту же саму сцену инициализируете в многопоточной среде. Например, вы можете сгенерировать массив точек в трёхмерной системе координат с помощью CUDA, или же вообще нарисовать картинку самостоятельно в своём ядре (kernel), и с помощью OpenGL просто отобразить результат на экране.

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

Когда начинаешь создавать шаблон проекта, очень важно не упустить ряд моментов:

1. Для некоторых функций, необходимых для привязки к буферу, необходим GLEW, то его заголовочный следует подключать над заголовочным файлом freeglut.

   #include <GL/glew.h>
   #include <GL/freeglut.h>
   ...
   #include <cuda_gl_interop.h>
   

   p.s. если что, вас предупредит компилятор.
2. Когда объявите нужные переменные указывающие на ресурс CUDA и на видео буфер, то случайно не очистите их где-нибудь в середине кода, думая, что они в текущем пробеге функции не нужны. Я долго голову ломал, пытаясь понять почему у меня возникает segfault ошибка:

   GLuint vbo;
   struct cudaGraphicsResource *cuda_vbo_resource;
   

3. Теперь мы можем использовать эти переменные в функциях отвечающих за связывание буфера кадра с CUDA:

   void createVBO(GLuint *vbo, struct cudaGraphicsResource **vbo_res, unsigned int vbo_res_flags) {
   	unsigned int size = width * height * sizeof(uchar4);
   
   	glGenBuffers(1, vbo);
   	glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, *vbo);
   	glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, size, NULL, GL_DYNAMIC_DRAW);
   
   	HANDLE_ERROR( cudaGraphicsGLRegisterBuffer(vbo_res, *vbo, vbo_res_flags) );
   }
   
   void deleteVBO(GLuint *vbo, struct cudaGraphicsResource *vbo_res) {
   	HANDLE_ERROR( cudaGraphicsUnregisterResource(cuda_vbo_resource) );
   
   	glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, *vbo);
   	glDeleteBuffers(1, vbo);
   
   	*vbo = 0;
   }
   

   Основная идея в том, что позже, в коде, можно получить указатель на созданный буфер в виде массива пикселей, и потом менять его в ядре. К сожалению, я не сильно сведущ в OpenGL, и потом мне не ясен момент, почему часто перед удалением, вызывают функцию glBindBuffer ещё раз, и передают 1 в качестве первого аргумента. Если я правильно понимаю, то если я использую только один буфер в своём приложении, то я сразу могу очистить его после завершения.
   
4. Вспомогательные функции созданы, теперь мы можем их использовать для создания и освобождения ресурса CUDA (и связывания пиксельного буфера):

   HANDLE_ERROR( cudaGLSetGLDevice(deviceId) );
   ...
   createVBO(&vbo, &cuda_vbo_resource, cudaGraphicsMapFlagsWriteDiscard);
   ...
   deleteVBO(&vbo, cuda_vbo_resource);
   

Вроде бы всё! Теперь мы можем обработать idle событие у приложения, где заниматься формированием кадра, а потом просто отрисовывать его. Следует напомнить, что мы рисуем кадр сразу же в буфере видеокарты, и потом просто просим её его отобразить.

void idle(void) {
	uchar4* dev_screen;
	size_t size;

	HANDLE_ERROR( cudaGraphicsMapResources(1, &cuda_vbo_resource, 0) );
	HANDLE_ERROR( cudaGraphicsResourceGetMappedPointer((void**) &dev_screen, &size, cuda_vbo_resource) );

	// Render Image
	renderFrame<<<blocks, threads>>>(dev_screen);
	HANDLE_ERROR( cudaDeviceSynchronize() );

	HANDLE_ERROR( cudaGraphicsUnmapResources(1, &cuda_vbo_resource, 0) );

	glutPostRedisplay();
}

p.s. Из кода я выкинул участок, отвечающий за вычисление времени создания кадра и отображение его в заголовке приложения.
Отрисовкой занимается функция draw, которая вызывается freeglut автоматически или по требованию glutPostRedisplay:

void draw(void) {
	glClearColor(0.0, 0.0, 0.0, 1.0);
	glClear(GL_COLOR_BUFFER_BIT);

	glDrawPixels(width, height, GL_RGBA, GL_UNSIGNED_BYTE, 0);

	glutSwapBuffers();
}

Обратите внимание, что я не указываю массив пикселей хранящих картинку, потому что буфер кадра уже изменён и нужно просто сообщить OpenGL о необходимости отражения изменений.

Когда я отошел от использования простых структур представляющих заряд, в сторону float4, то смог сделать выполнение кода быстрее. Так, например, кадр 1024*768 с 10ю зарядами создаётся теперь за 25-30 мс. И это позволило мне создавать плавную анимацию с помощью мыши — drag & drop:

Ну и как всегда, готовый код вы можете найти здесь. И если у вас возникнут вопрос или пожелания, пишите мне. Постараюсь исправиться.

Послесловие

В будущем планирую сделать режим отображающий напряжённость поля в виде опилок выброшенных на бумагу. Идея — рисовать короткие отрезки через равные промежутки. Ну и если уж совсем повезет, то попробую реализовать это в 3д :)