Используем OpenCL в Python

В последнее время параллельные вычисления прочно входят в жизнь, в частности, с использованием 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 %i\n", 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) в личку.

Полезные ссылки по теме




EDIT. Исправил опечатки.

Текст статьи распространяется под лицензией Creative Commons Attribution-ShareAlike 3.0, фрагменты кода — MIT.

  • +28
  • 14k
  • 6
Поделиться публикацией
AdBlock похитил этот баннер, но баннеры не зубы — отрастут

Подробнее
Реклама
Комментарии 6
  • 0
    Спасибо большое! Код выглядит гораздо более простым и понятным, чем аналогичный на С++. Пожалуй, начну изучение OpenCL в связке именно с Python
    • +2
      Спасибо за статью )Но, если честно, код выглядит для меня вообще не похоже на питон, скорее на Си или Дельфи.
      • +1
        Возможно, это из-за того, что я долгое время писал на С++. Но в питоне батарейки и краткость.
        • 0
          Скорее это из-за самой библиотеки
      • 0
        А PyOpenCL уже стабильна? Я задумывался не так давно о переходе на связку Python/OpenCL (сейчас Python/C++/OpenCL), но как-то страшно было, так как PyOpenCL весьма молода.
        • +1
          Да, видимо вполне. Библиотеке-то уже 3 года, да и автор активен. Судя по чейнджлогу, багфиксам он уделяет внимание.

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