В последнее время параллельные вычисления прочно входят в жизнь, в частности, с использованием GPU.
Здесь было много статей на эту тему, поэтому ограничусь лишь поверхностным описанием технологии. GPGPU — использование графических процессоров для задач общего назначения, т.е. не связанных напрямую с рендерингом. Пример — библиотека Nvidia PhysX для расчёта физики в некоторых современных играх. Эта технология выгодна тем, что GPU хороши на параллельном выполнении с множеством потоков. Правда, потоков должно быть много, иначе производительность сильно упадет. Ещё из-за особенностей работы с памятью приходится несколько хитрить с передачей данных из оперативной памяти в видеопамять. Известные реализации: CUDA (Nvidia, только для видеокарт), OpenCL (Khronos Group, для гетерогенных систем) и AMD FireStream. Здесь будет обсуждаться только OpenCL.
Итак, приступим к практике. В качестве языка основной программы выберем Python. Он, конечно, не очень быстр сам по себе, зато отлично работает как «клей» — во многих применениях основной расчёт идёт в OpenCL, а код на Python только «подносит патроны». Существует отличная библиотека PyOpenCL, которой и будем пользоваться.
Инсталляция
Но, прежде всего, надо установить всё необходимое для работы с OpenCL. Для карт Nvidia поддержка OpenCL идёт вместе с CUDA SDK, а также в официальном драйвере. В nouveau, насколько знаю, поддержки пока нет. Для Windows есть инсталлятор на официальном сайте, для GNU/Linux следует поставить нужный софт из репозиториев. Для ArchLinux это пакеты cuda-toolkit, cuda-sdk, nvidia, nvidia-utils, nvidia-opencl. Для процессоров Intel есть весьма неплохой Intel OpenCL SDK, для AMD также есть родной SDK.
В CUDA SDK есть отличный пример oclDeviceQuery, который показывает море информации по устройству OpenCL по умолчанию. Весьма удобная вещь для проверки общей работоспособности.
Финальный аккорд — установка библиотек для Python. Потребуются, во-первых, NumPy и, во-вторых, сам PyOpenCL.
Некоторые замечания по практике
Собственно, многие важные для практики вещи были рассказаны в недавней статье, поэтому очень советую с ней тоже ознакомиться. Здесь я остановлюсь на особенностях использования PyOpenCL.
В целом стоит упомянуть об основных особенностях GPU. Во-первых, они требовательны к наличию большого числа потоков, иначе производительность будет падать. Т.е., в отличие от CPU, они любят много легких потоков, нежели мало тяжелых. Во-вторых, следует быть очень осторожным при работе с памятью. Контроллер памяти выдаст обещанные Гб/с только при условии того, что за раз передается большой объем данных. Также следует иметь ввиду, что для при переносе алгоритма на GPU нужно как будто собираться в поход — всё нужное брать с собой, ибо с GPU ни достучаться до оперативной памяти или диска, ни выделить большой блок видеопамяти не получится. И в-третьих, условные переходы сейчас обрабатываются очень плохо, поэтому их следует использовать по минимуму.
Исходя из этого, следует помнить, что далеко не все алгоритмы работают хорошо на GPU. Видимо, обработка изображений и математическое моделирование наиболее подходят для видеокарт.
Инициализация
Работа с OpenCL начинается с инициализации устройства. На странице официальной документации есть пример, как можно использовать create_some_context() для выбора устройства по умолчанию, написав минимум кода. Копипастить его не буду, но весьма полезно ознакомиться сначала с ним.
Здесь я рассмотрю несколько более сложный случай. Я буду выбирать конкретное устройство, а также передавать разные типы параметров. Следует отметить, что массивы для передачи на устройство (здесь и далее «устройство» — это устройство OpenCL, т.е. то, на котором будет выполняться код на OpenCL) не могут быть встроенными списками или кортежами. Можно использовать только NumPy'евские массивы. Это происходит из-за того, что библиотека ожидает получить массив одного типа.
import pyopencl as cl
import numpy
set_simple_args(model_params)
set_device()
create_buffers(np.array(particles, dtype=np.float64))
program = cl.Program(self.context, open('worker.cl').read()).build()
def set_device():
device = cl.get_platforms()[settings.DEVICE_NUM].get_devices(getattr(cl.device_type, settings.CL_DEVICE))[0]
context = cl.Context(devices=[device], dev_type=None)
cqueue = cl.CommandQueue(self.context)
def set_simple_args(mp):
# полезно передавать константы преобразованными в типы numpy
# это избавит от головной боли с типами, ведь язык OpenCL статически типизирован
central_size = numpy.float64(mp['central size'])
dt = numpy.float64(mp['dt'])
skip = numpy.int32(mp['skip'])
q_e = numpy.float64(mp['e charge'])
q_i = numpy.float64(mp['ion charge'])
ion_m = numpy.float64(mp['ion mass'])
def create_buffers(particles):
# particles - большой массив координат и скоростей
mf = cl.mem_flags
buf_particles_1 = cl.Buffer(self.context,
mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=particles)
buf_particles_2 = cl.Buffer(self.context,
mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=particles)
buf_charge = cl.Buffer(self.context,
mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=np.float64(0.0))
buf_q_e = cl.Buffer(self.context,
mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=np.int32(0.0))
buf_q_i = cl.Buffer(self.context,
mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=np.int32(0.0))
Что же здесь происходит? Сначала выставляются «простые» аргументы — они передаются как константы. Далее конфигурируется желаемое устройство. settings.DEVICE_NUM — номер устройства, settings.CL_DEVICE — тип (GPU, CPU или что-то более экзотичное). Теперь загружаются буферы памяти видеокарты. Этот шаг можно производить только для массивов, а также переменных, значение которых нужно прочесть после выполения ядра. И, наконец, создается объект программы (в worker.cl исходный код ядра).
Следует заметить, что при ошибке компиляции выбрасывается исключение и соотв. сообщение на экран. Можно также включить показ предупреждений и другой чуши от компилятора. Это можно сделать, добавив export PYOPENCL_COMPILER_OUTPUT=1 в .bashrc. Тут можно упомянуть интеловский компилятор, который всегда явно рапортует о том, какой хорошо он всё векторизовал, в отличии от более юниксвейного от Nvidia.
Запуск ядра
Здесь всё относительно просто. Наши аргументы скармливаются как параметры для фунции запуска ядра (через __call__()). В данном коде каждая новая итерация — новый запуск ядра.
def iterate():
program.dust_plasma(cqueue, [global_size], None,
buf_particles_1, buf_particles_2,
buf_charge, buf_q_e, buf_q_i,
q_e, q_i, ion_m,
central_size, outer_boundary,
dt, skip)
output_particles, cur_charge = empty_particles, numpy.empty_like(numpy.float64(0.0))
# копируем в output_particles из буфера buf_particles_2
cl.enqueue_copy(cqueue, output_particles, buf_particles_2)
cl.enqueue_copy(cqueue, cur_charge, buf_charge)
return output_particles, cur_charge
Для удобства понимания, вот кусочек кода ядра OpenCL (как можно видеть, это что-то вроде C):
__kernel void dust_plasma(__global double4* particles_1, // эта штука передавалась как буфер
__global double4* particles_2, // и эта тоже
__global double* dust_charge,
__global int* charge_new_e, __global int* charge_new_i,
__const double charge_e, __const double charge_i, // а тут уже просто
__const double ion_mass, // как обычные
__const double central_size, // переменные, ибо
__const double outer_boundary, // константы
__const double dt, __const unsigned skip)
{
int id = get_global_id(0);
// ...
for(unsigned i = 0; i < skip; i++)
{
//printf("Iteration %in", i);
params.previous = &particles_1; params.next = &particles_2;
one_iteration_rk2(&params);
params.previous = &particles_2; params.next = &particles_1;
one_iteration_rk2(&params);
}
barrier(CLK_GLOBAL_MEM_FENCE);
// ...
}
Заключение
Надо отметить, что документация весьма адекватна и удобна. При возникновении вопросов по библиотеке можно писать в рассылку pyopencl@tiker.net, сообщество там довольно живое и отвечает быстро. Библиотека распространяется под либеральной лицензией MIT, что весьма приятно.
По поводу самого языка OpenCL — я считаю, что для его изучения не требуется особенных мануалов. Просто писать, как на С99, иногда вспоминая про ограничения (если что, компилятор подскажет) и листая спецификацию. Пожалуй, что следует заранее знать, так это встроенные типы данных, вроде float4 (переменная, состоящая из 4-х отдельных значений float, работать с ней быстрее, чем с 4 отдельными float).
Также хочется поблагодарить автора PyOpenCL по имени Andreas Klöckner за отличную библиотеку.
Просьба 1) сообщать об опечатках/ошибках в тексте 2) в личку.
Полезные ссылки по теме
- Википедия
- Спецификация по OpenCL
- Официальная страница PyOpenCL, документация
- Та самая недавняя статья, которую полезно почитать
Текст статьи распространяется под лицензией Creative Commons Attribution-ShareAlike 3.0, фрагменты кода — MIT.
Автор: ZaWarudo
