Как ускорить вычисления и повысить производительность программ с помощью принципов массивного параллелизма и OpenCL

Стандартный подход к написанию программ является линейным – операция b выполняется после завершения операции a. Но что делать в случае если таких операций десятки тысяч, а задача требует быстрого произведения данных операций?

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

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

В этой статье мы познакомимся с массивным параллелизмом и напишем программу для параллельных вычислений, используя фреймворк OpenCL.

OpenCL (Open Computing Language) — фреймворк разработанный Apple в 2008 году и поддерживаемый Khronos Group с 2009 года. Он позволяет создавать программы для параллельного исполнения на различных вычислительных девайсах (CPU и GPU), упакованные в "кернели"-ядра (kernels) — части кода, которые будут отправлены на вычислительный девайс для произведения каких-то операций.

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

Массивный параллелизм — топик для целого цикла лекций по информатике и инженерии, по этому давайте ограничимся небольшим обобщением, чтобы понимать с чем мы имеем дело.

В линейном программировании мы имеем чёткую последовательность действий: a, b, c, d; действие b не будет выполнено до того как завершиться a и c не будет выполнено пока не завершиться b. Но что делать, если нам, например, требуется найти суммы элементов из двух массивов (листов), и в каждом массиве по 100,000 элементов? Последовательное вычисление заняло бы достаточно долгое время, так-как нам пришлось бы совершить минимум 100,000 операций. А что если такая процедура требует многочисленного повторения и результат нужен в реальном времени с минимальной задержкой? Тут нам и приходит на помощь массивный параллелизм!

Допустим, мы хотим вычислить суммы 0 + 3, 1 + 2, 2 + 1, 3 + 0 и записать результаты в массив. В линейном программировании, мы воспользуемся циклом for или while, где операции будут выполняться последовательно, и схема вычислений будет выглядеть примерно так:

Через последовательную итерацию, мы будем выполнять первое действие, записывать результат в массив и переходить к следующему действию. Поскольку для выполнения данных вычислений нам требуется 4 итерации, для наглядности, давайте обозначим что время (t) потраченное на выполнение данных операций равняется 4.

Массивный параллелизм позволяет нам сформировать и отправить данную задачу (0 + 3 и т.д.) для выполнения используя ресурсы, например, видеокарты — она имеет десятки, сотни, тысячи вычислительных единиц (ядер), которые могут производить операции параллельно, независимо друг от друга.

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

Другими словами, если 0 + 3 — вычисление номер один, а 1 + 2 — вычисление номер два, мы можем посчитать 1 + 2 и записать ответ во второе место в массиве (res[1]) не зависимо от того записан ли ответ 0 + 3 в первое место (res[0]). При этом, если мы попытаемся одновременно записать информацию в ячейки res[0] (первое место) и res[1] (второе место), у нас не возникнет ошибки.

Таким образом, с помощью массивного параллелизма мы можем одновременно и независимо друг от друга выполнить все нужные нам операции, и записать все ответы в массив, содержащий результаты, всего за одно действие, тем самым, сократив t до 1. Мы только что сократили временную сложность алгоритма (время работы алгоритма) до константного значения 1 (одно действие).

Перед тем как мы перейдём к практике, важно прояснить специфику работы с OpenCL и разобраться как переписывать линейный код в параллельном формате. Не волнуйтесь — это совсем не сложно!

Кернель (kernel, вектор — функция, отправляемая на вычислительный девайс в контексте работы OpenCL) и хост (host — код, вызывающий OpenCL; ваш код) существуют, по большему счёту, изолированно друг от друга.

Компиляция и запуск кернеля OpenCL происходят внутри вашего кода, во время его исполнения (онлайн, или runtime execution). Так-же, важно понимать что кернель и хост не имеют общего буфера памяти и мы не можем динамически выделять память (с помощью malloc() и подобных) внутри кернеля. Обмен информацией между двумя системами происходит посредством отправления между ними заранее выделенных регионов памяти.

Другими словами, если у меня в хосте есть объект Х, для того чтобы обратится к нему из кернеля OpenCL, мне для начала нужно его туда отправить. Если вы всё ещё не до конца поняли о чём идёт речь — читайте дальше и всё прояснится!

Модель памяти OpenCL выглядит так:

При создании контекста OpenCL (среда в которой существует данный инстанс OpenCL), мы обозначаем нашу задачу как NDRange — общий размер вычислительной сетки — количество вычислений которые будут выполнены (в примере с суммами использованном выше, NDRange = 4). Информацию записанную в глобальную память (global memory) мы можем получить из любого элемента NDRange.

При отправке задачи на девайс OpenCL (например, видеокарту компьютера), наш NDRange разбивается на рабочие группы (work-groups) — локальная память (local memory), которые содержат в себе рабочие единицы (work-items, изолированные инстансы кернеля) — приватная память (private memory).

Мы не можем считать объект из локальной памяти в глобальную. Это значит что значение существующее внутри одной рабочей группы не будет доступно из другой группы.

Локальная память имеет доступ к объектам из глобальной памяти, но не может получить информацию из приватной памяти. Наконец, приватная память может считывать из глобальной и локальной памяти, но всё что существует внутри её, остаётся в ней.

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

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

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

В общих чертах, схема работы с OpenCL выглядит следующим образом:

Начнём с установки и настройки OpenCL на нашем компьютере.

Первым делом, скачайте и обновите драйвера вашей видеокарты. Это очень важно, так-как OpenCL не будет работать если ваши драйвера его не поддерживают.

Если вы пользуетесь Apple MacOS, всё что вам нужно сделать — убедиться что у вас установлена новейшая версия ОС и XCode. OpenCL поставляется с вашей системой.

В противном случае, вам требуется скачать и установить подходящую имплементацию OpenCL с сайта производителя вашей видеокарты.

Для AMD GPU/CPU скачайте AMD APP SDK

Для NVIDIA GPU установите последние драйвера вашей видеокарты, поддержка OpenCL включена в них

Для Intel GPU/CPU скачайте Intel OpenCL SDK

В качестве простого примера, давайте возьмём задачу векторного сложения.

Допустим, у нас есть два массива, A и B, равного размера. Наша цель — найти сумму элементов А и B и записать её в элемент нового массива, С, такого-же размера.

Стандартный метод решения данной задачи подразумевает использование цикла for или while, для последовательной итерации через массив, выполняя операцию:

Алгоритм очень простой, но имеет линейную временную сложность, O(n), где n — размер массива. Так-как каждая итерация этого цикла не зависит от других итераций, мы можем переформулировать наш алгоритм под параллелизм: каждая итерация этого цикла может быть выполнена параллельно и одновременно. Таким образом, если мы имеем n ядер в нашем вычислительном девайсе, временная сложность становится константной O(1). Это как-раз то, о чём мы уже говорили раньше.

Начать написание нашей программы стоит именно с кернеля — функции, которую мы отправим на видеокарту для параллельного вычисления. Давайте рассмотрим как это делается.

Функцию кернеля нужно прописывать в отдельном файле с расширением .cl

Давайте поместим этот файл в корневую директорию нашего проекта и назовём vector_add.cl

Так-как язык OpenCL это слегка модифицированный С, проблем с написанием у нас не возникнет, и данный код выполнит нашу задачу:

• Здесь, всё что мы делаем, это избавляемся от цикла for. Мы представляем что нужную нам операцию требуется выполнить всего один раз. Делается это потому что наш кернель будет отправлен на какое-то количество ядер видеокарты одновременно, и каждое ядро параллельно посчитает один элемент под индексом i из нашего массива. Стоит заметить что OpenCL не гарантирует выполнение операций по порядку — он может посчитать сначала двадцатый элемент, а потом третий.
• Ключевое слово __kernel даст компилятору понять что функция addition — именно функция кернеля. Каждая декларация основной функции кернеля должна с него начинаться. Поскольку кернели ничего не возвращают, мы всегда указываем void как тип return функции (речь идёт об основных функциях, не вспомогательных. С ними вы можете работать так-же как и с обычными функциями С)
• __global перед типом аргумента означает что данная переменная будет обработана в глобальной памяти (для обращения к локальной или приватной памяти, используются __local и __private соответственно)
• Указатели на массивы A и B имеют модификатор constant, так как в этих массивах мы будем передавать информацию в кернель. В массив С мы будем записывать результат, по этому такого модификатора нет.
• Функция get_global_id() позволяет получить уникальный идентификатор рабочей единицы (work-item) для данного измерения. Мы отправляем 0, потому-что у нас одномерный массив, соответственно, мы используем 0 для указания первого (и единственного) измерения. Да, с помощью OpenCL мы можем обрабатывать 2-х, 3-х, ...-мерные объекты.

Программа-хост контролирует запуск и выполнение кернелей на вычислительных девайсах. Хост пишется на языке С, но существуют различные библиотеки-обвязки для работы с другими языками, например С++ или Python.

Всё приведённое ниже находится в теле функции main(). Для наглядности, я не стал это прописывать.

• clGetPlatformIDs() — позволит нам узнать все доступные платформы для работы с OpenCL
• clGetDeviceIDs() — сообщит нам об OpenCL девайсах доступных на платформе, указанной в platform_id. Вторым аргументом мы отправляем CL_DEVICE_TYPE_GPU, так-как хотим использовать видеокарту для вычислений. Мы так-же можем использовать CL_DEVICE_TYPE_CPU для операций на процессоре или CL_DEVICE_TYPE_DEFAULT операций на дефолтном девайсе компьютера

• clCreateContext() — создаст нам контекст, внутри которого будет существовать данный инстанс OpenCL. В одной программе (хосте) может быть любое количество контекстов OpenCL. Не переживайте по поводу кучи, на первый взгляд, непонятных NULL — они сейчас совсем не важны.
• clCreateCommandQueue() — позволит нам создать очередь команд, согласно которой будет работать наша OpenCL программа.
  

• clCreateBuffer() позволит нам создать буфер памяти для работы с OpenCL
• Для объектов А и В, вторым аргументом мы отправляем дефайн (определён в хедере opencl.h) CL_MEM_READ_ONLY, потому-что в них будет только информация для считывания — мы не будем ничего туда записывать
• В объект С мы отправляем CL_MEM_WRITE_ONLY, так-как туда мы будем записывать ответы
• Обратите внимание, что третьим аргументом, где мы указываем размер выделяемой памяти, тип данных я указал как cl_int. Это сделано потому-что OpenCL по своему обрабатывает различные типы данных, и все типы данных используемые внутри кернеля OpenCL имеют префикс cl_. В реальности, это не критично, вы можете не добавлять префикс и ошибку вы не получите, но это хорошая практика.
• Говоря о типах данных, стоит добавить, что не все компьютеры поддерживают тип double для вычислений на OpenCL. Поддержку можно узнать на сайте производителя вашей видеокарты. Если ваша карта не поддерживает повышенную точность double, то перед отправкой в кернель, их стоит конвертировать (через каст) в float

• clEnqueueWriteBuffer() позволяет записывать информацию из памяти хоста в объекты памяти OpenCL. Это важная команда с кучей возможностей, по этому очень советую ознакомиться с её документацией. С её помощью мы можем передавать как индивидуальные данные, массивы, и даже целые структуры в кернель OpenCL. Так-же, она позволяет, вместо создания новой области памяти для OpenCL, копировать или обращаться (что-то вроде указателя) к уже существующей на хосте области памяти внутри кернеля
  
• Вторым аргументом мы передаём объект памяти OpenCL, в который информация будет записана
• Шестым аргументом мы передаём указатель на область памяти, из которой информацию мы будем записывать
• Остальные аргументы не так важны для данной статьи

• clCreateProgramWithSource() создаёт объект программы из исходного кода для данного контекста (первый аргумент) и загружает туда исходный код
• Для одного контекста мы можем создать неограниченное количество программ. Количество программ указывается вторым аргументом этой функции
• clBuildProgram() собирает программу из объекта, созданного на предыдущем шаге
  
• clCreateKernel() создаёт наш кернель. Эта команда найдёт функцию addition (второй аргумент) в нашем файле vector_add.cl и определит её как функцию, которую мы исполним при запуске нашего кернеля
  
• Опять, в нашем кернеле может быть неограниченное число функций, как основных (с ключевым словом __kernel), так и вспомогательных, которые мы можем использовать в контексте OpenCL

• clSetKernelArg() отправляет в наш кернель аргументы. Вторым аргументом этой функции мы указываем номер аргумента в кернеле (начиная считать с нуля)
• Второй аргумент функции — это номер аргумента в кернеле, третий — размер аргумента, а четвёртый — это то что туда нужно отправить
• Таким образом, 0 соответствует __global const int *A
• 1 соответствует __global const int *B
• 2 соответствует __global int *С

• Размер вычислительной сетки — это общее количество элементов которое требуется посчитать. В нашем случае, 1024
• Стоит добавить, что сейчас мы работаем в одном измерении, по этому у нас только 1 NDRange.
• В случае, если у нас несколько измерений, NDRange мы сделали бы массивом, в котором каждый элемент содержал бы размер сетки (элементов) в каждом из измерений (например 10 на 10 для 2-х измерений). Так-же в этом случае мы укажем количество измерений 3-м аргументом (для 2-х измерений, мы отправим 2)
• work_size — это размер нашей рабочей группы (work-group). Размер групп должен быть кратен NDRange'у(так-как мы делим его на рабочие группы). Размер групп подбирается экспериментально, нет универсального правила выбора размера, так-что выбирайте то, что подходит логически. Например, если у вас в сетке 100 элементов, нет смысла задавать размер групп по 50, так-как вы задействуете только 2 ядра вычислительного девайса, которые обработают по 50 элементов. Будет логичнее задать размер, например, 10, чтобы задействовать 10 ядер, каждое из которых выполнит по 10 операций параллельно
• clEnqueueNDRangeKernel() отправляет наш кернель для вычисления на OpenCL девайсе. Эта команда запускает наш код, видеокарта находит наши суммы и записывает результат в массив C (но, пока-что только в объект памяти OpenCL, c_mem_obj)
  

• clEnqueueReadBuffer() позволит нам записать данные из буфера OpenCL (c_mem_obj) в локальный буфер хоста (C)

Компилируется наша программа следующим образом (из командной строки)

Исходный код программы вы можете скачать здесь

Ура!

Только что мы с вами разобрались в достаточно трудной теме и написали первую программу с использованием OpenCL!

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

Когда я сам изучал OpenCL, я написал небольшую библиотеку-обёртку-шаблон для упрощения работы со спецификацией. Вы можете её изучить, и, надеюсь, найдёте её полезной в собственной работе. Найти её вы можете здесь.

Так-же, поскольку OpenCL не предоставляет поддержки комплексных чисел и операций с ними (как например стандартный хедер complex.h), если вам понадобится такой функционал, у меня так-же есть библиотека для OpenCL, которая предоставит вам данный функционал. Она находится тут.

Если вам интересно посмотреть на реальный пример использования OpenCL, могу предложить вам ознакомится с моим проектом, для которого мне и пришлось всё это выучить. Проект представляет собой графическую программу, написанную на языке С, для визуализации фракталов. OpenCL позволил мне совершать deep-zooms (глубокое увеличение во фрактал) без видимых задержек, а так-же рендер поверхности фракталов через эффект карты нормали. Найти его вы можете здесь.

В случае, если вам интересно изучить тему подробнее, предлагаю следующие ресурсы:

• Книга OpenCL — Открытый язык для параллельных программ (рус)

• Вводный экскурс в OpenCL Dr.Dobb'а (англ)
• Слайды лекции по OpenCL от университета Jülich Forschungszentrum (англ)
• Гайд Apple Developer по OpenCL (англ)

Спасибо большое за то что пришли на мой Ted Talk и желаю удачи в разработке! Всем пис.

P.S. Если вы захотите использовать иллюстрации из данной статьи в своих целях, прошу вас делать это с указанием авторства.

#selectel_инструкция