OpenCL под C# это просто
Хотя технология OpenCL появилась ещё в 2008 году, большого распространения она не получила до сих пор. Плюсы технологии несомненны: ускорение вычислений, кроссплатформенность, способность исполнять код как под GPU, так и под CPU, поддержка стандарта целым рядом компаний: Apple, AMD, Intel, nVidia и некоторыми другими. Минусов не так много, но и они есть: более медленная работа на nVidia, чем через CUDA, сложность использования. Первый из минусов влияет только при серьёзной разработке, где скорость программы важнее кроссплатформенности. Второй и является основным препятствием на пути разработчиков, делающих выбор в пользу того или иного метода разработки. Чтобы разобраться в куче хэдэров, драйверов и стандартов требуется куча времени. Но не всё так плохо. Темой этой статьи будет короткий guide по тому, как наиболее простым способом можно запустить OpenCL под C# и получить удовольствие от параллельного программирования.
Настройка драйверов
nVidia.
Intel.
В отличие от AMD и nVidia проблем с драйверами от Intel у меня не возникло ни разу. Взять их можно тут.
Установка драйверов на систему ещё не гарантирует что они будут у вас работать. Количество багов у AMD зашкаливает (проблемы возникали при установке на 2 из 3х компьютеров), у nVidia оно велико (1 из 3х). Поэтому, после установки рекомендую сначала проверить, подключилось ли OpenCL. Наиболее просто это можно сделать через программы показывающие параметры видеокарт. Я пользуюсь GpuCapsViewer, так же работают opencl-z и GPU-Z.
Если проблемы возникли… Удалите все старые драйвера, переставьте. Для AMD — убедитесь что ставите правильную версию драйверов. Драйвера для ноутбуков у них часто глючные. Если проблемы не исчезли — переустановка винды вас спасёт.
Врапперы
Так как нашей целью является максимально простое программирование на OpenCL под C#, мы не будем заниматься извращениями и подключать OpenCL хэдэры, а воспользуемся уже готовыми библиотеками, упрощающими разработку. Самой полной и безглючной версией, как мне кажется, на сегодняшний день является cloo.dll, входящая в OpenTK. Самой простой в применении, автоматизирующей множество операций является OpenCLTemplate, являющаяся надстройкой над cloo. Из минусов последней — некоторые глюки при работе с AMD, например с последней версией драйверов (11.6) могут отказаться инициализироваться устройства. Так как проект OpenSource, глюки которые у меня были я нашёл и поправил, но когда зарелизят новую версию библиотеки — не знаю. Так же есть несколько менее известных врапперов, которые можно найти на просторах интернета.
Первая программа
В качестве первой программы посчитаем через OpenCL сумму двух векторов, v1 и v2. Пример программы написанной с использованием cloo.dll:
Программа на Cloo
private void button4_Click( object sender, EventArgs e)
<
//Установка параметров, инициализирующих видеокарты при работе. В Platforms[1] должен стоять индекс
//указывающий на используемую платформу
ComputeContextPropertyList Properties = new ComputeContextPropertyList(ComputePlatform.Platforms[1]);
ComputeContext Context = new ComputeContext(ComputeDeviceTypes.All, Properties, null , IntPtr .Zero);
//Текст програмы, исполняющейся на устройстве (GPU или CPU). Именно эта программа будет выполнять паралельные
//вычисления и будет складывать вектора. Программа написанна на языке, основанном на C99 специально под OpenCL.
string vecSum = @»
__kernel void
floatVectorSum(__global float * v1,
__global float * v2)
<
int i = get_global_id(0);
v1[i] = v1[i] + v2[i];
>
» ;
//Список устройств, для которых мы будем компилировать написанную в vecSum программу
List Devs = new List ();
Devs.Add(ComputePlatform.Platforms[1].Devices[0]);
Devs.Add(ComputePlatform.Platforms[1].Devices[1]);
Devs.Add(ComputePlatform.Platforms[1].Devices[2]);
//Компиляция программы из vecSum
ComputeProgram prog = null ;
try
<
prog = new ComputeProgram(Context, vecSum); prog.Build(Devs, «» , null , IntPtr .Zero);
>
//Инициализация новой программы
ComputeKernel kernelVecSum = prog.CreateKernel( «floatVectorSum» );
//Инициализация и присвоение векторов, которые мы будем складывать.
float [] v1 = new float [100], v2 = new float [100];
for ( int i = 0; i //Загрузка данных в указатели для дальнейшего использования.
ComputeBuffer float > bufV1 = new ComputeBuffer float >(Context, ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.UseHostPointer, v1);
ComputeBuffer float > bufV2 = new ComputeBuffer float >(Context, ComputeMemoryFlags.ReadWrite | ComputeMemoryFlags.UseHostPointer, v2);
//Объявляем какие данные будут использоваться в программе vecSum
kernelVecSum.SetMemoryArgument(0, bufV1);
kernelVecSum.SetMemoryArgument(1, bufV2);
//Создание програмной очереди. Не забудте указать устройство, на котором будет исполняться программа!
ComputeCommandQueue Queue = new ComputeCommandQueue(Context, Cloo.ComputePlatform.Platforms[1].Devices[0], Cloo.ComputeCommandQueueFlags.None);
//Старт. Execute запускает программу-ядро vecSum указанное колличество раз (v1.Length)
Queue.Execute(kernelVecSum, null , new long [] < v1.Length >, null , null );
//Считывание данных из памяти устройства.
float [] arrC = new float [100];
GCHandle arrCHandle = GCHandle.Alloc(arrC, GCHandleType.Pinned);
Queue.Read float >(bufV1, true , 0, 100, arrCHandle.AddrOfPinnedObject(), null );
>
* This source code was highlighted with Source Code Highlighter .
А ниже та же программа, реализованная через OpenCLTemplate.DLL
private void btnOpenCL_Click( object sender, EventArgs e)
<
//Текст програмы, исполняющейся на устройстве (GPU или CPU). Именно эта программа будет выполнять паралельные
//вычисления и будет складывать вектора. Программа написанна на языке, основанном на C99 специально под OpenCL.
string vecSum = @»
__kernel void
floatVectorSum(__global float * v1,
__global float * v2)
<
int i = get_global_id(0);
v1[i] = v1[i] * v2[i];
>» ;
//Инициализация платформы. В скобках можно задавать параметры. По умолчанию инициализируются только GPU.
//OpenCLTemplate.CLCalc.InitCL(Cloo.ComputeDeviceTypes.All) позволяет инициализировать не только
//GPU но и CPU.
OpenCLTemplate.CLCalc.InitCL();
//Команда выдаёт список проинициализированных устройств.
List L = OpenCLTemplate.CLCalc.CLDevices;
//Команда устанавливает устройство с которым будет вестись работа
OpenCLTemplate.CLCalc.Program.DefaultCQ = 0;
//Компиляция программы vecSum
OpenCLTemplate.CLCalc.Program.Compile( new string [] < vecSum >);
//Присовоение названия скомпилированной программе, её загрузка.
OpenCLTemplate.CLCalc.Program.Kernel VectorSum = new OpenCLTemplate.CLCalc.Program.Kernel( «floatVectorSum» );
int n = 100;
float [] v1 = new float [n], v2 = new float [n], v3 = new float [n];
//Инициализация и присвоение векторов, которые мы будем складывать.
for ( int i = 0; i //Загружаем вектора в память устройства
OpenCLTemplate.CLCalc.Program.Variable varV1 = new OpenCLTemplate.CLCalc.Program.Variable(v1);
OpenCLTemplate.CLCalc.Program.Variable varV2 = new OpenCLTemplate.CLCalc.Program.Variable(v2);
//Объявление того, кто из векторов кем является
OpenCLTemplate.CLCalc.Program.Variable[] args = new OpenCLTemplate.CLCalc.Program.Variable[] < varV1, varV2 >;
//Сколько потоков будет запущенно
int [] workers = new int [1] < n >;
//Исполняем ядро VectorSum с аргументами args и колличеством потоков workers
VectorSum.Execute(args, workers);
//выгружаем из памяти
varV1.ReadFromDeviceTo(v3);
* This source code was highlighted with Source Code Highlighter .
Как видно, второй вариант куда проще и интуитивнее.
Ссылки напоследок
В первую очередь программированию на OpenCL через С# посвящён этот сайтик — www.cmsoft.com.br. К сожалению, его ведёт мало народу, поэтому примеры часто неадекватные, а OpenCLTemplate, созданный авторами сайта весьма глючный.
Полезным местом является сайт www.opentk.com где весьма оперативно отвечают на вопросы о программирование через cloo.dll
Примерно те же ответы можно получить и на сайте sourceforge.net/projects/cloo
Стандарт программирования OpenCl, основанный на C99 описан здесь — www.khronos.org/opencl
Продолжение статьи «Введение в OpenCl» рассказывает об особенностях языка программирования, которым мы программируем видеокарту.
Введение в 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, то драйвера будут автоматически запускать каждый процесс со следующей задачей после его завершения. Понимание физического уровня требуется только для того, чтобы иметь представление о возможностях взаимодействия между процессами и доступа процессов в память.
Базовые особенности
__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 и видеокарт понимали, что основная цель их детища — ускорить сложные расчёты. Для этого в язык был добавлен ряд специализированных особенностей, позволяющих при их использовании получить прирост в скорости на математических задачах.
Встроенные вектора
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. Эти функции выполняются быстрее, чем соответствующие им математические действия.
Пример
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 раз быстрее.
Разрешения
#pragma OPENCL EXTENSION extension name : behavior
* This source code was highlighted with Source Code Highlighter .
#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];». Они позволяют выделить глобальную переменную в группе процессов и во всех процессах уже во время их исполнения.
Единичные операции
#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 .
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 for ( int j=-10;j
* This source code was highlighted with Source Code Highlighter .
