Что такое opencl

Стандарт параллельного программирования OpenCL 3.0: новые функции и поддерживаемые устройства

27 апреля 2020 года на мероприятии International Workshop on OpenCL консорциум Khoronos Grop представил предварительные спецификации стандарта OpenCL™ 3.0.

OpenCL (Open Computing Language) — это открытый, не требующий лицензионных отчислений стандарт для кроссплатформенного параллельного программирования различных ускорителей, которые можно найти в суперкомпьютерах, облачных серверах, персональных компьютерах, мобильных устройствах и встроенных платформах. Предыдущая версия стандарта — OpenCL 2.1 — была представлена в марте 2015 года.

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

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

Разработкой стандарта занимается Khronos Group — промышленный консорциум, целью которого является выработка открытых стандартов интерфейсов программирования (API) в области создания и воспроизведения динамической графики и звука на широком спектре платформ и устройств с поддержкой аппаратного ускорения. В него входят более 100 компаний, в числе которых AMD, Apple, ARM, Google, Intel, Nvidia и другие.

OpenCL 3.0 создавался таким образом, чтобы производители «железа» могли обеспечить для разработчиков широкие функциональные возможности, гибкость развертывания, предоставляя OpenCL-совместимым решениям возможность быть ориентированными на функциональность, соответствующую их целевым рынкам.

пишет в Twitter Баладжи Калидас (Balaji Calidas), директор по разработкам в Qualcomm.

Изменения, которые включены в новую версию стандарта, главным образом служат тому, чтобы API мог использоваться для работы с новыми типа «железа», такими как FPGA, специализированные процессоры для встраиваемых систем, процессоры мобильных устройств, системы граничных вычислений (когда обработка данных происходит непосредственно на IoT устройствах), что может быть полезным, к примеру в робототехнике, автоматизированных производственных линиях, при массивно-распределённой аналитике. Более ранние версии стандарта в основном ориентировались на центральные процессоры и дискретные видеокарты.

Вычисления OpenCL могут выполняться на CPU, графических процессорах, большом спектре специализированных микросхем

Новыми решениями, включенными в OpenCL 3.0, стали поддержка асинхронных расширений DMA (Direct Memory Access), для передачи сложных структур памяти, например изображений и 2D/3D структур. Дополнительно введена поддержка языка SPIR-V 1.3 (промежуточный язык для параллельных вычислений и графики, разработанный Khronos Group), что позволит повысить взаимодействия между OpenCL и API Vulkan (интерфейс прикладного программирования для 2D, 3D графики от Khronos Group).

OpenCL 3.0 разработан таким образом, чтобы полностью совместимым только со спецификациями OpenCL 1.2, все остальные функции становятся необязательными. Все функции OpenCL 2.X согласованно определены в новой унифицированной спецификации, и текущие реализации OpenCL 2.X, которые обновляются до OpenCL 3.0, могут продолжать поставлять свои существующие функциональные возможности с полной обратной совместимостью.

сказал Винсент Хиндриксен (Vincent Hindriksen), основатель и генеральный директор Stream HPC.

Новое поколение стандарта имеет совместимость как с предыдущими поколениями, так и с другими стандартами и языками

сказал Нил Треветт, вице-президент NVIDIA, президент Khronos Group и председатель рабочей группы OpenCL в пресс-релизе.

OpenCL может применяться вместе с высокоуровневыми API, что может упростить разработку

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

сказал Джефф МакВей (Jeff McVeigh), вице-президент Intel по архитектуре, графике и программному обеспечению.

Введение в OpenCL

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

Компилятор

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

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

Прежде чем описывать сам язык я дам краткое описание физической модели устройства с которой он взаимодействует. Исполнение команд языка идёт на объектах, называемых «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 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 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 на русском. В них довольно мало инфы и нет связующего текста. Правда есть примеры неплохие. Первая. Вторая.

Технология параллельного программирования OpenCL

Эту статью можно считать продолжением статьи [1] о технологии CUDA, здесь мы поговорим о технологии параллельного программирования OpenCL.

1. Введение

OpenCL (Open Computing Language ) это спецификация, описывающая технологию параллельного программирования, которая в первую очередь ориентирована на GPGPU. Изначально она была разработана компанией Apple, в последствии для развития спецификаций OpenCL был образована группа разработчиков Khronos Compute [2], в неё вошли Apple, nVidia, AMD, IBM, Intel, ARM, Motorola и др. Первая версия стандарта была опубликована в конце 2008 года.

В отличии от nVidia CUDA, AMD Stream и т.п., в OpenCL изначально закладывалась мультиплатформенность, т.е. OpenCL программа должна без изменений в коде работать на GPU разных типов (разных производителей). Такая программа без изменений должна работать даже на CPU без GPU, хотя в этом случае она может выполняться существенно медленнее чем на GPU.

2. Схема работы с аппаратурой

Итак — хотим мультиплатформенность и желательно без существенных потерь в производительности. Достигается этот результат следующим образом [3,4].

OpenCL-программа работает с т.н. платформами (platform) . Платформа это программный пакет, который поставляется соответствующим разработчиком аппаратных средств. Например «AMD Accelerated Parallel Processing» или «Intel OpenCL». При этом несколько платформ могут работать одновременно на одной машине.

Каждая платформа включает в себя ICD (Installable Client Driver) — программный интерфейс OpenCL для работы с устройствами, которые эта платформа поддерживает.

В среде Linux список ссылок на ICD, присутствующих в системе, обычно хранится в каталоге /etc/OpenCL/vendors/ , а библиотека libOpenCL.so выполняет роль диспетчера (ICD loader) , т.е. она направляет вызовы OpenCL функций на устройства, через соответствующие ICD.

Рис.1: схема работы OpenCL программы с аппаратурой

[ Здесь ] можно увидеть пример конфигурации оборудования. На машине с процессором Intel Core2 CPU 6300 и графическим ускорителем nVidia Quadro FX1700, развернуты три платформы OpenCL: nVidia, AMD, Intel. При этом платформа nVidia поддерживает только GPU Quadro FX1700, а платформы AMD и Intel — только Intel Core2 CPU 6300 в качестве вычислительного устройства. Таким образом каждая платформа содержит одно вычислительное устройство.

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

Тут надо ещё сказать о версиях OpenCL SDK. При развёртывании платформы на компьютере необходимо уточнить версию пакета. Например, в процессе проведения экспериментов выяснилось, что Intel OpenCL SDK 2014 не поддерживает работу с процессором Intel Core2 CPU 6300, который был установлен на машине. Проблема решилась установкой более старого пакета Intel OpenCL SDK 2012.

3. Структура OpenCL-программы

В OpenCL (аналогично CUDA), программа разделяется на две части: первая часть — управляющая, вторая — вычислительная. В роли управляющего устройства ( host ) выступает центральный процессор (CPU), вычислительное устройство ( device ) выбираем из списка платформ и их устройств. Обычно используется GPU, но не обязательно, это может быть и тот же CPU.

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

Текст ядра (kernel) , т.е. части программы выполняемой на device, включается в основную часть программы (выполняемой на host) в «чистом» виде т.е. в виде текстовой строки. Этот исходник компилируется средствами OpenCL непосредственно в процессе работы программы (runtime) для выбранного в данный момент вычислительного устройства, это происходит каждый раз при запуске OpenCL-программы.

Так же есть возможность, скомпилированный таким образом, код ядра сохранить и при инициализации вычислительной программы загружать готовый бинарник. Но в этом случае мы теряем универсальность, этот код будет работать только на одном (данном) типе устройств.

1. получить информацию о платформах и устройствах
   clGetPlatformIDs(),clGetPlatformInfo(), clGetDeviceIDs(),clGetDeviceInfo()
2. выбрать устройства и создать для них контекст
   clCreateContext()
3. создать ядро из текста программы
   clCreateProgramWithSource(), clBuildProgram(), clCreateKernel()
4. выделить память для данных на устройствах
   clCreateBuffer()
5. создать очередь комманд для устройтва
   clCreateCommandQueue(), clCreateCommandQueueWithProperties( )
6. скопировать данные с host на device
   clEnqueueWriteBuffer()
7. назначить параметры выполнения ядра
   clSetKernelArg()
8. запуск ядра
   clEnqueueNDRangeKernel()
9. скопировать результат с device на host
   clEnqueueReadBuffer()
10. обработка результата
11. завершение работы, освобождение ресурсов
    clReleaseMemObject(), clReleaseKernel(), clReleaseProgram(), clReleaseCommandQueue(), clReleaseContext()

4. Пример OpenCL-программы

В качестве первого примера рассмотрим простую вычислительную задачу: С := d * A + B, где d — константа, А,В и С векторы заданного размера.

__kernel void kernel1(const float alpha, __global float *A, __global float *B, __global float *C) < int idx = get_global_id(0); C[idx] = alpha* A[idx] + B[idx]; >
Листинг 1: код ядра для задачи сложения векторов

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

Код программы можно скачать [ здесь ]. Для сравнения напишем ещё простую (последовательную) программу и сравним время затраченное на вычисления. Ниже в таблице представлены результаты скорости выполнения программы на разных платформах.

платформа	устройство	время(ms)
NVIDIA	Quadro FX 1700	3.2
AMD	Intel Core2 CPU 6300	19.1
Intel	Intel Core2 CPU 6300	11.0
—	Intel Core2 CPU 6300	9.3

Таблица 1: результаты работы программы saxpy для разных платформ

По результатам, представленным в таб.1, видно, что выполнение примера на GPU почти в три раза быстрее чем на CPU. Для CPU платформа Intel показала лучший результат чем AMD, что очевидно для Intel Core2 CPU 6300. Наилучший результат для CPU показала простая программа, что можно объяснить эффективной работой компилятора и отсутствием необходимости выполнять дополнительные процедуры OpenCL.

5. Группы процессов и их конфигурация

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

Иногда треды удобно формировать в виде решетки. Рассмотрим пример задачи умножения матриц.

Листинг 2: код ядра для задачи умножения матриц

Имеем на входе матрицы A[M*K], B[K*N] и соответственно буфер для результата C[M*N]. Программа создаёт M*N тредов в виде решетки MxN, каждый тред [r,c] отрабатывает одну ячейку в матрице результата.

Код программы можно скачать [ здесь ].

Для сравнения напишем ещё простую (последовательную) программу и сравним время затраченное на вычисления. Ниже в таблице представлены результаты скорости выполнения программы на разных платформах.

платформа	устройство	время(ms)
NVIDIA	Quadro FX 1700	5169.5
AMD	Intel Core2 CPU 6300	2279.0
Intel	Intel Core2 CPU 6300	920.1
—	Intel Core2 CPU 6300	1340.6

Таблица 2: результаты работы программы gemm1 для разных платформ

Результаты несколько расстраивают, поскольку из таб.2 видно, что GPU показала наихудший результат. Проблема в крайне неэффективном использовании памяти GPU. Далее мы исправим это затруднение.

6. Модели памяти и синхронизация процессов

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

• Память global — основная память уcтройства, самая большая по размеру (512MB для FX1700) и самая медленная, она является общей для всех тредов.
• Память local — общая память для одной группы тредов (shared в терминах CUDA), этот тип быстрее global но существенно меньше по размеру (16KB для FX1700)
• Память private — память треда, быстрая но маленькая (8KB для FX1700)

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

Листинг 3: код ядра для задачи умножения матриц с кэшированием

Имеем на входе матрицы A[M*K], B[K*N] и соответственно результат C[M*N]. Программа создаёт M*N тредов в виде решетки MxN, группами размера TSxTS, каждый тред отрабатывает одну ячейку в матрице результата, группа тредов кэширует блоки исходных матриц и далее выполняет операции с этим кэшем.

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

Код программы можно скачать [ здесь ].

Ниже в таблице представлены результаты скорости выполнения программы на разных платформах.

платформа	устройство	время(ms)
NVIDIA	Quadro FX 1700	169.5
AMD	Intel Core2 CPU 6300	2301.1
Intel	Intel Core2 CPU 6300	1402.7
—	Intel Core2 CPU 6300	1340.6

Таблица 3: результаты работы программы gemm2 для разных платформ

Из таблицы 3 видно, что модифицированная программа умножения матриц (gemm2) показывает вполне удовлетворительный результат производительности, работая гораздо быстрее первого варианта (gemm1).

7. Заключение

Хотя программы OpenCL могут выполняться с меньшей скоростью в сравнении с CUDA, но они обладает важным свойством — переносимость, этот стандарт имеет хорошие перспективы развития.

В заключении можно ещё привести ссылку на список библиотек, основанных на OpenCL.
https://www.khronos.org/opencl/resources/opencl-libraries-and-frameworks-with-opencl-acceleration

Список литературы

1. Е.С.Борисов Технология параллельного программирования CUDA — http://mechanoid.kiev.ua/parallel-cuda.html
2. Khronos group. OpenCL. — https://www.khronos.org/opencl/
3. Janusz Kowalik, Tadeusz Puzniakowski Using OpenCL. Programming Maassively Paarallel Commputers. — IOS Press — 2012
4. Benedict R. Gaster, Lee Howes, David R. Kaeli, Perhaad Mistry, Dana Schaa Heterogeneous Computing with OpenCL Revised OpenCL. — Advanced Micro Devices, Inc. Published by Elsevier Inc. — 2013
5. Cedric Nugteren Tutorial: OpenCL SGEMM tuning for Kepler — http://www.cedricnugteren.nl/tutorial.php

OpenCL для чайников

Хочу чисто для интереса посмотреть, что такое OpenCL и как при помощи него программировать. Для начала надо разобраться с необходимыми средствами разработки. Гугл выдаёт не особо много информации (как я понимаю, из-за новизны OpenCL), поэтому задаю вопросы, на которые не нашёл ответа, здесь.

1) OpenCL ведь предназначен для вычислений не только на GPU, но и на CPU. Поэтому интересует вопрос, можно ли сейчас в Linux использовать OpenCL без задействования видеокарты и проприетарных средств разработки от её производителя. Если да, то как?