26 июля 2011 в 09:47

Введение в OpenCL

Эта статья посвящена основам программирования на OpenCl. OpenCl -это язык программирования на GPU/CPU, по своей структуре близкий к стандарту c99. Его развитием занимается Khronos Group, где на их сайте доступна полная документация. Во избежание полемики на тему «ну это же всё тривиально, достаточно покопаться в инете» сразу оговорюсь: в рунете информация на эту тематику практически полностью отсутствует, а в западном инете доступна весьма в разрозненном состоянии на десятке сайтов. Здесь будет приведена некоторая компиляция базовых принципов, максимально упрощающая начинающему программисту жизнь, а так же позволяющая с самого первого проекта максимально задействовать вычислительные мощности видеокарты. Людям написавшим 2-3 серьёзных программы на OpenCl это будет уже неинтересно. Статья в некотором смысле является продолжением моей прошлой статьи.

Компилятор


В первую очередь вопрос: где писать сам код. Насколько мне известно под .NET пока что нет никакой свистелки, позволяющей обрабатывать код ядра непосредственно в студии. Поэтому приходиться использовать сторонние редакторы. AMD, nVidia и Intel прилагают их к своим пакетам SDK. Мне почему-то больше нравиться именно Интеловский. Так же, как вариант, есть несколько редакторов, написанных фанатами. Из них мне больше всего нравиться редактор, прилагающийся к OpenCLTemplate. Стоит отметить, что это именно редакторы. Компиляция кода происходит непосредственно перед запуском на GPU/CPU.

Модель памяти устройства


image
Прежде чем описывать сам язык я дам краткое описание физической модели устройства с которой он взаимодействует. Исполнение команд языка идёт на объектах, называемых «work-item». Каждый «work-item» не зависим от другого и может исполнять код параллельно с остальными. Если же процесс из одного work-item хочет получить данные, используемые или уже обработанные любым другим work-item он может это сделать через общую память. Общая память весьма медленная, зато имеет большой объём. Чтобы ускорить вычисления имеется локальная память. Если вы знакомы с CUDA, то там она называется «разделяемая память». Она значительно быстрее общей, но не любой процесс может получить к ней доступ. К локальной памяти могут обращаться только work-item одной группы. Эти группы называются «Compute Unit» или «Workgroup» (первое название относится к физическому разбиению на уровне железа, а второе к логическому на уровне программы). В зависимости от устройства в каждой из этих групп различное количество work-item (например 240 для NVIDIA GT200 или 256 для Radeon 5700 Series). Количество этих юнитов ограниченно достаточно маленьким числом (30 для NVIDIA GT200 или 9-10 для Radeon 5700 Series). Так же существует сверхбыстрая «private memory» к которой work-item может обращаться единолично.
Драйвера OpenCL устройств автоматизируют старт и работу work-item и workgroup. Например если нам нужно выполнить миллион процессов, а у нас в распоряжении всего тысяча work-item, то драйвера будут автоматически запускать каждый процесс со следующей задачей после его завершения. Понимание физического уровня требуется только для того, чтобы иметь представление о возможностях взаимодействия между процессами и доступа процессов в память.

Базовые особенности


Так как в основе языка лежит практически стандартный с++ я буду рассматривать только те особенности, которые отличают OpenCL от него. Рассмотрим код простейшей программы ядра, который я приводил в прошлой статье. Этот код складывает два вектора, v1 и v2, кладя результат в первый из них.
__kernel void
          floatVectorSum(__global    float * v1,
                  __global    float * v2)
          {
            int i = get_global_id(0);
            v1[i] = v1[i] + v2[i];
          }

* This source code was highlighted with Source Code Highlighter.


Объявление процедур

В первую очередь в глаза бросается загадочный "__kernel ". Этой директивой должна быть помечена любая процедура, которую мы хотим вызвать извне. Если процедура не нужна при работе извне, её можно не отмечать.

Типы памяти

Тип данных "__global " обозначает память, которая выделяется из глобального адресного пространства работающего устройства. Она достаточна медленная, зато вместительная. Для современных видеокарт измеряется гигобайтами. Если вы работаете на процессоре — под global подразумевается оперативная память.
Кроме global есть "__local ". К ней может обращаться только рабочая группа(workgroup). На каждую такую группу выделяется примерно 8 килобайт.
Так же быстрой памятью является "__privat ". Это память к которой имеет доступ только отдельный поток (work-item). Всего на поток выделяется 32 регистра этой памяти.
Остальные типы памяти, которые можно объявлять при создании ядра основаны на типе "__global ". Во-первых, это "__constant ", который может использоваться только для чтения. Во-вторых, это "__read_only", "__write_only" и "__read_write" — структуры, использование которых разрешено только для изображений.

Идентификаторы процессов

После запуска на видеокарте все процессы равнозначны и исполняют равнозначный код. Но, очевидно, нам не нужно многократное повторение одного и того же действия — каждый процесс должен делать свой кусок задачи. Для осознания своего места в окружающем мире служат идентификаторы процессов. Самый простой идентификатор — " get_global_id(0)". В случае приведённого примера он указывает на i номер вектора, который должен сложить этот процесс. Если же мы обрабатываем не одномерный вектор, а двухмерное изображение — нам нужно знать положение процесса по двум осям. Конечно, это значение можно вычислить. Но это лишние операции. Поэтому для удобства при запуске можно указать, что нам нужно пространство двухмерной размерности. Тогда в процессе можно получить оба идентификатора положения: «get_global_id(0)», «get_global_id(1)». Так же можно сделать и для трёхмерного пространства. Часто может потребоваться и размерности пространства в которой мы работаем. Например для изображения практически при любой его обработке нам нужны его ширина и высота. Для получения размерности пространства используется идентификатор «get_global_size(i)». Кроме этого есть идентификаторы процессов внутри рабочей группы — «get_local_id(i)», «get_local_size(i)» и идентификатор самой группы -" get_group_id(i)", «get_num_groups(i)». Большая часть этих соотношений связанна друг с другом: num_groups * local_size = global_size, local_id + group_id * local_size = global_id, global_size % local_size = 0.

Оптимизация расчётов


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

Встроенные вектора

Самой первой особенностью являются вектора и векторная математика. В OpenCl можно объявлять в качестве переменных 2, 4, 8 и 16-мерные вектора. Это делается соответственно: int2, int4, int8, int16. Так же можно объявлять double, byte и все остальные типы. Вектора соответствующей размерности можно складывать/вычитать/делить/умножать, а так же любой вектор можно делить/умножать на число:
uint4 sumall = (uint4)(1,1,1,1);
small += (uint4)(1,1,1,1);
sumall = sumall/2;

* This source code was highlighted with Source Code Highlighter.

Кроме того есть ряд функций, оптимизированных под вектора и позволяющих работать непосредственно с ними. К таким функциям относятся функции вычисления расстояния, функции векторного произведения. Например:
float4 dir1 = (float4)(1, 1, 1, 0);
float4 dir2 = (float4)(1, 2, 3, 0);
float4 normal = cross(dir1, dir2);

* This source code was highlighted with Source Code Highlighter.

Ещё можно сливать вектора, беря части от одного и другого, а так же склеивать в более большие:
int4 vi0 = (int4) -7 ;
int4 vi1 = (int4) ( 0, 1, 2, 3 ) ;
vi0.lo = vi1.hi; // слияние
int8 v8 = (int8)(vi0.s0123, vi1.s0123); //склейка

* This source code was highlighted with Source Code Highlighter.


Простые функции

Следующей особенностью OpenCl является встроенная библиотека функций. Кроме стандартного набора math.lib в OpenCl имеются так называемые native функции. Это функции, основаны непосредственно на использовании некоторых функций видеокарт и на загрублённой математике. Не советуется применять их при сверхточных расчётах, но в случае фильтрации изображений разницу невозможно заметить. К таким функциям, например, относятся: «native_sin», «native_cos», «native_powr». Я не буду приводить более подробное объяснение этих функций, их очень много, да и принципы разные. Если они вам понадобятся — смотрите документацию.

Часто встречающиеся функции

Кроме «простых функций» разработчики создали целый ряд называемый common function. Это функции, часто встречающиеся при обработке изображений. Например: mad(a,b,c) = a*b + c, mix(a,b,c) = a + (b-a)*c. Эти функции выполняются быстрее, чем соответствующие им математические действия.

Пример

На сайте www.cmsoft.com.br есть замечательный пример, показывающий возможности при оптимизации кода средствами native и common функций:
kernel void regularFuncs()
{
  for (int i=0; i<5000; i++)
  {
    float a=1, b=2, c=3, d=4;
    float e = a*b+c;
    e = a*b+c*d;
    e = sin(a);
    e = cos(b);
    e = a*b+c*d;
    e = sin(a);
    e = cos(b);
    e = a*b+c*d;
    e = sin(a);
    e = cos(b);
    float4 vec1 = (float4)(1, 2, 3, 0);
    float4 vec2 = (float4)(-1, 3, 1, 0);
    float4 vec = distance(vec1, vec2);
    double x=1, y=2, z=3;
    double resp = x*y+z;
  }
}
kernel void nativeFuncs()
{
  for (int i=0; i<5000; i++)
  {
    float a=1, b=2, c=3, d=4;
    float e = mad(a,b,c);
    e = mad(a,b,c*d);
    e = native_sin(a);
    e = native_cos(b);
    e = mad(a,b,c*d);
    e = native_sin(a);
    e = native_cos(b);
    e = mad(a,b,c*d);
    e = native_sin(a);
    e = native_cos(b);
    float4 vec1 = (float4)(1, 2, 3, 0);
    float4 vec2 = (float4)(-1, 3, 1, 0);
    float4 vec = fast_distance(vec1, vec2);
    double x=1, y=2, z=3;
    double resp = mad(x,y,z);
  }
}

* This source code was highlighted with Source Code Highlighter.

Вторая процедура (использующая оптимизацию) выполняется в 35 раз быстрее.

Разрешения


Стоит отметить, что в OpenCl есть ряд директив, позволяющих включать различные дополнительные функционалы. Причин к этому две. Первая — исторически не все из этих функционалов поддерживались. Вторая — эти функционалы могут влиять на производительность. Обычно функционал включается следующей командой:
#pragma OPENCL EXTENSION extension name : behavior

* This source code was highlighted with Source Code Highlighter.

Для примера. Слудющие команды включают: возможность использования типа byte, двойную точность рассчётов и все математические функции
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#pragma OPENCL EXTENSION cl_khr_fp64 : enable

* This source code was highlighted with Source Code Highlighter.


Синхронизация



Барьеры

Часто в вычислениях нужна синхронизация. Это достигается несколькими способами. Первый — барьеры. Барьер это такая команда, на которой остановится процесс пока все прочие процессы или процессы его рабочей группы не достигнут. Приведём два примера:
kernel void localVarExample()
{
  int i = get_global_id(0);
  __local int x[10];
  x[i] = i;
  barrier(CLK_LOCAL_MEM_FENCE);
  if (i>0) int y = x[i-1];
}
kernel void globalVarExample()
{
  int i = get_global_id(0);
  __global int x[10];
  x[i] = i;
  barrier(CLK_GLOBAL_MEM_FENCE);
  if (i>0) int y = x[i-1];
}

* This source code was highlighted with Source Code Highlighter.


В первом примере на команде barrier ожидают все процессы рабочей группы, во втором — все процессы OpenCL устройства.
Стоит отметить особенность этого примера, команды "__local int x[10];" и "__global int x[10];". Они позволяют выделить глобальную переменную в группе процессов и во всех процессах уже во время их исполнения.

Единичные операции

Второй вариант синхронизации между потоками — atomic. Это функции, предотвращающие одновременное обращение к памяти. Перед их использованием нужно включить следующие директивы:
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable

* This source code was highlighted with Source Code Highlighter.

Простейший пример работы этой функции:
__kernel void test(global int * num)
{
 atom_inc(&num[0]);
}

* This source code was highlighted with Source Code Highlighter.

Если бы вместо «atom_inc(&num[0]);» было написано num++, то результат исполнения программы был непредсказуем, так как все процессы бы одновременно обратились к памяти и считали там одно и то же значение. Всего есть 11 функций еденичных операций: «add, sub, xchg, inc, dec, cmp_xchg, min, max, and, or, xor».
При помощи этих функций не сложно создать семафор.
void GetSemaphor(__global int * semaphor) {
  int occupied = atom_xchg(semaphor, 1);
  while(occupied > 0)
  {
   occupied = atom_xchg(semaphor, 1);
  }
}

void ReleaseSemaphor(__global int * semaphor)
{
  int prevVal = atom_xchg(semaphor, 0);
}

* This source code was highlighted with Source Code Highlighter.


Работа с изображениями



Последней вещью, которую я хочу включить в этот guide является работа с изображениями через OpenCL. Создатели попробовали сделать так, чтобы работа с изображениями требовала минимума мозга пользователя. Это очень приятно. Загрузка изображений возможна в типы image2d_t и image3d_t. Первые — это обычные изображения, вторые — трёхмерные. Так же загружаемое изображение должно быть одного из форматов: " __read_only", " __write_only", "__read_write". Чтение и запись данных из изображения возможны только специальными процедурами: значение = read_imageui(изображение, сэмплер, положение), write_imageui(изображение, положение, значение).
На мой взгляд здесь всё понятно кроме понятия «сэмплер». Сэмплер — это штука, которая будет оптимизировать вашу работу с изображением. У него есть три параметра: «normalized coords», «address mode», «filter mode». Первый имеет два значения: «CLK_NORMALIZED_COORDS_TRUE, CLK_NORMALIZED_COORDS_FALSE». В соответствии с названием он должен показывать, нормализованы ли входные координаты или нет. Второй показывает, что делать в случае, если вы пробуете прочитать координаты из-за пределов границ изображения. Возможные варианты: зеркально продолжить изображение(CLK_ADDRESS_MIRRORED_REPEAT), взять ближайшее граничное значение (CLK_ADDRESS_CLAMP_TO_EDGE), взять базовый цвет (CLK_ADDRESS_CLAMP), ничего не делать (пользователь гарантирует что такого не произойдёт CLK_ADDRESS_NONE). Третий показывает, что делать, если на входе не целые координаты. Возможные варианты: приблизить ближайшим значением (CLK_FILTER_NEAREST), линейно проинтерполировать (CLK_FILTER_LINEAR).
Краткий пример. Замыливаем изображение по среднему значению в области:
__kernel void ImageDiff(__read_only image2d_t bmp1, __write_only image2d_t bmpOut)

            {
             const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE |
                       CLK_ADDRESS_CLAMP_TO_EDGE |
                       CLK_FILTER_NEAREST;
             int2 coords = (int2)(get_global_id(0), get_global_id(1));
             uint4 sumall = (uint4)(0,0,0,0);
             int sum = 0;
             for(int i=-10;i<11;i++)
              for(int j=-10;j<11;j++)
              {
                int2 newpol = (int2)(i,j)+coords;
                sumall+= read_imageui(bmp1, smp, newpol);
                sum++;
              }
             sumall = sumall/sum;
             write_imageui(bmpOut, coords, sumall);
            }

* This source code was highlighted with Source Code Highlighter.


Полезности


Ну, думаю, с кратким описанием я справился. Теперь несколько ссылок для более подробного изучения, ежели кому понадобиться.
Официальный сайт с документацией.
Сайт с примерами и понятными описаниями.
Неплохой пдфничек, там структура OpenCl устройств неплохо прорисованна.
Так же есть 2 презентации про OpenCL на русском. В них довольно мало инфы и нет связующего текста. Правда есть примеры неплохие. Первая. Вторая.
Мальцев Антон @ZlodeiBaal
карма
293,7
рейтинг 0,0
Computer Vision, Machine Learning
Самое читаемое Разработка

Комментарии (26)

  • +3
    Давно хотел такую статью…
    • 0
      поддерживаю, вменяемого текста, который может дать отправную точку изучения технологии, очень мало. а собирать общую картину из разрозненного материала тяжело, я пока не осилил.
      за статью спасибо, она значительно больше похожа на желаемую отправную точку) чем всё остальное что я читал.
  • 0
    Вот бы по OpenCL сделать такую же документацию, как делает noonv по OpenCV.
  • +1
    This source code was highlighted with Source Code Highlighter.

    Хабр уже достаточно давно поддерживает тег source
    • +1
      О, спасибо. Я просто посмотрел на 2-3 статьи и сделал по аналогии с ними. В будующем исправлюсь.
  • –2
    «Так же быстрой памятью является „__privatE “. Это память к которой имеет доступ только отдельный поток (work-item). Всего на поток выделяется 32 регистра этой памяти.»…
    Одкуда вы взяли 32 регистра? Приватная и локальная память очень совместимы, часто локальная это сумма всех приватных. Причём и это тоже не всегда правда. На каждом HW поразному…

    «Вторая процедура (использующая оптимизацию) выполняется в 35 раз быстрее.»
    На каком HW? ctrl+c & ctrl+v?

    А где синхронизация с хоста?

    «Сэмплер — это штука, которая будет оптимизировать вашу работу с изображением. У него есть три параметра:,, .»… запятая запятая и точка?

    Это не введение, это что-то не продуманное с кучей скопированных примеров с сайта с которого вы учились :-).

    • +2
      " На каждом HW поразному…"
      Ну, про тот же local везде пишут что он разный. Про private я в 2х местах натыкался про 32 регистра. Хотя на официальном сайте этого пожалуй нет. Что он любой может быть — спорить не буду, но что 32 — это очень часто на сегодняшний день встречающаяся конфигурация, всё-таки возьмусь утверждать.

      «На каком HW? ctrl+c & ctrl+v?»
      Откуда взят пример — я писал. Там у автора nVidia. У меня две разных карточки ADM 5700 серии. Примерно один и тот же результат на всём. Естественное, говорить строго «в 35 раз» может не правильно, но на мой взгляд было нужно показать порядок ускорения при использовании native кода.

      «А где синхронизация с хоста?»
      Показать ещё 5 вещей которые я тут не включил в описание? :-) Мне показалось что это менее нужная вещь чем те что я включил.

      «запятая запятая и точка?»
      Спасибо, сейчас поправлю. Какой-то глюк при копировании из ворда…

      «Это не введение, это что-то не продуманное с кучей скопированных примеров с сайта с которого вы учились»
      Во-первых, не с одного сайта а с 5-6 разных, разбавленных авторским текстом. Во-вторых, у каждого свой взгляд на мир, мне кажется что тут более-менее продуманно и структурировано для любого начинающего программиста. В принципе во избежание таких дискуссий я в первым абзаце написал что собой представляет эта статья.
      • 0
        Спасибо за ответ.

        Если бы вы рассказали как кернелы спускаются, тогда бы наверное упомянули и об синхронизации на стороне хоста. Потому, что часто кернелы надо пускать несколько (тысяч) раз подряд с разными размерами локальной / глобальной группы.

        Меня все-равно радует, что хоть кто-то об и на OpenCL пишет! Я бы тоже писал, да меня с моей критикой слили в унитаз.
        Хотелось бы следующую статью и потом следующую!
        Я может и немного строго, но всетаки душа хочет качества, когда разговор идёт об OpenCL.
        • 0
          Я сам на OpenCL только с весны начал писать. Да и по работе лишь к одному проекту и то по мелочи его подключал. В основном для себя баловался. Так что что-то сильно глубокое по OpenCl выдать не смогу…
          Интересно, а вы его где-нибудь использовали? Я до сих пор только несколько проектов серьёзных на CUDA видел и ни одного на CL.
          • 0
            Я в университете писал на OpenCL. А так в практике мне не разрешают такие «бешенства» использовать, потому, что если я уйду, никто не сможет удерживать продукт. Планирую себе найти работу в будующем с OpenCL.
          • 0
            Кстати, предлагаю следующую статью сделать по мотивам www.cmsoft.com.br/index.php?option=com_content&view=category&layout=blog&id=99&Itemid=150.

            ПС: изнасиловать в большое анальное отверстие карму! 2 комента посылаю 10 минут… ппц!
            • 0
              Ага, я читал, оно забавно. Но сам пока не сподобился написать код и помучать его. Как будет время — поиграюсь напишу что-нибудь.

              ЗЮ, Я плюсанул:)
              • 0
                Я когда-то мучал мозг этой статьёй. Очень приятное занятие! :-)
  • НЛО прилетело и опубликовало эту надпись здесь
    • 0
      Хотите смотреть на граф с разницами в один или два процента? Google.
      OpenMP будет немного быстрее, у него немножко другая архитектура. Но не на десятки процентов.
      • НЛО прилетело и опубликовало эту надпись здесь
        • 0
          ИМХО, написание прогаммы в ОпенЦЛ займёт примерно в 2-3 раза дольше времени…
          • 0
            Да, это точно. На OpenMP всё уже оптимизировано к тому же. Не факт, что удастся получить код того же качества для сложного проекта.
    • 0
      На хабре было уже: habrahabr.ru/blogs/hi/96122/
      Правда там использовался Core 2 Duo
      • 0
        Это смешные тесты! Потому, что OpenCL вы должны >>сами
        • 0
          парсер лох…
          Вы должны сами оптимализировать код (векторы, функции, итд..), у OpenMP за вас это делает компилятор. Когда кернел подргужается, компилируется или подгружается сразу скомпилированным, тогда это тоже время жрёт.
          У ОпенМП всё уже в бинарном виде.
      • НЛО прилетело и опубликовало эту надпись здесь
  • 0
    Загрузка изображений возможна в типы image2d_t и image3d_t. Первые — это обычные изображения, вторые — двухмерные.

    Может трехмерные? :-)
    • 0
      Да))) Спасибо, поправлю))
  • 0
    Блин… Поймал себя на мысли что: читаю — OpenCL, думаю — биткоины…
    А за статью спасибо )
  • 0
    >> Так как в основе языка лежит практически стандартный с++…

    Вы уверены?

Только полноправные пользователи могут оставлять комментарии. Войдите, пожалуйста.