OpenCL. Как начать04.07.2015 23:48

Тяжелый старт


Всем привет! Какое-то время назад я начал копать тему с OpenCL под C#. Но наткнулся на трудности, связанные с тем, что не то, что под C#, а вообще по этой теме очень мало материала. Какую-то вводную по OpenCL можно почерпнуть здесь. Так же простой, но работающей старт OpenCL описан вот тут. Ни на йоту не хочу обидеть авторов, но все статьи, что я находил на русском (и на хабре в том числе) страдают одной и той же проблемой — очень мало примеров. Документация есть, её много и как принято для хорошей документации читается сложно. В своей статье (а если всё будет нормально, то и в цикле статей), я постараюсь поподробней описать эту область, с точки зрения человека, который начал её копать с нуля. Думаю такой подход будет полезен тем кто хочет быстро стартовать в высоко производительных вычислениях.
Первоначально я хотел написать статью-минисамоучитель OpenCL, которая содержала в себе информацию, о там что это, как устроено, как писать код и какие-то рекомендации, основанные на моем опыте. Но в процессе понял, что если даже быть кратким, то уткнусь в ограничения объема статьи. Потому что, имхо, статья должна быть такого объема, чтобы усвоить её объем было не сложно. По этому в данной статье (которая станет первой) я планирую описать, то как стартовать в OpenCL, проверить что локально все корректно законфигурино, и написать простейшую программу. Вопросы с устройством памяти, архитектурой и прочим будут описаны в следующих статьях.

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

Что. Где. Как.


OpenCL это технология связанная с параллельными компьютерными вычислениями на различных типах графических и центральных процессоров. Тема с параллельным вычислениями на GPU совсем недавно широко продвигалась вместе с технологией CUDA. Данное продвижение в основном обеспечивалось усилиями компании Nvidia. Отличия OpenGL и CUDA уже широко обсуждались.

OpenGL позволяет работать как с CPU так и с GPU, но думаю нам более интересно будет сосредоточиться на работе с GPU. Для использования данной технологии понадобиться мало мальски современная видеокарта. Главное это проверить, что устройство функционирует нормально. На всякий случай напоминаю что это можно сделать в диспетчере устройств.

231b54fab6f04c9583f6ac0f92da8119.png

Если в данном окне вы видите какие-то фейлы или ворнинги, то вам прямая дорога на сайт производителя вашей видеокарты. Свежие драйвера должны решить проблему с функционированием железа и как следствие дать доступ к мощностям OpenCL.

Первоначально я планировал использовать OpenCL по C#. Но наткнулся на проблему, что все существующие фреймворки типа Cloo или OpenCLNet являются самописными и Khronus не имеет к ним никакого отношения и следовательно не гарантирует их стабильную работу. Ну и все мы помним главную проблему — очень мало примеров. Исходя из этого вначале я бы хотел представить примеры написанные на C++, а уже потом, получив подтверждение того что OpenCL ведет себя так как мы ожидаем, привинтить прокси в виде C# фреймворка.
Итак, чтобы использовать OpenCL через С++ необходимо найти его API. Для этого открывайте переменные среды, и ищете там переменную со страшным названием, намекающую своим названием на производителя вашей видеокарты. У меня данная переменная называется «AMDAPPSDKROOT». После этого можете посмотреть что лежит оп указанному пути. Там ищите папочку include\CL.

614ab39e8e79414090aeebd1ab2b34fc.png

Кстати, обычно в папочки include, рядом с папкой CL лежит и папка GL, предоставляющая доступ к знаменитой графической библиотеки.

Теперь создаем проект в Visual Studio, подключаем в свойствах проекта папку include (в моем случае $(AMDAPPSDKROOT)\include\) и в бой!

Инфраструктура


Мы должны помнить, что мы будем работать с OpenCL не через API, а при помощи API. Вроде бы кажется, что эти две фразы практически идентичны, но это не так. К примеру, вспомните OpenGL. Как происходит там работа (урощённый вариант) — вначале мы настраиваем какие-то общие параметры, а потом прямо из кода вызываем методы типа «нарисовать сферу», «изменить параметры источника света» и т.д.
Так вот в OpenCL сценарий другой:

  1. При помощи API получаем доступ к устройствам, которые поддерживают OpenCL. Это часть приложения обычно называется хостом;
  2. Пишем код который будет выполняться на устройстве. Этот код называется kernel. Данный код о хосте вообще ничего не знает. Его может дернуть любой хост
  3. При помощи API прогружаем код kernel и запускаем его выполнение на выбранном устройстве.


Как видите наше приложение будет иметь комплексную инфраструктуру. Давайте займемся ее настройкой!

Так как на предыдущем шаге мы предусмотрительно подсоединили папочку include, то теперь вы можем просто добавить ссылку на заголовочный файл cl.h, который даст доступ к API. При добавление cl.h, стоит добавить проверку выбора платформы: 

#ifdef __APPLE__
#include 
#else
#include 
#endif


Теперь необходимо выбрать устройство на котором будет отрабатывать наш код и создать контекст в котором будут жить наши переменные. Как это сделать показано ниже: 

/* получить доступные платформы */
ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);

/* получить доступные устройства */
ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);

/* создать контекст */
context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);

/* создаем команду */
command_queue = clCreateCommandQueue(context, device_id, 0, &ret);


Обращаю внимание на переменную ret. Это переменная, которая содержит числовое значение которое возвращает та или иная функция. Если ret== 0, то функция выполнилась корректно, если нет, то значит произошла ошибка.
Так же заслуживает внимание константа CL_DEVICE_TYPE_DEFAULT, она запрашивает устройство, которое используется для вычислений на OpenCL по умолчанию. Вместо данной константы могут быть использованы другие. К примеру:

  • CL_DEVICE_TYPE_CPU — запросит существующие CPU.
  • CL_DEVICE_TYPE_GPU — запросит существующие GPU.


Kernel


Отлично. Настроили инфраструктуру. Теперь возьмемся за kernel. Kernel — это просто функция объявление, которой начинается с ключевого слова __kernel. Синтаксис языка программирования OpenCL базируется на стандарте C99, но имеет ряд специфических и очень важных изменений. Об этом будет (я очень надеюсь) отдельная статья. Пока базовая информация:

  1. Код который, будет дергаться с хостовой части, для исполнения, должен начинаться с ключевого слова __kernel;
  2. Функция с ключевым словом __kernel всегда возвращает void;
  3. Существуют квалификаторы типов памяти: __global, __local, __constant, __private, которые будут определять, в какой памяти будут храниться переменные. Если квалификатора перед переменной нет, то она является __private;
  4. «Общение» между хостом и kernel будет через параметры kernel. Чтобы kernel мог что-то передать хосту через параметр, параметр должен быть с квалификатором __global (пока будем использовать только __global);
  5. Код kernel принято хранить в файле с расширением cl. Но по сути подобный код может генерироваться и на лету. Это позволяет обойти некоторые ограничения. Но об этом в другой раз :)


Простейший пример kernel приведен ниже: 

__kernel void test(__global int* message)
{
        // получаем текущий id.
        int gid = get_global_id(0);

        message[gid] += gid;
}

Что делает данный код. Первое — получает глобальный id work-item который сейчас выполняется. Work-item — это то что и выполняет наш kernel. Так как мы имеем дела с параллельными вычислениями, то для каждого work-item создается свой kernel который ничего не знает о других. И никто не может гарантировать в каком порядке все work-item отработают. Но об этом подробней будет в отдельной статье (уже утал это повторять). В нашем примере это по сути индекс элемента в массиве, потому что мы будем каждый элемент массива обрабатывать в отдельном work-item. Думаю вторую строчку строчку в kernel комментировать излишни :)

Формируем kernel


Следующий шаг скомпилировать, то что лежит в файле *.cl. Делается это следующим образом: 

cl_program program = NULL;
cl_kernel kernel = NULL;

FILE *fp;
const char fileName[] = "../forTest.cl";
size_t source_size;
char *source_str;
int i;

try {
        fp = fopen(fileName, "r");
        if (!fp) {
                fprintf(stderr, "Failed to load kernel.\n");
                exit(1);
        }
        source_str = (char *)malloc(MAX_SOURCE_SIZE);
        source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
        fclose(fp);
} 
catch (int a) {
        printf("%f", a);
}

/* создать бинарник из кода программы */
program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);

/* скомпилировать программу */
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

/* создать кернел */
kernel = clCreateKernel(program, "test", &ret);


Типы cl_program и cl_kernel определены в cl.h. Сам сценарий довольно прост — загружаем файл, создаем бинарник (clCreateProgramWithSource) и компилируем. Если переменная ret по прежнему содержит 0, то вы все сделали правильно. И останется только создать сам kernel. Важно, чтобы имя передаваемое в команду clCreateKernel, совпадало с именем kernel в файле cl. В нашем случае это «test».

Параметры


Я уже упоминал, что «общения» kernel с хостом происходит за счет записи/чтения в параметры, которые передаются в kernel. В нашем случае это параметр message. Параметры, которые позволяют вот так общаться хосту с kernel, называются буферами (buffer). Давайте создадим такой буфер на стороне хоста и передадим в kernel через API: 

cl_mem memobj = NULL;
int memLenth = 10;
cl_int* mem = (cl_int *)malloc(sizeof(cl_int) * memLenth);

/* создать буфер */
memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, memLenth * sizeof(cl_int), NULL, &ret);

/* записать данные в буфер */
ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, memLenth * sizeof(cl_int), mem, 0, NULL, NULL);

/* устанавливаем параметр */
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);

Важно отметить константу CL_MEM_READ_WRITE, она означает, что мы у нас есть права для буфера на чтение и запись, на стороне kernel. Так же могут быть использованы константы типа CL_MEM_WRITE_ONLY, CL_MEM_READ_ONLY и др. Так же в методе clSetKernelArg, важен второй аргумент, он содержит индекс параметра. В данном случае 0, так как параметр message идет первым в сигнатуре kernel. Если бы он шел вторым, то мы бы написали: 

/* устанавливаем параметр */
ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobj);

clEnqueueWriteBuffer записывает данные из массива mem в буфер memobj.
Ну что в целом все готово. Осталось только выполнить kernel.

Исполняем kernel


Погнали, отправляем код на исполнение: 

size_t global_work_size[1] = { 10 };

/* выполнить кернел */
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);

/* считать данные из буфера */
ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, memLenth * sizeof(float), mem, 0, NULL, NULL);

global_work_size содержит число work-item которые будут созданы. Я уже говорил, что на обработку каждого элемента массива у нас будет свой work-item. Элементов в массиве у нас 10, следовательно work-item содержит 10. clEnqueueNDRangeKernel особых вопросов порождать не должна — просто запускает указанный kernel заданное число раз. clEnqueueReadBuffer считывает данные из буфера с именем memobj и помещает данные в массив mem. Данные в mem и есть наш результат!

Итоги и выводы


Друзья, вот так я представляю старт в OpenCL для новичка. Надеюсь на ваши конструктивные замечания в комментариях, чтобы можно было внести апдейты в будущем. Я пытался быть кратким, но все равно объем вышел не маленький. Так что могу сказать, что материала для 2–3 статей найти еще смогу.

Спасибо, всем кто дочитал до конца!

© Habrahabr.ru