Общий доступ к поверхностям в OpenCL™ и OpenGL* 4.3 в на ГП Intel® Processor Graphics с помощью неявной синхронизации

Download Sharing Surfaces Code Sample Zipfile

Введение

В этом примере демонстрируется создание текстуры в OpenGL* 4.3, подчиненная область которой обновляется ядром С OpenCL™, выполняющимся на ГП Intel® Processor Graphics под управлением Microsoft Windows*. Одним из назначений такой технологии могут быть приложения компьютерного зрения в реальном времени, где необходимо запускать детектор определенных элементов изображения в OpenCL, но в реальном времени выводить готовое изображение с четко отмеченными детекторами на экран. В этом случае нужен доступ ко всем возможностям языка С ядра OpenCL, а также возможности рендеринга API OpenGL для совместимости с существующим конвейером рендеринга. Еще один пример использования такой технологии: если динамически создаваемые в OpenCL процедурные текстуры используются для рендеринга трехмерных объектов на сцене. И наконец, представьте себе постобработку изображения в OpenCL после рендеринга сцены с помощью 3D конвейера. Это может быть полезно для преобразования цветов, изменения разрешения или выполнения сжатия в определенных сценариях.

В этом примере показано обновление в OpenCL текстуры, созданной в OpenGL. Такие же рекомендации применяются для обновления объекта вертексного буфера или внеэкранного кадрового буфера, который может использоваться в автономном конвейере обработки изображений.

Расширение общего доступа к поверхностям определяется в спецификации расширений OpenCL строкой cl_khr_gl_sharing. Мы также используем расширение cl_khr_gl_event, которое поддерживается ГП Intel.

Мотивация

Назначение этого учебного руководства в том, чтобы ознакомить читателей с возможностью создания поверхностей, общих для OpenCL и OpenGL. Также вы сможете лучше понять работу API, соображения производительности различных путей создания текстур в API OpenGL, в частности на ГП Intel, а также разницу между таким подходом и использованием дискретных ГП.

Основной принцип

Для создания текстур OpenGL и доступа к ним как к изображениям OpenCL с наивысшей производительностью ГП Intel не следует создавать объект пиксельного буфера (РВО) OpenGL. Объекты PBO не обладают преимуществами производительности на ГП Intel. Кроме того, они создают по крайней мере одну дополнительную линейную копию данных, которые затем копируются в формат текстур, используемый в ГП для рендеринга. Во-вторых, вместо использования glFinish() для синхронизации между OpenCL и OpenGL мы можем использовать механизм неявной синхронизации, поскольку ГП Intel поддерживает расширение cl_khr_gl_event.

ГП Intel® с общей физической памятью

ГП Intel® и ЦП вместе используют общую память. Их взаимоотношение показано на рисунке 1. Существует несколько архитектурных механизмов (не показанных на этом рисунке), расширяющих возможности подсистемы памяти. Например, для повышения производительности подсистемы памяти применяются иерархии кэша, сэмплеры, элементарные операции, очереди чтения и записи.

Intel® processor graphics relationship
Рисунок 1. Взаимоотношения между ЦП, ГП Intel® и основной памятью. Обратите внимание, что ЦП и ГП используют общий пул памяти (в отличие от дискретных ГП с собственной выделенной памятью, управление которой осуществляет драйвер)

Почему не следует использовать объекты пиксельного буфера (РВО) с ГП Intel

Спецификация OpenGL рекомендует использовать объекты пиксельного буфера при общем доступе к ресурсам со стороны ЦП и ГП.

«Основное преимущество использования объекта буфера для промежуточного хранения данных текстуры состоит в том, что передача из объекта буфера в текстуру не должна обязательно происходить немедленно, если она происходит до момента, когда данные требуются шейдеру. Это позволяет осуществлять передачу параллельно с выполнением приложения. Если же данные находятся в памяти приложения, то семантика glTexSubImage2D() требует, чтобы перед возвратом функции была создана копия данных, благодаря чему исключается параллельная передача. Преимущество такого подхода состоит в том, что приложение может свободно изменять данные, переданные в функцию, сразу после возврата функции».

Обратите внимание, что смысл этого вызова API заключается в общем доступе между памятью приложения (т. е. памятью ЦП) и ГП, а не в общем доступе между двумя API, каждый из которых выполняет свой поток команд на одном и том же устройстве и одной и той же физической памяти, как показано на рисунке 1.

Использование объектов PBO на самом деле приводит к снижению производительности на устройствах, где используется общая физическая память. Во-первых, объект РВО — это дополнительная промежуточная область, что означает увеличение объема памяти, потребляемого приложением. Во-вторых, данные в РВО хранятся в линейном виде, а если данные требуются в сегментированном виде, как, например, в текстурах OpenGL или в изображениях OpenCL, то приходится преобразовывать данные в нужный формат. И наконец, копирование между двумя API занимает определенное время, что также отрицательно сказывается на производительности приложения.

В случае общего доступа с дискретным ГП использование объектов РВО вполне целесообразно: можно запустить передачу DMA, работающую асинхронно по отношению к ЦП. Без РВО семантика OpenGL требует синхронной записи и дожидается возвращения результата, что также снижает производительность. В нашем случае нет передачи данных из ЦП в подсистему памяти ГП.

В каких случаях можно использовать РВО при общем доступе к поверхностям?

Существуют сценарии, когда имеет смысл применять объекты РВО. Например, если не существует подходящего формата поверхностей, совместимого с OpenGL и OpenCL согласно таблице 9.4 в спецификации расширений OpenCL. В этом случае можно создать РВО и предоставить к нему общий доступ для API, связанных с общим доступом к буферу. Тем не менее старайтесь избегать таких сценариев, чтобы не допустить снижения производительности, о котором было сказано выше. Если это необходимо, см. пример Максима Шевцова, ссылка на который приводится в разделе справочных материалов.

Синхронизация между OpenCL™ и OpenGL*

Во время выполнения важно добиться наивысшей производительности OpenCL и OpenGL. В спецификации сказано следующее:

«Перед вызовом объектов clEnqueueAcquireGLObjects приложение должно убедиться в завершении всех отложенных операций GL, располагающих доступом к объектам, указанным в mem_objects. Чтобы сделать это с сохранением переносимости, можно выполнить и дождаться завершения команды glFinish для всех контекстов GL с отложенными ссылками на эти объекты. В разных реализациях могут быть доступны более эффективные методы синхронизации. Например, на некоторых платформах может оказаться достаточно вызвать glFlush, или же синхронизация может быть неявной внутри потока, или могут быть поддерживаемые данным поставщиком расширения, позволяющие разграничивать поток команд GL и дожидаться завершения каждой части в очереди команд CL. Обратите внимание, что в данный момент единственным методом синхронизации, поддерживающим перенос между различными реализациями OpenGL, является glFinish».

Для наибольшей переносимости, согласно спецификации, нужно вызывать glFinish(), но это блокирующий вызов! На ГП Intel будет эффективнее использовать неявную синхронизацию или объекты синхронизации между OpenCL и OpenGL с расширением cl_khr_gl_events. Подробнее это будет описано ниже. Использование неявной синхронизации не является обязательным. В образце кода содержатся закомментированные фрагменты, которые можно задействовать, если нужно использовать неявную синхронизацию.

Обзор общего доступа к поверхностям для OpenCL и OpenGL

Сначала опишем этапы, необходимые для поддержки общего доступа к поверхностям при инициализации, выполнении и завершении работы. Затем более подробно опишем API и синтаксис языка. И наконец, мы расскажем, как можно развить эти идеи, чтобы охватить другие форматы текстур, выходящие за рамки данного примера. Мы используем общедоступную библиотеку freeglut для управления окнами, а также библиотеку glew. Использование этих библиотек является стандартной практикой в образцах приложений OpenGL, поэтому мы не будем описывать их подробнее.

Инициализация

  1. OpenCL:
    1. Выдайте запрос, чтобы определить, поддерживаются ли расширения; завершение и выход, если не поддерживаются.
    2. Создайте контекст, передающий соответствующие параметры устройства.
    3. Создайте очередь на устройстве и контекст, поддерживающий обмен данными между OpenGL и OpenCL.
  2. OpenGL: Создайте текстуру OpenGL, доступ к которой нужно предоставить для OpenCL.
  3. OpenCL: С помощью дескриптора OpenGL, созданного на шаге 2, создайте общую поверхность посредством расширения OpenCL.

Шаги 1 и 2 можно поменять местами. Шаг 3 должен следовать за шагами 1 и 2.

Запись на общую поверхность в OpenCL

  1. Заблокируйте поверхность для монопольного доступа OpenCL.
  2. Запишите на эту поверхность через ядро C OpenCL. При работе с данными текстур необходимо использовать функции чтения или записи изображения и соответствующим образом передавать изображение.
  3. Разблокируйте поверхность, чтобы предоставить OpenGL доступ к ней на чтение или запись.

Шаги 1, 2 и 3 должны следовать в указанном порядке.

Цикл

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

Завершение работы

  1. Очистка состояния OpenCL
  2. Очистка состояния OpenGL

Подробные сведения об общем доступе к поверхностям OpenGL и OpenCL

В этом разделе приводятся подробные сведения об этапах, описанных в предыдущем разделе.

Инициализация

  1. OpenCL:
    1. Выдайте запрос, чтобы определить, поддерживаются ли расширения; завершение и выход, если не поддерживаются.

      Не все реализации OpenCL поддерживают общий доступ к поверхностям OpenCL и OpenGL, поэтому сначала нужно определить, есть ли вообще в системе нужное расширение. Мы последовательно перебираем платформы, чтобы найти строку расширения для платформы, поддерживающей общий доступ к поверхностям. Внимательное изучение спецификации показывает, что это расширение платформы, а не устройства. Затем мы создаем контекст, который нужно будет опросить, чтобы определить, какие из наших устройств в контексте поддерживают общий доступ к контексту OpenGL.

      Этот пример поддерживается только на ГП Intel, но можно без особых усилий реализовать поддержку и других ГП. Нужное нам расширение — cl_khr_gl_sharing. Вот соответствующий фрагмент кода.

       

      char extension_string[1024];
      memset(extension_string, '', 1024);
      status = clGetPlatformInfo( platforms[i], 
      							CL_PLATFORM_EXTENSIONS,
      							sizeof(extension_string), 
      							extension_string, 
      							NULL);
      char *extStringStart = NULL;
      extStringStart = strstr(extension_string, "cl_khr_gl_sharing");
      if(extStringStart != 0){
      printf("Platform does support cl_khr_gl_sharingn");
      …
      }
      
    2. Если это поддерживается, создайте контекст, передающий соответствующие параметры устройства.

      Если OpenCL поддерживает общий доступ к поверхностям вместе с OpenGL, то нужно создать контекст OpenCL с поддержкой этой возможности. В Windows мы передаем обработчик текущему контексту отрисовки GL и текущему контексту устройства. Обратите внимание, что на разных платформах нужно передавать разные флаги среде выполнения. В таблице 4.5 спецификации расширений OpenCL содержится описание флагов контекста отрисовки, которые нужно передавать API clCreateContext(). CL_WGL_HDC_KHR используется в Windows 7 и Windows 8, а в MacOS — CL_CGL_SHAREGROUP_KHR. Получить эти значения можно несколькими способами, нужно использовать API управления окнами из документации к ОС.

      В примере для Windows мы используем:

       

      //get the GL rendering context
      HGLRC hGLRC = wglGetCurrentContext();
      //get the device context
      HDC hDC = wglGetCurrentDC(); 
      cl_context_properties cps[] = 
      { 
      	CL_CONTEXT_PLATFORM, (cl_context_properties)platformToUse, 
      	CL_GL_CONTEXT_KHR, (cl_context_properties)hGLRC,
      	CL_WGL_HDC_KHR, (cl_context_properties)hDC,
      	0 
      };
      
      //create an OCL context using the context properties 
      g_clContext = clCreateContext(cps, 1, g_clDevices, NULL, NULL, &status);
    3. Создайте очередь на устройстве и контекст, поддерживающий обмен данными между OpenGL и OpenCL

      Мы запрашиваем контекст для устройства с определенного устройства, для которого планируется общий доступ между OpenCL и OpenGL. Поскольку мы уже убедились, что расширение поддерживается, можно получить указатель на нужное расширение:

      clGetGLContextInfoKHR_fn pclGetGLContextInfoKHR = (clGetGLContextInfoKHR_fn)
      clGetExtensionFunctionAddressForPlatform(g_platformToUse, "clGetGLContextInfoKHR");

      Затем запрашиваем идентификатор устройства, поддерживающего общий доступ к поверхностям для OpenCL и OpenGL:

      devID = pclGetGLContextInfoKHR(cps, CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR, bytes, g_clDevices, NULL);

      И наконец, создайте очередь команд для приложения на этом устройстве:

      //create an openCL commandqueue
      g_clCommandQueue = clCreateCommandQueue(g_clContext, devID, 0, &status);
      testStatus(status, "clCreateCommandQueue error");

      Пример кода несколько отличается: здесь код немного изменен, чтобы можно было использовать его как шаблон для передачи CL_DEVICES_FOR_GL_CONTEXT_KHR. Дополнительные сведения см. в спецификации в разделах 9.5.5 и 9.5.6, вопросы 7 и 8, а также в другом примере Intel, указанном в разделе справочных материалов [Шевцов 2014]

  2. OpenGL: создайте текстуру OpenGL, которая будет общей с OpenCL

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

    //create a texture object and assign data to it
    GLuint texture_object;
    glGenTextures(1, &texture_object);
    //bind the texture
    glBindTexture(GL_TEXTURE_2D, texture_object);
    //allocate storage for texture data
    glTexStorage2D(GL_TEXTURE_2D, 1, GL_RGBA8, tex_width, tex_height);
    	
    //specify data	
    glTexSubImage2D(GL_TEXTURE_2D, 
    				0, //mip map level, in this case we only have 1 level==0
    				0, 0, //subregion offset
    				tex_width, tex_height, //width and height of subregion						
    				GL_RGBA, GL_UNSIGNED_BYTE,
    				texture); //works! We specify initial data

    Пропускаем стандартные вызовы API OpenGL для фильтрации, оболочки и множественные отображенияMIP-карты. Они показаны в примере кода. Общий доступ к поверхностям для OpenCL и OpenGL в настоящее время не поддерживает множественные отображенияMIP-карты. Для поддержки множественных отображенийMIP-карт используется отдельное расширение OpenCL cl_khr_mipmap_image.

    В этом случае нужно быломы хотели создать как можно более универсальный код и указать полную поверхность RGBA для общего доступа OpenGL и OpenCL, поскольку в этом случае обеспечивается наибольшая функциональность API. Кроме того, мы используем glTexSubImage2D() в силу универсальности на случай, если нужно будет обновить часть текстуры или текстуру целиком. Мы используем этот API, чтобы передать инициализированный буфер в OpenCL для имитации сценария с поверхностью, в которую записывают и OpenCL, и OpenGL, и для демонстрации того, что записываемые данные OpenGL сохраняются после того, как OpenCL записывает часть подмножества пикселей.

    Чтобы снизить нагрузку на память для общих алгоритмов, можно использовать только один канал для текстур, например, GL_R. Этот сценарий был проверен в примере кода, он работает. Важно убедиться в том, чтобы совпадали форматы, а также совпадала высота и ширина с размером группы обработки, ожидаемой ядром OpenCL. Для этого проще всего использовать только одно значение для каждого из глобальных измерений (высота и ширина), используемых для создания текстур в OpenGL, и в размере группы обработки при вызове clEnqueueNDRangeKernel().

    size_t global_dim[2];
    global_dim[0] = CL_GL_SHARED_TEXTURE_HEIGHT;
    global_dim[1] = CL_GL_SHARED_TEXTURE_WIDTH;
    
    status = clEnqueueNDRangeKernel(g_clCommandQueue, cl_kernel_drawBox, 2, NULL, global_dim, NULL, 0, NULL, NULL);
  3. OpenCL: с помощью дескриптора OpenGL, созданного на шаге 2, создайте поверхность OpenCL с вызовом clCreateFromGLTexture().

    void ShareGLBufferWithCL()
    {
    int status = 0;
    g_SharedRGBAimageCLMemObject = clCreateFromGLTexture(	g_clContext, 
    														CL_MEM_WRITE_ONLY,
     														GL_TEXTURE_2D, 
    														0, 
    														g_RGBAbufferGLBindName, 
    														&status);
    	if(status == 0)
    	{
    		printf("Successfully shared!\n");
    	}
    	else
    	{
    		printf("Sharing failed\n");
    	}
    }

    Это главный вызов API для общего доступа к поверхностям OpenGL и OpenCL. Мы передаем созданный ранее контекст OpenCL, свойства чтения и записи, чтобы описать наши намерения (чтение текстуры OpenGL, запись в нее или чтение с записью) и имя, созданное предыдущим вызовом API OpenGL API для glGenTextureName(). На выходе получаем объект cl_mem, который обрабатывается в ядре OpenCL, как любое обычное изображение. Если ядро работает для изображений, которые находились только по пути OpenCL, это же ядро будет работать и с текстурами OpenGL, если указаны параметр, совпадающие с параметрами, указанными при задании текстуры OpenGL.

Запись на общую поверхность в OpenCL

  1. Заблокируйте поверхность для монопольного доступа OpenCL

    Нужно заблокировать поверхность с помощью вызова API Map/Unmap API к OpenCL при записи OpenCL на эту поверхность. Благодаря этому OpenGL не будет пытаться изменять содержимое или использовать поверхность каким-либо образом, пока OpenCL записывает на эту поверхность. Кроме того, если расширение cl_khr_gl_event не поддерживается, то перед получением объектов OpenGL нужно убедиться в завершении всех операций OpenGL, работающих с объектами OpenGL. В спецификации указано, что единственным переносимым способом решения этой задачи без cl_khr_gl_event является вызов glFinish(). Необходимо убедиться в том, что не выдаются никакие дополнительные команды, влияющие на поверхность GL.

    Но поскольку ГП Intel поддерживает cl_khr_gl_event, можно воспользоваться следующим:

    «Кроме того, это расширение изменяет поведение clEnqueueAcquireGLObjects и clEnqueueReleaseGLObjects; неявно обеспечивается синхронизация с контекстом OpenGL, привязанным к тому же потоку, что и контекст OpenCL».

    Это означает, что даже не нужно создавать сами объекты синхронизации. С помощью clEnqueueAcquireGLObjects() и clEnqueueReleaseGLObjects() синхронизация запускается неявным образом. Это очень удобно и упрощает код, относящийся к созданию и управлению синхронизацией объектов между API, при этом повышается производительность по сравнению с использованием вызовов glFinish() или clFinish().Обратите внимание, что я сначала создал эти объекты и реализовал управление ими в примере кода, а уже потом понял, что они не нужны, и оценил преимущества семантики неявной синхронизации. Кроме того, если все же использовать объекты явной синхронизации, то важно правильно обрабатывать удержание, высвобождение и удаление события и объекта синхронизации, иначе приложение может работать нестабильно.

    Вызов API — clEnqueueAcquireGLObjects(), мы передаем очередь команд OpenCL и объект cl_mem, созданный ранее на шаге 3.

    status = clEnqueueAcquireGLObjects(
    			g_clCommandQueue, 
    			1, 
    			&g_SharedRGBAimageCLMemObject, 
    			0, 0, 0);
  2. Запишите на эту поверхность через ядро C OpenCL. При работе с данными текстур необходимо использовать функции чтения или записи изображения и соответствующим образом передавать изображение

    Ядро в этом примере обновляет только подмножество пикселей в ядре OpenCL C drawRect(). Есть два важных аспекта. Во-первых, подпись ядра:

    kernel void drawBox(__write_only image2d_t output)

    Мы объявили двухмерное изображение, которое передается в ядро с атрибутом __write_only.

    Во-вторых, запись в выходное изображение:  

    write_imagef(output, coord, color);

    При этом происходит запись в выходное изображение в положении (u,v) с указанным значением цвета. Для записи изображений не нужны сэмплеры, но если мы читаем данные из изображения, то есть возможность включения сэмплера.

  3. Отмените монопольный доступ к поверхности, чтобы теперь OpenGL мог использовать ее содержимое.

    Это последний интересующий нас вызов API перед отрисовкойрендерингом: нужно разблокировать поверхность или высвободить ее, чтобы можно было использовать обновленное содержимое в OpenGL. Поскольку ГП Intel поддерживают расширение для синхронизации объектов, можно воспользоваться неявной синхронизацией, нет необходимости вызывать clFinish() перед выполнением команд OpenGL для текстуры.

    status = clEnqueueReleaseGLObjects(g_clCommandQueue, 
    								   1, 
      								   &g_SharedRGBAimageCLMemObject, 
      								   0, NULL, NULL);

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

    void simulateCL()
    {
    	cl_int status;
    	static float fDimmerSwitch = 0.0f;
    
    	status = clEnqueueAcquireGLObjects(g_clCommandQueue, 1, 
    &g_SharedRGBAimageCLMemObject, 0, 0, 0);
    
    	status = clSetKernelArg(cl_kernel_drawBox, 0, sizeof(cl_mem), 
    &g_SharedRGBAimageCLMemObject);
    	testStatus(status, "clSetKernelArg");
    
    	size_t global_dim[2];
    	global_dim[0] = CL_GL_SHARED_TEXTURE_HEIGHT;
    	global_dim[1] = CL_GL_SHARED_TEXTURE_WIDTH;
    
    	status = clEnqueueNDRangeKernel(g_clCommandQueue, cl_kernel_drawBox, 2, NULL, 
    global_dim, NULL, 0, NULL, NULL);
    
    	status = clEnqueueReleaseGLObjects(g_clCommandQueue, 1, 
    &g_SharedRGBAimageCLMemObject, 0, NULL, NULL);
    }
    Вся функция ядра OpenCL C тривиальна, она просто обновляет подмножество значений изображения.
    kernel void drawBox(__write_only image2d_t output, float fDimmerSwitch)
    {
        int x = get_global_id(0);
        int y = get_global_id(1);
    
        int xMin = 0, xMax = 1, yMin = 0, yMax = 1;
    
        if((x >= xMin) && (x <= xMax) && (y >= yMin) && (y <= yMax))
        {      
            write_imagef(output, (int2)(x, y), (float4)(0.f, 0.f, fDimmerSwitch, 1.f));
        }
    }

Завершение работы

  1. Очистка объекта памяти CL

    Объектам OpenCL требуется очистка. Здесь нас интересует только объект cl_mem связанный с текстурой OpenGL.

    //cleanup all CL queues, contexts, programs, mem_objs
    status = clReleaseMemObject(g_SharedRGBAimageCLMemObject);
  2. Очистка поверхностей GL

    glDeleteTextures(1, &g_RGBAbufferGLBindName);

    В примере кода показаны другие объекты, которые необходимо очистить для OpenGL и OpenCL.

Был бы код быстрее при синхронизации объектов вместо использования неявной синхронизации?

В данное время ГП Intel поддерживают переключение между контекстами очистки OpenCL и OpenGL. Поэтому использование объектов явной синхронизации не приведет к повышению производительности. Кроме того, при использовании неявной синхронизации разработчикам не требуется вставлять собственные функции очистки или завершения.

Дополнительные сведения

В примере кода нужно задать для флага glewExperimental значение GL_TRUE перед вызовом glewInit() чтобы использовать объекты массива вертексов.

Задачи на будущее

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

События явной синхронизации между OpenCL и OpenGL

В OpenCL можно создавать объекты событий OpenCL объектами разграничения GLsync с помощью расширения cl_khr_gl_event. Существует и расширение OpenGL для общего доступа к событиям вместе с OpenCL: GL_ARB_cl_event В статье заявлено, что использование этих расширений не дает прироста производительности, но неплохо бы подтвердить это заявление настоящим кодом. Кроме того, в более сложных сценариях использования может потребоваться применение явной синхронизации, поэтому будет полезно понимать, что делать в таких случаях.

Общий доступ к буферам, данным глубины, формата и сглаживания с многоступенчатой выборкой (MSAA)

Расширения общего доступа к поверхностям OpenCL с поддержкой OpenGL допускают общий доступ ко всем типам поверхностей, которые могут быть созданы в OpenGL (буферы, глубина и пр.), но существуют некоторые ограничения по форматам этих поверхностей. Ограничения описаны в спецификации расширения. Это учебное руководство посвящено общему доступу к текстуре, которая отображается или используется в качестве карты текстур в OpenGL, но общий доступ к другим поверхностям работает точно так же. Рекомендации по общему доступу к текстурам и буферам OpenGL аналогичны для ГП Intel: в OpenGL не нужно создавать объекты пиксельного буфера и буфера вертексов при общем доступе к поверхностям между API.

Двойная Буферизация

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

Что делать, если не поддерживается общий доступ к поверхностям?

В разделе справочных материалов Максим Шевцов приводит примеры кода, касающиеся случаев, когда необходимо копирование между OpenCL и OpenGL. Рекомендую ознакомиться с этим примером кода для данного сценария использования. В настоящее время Intel не поддерживает, например, общий доступ к поверхностям в Linux*. Это решение может быть изменено в зависимости от пожеланий заказчиков.

Пример общего доступа к поверхностям

Мы написали это учебное руководство для демонстрации общего доступа к поверхностям. Мы проверили его на ГП Intel с драйверами OpenCL 2.0 и OpenGL 4.3. Но этот же код с незначительным портированием, по всей вероятности, будет вполне работоспособным на гораздо большем количестве платформ и устройств. Программируемые шейдеры вертексов и пикселей в OpenGL тривиальны и также должны работать в гораздо более ранних версиях OpenGL. Ядро OpenCL C очень простое, но ясно демонстрирует принципы, описанные в этой статье.

Зависимости

Для сборки программ, описанных в этой статье, нужно загрузить определенные библиотеки и задать соответствующие пути include и library. URL-адрес этих библиотек приводятся ниже в разделе справочных материалов.

  • freeglut.h, freeglut.dll- в путь следует добавить расположение этой библиотеки. Путь к двоичному файлу freeglut в Интернете указан в разделе справочных материалов. После загрузки можно распаковать файл freeglut.dll, находящийся в папке freeglut\bin
  • glew.h, glew32s.lib-это библиотека из GLEW 1.11.0, она обрабатывает все управление API OpenGL и расширением за сценой. Также обратите внимание на #define GLEW_STATIC перед включением glew.h.
  • cl.h, cl_gl.h, openCL.lib-из Intel® OpenCL™ SDK
  • gl.h - входит в состав каталога Windows Kits 8.0

Я использовал следующие параметры, хотя в вашем случае они могут несколько отличаться:

  • Библиотека freeglut.dll скопирована в папку debug или release учебного руководства на уровне решения (а не на уровне проекта). Вы можете обрабатывать DLL иначе, но этот способ тоже работает.
  • Задайте путь к библиотеке glew32s.lib:
    • C:\src\glew-1.11.0\lib\Release\Win32
  • Добавьте расположение cl.h и cl_gl.h в путь к включаемым файлам. Например:
    • C:\Program Files (x86)\Intel\OpenCL SDK\3.0\include\CL
  • Добавьте расположение библиотеки OpenCL в путь к библиотекам: Например:
    • C:\Program Files (x86)\Intel\OpenCL SDK\3.0\lib\x86
  • Добавьте OpenCL.lib в набор статически компонуемых библиотек.

Файл примера и структура папок

В этом разделе содержатся сведения о том, как разделяется этот код. При этом упор делается на создание простого примера на С, а не на высококачественной реализации продукта.

Проект находится в одной папке: CL_20_GL_43_surface_sharing. Название указывает, что проект был протестирован с OpenCL 2.0 и OpenGL 4.3. Но здесь не используется ничего, относящегося именно к версии OpenCL 2.0; мы даже не включаем компилятор OpenCL 2.0 при компиляции ядра OpenCL. В OpenGL мы используем очень простые программируемые шейдеры вертексов и пикселей. Using #version we requested at least version 3.3 for the shaders.

Используются следующие файлы:

  • main.cpp: OpenCL, OpenGL и систему управления окнами. Здесь также содержится обработчик клавиатуры для обработки событий клавиши Escape.
  • API OpenGL и OpenCL находятся в отдельных файлах, соответственно OGL.h, OGL.cpp, OCL.h и OCL.cpp. Некоторые общие используемые флаги и переменные находятся в commonCLGL.h.
  • Шейдер OpenCL, просто рисующий прямоугольник на общем изображении, находится в файле OpenCLRGBAFile.cl , а шейдеры OpenGL находятся в файлах triangles.frag and triangles.vert.

Сборка и запуск примера

Соберите это решение, выбрав в главном меню Сборка ->Собрать решение. Должны быть созданы все исполняемые файлы. Их можно запускать напрямую в Visual Studio или перейти в папки Debug или Release, расположенные там же, где и файл решения CL_20_GL_43_surface_sharing solution file.

Для запуска нажмите клавишу F5 в среде разработки Visual Studio. Для запуска из командной строки нужно скопировать шейдеры из папки проекта в папку с выполняемым файлом. Три нужных нам файла ядра и шейдеров: OpenCLRGBAFile.cl, triangles.frag, and triangles.vert.

Дополнительные сведения

У Максима Шевцова также есть учебное руководство по общему доступу к поверхностям с действующим примером кода: https://software.intel.com/
en-us/articles/opencl-and-opengl-interoperability-tutorial
. Он обсуждает недостатки использования объектов пиксельного буфера, использование glMapBuffer() и предоставляет образец кода. Также показано, как действовать в случае, если расширение не поддерживается, например, в реализации OpenCL на Linux. Пример кода создает окно с помощью традиционных API win32 и использует фиксированный графический конвейер OpenGL, записывая последовательно изменяющиеся цвета на поверхность. Также делаются попытки запуска на других платформах, включая ЦП Intel, и предоставляются очень полезные ссылки на дополнительные материалы.

В исходном коде в этой статье мы используем freeglut, glew, программируемые шейдеры вертексов и пикселей OpenGL 4.3, наложение текстур на ориентированную по экрану многоугольную поверхность, которая отрисовывается на экране. Запись производится только в часть текстуры в ядре OpenCL, чтобы продемонстрировать, что получающийся цвет пикселей является сочетанием всех попиксельных операций OpenGL и конвейера OpenCL, работающих вместе.

sample execution
Рисунок 2. Ожидаемый результат выполнения примера. Цвет фона — зеленый. Изображение представляет собой прямоугольник, ориентированный согласно положению экрана и состоящий из двух треугольников с нанесенными на них текстурами. Текстура представляет собой небольшую красную карту текстур 4x4, при этом левая нижняя часть текселей записывается с помощью OpenCL после первоначального заполнения с помощью OpenGL. OpenCL записывает в канал синего цвета значение, циклически изменяющееся от черного до синего (то есть от 0 до 255 в канале синего).

Благодарности

В подготовке этого материала мне помогли: Мурали Сундаресан, Аарон Кунце, Аллен Хакс, Паван Ланка, Максим Шевцов, Михаль Мроцек, Петр Умински, Стивен Джанкинс, Дэн Петр и Бен Эшбоу. Все они участвовали в технических обсуждениях, разъясняли нужные вопросы и всегда с готовностью проверяли написанные материалы.

Об авторах

Адам Лейк (Adam Lake) работает на должности старшего графического архитектора в компании Visual Products Group и является представителем с правом голоса в организации по стандартам Khronos OpenCL Он занимается программированием ГП уже свыше 12 лет. Ранее он работал в области виртуальной реальности, трехмерной и обычной графики, занимался компиляторами языков поточного программирования

Роберт Иоффе (Robert Ioffe) работает на должности технического инженера-консультанта в отделе Software and Solutions Group корпорации Intel.  Роберт принимал активное участие в разработке стандартов Khronos, занимался прототипированием самых последних функций и обеспечением их работоспособности в архитектуре Intel. 

Справочные материалы

  1. Спецификация OpenCL 1.2: https://www.khronos.org/registry/cl/
  2. Спецификация OpenCL 2.0, состоящая из трех томов: спецификация OpenCL для языка C, интерфейс API среды выполнения OpenCL и расширения OpenCL: https://www.khronos.org/registry/cl/
  3. Документ Стивена Джанкинса: Вычислительная архитектура Intel® поколения 7.5: https://software.intel.com/sites/default/files/managed/f3/13/Compute_Architecture_of_Intel_Processor_Graphics_Gen7dot5_Aug2014.pdf. Это обязательный к ознакомлению материал для всех, кто использует OpenCL на платформах с ГП Intel.
  4. Учебное руководство Адама Лейка по общему доступу к поверхностям с нулевым копированием: https://software.intel.com/ru-ru/articles/getting-the-most-from-opencl-12-how-to-increase-performance-by-minimizing-buffer-copies-on-intel-processor-graphics
  5. Руководство Максима Шевцова по взаимодействию: https://software.intel.com/ru-ru/articles/opencl-and-opengl-interoperability-tutorial
  6. Исходный код Freeglut: http://freeglut.sourceforge.net/
  7. Библиотека Freeglut .dll: www.transmissionzero.co.uk/software/freeglut-devel/
  8. Библиотека GLEW: http://glew.sourceforge.net/
    1. Также обратите внимание на этот материал: http://stackoverflow.com/questions/13558073/program-crash-on-glgenvertexarrays-call

Определения

Ниже приводятся определения некоторых терминов, использованных в этой статье. Подробные сведения см. в разделе справочных материалов.

  • Буферы: в OpenCL есть различие между буферами и изображениями. В OpenCL буферы располагаются линейно в памяти, поэтому буфер можно считать массивом.
  • Текстуры: буферы данных в сегментированном формате, прочитываемых сэмплерами OpenGL. Такой способ работы с памятью повышает производительность за счет применения сэмплеров текстур, отфильтровывающих прочитываемые в памяти пиксели с помощью заранее заданных ядер фильтров. 
  • Поверхность: буферы, текстуры или изображения. Это общий термин для данных в памяти, которые могут быть линейными или прямоугольными.  В некоторых случаях у поверхности могут быть дополнительные данные, например, измерение, высота, ширина и атрибуты разметки. Для управления этими атрибутами используется API (OpenCL, OpenGL, DirectX и пр.).
  • Сэмплеры: механизмы выборки, используются для чтения изображений в OpenCL и текстур в OpenGL. Сэмплер использует внутренние кэши и сегментированный формат находящихся в памяти изображений или текстур для повышения производительности при фильтрации. Сэмплеры включают кэши и логику для одновременной выборки на разных уровнях (а также на уровнях множественных отображений) и вывода одного значения текселя для одного запроса.
  • Изображения: буферы данных в сегментированном формате, прочитываемых сэмплерами OpenCL. Они равноценны текстурам OpenGL. Конкретный набор поддерживаемых и пригодных для общего доступа форматов зависит от конкретной реализации.
  • Общий доступ к поверхностям: имеется в виду общий доступ разных API к поверхностям, этот термин подразумевает создание поверхности в одном API и использование данных в другом. Цель состоит в снижении количества создаваемых копий одной и той же поверхности, но для достижения этого результата необходимо соблюдать набор ограничений, зависящих от устройства. В этом учебном руководстве описываются такие ограничения для ГП Intel..
  • Наложение текстур: сопоставление пикселей, находящихся в памяти, с многоугольником в графическом конвейере. В этом примере мы накладываем текстуру OpenGL на два ориентированных по экрану многоугольника для отображения.
  • Нулевое копирование: это технический жаргон, не вполне точно применяемый к главной системе (ЦП) и устройству (ГП). В этой статье под нулевым копированием понимается отсутствие необходимости копирования текстуры (изображения, буфера и т. д.) между потоками команд OpenGL и OpenCL. Это возможно, поскольку в обоих случаях используется одно и то же расположение хранилища и совместимая параметризация поверхности. При нулевом копировании объем хранилища снижается пропорционально размеру буфера, а также повышается производительность, поскольку отсутствует копирование в системе, где к фактическому хранилищу поверхности нет общего доступа.
  • Общая физическая память: система и устройство вместе используют одну и ту же физическую память DRAM. Этот сценарий отличается от общей виртуальной памяти, когда система и устройство вместе используют одни и те же виртуальные адреса (такое решение не рассматривается в этом документе).  Важная особенность оборудования, делающая возможным нулевое копирование, заключается в том, что у ЦП и ГП общая физическая память. Применение общей физической памяти и общей виртуальной памяти не исключают друг друга. Устройства, поддерживающие общую виртуальную память, могут не «видеть» всю физическую память.
  • Intelprocessorgraphics: этот термин обозначает текущие графические решения Intel. Наименования продуктов, где ГП Intel интегрированы в платформу с архитектурой «система на кристалле», включают Intel® Iris™, Intel® Iris™ Pro и Intel® HD Graphics (в зависимости от конкретной платформы). Дополнительные сведения об архитектуре оборудования см. в документе по вычислительной архитектуре Intel® поколения 7.5, ссылка на который приводится в разделе справочных материалов конце этой статьи, или по адресу http://ark.intel.com/ru/.

Дополнительные сведения об оптимизации компиляторов см. в нашем уведомлении об оптимизации.

AnexoTamanho
Image icon 1.png14.26 KB
Para obter informações mais completas sobre otimizações do compilador, consulte nosso aviso de otimização.