Встроенная поддержка параллельных вычислений: OpenCL

OpenCL — это открытый стандарт параллельного программирования, который позволяет создавать приложения для одновременного выполнения на множестве ядер современных процессоров, различных по архитектуре, в частности, графических (GPU) или центральных (CPU).

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

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

Вместе с тем, у начинающих часто возникает вопрос, можно ли с помощью OpenCL ускорить процессы тестирования и оптимизации советников. Ответы на оба вопроса: нет. Тестирование воспроизводит реальный процесс последовательной торговли и потому каждый следующий бар или тик зависит от результатов предыдущих, что делает невозможным распараллеливание расчетов одного прохода. Что же касается оптимизации, то агенты тестера поддерживают только ядра центрального процессора. Это связано со сложностью полноценного анализа котировок или тиков, отслеживания позиций и подсчета баланса и эквити. Однако, если сложность вас не пугает, вы можете реализовать собственный "движок" оптимизации на ядрах графических карт, перенеся все вычисления, эмулирующие торговое окружение с требуемой достоверностью, в OpenCL.

OpenCL "расшифровывается" как Open Computing Language — открытый язык вычислений. Он похож на языки C и C++, а стало быть, и на MQL5. Однако для того, чтобы подготовить ("откомпилировать") программу на OpenCL, передать в неё входные данные, запустить параллельно на нескольких ядрах и получить результаты вычислений, применяется специальный программный интерфейс (набор функций). Это OpenCL API доступно и для MQL-программ, желающих организовать параллельное исполнение.

Для использования OpenCL совсем не обязательно иметь видеокарту на Вашем ПК — вполне достаточно и наличия центрального процессора, но в любом случае требуется наличие специальных драйверов от производителя (требуется версия OpenCL 1.1 и выше). Если на вашем компьютере установлены игры или другой софт (например, научный, редактор видео и пр.), работающий напрямую с видеокартами, то необходимая программная "прослойка", скорее всего, уже имеется. Это можно проверить, попытавшись запустить в терминале MQL-программу с обращением к OpenCL (хотя бы простой пример из поставки терминала, см. далее).

При отсутствии поддержки OpenCL вы увидите в журнале ошибку.

OpenCL OpenCL not found, please install OpenCL drivers

Если же на вашем компьютере есть подходящее устройство и для него была включена поддержка OpenCL, терминал выведет сообщение с названием и типом этого устройства (их может быть несколько). Например:

OpenCL Device #0: CPU GenuineIntel Intel(R) Core(TM) i7-2700K CPU @ 3.50GHz with OpenCL 1.1 (8 units, 3510 MHz, 16301 Mb, version 2.0, rating 25)
OpenCL Device #1: GPU Intel(R) Corporation Intel(R) UHD Graphics 630 with OpenCL 2.1 (24 units, 1200 MHz, 13014 Mb, version 26.20.100.7985, rating 73)

Процедура установки драйверов описана для разных устройств в статье на сайте mql5.com. Разумеется, поддержка распространяется на наиболее популярные устройства от Intel, AMD, ATI, Nvidia.

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

Важно: Если на компьютере имеется видеокарта с поддержкой OpenCL, то ставить программную эмуляцию OpenCL на центральном процессоре не нужно!

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

Подготовительные операции по настройке среды исполнения OpenCL в MQL-программе выполняются однократно с помощью функций вышеупомянутого OpenCL API.

  1. Создание контекста для программы OpenCL (выбор устройства, например, видеокарты, ЦП или любого доступного): CLContextCreate(CL_USE_ANY). Функция вернет дескриптор контекста (целое число, обозначим его условно ContextHandle).
  2. Создание в полученном контексте OpenCL-программы: она "компилируется" на основе исходного кода на языке OpenCL с помощью вызова функции CLProgramCreate, в которую текст кода передается через параметр Source: CLProgramCreate(ContextHandle, Source, BuildLog). Функция вернет дескриптор программы (целое число ProgramHandle). Здесь важно отметить, что внутри исходного кода этой программы должны присутствовать функции (хотя бы одна), помеченные специальным ключевым словом __kernel (или просто kernel): именно они содержат части алгоритма, подлежащие распараллеливанию (см. пример ниже). Разумеется, программист может для упрощения (декомпозиции исходного кода) разнести логические подзадачи функции-кернела на другие вспомогательные функции и вызывать их из кернела: при этом помечать вспомогательные функции словом kernel не надо.
  3. Регистрация кернела для выполнения по имени одной из тех функций, что помечены в коде OpenCL-программы как "кернелобразующие": CLKernelCreate(ProgramHandle, KernelName). Вызов этой функции вернет дескриптор кернела (целое число, допустим — KernelHandle). Вы можете подготовить в OpenCL-коде много разных функций и зарегистрировать их как разные кернелы.
  4. При необходимости, создание буферов для массивов данных, передаваемых по ссылке в кернел и для возвращаемых значений/массивов: CLBufferCreate(ContextHandle, Size * sizeof(double), CL_MEM_READ_WRITE) и др. Буфера тоже идентифицируются и управляются с помощью дескрипторов.

Далее, однократно или, при необходимости, многократно (например, в обработчиках событий индикатора или эксперта), производятся непосредственно вычисления по следующей схеме:

  1. Передача входных данных и/или привязка входных/выходных буферов с помощью CLSetKernelArg(KernelHandle,...) и/или CLSetKernelArgMem(KernelHandle,..., BufferHandle). Первая функция обеспечивает установку скалярного значения, а вторая эквивалента передаче или получению значения (или массива значений) по ссылке. На этом этапе происходит перемещение данных из MQL5 в исполнительное ядро OpenCL. Для записи данных в буфер используется CLBufferWrite(BufferHandle,...). Параметры и буфера станут доступны OpenCL-программе во время выполнения кернела.
  2. Выполнение параллельных вычислений вызовом конкретного кернела CLExecute(KernelHandle,...). Функция-кернел сможет записать результаты своей работы в выходной буфер.
  3. Получение результата с помощью CLBufferRead(BufferHandle). На этом этапе происходит обратное перемещение данных из OpenCL в MQL5.

После завершения вычислений следует освободить все дескрипторы: CLBufferFree(BufferHandle), CLKernelFree(KernelHandle), CLProgramFree(ProgramHandle), CLContextFree(ContextHandle).

Данная последовательность условно обозначена на следующей схеме.

Схема взаимодействия MQL-программы и OpenCL-вложения

Схема взаимодействия MQL-программы и OpenCL-вложения

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

Стандартная библиотека заголовочных файлов, поставляемая с терминалом, содержит класс-обертку для работы с OpenCL: MQL5/Include/OpenCL/OpenCL.mqh.

А примеры использования OpenCL можно найти в каталоге MQL5/Scripts/Examples/OpenCL/. В частности, там имеется скрипт MQL5/Scripts/Examples/OpenCL/Double/Wavelet.mq5, производящий вейвлет-преобразование временного ряда (можно взять искусственную кривую по стохастической модели Вейерштрасса или приращения цен текущего финансового инструмента). В любом случае исходными данными для алгоритма выступает массив — двумерное изображение ряда.

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

Scripts script Wavelet (EURUSD,H1) loaded successfully
OpenCL  device #0: GPU NVIDIA Corporation NVIDIA GeForce GTX 1650 with OpenCL 3.0 (16 units, 1560 MHz, 4095 Mb, version 512.72)
OpenCL  device #1: GPU Intel(R) Corporation Intel(R) UHD Graphics 630 with OpenCL 3.0 (24 units, 1150 MHz, 6491 Mb, version 27.20.100.8935)
OpenCL  device performance test started
OpenCL  device performance test successfully finished
OpenCL  device #0: GPU NVIDIA Corporation NVIDIA GeForce GTX 1650 with OpenCL 3.0 (16 units, 1560 MHz, 4095 Mb, version 512.72, rating 129)
OpenCL  device #1: GPU Intel(R) Corporation Intel(R) UHD Graphics 630 with OpenCL 3.0 (24 units, 1150 MHz, 6491 Mb, version 27.20.100.8935, rating 136)
Scripts script Wavelet (EURUSD,H1) removed

В результате выполнения скрипт выводит в закладку Эксперты записи с измерениями скорости расчетов обычным способом (последовательно, на ЦПУ) и параллельным (на ядрах OpenCL).

OpenCL: GPU device 'Intel(R) UHD Graphics 630' selected
time CPU=5235 ms, time GPU=125 ms, CPU/GPU ratio: 41.880000

Соотношение скоростей, в зависимости от специфики задачи, может достигать десятков.

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

Исходный смоделированный ряд, его приращения и вейвлет-преобразование

Исходный смоделированный ряд, его приращения и вейвлет-преобразование

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

А вот как выглядит исходный OpenCL-код вейвлет-преобразования, вынесенный в отдельный файл MQL5/Scripts/Examples/OpenCL/Double/Kernels/wavelet.cl.

// требуется повышенная точность расчетов double
// (по умолчанию, без этой директивы получим float)
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
   
// Вспомогательная функция Morlet
double Morlet(const double t)
{
   return exp(-t * t * 0.5) * cos(M_2_PI * t);
}
   
// OpenCL kernel function
__kernel void Wavelet_GPU(__global double *data, int datacount,
   int x_size, int y_size, __global double *result)
{
   size_t i = get_global_id(0);
   size_t j = get_global_id(1);
   double a1 = (double)10e-10;
   double a2 = (double)15.0;
   double da = (a2 - a1) / (double)y_size;
   double db = ((double)datacount - (double)0.0) / x_size;
   double a = a1 + j * da;
   double b = 0 + i * db;
   double B = (double)1.0;
   double B_inv = (double)1.0 / B;
   double a_inv = (double)1.0 / a;
   double dt = (double)1.0;
   double coef = (double)0.0;
   
   for(int k = 0; k < datacount; k++)
   {
      double arg = (dt * k - b) * a_inv;
      arg = -B_inv * arg * arg;
      coef = coef + exp(arg);
   }
   
   double sum = (float)0.0;
   for(int k = 0; k < datacount; k++)
   {
      double arg = (dt * k - b) * a_inv;
      sum += data[k] * Morlet(arg);
   }
   sum = sum / coef;
   uint pos = (int)(j * x_size + i);
   result[pos] = sum;
}

Полную информацию о синтаксисе, встроенных функциях и принципах функционирования OpenCL можно найти на официальном сайте разработчика Khronos Group.

В частности, интересно отметить, что OpenCL поддерживает не только привычные скалярные числовые типы данных (начиная от char и заканчивая double), но и векторные (u)charN, (u)shortN, (u)intN, (u)longN, floatN, doubleN, где N = {2|3|4|8|16} и обозначает длину вектора. В данном примере это не используется.

Помимо упомянутого ключевого слова kernel важную роль в организации параллельных вычислений играет функция get_global_id: она позволяет узнать в коде номер вычислительной подзадачи, выполняющейся в данный момент. Очевидно, что расчеты в разных подзадачах должны отличаться (иначе не было бы смысла задействовать много ядер). И в данном примере, поскольку задача подразумевает анализ двумерного изображения, её фрагменты удобнее идентифицировать с помощью двух ортогональных координат — именно их в вышеприведенном коде мы получаем с помощью двух вызовов get_global_id(0) и get_global_id(1).

В принципе, размерность данных для задачи мы сами задаем при вызове в MQL5 функции CLExecute (см. далее).

В файле Wavelet.mq5 исходный код OpenCL подключен с помощью директивы:

#resource "Kernels/wavelet.cl" as string cl_program

Размер изображения задан макросами:

#define SIZE_X 600
#define SIZE_Y 200

Для управления OpenCL используется стандартная библиотека с классом COpenCL. Его методы имеют похожие названия и внутри используют соответствующие встроенные функции OpenCL из MQL5 API. С ним предлагается ознакомиться самостоятельно.

#include <OpenCL/OpenCL.mqh>

В упрощенном виде (без проверок на ошибки и визуализации) MQL-код, обеспечивающий запуск преобразования, представлен ниже. Связанные с вейвлет-преобразованием действия сведены в класс CWavelet.

class CWavelet
{
protected:
   ...
   int        m_xsize;              // размеры изображения по осям
   int        m_ysize;
   double     m_wavelet_data_GPU[]; // результат попадет сюда
   COpenCL    m_OpenCL;             // объект-обертка
   ...
};

Основные "параллельные" вычисления организует его метод CalculateWavelet_GPU.

bool CWavelet::CalculateWavelet_GPU(double &data[], uint &time)
{
   int datacount = ArraySize(data); // размер изображения (количество точек)
   
   // компилируем cl-программу по её исходному коду
   m_OpenCL.Initialize(cl_programtrue);
   
   // регистрируем единственную функцию-кернел из cl-файла
   m_OpenCL.SetKernelsCount(1);
   m_OpenCL.KernelCreate(0"Wavelet_GPU");
   
   // регистрируем 2 буфера для ввода и вывода данных, записываем входной массив
   m_OpenCL.SetBuffersCount(2);
   m_OpenCL.BufferFromArray(0data0datacountCL_MEM_READ_ONLY);
   m_OpenCL.BufferCreate(1m_xsize * m_ysize * sizeof(double), CL_MEM_READ_WRITE);
   m_OpenCL.SetArgumentBuffer(000);
   m_OpenCL.SetArgumentBuffer(041);
   
   ArrayResize(m_wavelet_data_GPUm_xsize * m_ysize);
   uint work[2];              // задача анализа двумерного изображения - отсюда размерность 2
   uint offset[2] = {00};   // начинаем с самого начала (а можно что-то пропустить)
   work[0] = m_xsize;
   work[1] = m_ysize;
   
   // задаем входные данные   
   m_OpenCL.SetArgument(01datacount);
   m_OpenCL.SetArgument(02m_xsize);
   m_OpenCL.SetArgument(03m_ysize);
   
   time = GetTickCount();     // отсечка времени для замера скорости
   // запуск вычислений на GPU, двумерная задача
   m_OpenCL.Execute(02offsetwork);
   
   // получение результатов в выходной буфер
   m_OpenCL.BufferRead(1m_wavelet_data_GPU00m_xsize * m_ysize);
   
   time = GetTickCount() - time;
   
   m_OpenCL.Shutdown(); // освобождаем все ресурсы - вызов всех нужных функций CL***Free
   return true;
}

В исходном коде примера присутствует закомментированная строка с вызовом PreparePriceData для подготовки входного массива на основе реальных цен: вы можете активировать её вместо предыдущей строки с вызовом PrepareModelData (которая генерирует искусственный ряд).

void OnStart()
{
   int momentum_period = 8;
   double price_data[];
   double momentum_data[];
   PrepareModelData(price_dataSIZE_X + momentum_period);
   
   // PreparePriceData("EURUSD", PERIOD_M1, price_data, SIZE_X + momentum_period);
   
   PrepareMomentumData(price_datamomentum_datamomentum_period);
   ... // визуализация ряда и приращений
   CWavelet wavelet;
   uint time_gpu = 0;
   wavelet.CalculateWavelet_GPU(momentum_datatime_gpu);
   ... // визуализация результата вейвлет-преобразования
}

Для работы с OpenCL выделен специальный набор кодов ошибок (с префиксом ERR_OPENCL_, начиная с кода 5100, ERR_OPENCL_NOT_SUPPORTED), доступный в справке. При возникновении проблем с исполнением OpenCL-программ, терминал выводит подробную диагностику в журнал с указанием кодов ошибок.