opencl с чего начать
Введение в OpenCL
Компилятор
Модель памяти устройства
Прежде чем описывать сам язык я дам краткое описание физической модели устройства с которой он взаимодействует. Исполнение команд языка идёт на объектах, называемых «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 «. Этой директивой должна быть помечена любая процедура, которую мы хотим вызвать извне. Если процедура не нужна при работе извне, её можно не отмечать.
Типы памяти
Тип данных «__global » обозначает память, которая выделяется из глобального адресного пространства работающего устройства. Она достаточна медленная, зато вместительная. Для современных видеокарт измеряется гигобайтами. Если вы работаете на процессоре — под global подразумевается оперативная память.
Кроме global есть «__local «. К ней может обращаться только рабочая группа(workgroup). На каждую такую группу выделяется примерно 8 килобайт.
Так же быстрой памятью является «__privat «. Это память к которой имеет доступ только отдельный поток (work-item). Всего на поток выделяется 32 регистра этой памяти.
Остальные типы памяти, которые можно объявлять при создании ядра основаны на типе «__global «. Во-первых, это «__constant «, который может использоваться только для чтения. Во-вторых, это «__read_only», «__write_only» и «__read_write» — структуры, использование которых разрешено только для изображений.
Идентификаторы процессов
Оптимизация расчётов
Разработчики OpenCL и видеокарт понимали, что основная цель их детища — ускорить сложные расчёты. Для этого в язык был добавлен ряд специализированных особенностей, позволяющих при их использовании получить прирост в скорости на математических задачах.
Встроенные вектора
Простые функции
Следующей особенностью 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. Эти функции выполняются быстрее, чем соответствующие им математические действия.
Пример
Вторая процедура (использующая оптимизацию) выполняется в 35 раз быстрее.
Разрешения
#pragma OPENCL EXTENSION extension name : behavior
Синхронизация
Барьеры
В первом примере на команде barrier ожидают все процессы рабочей группы, во втором — все процессы OpenCL устройства.
Стоит отметить особенность этого примера, команды «__local int x[10];» и «__global int x[10];». Они позволяют выделить глобальную переменную в группе процессов и во всех процессах уже во время их исполнения.
Единичные операции
void GetSemaphor(__global int * semaphor) <
int occupied = atom_xchg(semaphor, 1);
while (occupied > 0)
<
occupied = atom_xchg(semaphor, 1);
>
>
Работа с изображениями
Последней вещью, которую я хочу включить в этот 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)
OpenCL. Как начать
Тяжелый старт
Всем привет! Какое-то время назад я начал копать тему с OpenCL под C#. Но наткнулся на трудности, связанные с тем, что не то, что под C#, а вообще по этой теме очень мало материала. Какую-то вводную по OpenCL можно почерпнуть здесь. Так же простой, но работающей старт OpenCL описан вот тут. Ни на йоту не хочу обидеть авторов, но все статьи, что я находил на русском (и на хабре в том числе) страдают одной и той же проблемой — очень мало примеров. Документация есть, её много и как принято для хорошей документации читается сложно. В своей статье (а если всё будет нормально, то и в цикле статей), я постараюсь поподробней описать эту область, с точки зрения человека, который начал её копать с нуля. Думаю такой подход будет полезен тем кто хочет быстро стартовать в высоко производительных вычислениях.
Первоначально я хотел написать статью-минисамоучитель OpenCL, которая содержала в себе информацию, о там что это, как устроено, как писать код и какие-то рекомендации, основанные на моем опыте. Но в процессе понял, что если даже быть кратким, то уткнусь в ограничения объема статьи. Потому что, имхо, статья должна быть такого объема, чтобы усвоить её объем было не сложно. По этому в данной статье (которая станет первой) я планирую описать, то как стартовать в OpenCL, проверить что локально все корректно законфигурино, и написать простейшую программу. Вопросы с устройством памяти, архитектурой и прочим будут описаны в следующих статьях.
Друзья, сразу хотел бы сказать, что мой опыт в OpenCL пока, к сожалению, далек от гуру/йода уровня, но на вопросы постараюсь отвечать изо всех сил. А если что-то не знаю, то буду делиться ресурсами и видением того как это на самом деле должно работать.
Что. Где. Как.
OpenCL это технология связанная с параллельными компьютерными вычислениями на различных типах графических и центральных процессоров. Тема с параллельным вычислениями на GPU совсем недавно широко продвигалась вместе с технологией CUDA. Данное продвижение в основном обеспечивалось усилиями компании Nvidia. Отличия OpenGL и CUDA уже широко обсуждались.
OpenСL позволяет работать как с CPU так и с GPU, но думаю нам более интересно будет сосредоточиться на работе с GPU. Для использования данной технологии понадобиться мало мальски современная видеокарта. Главное это проверить, что устройство функционирует нормально. На всякий случай напоминаю что это можно сделать в диспетчере устройств.
Если в данном окне вы видите какие-то фейлы или ворнинги, то вам прямая дорога на сайт производителя вашей видеокарты. Свежие драйвера должны решить проблему с функционированием железа и как следствие дать доступ к мощностям OpenCL.
Первоначально я планировал использовать OpenCL по C#. Но наткнулся на проблему, что все существующие фреймворки типа Cloo или OpenCLNet являются самописными и Khronos не имеет к ним никакого отношения и следовательно не гарантирует их стабильную работу. Ну и все мы помним главную проблему — очень мало примеров. Исходя из этого вначале я бы хотел представить примеры написанные на C++, а уже потом, получив подтверждение того что OpenCL ведет себя так как мы ожидаем, привинтить прокси в виде C# фреймворка.
Итак, чтобы использовать OpenCL через С++ необходимо найти его API. Для этого открывайте переменные среды, и ищете там переменную со страшным названием, намекающую своим названием на производителя вашей видеокарты. У меня данная переменная называется «AMDAPPSDKROOT». После этого можете посмотреть что лежит по указанному пути. Там ищите папочку include\CL.
Кстати, обычно в папочки include, рядом с папкой CL лежит и папка GL, предоставляющая доступ к знаменитой графической библиотеки.
Инфраструктура
Так как на предыдущем шаге мы предусмотрительно подсоединили папочку include, то теперь вы можем просто добавить ссылку на заголовочный файл cl.h, который даст доступ к API. При добавление cl.h, стоит добавить проверку выбора платформы:
Теперь необходимо выбрать устройство на котором будет отрабатывать наш код и создать контекст в котором будут жить наши переменные. Как это сделать показано ниже:
Kernel
Что делает данный код. Первое — получает глобальный id work-item который сейчас выполняется. Work-item — это то что и выполняет наш kernel. Так как мы имеем дела с параллельными вычислениями, то для каждого work-item создается свой kernel который ничего не знает о других. И никто не может гарантировать в каком порядке все work-item отработают. Но об этом подробней будет в отдельной статье (уже утал это повторять). В нашем примере это по сути индекс элемента в массиве, потому что мы будем каждый элемент массива обрабатывать в отдельном work-item. Думаю вторую строчку строчку в kernel комментировать излишни 🙂
Формируем kernel
Следующий шаг скомпилировать, то что лежит в файле *.cl. Делается это следующим образом:
Типы cl_program и cl_kernel определены в cl.h. Сам сценарий довольно прост — загружаем файл, создаем бинарник (clCreateProgramWithSource) и компилируем. Если переменная ret по прежнему содержит 0, то вы все сделали правильно. И останется только создать сам kernel. Важно, чтобы имя передаваемое в команду clCreateKernel, совпадало с именем kernel в файле cl. В нашем случае это «test».
Параметры
Я уже упоминал, что «общения» kernel с хостом происходит за счет записи/чтения в параметры, которые передаются в kernel. В нашем случае это параметр message. Параметры, которые позволяют вот так общаться хосту с kernel, называются буферами(buffer). Давайте создадим такой буфер на стороне хоста и передадим в kernel через API:
Важно отметить константу CL_MEM_READ_WRITE, она означает, что мы у нас есть права для буфера на чтение и запись, на стороне kernel. Так же могут быть использованы константы типа CL_MEM_WRITE_ONLY, CL_MEM_READ_ONLY и др. Так же в методе clSetKernelArg, важен второй аргумент, он содержит индекс параметра. В данном случае 0, так как параметр message идет первым в сигнатуре kernel. Если бы он шел вторым, то мы бы написали:
clEnqueueWriteBuffer записывает данные из массива mem в буфер memobj.
Ну что в целом все готово. Осталось только выполнить kernel.
Исполняем kernel
Погнали, отправляем код на исполнение:
global_work_size содержит число work-item которые будут созданы. Я уже говорил, что на обработку каждого элемента массива у нас будет свой work-item. Элементов в массиве у нас 10, следовательно work-item содержит 10. clEnqueueNDRangeKernel особых вопросов порождать не должна — просто запускает указанный kernel заданное число раз. clEnqueueReadBuffer считывает данные из буфера с именем memobj и помещает данные в массив mem. Данные в mem и есть наш результат!
Итоги и выводы
Друзья, вот так я представляю старт в OpenCL для новичка. Надеюсь на ваши конструктивные замечания в комментариях, чтобы можно было внести апдейты в будущем. Я пытался быть кратким, но все равно объем вышел не маленький. Так что могу сказать, что материала для 2-3 статей найти еще смогу.
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
//Текст програмы, исполняющейся на устройстве (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
<
//Инициализация новой программы
ComputeKernel kernelVecSum = prog.CreateKernel( «floatVectorSum» );
А ниже та же программа, реализованная через 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);
Как видно, второй вариант куда проще и интуитивнее.
Ссылки напоследок
В первую очередь программированию на OpenCL через С# посвящён этот сайтик — www.cmsoft.com.br. К сожалению, его ведёт мало народу, поэтому примеры часто неадекватные, а OpenCLTemplate, созданный авторами сайта весьма глючный.
Полезным местом является сайт www.opentk.com где весьма оперативно отвечают на вопросы о программирование через cloo.dll
Примерно те же ответы можно получить и на сайте sourceforge.net/projects/cloo
Стандарт программирования OpenCl, основанный на C99 описан здесь — www.khronos.org/opencl
Продолжение статьи «Введение в OpenCl» рассказывает об особенностях языка программирования, которым мы программируем видеокарту.
OpenCL. Практика
Здравствуй, уважаемое хабрасообщество.
В предыдущих статьях мы рассмотрели OpenCL в целом, потом подробно вникли в суть стандарта и разобрали на каких идеях базируется эта технология.
OpenCL. Что это такое и зачем он нужен? (если есть CUDA)
OpenCL. Подробности технологии
Теперь настало время пощупать эту технологию живьем.
Приготовления
Итак, для работы нам понадобится: спецификация стандарта, SDK (AMD или NVidia) и, опционально, литература по OpenCL, например, отсюда.
Если вы устанавливаете Nvidia Computing SDK – вы автоматически получите все нужные документы. Кроме того бонусом Вы получите много интересных примеров программ (30 штук в последнем релизе SDK). Благодаря этим примерам легко научиться правильно использовать OpenCL, использовать несколько OpenCL устройств одновременно, пользоваться связкой OpenCL-OpenGL (это взаимодействие оговорено стандартом!) итд.
Компилятор OpenCL встроен в драйвер, поэтому выбор IDE для разработки никак не ограничен, посему не буду описывать процесс настройки какой-то определенной IDE. Все что Вам надо сделать – это прописать пути до заголовков и библиотек, которые установит SDK.
Поехали.
Напишем простую программу для суммирования двух векторов. Такая программа есть в примерах SDK для CUDA и OpenCL, но наша программа будет немного отличаться (уберем проверки кодов ошибок на каждом шаге, и немного упростим программу, оставив только самую суть).
GPU часть
Начнем с самого интересного, и, пожалуй, самого простого в данном примере – с кода, который будет исполнен на GPU.
Синтаксис OpenCL для написания kernel’ов сам по себе не представляет собой ничего особенного и слабо отличается от синтаксиса той же CUDA – это все старый добрый С с небольшими модификациями.
Создадим файл vectorAdd.cl – здесь будет располагаться наш kernel.
__kernel void VectorAdd(__global const float * a, __global const float * b, __global float * c, int iNumElements)
<
// get index into global data array
int iGID = get_global_id(0);
// bound check (equivalent to the limit on a ‘for’ loop for standard/serial C code
if (iGID >= iNumElements)
<
return ;
>
// add the vector elements
c[iGID] = a[iGID] + b[iGID];
>
Получили свой номер в глобальном пространстве индексов и сложили элементы векторов с соответствующим индексом, а если наш номер больше размера вектора – ничего не делам.
Все выглядит легко и просто: kernel – простая функция, объявление которой предваряется ключевым словом __kernel (два подчеркивания) а дальше все как в С – возвращаемый тип, название функции, параметры (при определении параметров так же необходимо указать модификаторы __global, __local, __private).
kernel пишется на языке С. Существует ряд расширений (кроме синтаксиса) и ограничений. Кратко об ограничениях можно прочитать тут. Более полно в стандарте, так же полезен может быть OpenCL Programming Guide
Расширениями языка являются: тип данных «изображение» 2d и 3d, типы данных вектор интов, флоатов итп. размерности от 2 до 4.
При объявлении переменных надо указывать область памяти, где они должны располагаться: __global, __local, __private. Если область памяти не указана – будет использована private-память.
Если в kernel необходимо использовать другие функции, скрытые для CPU то можно определить их в том же файле, но без указания модификатора __kernel.
Хостовая часть
Простейший kernel мы создали. Теперь давайте разберемся, как запустить этот kernel на видеокарте.
Хостовая часть программы тоже будет проста и ограничится запуском kernel.
Функции для работы с kernel предоставляет OpenCL API. Это С-функции. Тут можно скачать C++-bindings и документацию к ним.
Мануал по всем функциям API находится в том же документе, где описывается стандарт OpenCL.
Для работы любого kernel необходим контекст, в котором он будет исполняться. Создадим объект «контекст».
cl_context cxGPUContext = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, &ciErr1);
первый параметр – список свойств контекста и их значений. NULL означает использование дефолтных implementation-defined свойств.
Далее объясняем системе, что собираемся работать с GPU – это означает, что устройство с которым мы будем работать, может быть так же использовано и для 3d API, например OpenGL. (список возможных значений этого параметра есть так же в спецификации стандарта)
Следующие два параметра нужны для регистрации call-back-функции, которая будет вызвана OpenCL в случае появления ошибок в контексте.
Последний параметр – для возврата кода ошибки. Может быть NULL.
На самом деле это не единственный способ создания контекста. Но статья и не претендует на описание всех функций OpenCL API. Просто такой создания контекста способ нам более удобен.
Далее выберем устройство (у меня в системе оно всего одно, но на будущее пусть в нашей программе используется устройство с максимальным числом FLOPS).
cl_device_id cdDevice = oclGetMaxFlopsDev(cxGPUContext);
Отмечу, что в различных примерах из SDK весь процесс инициализации порой различается, может быть это сделано намеренно, дабы показать, что существует не один способ выполнять данные действия и заставить разработчика покопаться в спецификациях. К примеру, тут мы выбрали устройство с максимальными FLOPS, но мы могли бы воспользоваться функцией clGetContextInfo для получения списка всех устройств, ассоциированных с контекстом (см. оригинальный пример VectorAdd).
Выбрали и инициализировали устройство.
Теперь свяжем с нашим устройством очередь команд.
cl_command_queue cqCommandQue = clCreateCommandQueue(cxGPUContext, cdDevice, 0, &ciErrNum);
Из интересных параметров только 0. На самом деле это список свойств очереди команд: можно ли выполнять команды не последовательно и разрешено ли профилирование команд.
Для работы с устройством все готово, мы можем отправлять команды в очередь для исполнения.
Создадим объекты памяти, через которые будут связаны области памяти на устройстве и хосте.
cl_mem cmDevSrcA = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof (cl_float) * szGlobalWorkSize, NULL, &ciErr1);
cl_mem cmDevSrcB = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof (cl_float) * szGlobalWorkSize, NULL, &ciErr2);
cl_mem cmDevDst = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof (cl_float) * szGlobalWorkSize, NULL, &ciErr2);
указываем тип доступа к памяти объектов (для устройства), размер области памяти и область памяти хоста с которой связан объект памяти (тут NULL).
Инициализировать входные данные можно было бы и сразу, если предпоследним параметром передать указатель на область памяти хоста, которую надо скопировать на устройство.
Но мы сделаем это позже, перед самым запуском kernel, чтобы не занимать место на устройстве раньше времени.
Все подготовительные работы завершены, теперь мы примемся за сам kernel. Как Вы помните, компилятор OpenCL встроен прямо в платформу. По этому причине сборка OpenCL-kernel должна осуществляться во время исполнения (собирать kernel можно как из исходников так и из бинарников).
Приступим.
Получили исходник программы в строке char*. Source_path – полный путь до файла vectorAdd.cl, далее следует «преамбула» — обычно это header или список define’ов. Последний параметр – размер выходной строки.
ВНИМАНИЕ. oclLoadProgSource — не является функцией OpenCL API, а находятся в вспомогательной библиотеке, поставляемой вместе с Nvidia Computing SDK.
Создаем объект программы из полученных исходников, последующие функции – это OpenCL API.
cl_program cpProgram = clCreateProgramWithSource(cxGPUContext, 1,( const char **)&source, &program_length, &ciErrNum);
программа может состоять из нескольких файлов, каждый из которых необходимо загрузить в отдельную строку char*, массив таких строк мы передаем для создания программы. Второй параметр тут означает размер этого массива. В нашем случае – 1.
После массива строк передается массив длин этих строк.
Все остальные параметры не заслуживают внимания.
Слепили программу из кучи файлов, теперь давайте ее соберем (компиляция и линковка)
сlBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);
clSetKernelArg(ckKernel, 0, sizeof (cl_mem), ( void *)&cmDevSrcA);
clSetKernelArg(ckKernel, 1, sizeof (cl_mem), ( void *)&cmDevSrcB);
clSetKernelArg(ckKernel, 2, sizeof (cl_mem), ( void *)&cmDevDst);
clSetKernelArg(ckKernel, 3, sizeof (cl_int), ( void *)&iNumElements);
указываем порядковый номер параметра, размер и объект памяти.
Вот теперь точно все. Начинается работа с очередью:
Скопируем (асинхронно; за это отвечает третий аргумент) данные на устройство.
clEnqueueWriteBuffer(cqCommandQue, cmDevSrcA, CL_FALSE, 0, sizeof (cl_float) * szGlobalWorkSize, srcA, 0, NULL, NULL);
clEnqueueWriteBuffer(cqCommandQue, cmDevSrcB, CL_FALSE, 0, sizeof (cl_float) * szGlobalWorkSize, srcB, 0, NULL, NULL);
С первыми двумя параметрами все понятно.
Третий параметр – размерность пространства индексов. Вектор – одномерный.
За ним следует аргумент, который означает размер сдвига в пространстве индексов и в текущей версии стандарта должен быть всегда NULL.
szGlobalWorkSize указывает размер пространства индексов — это общее количество work-item’ов, которые будут выполняться.
Размер группы оставляем на усмотрение драйверу (NULL).
Следующие два парметра используются для синхронизации при использовании out-of-order исполнения команд. Это список событий, которые должны завершиться перед запуском этой команды (сначала идет размер списка, потом сам список).
Через последний параметр возвращается объект-событие, сигнализирующее о завершении команды.
Осталось только прочитать результат. Сделаем это синхронно:
clEnqueueReadBuffer(cqCommandQue, cmDevDst, CL_TRUE, 0, sizeof (cl_float) * szGlobalWorkSize, dst, 0, NULL, NULL);
Теперь осталось очистить память, удалив созданные объекты памяти и программ. Это не сложно, и легко найти в любом примере из SDK, поэтмоу я не буду приводить здесь этот код.
Заключение
Простейшая программа готова.
Дальше можно изучать спецификации вглубь, изучать особенности конкретных реализаций OpenCL, пробовать различные опции оптимизации в компиляторах, читать документы, выпускаемые поставщиками этих реализаций и, конечно же, писать более сложные программы.
Существует множество нюансов и подводных камней при написании OpenCL программ, так же на данный момент доступен профилировщик OpenCL приложений раработанный Nvidia, предоставляющий ряд интересных возможностей. Но это все уже выходит за рамки данной статьи и, если уважаемые читатели проявят интерес, про различные тонкости и особенности OpenCL приложений и OpenCL для Nvidia GPU можно будет написать отдельную статью.