Знакомство с 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. Из кода я выкинул участок, отвечающий за вычисление времени создания кадра и отображение его в заголовке приложения.
Когда я отошел от использования простых структур представляющих заряд, в сторону float4, то смог сделать выполнение кода быстрее. Так, например, кадр 1024*768 с 10ю зарядами создаётся теперь за 25-30 мс. И это позволило мне создавать плавную анимацию с помощью мыши — drag & drop:
Ну и как всегда, готовый код вы можете найти здесь. И если у вас возникнут вопрос или пожелания, пишите мне. Постараюсь исправиться.

Послесловие

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