Здравствуй, уважаемое хабрасообщество.
В предыдущих статьях мы рассмотрели 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];
}
* This source code was highlighted with Source Code Highlighter.Получили свой номер в глобальном пространстве индексов и сложили элементы векторов с соответствующим индексом, а если наш номер больше размера вектора – ничего не делам.
Все выглядит легко и просто: 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);
* This source code was highlighted with Source Code Highlighter.первый параметр – список свойств контекста и их значений. NULL означает использование дефолтных implementation-defined свойств.
Далее объясняем системе, что собираемся работать с GPU – это означает, что устройство с которым мы будем работать, может быть так же использовано и для 3d API, например OpenGL. (список возможных значений этого параметра есть так же в спецификации стандарта)
Следующие два параметра нужны для регистрации call-back-функции, которая будет вызвана OpenCL в случае появления ошибок в контексте.
Последний параметр – для возврата кода ошибки. Может быть NULL.
На самом деле это не единственный способ создания контекста. Но статья и не претендует на описание всех функций OpenCL API. Просто такой создания контекста способ нам более удобен.
Далее выберем устройство (у меня в системе оно всего одно, но на будущее пусть в нашей программе используется устройство с максимальным числом FLOPS).
cl_device_id cdDevice = oclGetMaxFlopsDev(cxGPUContext);
* This source code was highlighted with Source Code Highlighter.Отмечу, что в различных примерах из SDK весь процесс инициализации порой различается, может быть это сделано намеренно, дабы показать, что существует не один способ выполнять данные действия и заставить разработчика покопаться в спецификациях. К примеру, тут мы выбрали устройство с максимальными FLOPS, но мы могли бы воспользоваться функцией clGetContextInfo для получения списка всех устройств, ассоциированных с контекстом (см. оригинальный пример VectorAdd).
Выбрали и инициализировали устройство.
Теперь свяжем с нашим устройством очередь команд.
cl_command_queue cqCommandQue = clCreateCommandQueue(cxGPUContext, cdDevice, 0, &ciErrNum);
* This source code was highlighted with Source Code Highlighter.Из интересных параметров только 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);
* This source code was highlighted with Source Code Highlighter.указываем тип доступа к памяти объектов (для устройства), размер области памяти и область памяти хоста с которой связан объект памяти (тут NULL).
Инициализировать входные данные можно было бы и сразу, если предпоследним параметром передать указатель на область памяти хоста, которую надо скопировать на устройство.
Но мы сделаем это позже, перед самым запуском kernel, чтобы не занимать место на устройстве раньше времени.
Все подготовительные работы завершены, теперь мы примемся за сам kernel. Как Вы помните, компилятор OpenCL встроен прямо в платформу. По этому причине сборка OpenCL-kernel должна осуществляться во время исполнения (собирать kernel можно как из исходников так и из бинарников).
Приступим.
Создадим kernel из того файла, который мы написали ранее.
char *source = oclLoadProgSource(source_path, "", &program_length);
* This source code was highlighted with Source Code Highlighter.Получили исходник программы в строке 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);
* This source code was highlighted with Source Code Highlighter.программа может состоять из нескольких файлов, каждый из которых необходимо загрузить в отдельную строку char*, массив таких строк мы передаем для создания программы. Второй параметр тут означает размер этого массива. В нашем случае – 1.
После массива строк передается массив длин этих строк.
Все остальные параметры не заслуживают внимания.
Слепили программу из кучи файлов, теперь давайте ее соберем (компиляция и линковка)
сlBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);
* This source code was highlighted with Source Code Highlighter.Собирает программу для выбранных устройств из списка устройств связанных с контекстом (контекст не передаем, так как он связан с объектом программы).
Тут можно указать для каких из устройств связанных с контекстом надо собрать программу. Так же можно установить опции компилятора и настроить call-back-функцию для возможности асинхронной компиляции.
У нас пока всего одно устройство, и мы можем позволить себе синхронную компиляцию. И опциями компилятора пользоваться пока не станем.
Все это подробно описано в спецификации стандарта.После сборки с нашим объектом программы уже связан готовый исполняемый файл. Теперь из функции (а это ведь пока просто функция) с идентификатором __kernel надо создать kernel.
cl_kernel ckKernel = clCreateKernel(cpProgram, "VectorAdd", &ciErr1);
* This source code was highlighted with Source Code Highlighter.Все готово, все приготовления завершены и настал момент истины: запуск kernel на исполнение.
НО! Предварительно надо установить входные параметры, с которыми будет вызван наш kernel.
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);
* This source code was highlighted with Source Code Highlighter.указываем порядковый номер параметра, размер и объект памяти.
Вот теперь точно все. Начинается работа с очередью:
Скопируем (асинхронно; за это отвечает третий аргумент) данные на устройство.
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);
* This source code was highlighted with Source Code Highlighter.Самое главное – поставить kernel на исполнение.
clEnqueueNDRangeKernel(cqCommandQue, ckKernel, 1, NULL, &szGlobalWorkSize, NULL, 0, NULL, NULL);
* This source code was highlighted with Source Code Highlighter.С первыми двумя параметрами все понятно.
Третий параметр – размерность пространства индексов. Вектор – одномерный.
За ним следует аргумент, который означает размер сдвига в пространстве индексов и в текущей версии стандарта должен быть всегда NULL.
szGlobalWorkSize указывает размер пространства индексов — это общее количество work-item'ов, которые будут выполняться.
Размер группы оставляем на усмотрение драйверу (NULL).
Следующие два парметра используются для синхронизации при использовании out-of-order исполнения команд. Это список событий, которые должны завершиться перед запуском этой команды (сначала идет размер списка, потом сам список).
Через последний параметр возвращается объект-событие, сигнализирующее о завершении команды.
Осталось только прочитать результат. Сделаем это синхронно:
clEnqueueReadBuffer(cqCommandQue, cmDevDst, CL_TRUE, 0, sizeof(cl_float) * szGlobalWorkSize, dst, 0, NULL, NULL);
* This source code was highlighted with Source Code Highlighter.Теперь осталось очистить память, удалив созданные объекты памяти и программ. Это не сложно, и легко найти в любом примере из SDK, поэтмоу я не буду приводить здесь этот код.
Заключение
Простейшая программа готова.
Дальше можно изучать спецификации вглубь, изучать особенности конкретных реализаций OpenCL, пробовать различные опции оптимизации в компиляторах, читать документы, выпускаемые поставщиками этих реализаций и, конечно же, писать более сложные программы.
Существует множество нюансов и подводных камней при написании OpenCL программ, так же на данный момент доступен профилировщик OpenCL приложений раработанный Nvidia, предоставляющий ряд интересных возможностей. Но это все уже выходит за рамки данной статьи и, если уважаемые читатели проявят интерес, про различные тонкости и особенности OpenCL приложений и OpenCL для Nvidia GPU можно будет написать отдельную статью.
Cписок полезных ссылок
www.nvidia.com/object/cuda_opencl.html — тут можно зарегистрироваться для получения доступа к Nvidia Computing SDK и скачать полезные документы
www.khronos.org/registry/cl — страница на сайде Khronos Group. Спецификации, заголовочные файлы итд.
developer.amd.com/GPU/ATISTREAMSDKBETAPROGRAM/Pages/default.aspx — AMD Stream SDK с поддержкой OpenCL.
Настоятельно рекомендую ознакомиться с двумя документами Nvidia: OpenCL Programming Guide и OpenCL Best Practices Guide.