на главную ] 

Технология параллельного программирования OpenCL

Е.С.Борисов

среда, 11 ноября 2015 г.

Эту статью можно считать продолжением статьи [1] о технологии CUDA, здесь мы поговорим о технологии параллельного программирования OpenCL.

1. Введение

OpenCL (Open Computing Language ) это спецификация, описывающая технологию параллельного программирования, которая в первую очередь ориентирована на GPGPU. Изначально она была разработана компанией Apple, в последствии для развития спецификаций OpenCL был образована группа разработчиков Khronos Compute [2], в неё вошли Apple, nVidia, AMD, IBM, Intel, ARM, Motorola и др. Первая версия стандарта была опубликована в конце 2008 года.

В отличии от nVidia CUDA, AMD Stream и т.п., в OpenCL изначально закладывалась мультиплатформенность, т.е. OpenCL программа должна без изменений в коде работать на GPU разных типов (разных производителей). Такая программа без изменений должна работать даже на CPU без GPU, хотя в этом случае она может выполняться существенно медленнее чем на GPU.

2. Схема работы с аппаратурой

Итак - хотим мультиплатформенность и желательно без существенных потерь в производительности. Достигается этот результат следующим образом [3,4].

OpenCL-программа работает с т.н. платформами (platform). Платформа это программный пакет, который поставляется соответствующим разработчиком аппаратных средств. Например "AMD Accelerated Parallel Processing" или "Intel OpenCL". При этом несколько платформ могут работать одновременно на одной машине.

Каждая платформа включает в себя ICD (Installable Client Driver) -- программный интерфейс OpenCL для работы с устройствами, которые эта платформа поддерживает.

В среде Linux список ссылок на ICD, присутствующих в системе, обычно хранится в каталоге /etc/OpenCL/vendors/, а библиотека libOpenCL.so выполняет роль диспетчера (ICD loader), т.е. она направляет вызовы OpenCL функций на устройства, через соответствующие ICD.


Рис.1: схема работы OpenCL программы с аппаратурой

[ Здесь ] можно увидеть пример конфигурации оборудования. На машине с процессором Intel Core2 CPU 6300 и графическим ускорителем nVidia Quadro FX1700, развернуты три платформы OpenCL: nVidia, AMD, Intel. При этом платформа nVidia поддерживает только GPU Quadro FX1700, а платформы AMD и Intel - только Intel Core2 CPU 6300 в качестве вычислительного устройства. Таким образом каждая платформа содержит одно вычислительное устройство.

Для определения возможностей OpenCL можно воспользоваться утилитой clinfo, также [ здесь ] можно скачать исходник программы, считывающую эту информацию.

Тут надо ещё сказать о версиях OpenCL SDK. При развёртывании платформы на компьютере необходимо уточнить версию пакета. Например, в процессе проведения экспериментов выяснилось, что Intel OpenCL SDK 2014 не поддерживает работу с процессором Intel Core2 CPU 6300, который был установлен на машине. Проблема решилась установкой более старого пакета Intel OpenCL SDK 2012.

3. Структура OpenCL-программы

В OpenCL (аналогично CUDA), программа разделяется на две части: первая часть - управляющая, вторая - вычислительная. В роли управляющего устройства ( host) выступает центральный процессор (CPU), вычислительное устройство ( device) выбираем из списка платформ и их устройств. Обычно используется GPU, но не обязательно, это может быть и тот же CPU.

Код, который должен выполняться на device, оформляется специальным способом. Поскольку device могут быть от разных производителей и разных типов, то скомпилировать один бинарник для всех не получится. Решается эта проблема следующим образом.

Текст ядра (kernel), т.е. части программы выполняемой на device, включается в основную часть программы (выполняемой на host) в "чистом" виде т.е. в виде текстовой строки. Этот исходник компилируется средствами OpenCL непосредственно в процессе работы программы (runtime) для выбранного в данный момент вычислительного устройства, это происходит каждый раз при запуске OpenCL-программы.

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

Далее рассмотрим общую структуру OpenCL-программы. Она выглядит более сложной чем аналогичная программой CUDA, поскольку необходимо выполнять дополнительные действия по конфигурированию среды выполнения и подготовке кода для device. Эта среда в терминах OpenCL именуется контекстом (context), она включает в себя платформу (platform), вычислительные устройства (device) и буферы памяти для них. Для device в рамках context создаётся очередь команд на исполнение (command queue). Операции с device, такие как чтение/запись данных и запуск ядра, помещаются в эту очередь и последовательно исполняются. Таким образом общая схема OpenCL-программы выглядит так:

  1. получить информацию о платформах и устройствах
    clGetPlatformIDs(),clGetPlatformInfo(), clGetDeviceIDs(),clGetDeviceInfo()
  2. выбрать устройства и создать для них контекст
    clCreateContext()
  3. создать ядро из текста программы
    clCreateProgramWithSource(), clBuildProgram(), clCreateKernel()
  4. выделить память для данных на устройствах
    clCreateBuffer()
  5. создать очередь комманд для устройтва
    clCreateCommandQueue(), clCreateCommandQueueWithProperties()
  6. скопировать данные с host на device
    clEnqueueWriteBuffer()
  7. назначить параметры выполнения ядра
    clSetKernelArg()
  8. запуск ядра
    clEnqueueNDRangeKernel()
  9. скопировать результат с device на host
    clEnqueueReadBuffer()
  10. обработка результата
  11. завершение работы, освобождение ресурсов
    clReleaseMemObject(), clReleaseKernel(), clReleaseProgram(), clReleaseCommandQueue(), clReleaseContext()

4. Пример OpenCL-программы

В качестве первого примера рассмотрим простую вычислительную задачу: С := d * A + B, где d - константа, А,В и С векторы заданного размера.

__kernel void kernel1(const float alpha, __global float *A, __global float *B, __global float *C) { int idx = get_global_id(0); C[idx] = alpha* A[idx] + B[idx]; }
Листинг 1: код ядра для задачи сложения векторов

Программа запускает много (по количеству элементов векторов) параллельных процессов (тредов) на device, каждый тред получает свой номер (get_global_id), исходя из него считывает свою часть данных из глобальной (__global) памяти device, выполняет вычисления и записывает свою часть результата обратно в память.

Код программы можно скачать [ здесь ]. Для сравнения напишем ещё простую (последовательную) программу и сравним время затраченное на вычисления. Ниже в таблице представлены результаты скорости выполнения программы на разных платформах.

платформаустройствовремя(ms)
NVIDIAQuadro FX 17003.2
AMDIntel Core2 CPU 6300 19.1
IntelIntel Core2 CPU 6300 11.0
- Intel Core2 CPU 6300 9.3
Таблица 1: результаты работы программы saxpy для разных платформ

По результатам, представленным в таб.1, видно, что выполнение примера на GPU почти в три раза быстрее чем на CPU. Для CPU платформа Intel показала лучший результат чем AMD, что очевидно для Intel Core2 CPU 6300. Наилучший результат для CPU показала простая программа, что можно объяснить эффективной работой компилятора и отсутствием необходимости выполнять дополнительные процедуры OpenCL.

5. Группы процессов и их конфигурация

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

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

// конвертер индексов матрицы в линейный адрес #define IDX2LIN(i,j,l) (i+j*l) __kernel void myGEMM1(const int M, const int N, const int K, const __global float* A, const __global float* B, __global float* C) { // номер треда (2D решетка) const int r = get_global_id(0); // строка 0..M const int c = get_global_id(1); // столбец 0..N // вычисляем элемент [r,c] результирующей матрицы C float acc = 0.0f; for (int i=0; i<K; i++) { acc += A[ IDX2LIN(r,i,M) ] * B[ IDX2LIN(i,c,K) ]; } C[ IDX2LIN(r,c,M) ] = acc; // сохраняем результат }
Листинг 2: код ядра для задачи умножения матриц

Имеем на входе матрицы A[M*K], B[K*N] и соответственно буфер для результата C[M*N]. Программа создаёт M*N тредов в виде решетки MxN, каждый тред [r,c] отрабатывает одну ячейку в матрице результата.

Код программы можно скачать [ здесь ].

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

платформаустройствовремя(ms)
NVIDIAQuadro FX 1700 5169.5
AMDIntel Core2 CPU 6300 2279.0
IntelIntel Core2 CPU 6300 920.1
- Intel Core2 CPU 6300 1340.6
Таблица 2: результаты работы программы gemm1 для разных платформ

Результаты несколько расстраивают, поскольку из таб.2 видно, что GPU показала наихудший результат. Проблема в крайне неэффективном использовании памяти GPU. Далее мы исправим это затруднение.

6. Модели памяти и синхронизация процессов

В статье [1] была приведена схема организации GPU. Из этой схемы видно, что этот тип устройств обладает сложно организованной памятью, которая может работать с разной скоростью.

Модель памяти OpenCL имеет несколько типов. Рассмотрим их подробней : global, local, private.

Вернёмся к предыдущему примеру с умножением матриц. В процессе работы ядро выполняет много повторных чтений из памяти global. Попробуем сократить количество чтений с помощью организации быстрого кэша [5].

// размер кэша тредов #define TS 16 // конвертер индексов матрицы в линейный адрес #define IDX2LIN(i,j,l) (i+j*l) __kernel void myGEMM2(const int M, const int N, const int K, const __global float* A, const __global float* B, __global float* C) { // 2D номер треда в группе const int r = get_local_id(0); const int c = get_local_id(1); // номер ячейки в матрице результата const int gr = get_group_id(0)*TS + r; // 0..M const int gc = get_group_id(1)*TS + c; // 0..N // общий кэш для тредов группы __local float Asub[TS][TS]; __local float Bsub[TS][TS]; float acc = 0.0f; // результат работы треда for (int t=0; t<K/TS; t++) { // цикл по всем блокам матриц const int tr = t*TS + r; const int tc = t*TS + c; // загружаем блоки в кэш Asub[c][r] = A[ IDX2LIN(gr,tc,M) ]; Bsub[c][r] = B[ IDX2LIN(tr,gc,K) ]; barrier(CLK_LOCAL_MEM_FENCE); // ждём пока треды группы заполнят общий кэш // вычисляем for (int k=0; k<TS; k++) { acc += Asub[k][r] * Bsub[c][k]; } barrier(CLK_LOCAL_MEM_FENCE); // ждём пока треды группы завершат вычисления } C[ IDX2LIN(gr,gc,M) ] = acc; // сохраняем результат }
Листинг 3: код ядра для задачи умножения матриц с кэшированием

Имеем на входе матрицы A[M*K], B[K*N] и соответственно результат C[M*N]. Программа создаёт M*N тредов в виде решетки MxN, группами размера TSxTS, каждый тред отрабатывает одну ячейку в матрице результата, группа тредов кэширует блоки исходных матриц и далее выполняет операции с этим кэшем.

Для того, что бы кэш корректно был заполнен необходима синхронизация группы тредов (barrier), достигая barrier программа ждет, пока все треды группы соберутся в этой точке и только после этого продолжает выполнение.

Код программы можно скачать [ здесь ].

Ниже в таблице представлены результаты скорости выполнения программы на разных платформах.

платформаустройствовремя(ms)
NVIDIAQuadro FX 1700 169.5
AMDIntel Core2 CPU 6300 2301.1
IntelIntel Core2 CPU 6300 1402.7
- Intel Core2 CPU 6300 1340.6
Таблица 3: результаты работы программы gemm2 для разных платформ

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

7. Заключение

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

В заключении можно ещё привести ссылку на список библиотек, основанных на OpenCL.
https://www.khronos.org/opencl/resources/opencl-libraries-and-frameworks-with-opencl-acceleration

Список литературы

  1. Е.С.Борисов Технология параллельного программирования CUDA -- http://mechanoid.kiev.ua/parallel-cuda.html
  2. Khronos group. OpenCL. -- https://www.khronos.org/opencl/
  3. Janusz Kowalik, Tadeusz Puzniakowski Using OpenCL. Programming Maassively Paarallel Commputers. -- IOS Press -- 2012
  4. Benedict R. Gaster, Lee Howes, David R. Kaeli, Perhaad Mistry, Dana Schaa Heterogeneous Computing with OpenCL Revised OpenCL. -- Advanced Micro Devices, Inc. Published by Elsevier Inc. -- 2013
  5. Cedric Nugteren Tutorial: OpenCL SGEMM tuning for Kepler -- http://www.cedricnugteren.nl/tutorial.php
Яндекс.Метрика
TOP.zp.ua PR-CY.ru
При использовании материалов этого сайта, пожалуйста вставляйте в свой текст ссылку на мою статью.