MetaTrader 5 / Примеры
English
preview
Понимание и эффективное использование OpenCL API путем воссоздания встроенной поддержки в виде DLL в Linux (Часть 1): Мотивация и проверка

Понимание и эффективное использование OpenCL API путем воссоздания встроенной поддержки в виде DLL в Linux (Часть 1): Мотивация и проверка

MetaTrader 5 — Примеры | 11 апреля 2023, 16:29
1 072 0
Wasin Thonkaew
Wasin Thonkaew

Содержание

  • Введение
  • Проблема (мотивация)
  • Обходной путь
  • План работы
  • Определение OpenCL
  • Терминология OpenCL (кратко)
  • Диаграмма классов OpenCL
  • Этап I - Проверка с использованием устройства GPU с OpenCL путем разработки простой программы тестирования OpenCL

    Темы будущих статей

    • Этап II — Разработка простой экспериментальной версии поддержки OpenCL в виде DLL
    • Этап III - Разработка полной поддержки OpenCL в виде DLL
    • Port Includes/OpenCL/OpenCL.mqh
    • Портирование примеров OpenCL (MQL5) для использования нового решения OpenCL
      • BitonicSort
      • FFT
      • MatrixMult
      • Wavelet
    • Тест производительности как встроенного OpenCL, так и нашего решения OpenCL (исключительно ради интереса)
    • Заключение
    Моя предыдущая статья "Разработка экспериментальной DLL с поддержкой многопоточности в C++ для MetaTrader 5 на Linux" помогает составить представление о том, как разрабатывать инструменты/решения для MetaTrader 5 в Linux. Эти знания также будут использованы и в дальнейшем для моих будущих статей.


    Введение

    OpenCL (Open Computing Language) — это фреймворк, который позволяет создавать программы для выполнения на графическом и центральном процессорах или выделенном устройстве-ускорителе с тем преимуществом, что оно может ускорить тяжелые вычисления, необходимые для той или иной предметной области. Использование OpenCL особенно актуально при работе с графическим процессором для параллельной обработки больших объемов данных с высокой пропускной способностью памяти и выделенным набором инструкций, оптимизированным для математических вычислений. Он имеет множество процессорных ядер, называемых вычислительными блоками, каждый из которых может независимо выполнять вычисления. Центральный же процессор предназначен для выполнения более общих задач. У него меньше процессорных ядер. Графический процессор подходит для тяжелых вычислений, особенно в графической области.

    MetaTrader 5 поддерживает OpenCL версии 1.2. Платформа имеет несколько встроенных функций, полезных для пользователей.


    Проблема (мотивация)

    Проблема заключается в том, что встроенная поддержка OpenCL в MetaTrader 5 хотя и может обнаруживать графический процессор, но не может выбрать его для использования с OpenCL при вызове CLContextCreate() с CL_USE_GPU_ONLY или CL_USE_GPU_DOUBLE_ONLY. Она всегда возвращает код ошибки 5114.

    MetaTrader 5 может обнаружить графический процессор

    Встроенная поддержка OpenCL в MetaTrader 5 способна обнаруживать графический процессор


    Ошибка выбора устройства встроенной поддержки MetaTrader 5 OpenCL

    Встроенная поддержка OpenCL в MetaTrader 5 всегда выдает ошибку 5114 (ошибка выбора устройства) при тестировании в моей системе Linux

    В разделе Ошибки времени выполнения такой код ошибки соответствует следующему объяснению.

    ERR_OPENCL_SELECTDEVICE

    5114

    Ошибка выбора OpenCL устройства

    Для справки, ниже приведены характеристики моей машины

    • Ubuntu 20.04.3 LTS с ядром 5.16.0
    • Процессор: 6-ядерный процессор AMD Ryzen 5 3600 (2 потока на ядро)
    • Графический процессор: AMD Radeon RX 570 Series (поддержка OpenCL 1.2)
    • Поддержка графического драйвера осуществляется через драйвер с открытым исходным кодом, доступный в Ubuntu (как OpenGL, так и Vulkan), а не через проприетарный AMDGPU-Pro.

    Как я уже отметил, MetaTrader 5 может отображать список всех устройств, включая графический процессор, как показано на вкладке "Журнал", поэтому проблем с графическим драйвером, установленным на моем компьютере с Linux, нет. Проблема не в открытом исходном коде и не в проприетарном графическом драйвере. Весьма вероятно наличие ошибки в коде выбора устройства всякий раз, когда эти два упомянутых флага используются с CLContextCreate(). Мы используем эту проблему в качестве мотивации для проверки нашего предположения об ошибке, а затем приступим к разработке полноценного решения в одной из последующих статей серии. Наши усилия направлены на то, чтобы лучше понять концепции OpenCL, его терминологию и, что наиболее важно, как мы можем эффективно использовать его API, особенно когда мы используем его с MQL5 для разработки связанных инструментов на платформе MetaTrader 5.

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


    Обходной путь

    Мне удалось найти обходной путь для решения этой проблемы. Он связан с использованием порядкового номера графического процессора  (как он указан на вкладке "Журнал") непосредственно в виде функции.

    В моем случае графический процессор указан в первой записи на вкладке "Журнал" (первый рисунок в верхней части статьи). Назовем функцию

    CLContextCreate(0)

    Работает вполне сносно.


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

    Замените код

    ...
bool COpenCL::Initialize(const string program,const bool show_log)
  {
   if(!ContextCreate(CL_USE_ANY))
      return(false);
   return(ProgramCreate(program,show_log));
  }
...

    на

    ...
bool COpenCL::Initialize(const string program,const bool show_log)
  {
   if(!ContextCreate(0))
      return(false);
   return(ProgramCreate(program,show_log));
  }
...

    где значение параметра должно быть порядковым номером графического процессора на вкладке "Журнал" при запуске MetaTrader 5. В моем случае это 0.

    Это не идеальное решение. Мы можем улучшить его, изменив следующим образом:

    bool COpenCL::Initialize(const string program, const int device, const bool show_log)
  {
   if(!ContextCreate(device))
      return(false);
   return(ProgramCreate(program,show_log));
  }

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

    COpenCL::Initialize() повсеместно используется встроенными примерами OpenCL (Scripts/Examples/OpenCL), в частности

    • BitonicSort - Float и Double
    • FFT - Float и Double
    • MatrixMult - Float и Double
    • Wavelet - Float и Double
    • Seascape

    Float/BitonicSort.mq5,

    void OnStart()
  {
//--- OpenCL
   COpenCL OpenCL;
   if(!OpenCL.Initialize(cl_program,true))
     {
      PrintFormat("Error in OpenCL initialization. Error code=%d",GetLastError());
      return;
     }
     ...

    заменяется на

    void OnStart()
  {
//--- OpenCL
   COpenCL OpenCL;
   if(!OpenCL.Initialize(cl_program, 0, true))
     {
      PrintFormat("Error in OpenCL initialization. Error code=%d",GetLastError());
      return;
     }
     ...


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


    План работы

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

    Затем мы разработаем экспериментальную поддержку OpenCL в виде разделяемой библиотеки (DLL) на C++, которая пытается подключиться к графическому процессору как обычно, а затем использовать ее с помощью MQL5 на платформе MetaTrader 5. Если эксперимент будет успешен, мы можем перейти к дальнейшей реализации эквивалентной поддержки OpenCL, как во встроенных API.

    Более того, мы портируем Includes/OpenCL/OpenCL.mqh , чтобы он основывался на нашем новом решении. Также портируем несколько примеров OpenCL, чтобы тщательно протестировать наше решение, а затем, наконец, провести тест встроенной и разработанной поддержкой OpenCL (исключительно ради интереса).

    Ниже представлен краткий план работы.

    1. Разработка простой программы тестирования OpenCL в виде отдельного исполняемого файла (сосредоточимся исключительно на проверке использования графического процессора для выполнения функций ядра).
    2. Разработка простой поддержки OpenCL в виде DLL для тестирования скрипта в MetaTrader 5 с помощью MQL5.
    3. Разработка поддержки OpenCL в виде DLL, которая имеет функции, схожие с функциями встроенного OpenCL в MetaTrader 5.
    4. Портирование Includes/OpenCL/OpenCL.mqh
    5. Портирование примеров OpenCL
    6. Сравнительный анализ встроенного решения и нашего решения OpenCL

    Цель здесь состоит в том, чтобы сделать всё самостоятельно, чтобы найти эффективное решение и освоить OpenCL. Наше решение не призвано заменить встроенную поддержку OpenCL в MetaTrader 5.


    Подготовка

      Я использую для работы Ubuntu 20.04. Если вы используете другие дистрибутивы, адаптируйте их в соответствии с вашими потребностями.

      • Установите пакет mingw64
      • Установите пакет winehq-devel
      • При необходимости установите поддержку драйвера OpenCL, которая зависит от модели вашей видеокарты.
        Пожалуйста, ознакомьтесь со статьей "Как установить и использовать в расчетах OpenCL". Если вы предпочитаете графический драйвер с открытым исходным кодом, то я предлагаю найти ROCm только для поддержки OpenCL, но вы можете установить графический драйвер с открытым исходным кодом, поскольку он обычно доступен в Ubuntu.


      Определение OpenCL

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

      Ниже представлен обзор OpenCL. Далее мы рассмотрим каждый подраздел.

      Модель платформы OpenCL

      Модель платформы OpenCL. Взято из Практического введения в OpenCL - OpenCL: A Hands-on Introduction (Tim Mattson, Alice Koniges, and Simon McIntosh-Smith)


      Как видно из рисунка, OpenCL состоит из множества вычислительных единиц. Каждая из них состоит из десятка процессорных элементов (Processing Element, PE). В самом общем виде память делится на память хоста и память устройства. Рассмотрим всё более детально. Взглянем на архитектуру устройства OpenCL.

      Архитектура устройства OpenCL

      Архитектура устройства OpenCL. Взято из Справочного руководства по OpenCL API 1.2 (OpenCL API 1.2 Reference Guide by Khronos)

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

      На следующем рисунке показан еще один более четкий обзор модели памяти с некоторыми взаимозаменяемыми терминами.


      OpenCL Memory Model

      Модель памяти OpenCL. Взято из Практического введения в OpenCL - OpenCL: A Hands-on Introduction (Tim Mattson, Alice Koniges, and Simon McIntosh-Smith)


      Здесь мы видим ключевые термины work-item (рабочий элемент) и work-group (рабочая группа). Мы можем рассматривать процессорный элемент как рабочий элемент. В одной рабочей группе может быть много рабочих элементов. Как упоминалось ранее, модель памяти рассредоточена по всей архитектуре, включая рабочие элементы и группы, а также устройство и хост (например, ПК). В нижней части рисунка видно, что мы, как пользователи OpenCL, будем нести ответственность за перемещение данных туда и обратно между хостом и устройством. Мы рассмотрим этот момент более пристально, когда займемся кодом. Если коротко, данные на обоих концах должны быть синхронизированы для согласованности при использовании результатов вычислений или подаче данных для выполнения ядра.

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

      N-мерный домен рабочих элементов (связанных с рабочей группой)

      Отношения между рабочими элементами и рабочими группами в контексте ядра для выполнения - Взято из Практического введения в OpenCL - OpenCL: A Hands-on Introduction (Tim Mattson, Alice Koniges, and Simon McIntosh-Smith)


      На приведенном выше рисунке показана воображаемая работа, которую необходимо выполнить с точки зрения проблемного пространства. Понимание проблемного пространства позволит нам освоить глобальные размерности и потенциально приведет нас к пониманию местных размерностей. Размерность является одной из важных настроек для эффективной работы с OpenCL API. Сам API ограничен лишь тремя измерениями. Если нашему проблемному пространству нужно больше, мы должны перевести такую проблему в трехмерную задачу. В качестве примера можно привести обработку изображений. В этой задаче имеются два измерения - ширина и высота. Мы можем вывести локальные измерения, чтобы наилучшим образом соответствовать возможностям устройства, которое мы будем использовать для выполнения такой задачи, но в большинстве случаев значение каждого глобального измерения делится на соответствующее значение в локальном измерении. Но как узнать такое значение для локальной размерности? Мы можем запросить возможности устройства с помощью OpenCL API. Нам не нужно угадывать. Мы можем реализовать универсальное решение, которое работает, адаптируясь ко всему диапазону возможностей графического процессора.

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



      Терминология OpenCL (кратко)

      Здесь рассмотрены ключевые термины OpenCL (отмеченные подчеркиванием в предыдущем разделе).

      Термин Пояснение
      Хост (host) Система, которая запускает приложение OpenCL и обменивается данными с устройством OpenCL. Хост отвечает за подготовку и работу по управлению перед выполнением ядра. Это может быть ПК, рабочая станция или кластер.
      Устройство (device) Компонент OpenCL, который запускает ядро OpenCL. Это может быть центральный или графический процессор, ускоритель или пользовательское устройство.
      Ядро (kernel) Функция, написанная на языке OpenCL C, в которой она будет выполнять математические вычисления на устройствах, которые поддерживает OpenCL.
      Рабочий элемент (work-item) Функциональный модуль внутри ядра. Коллекция рабочих элементов — это отдельные экземпляры ядра, которые выполняются параллельно. Несколько рабочих элементов называются рабочей группой.
      Рабочая группа (work-group) Группа рабочих элементов, которые выполняются параллельно. Это логическая группа, внутри которой можно обмениваться данными и синхронизировать их выполнение.
      Собственная память (private memory) Область памяти, специфичная для рабочего элемента. Каждый рабочий элемент имеет собственную память, которая не используется другими рабочими элементами. Она предназначена для хранения локальных переменных и временных данных, которые нужны только одному рабочему элементу.
      Локальная память (local memory) Область памяти, совместно используемая рабочими элементами в одной рабочей группе. Локальная память позволяет рабочим элементам взаимодействовать друг с другом. Доступ к ней происходит быстрее, чем к глобальной/постоянной памяти.
      Глобальная память/кэш (global memory/cache) Область общей памяти, к которой может обращаться рабочая группа (а значит, и рабочие элементы). Кэш расположен рядом с процессором, при этом память находится далеко (аналогично кэшу центрального процессора и ОЗУ соответственно).
      Постоянная память/кэш (constant memory/cache) Область общей памяти, предназначенная только для чтения (данные не должны изменяться во время выполнения), к которой может получить доступ рабочая группа (а значит, и рабочие элементы). Кэш расположен рядом с процессором, при этом память находится далеко (аналогично кэшу команд центрального процессора и ПЗУ).
      Глобальное измерение (global dimension) Конфигурация в виде значения и количества измерений, например 1024x1024 (два измерения по 1024 в каждом) для описания проблемного пространства. Максимальное значение - три измерения.
      Локальные измерения (local dimensions) Конфигурация в виде значения и количества измерений, например 128x128 (два измерения по 128 в каждом) для описания количества рабочих элементов и рабочих групп для выполнения ядра. Например, в случае глобальных размеров 1024x1024 и локальных размеров 128x128 количество рабочих элементов для одной рабочей группы составляет 128*128 = 16 384, необходимое количество рабочих групп составляет 1024/128 * 1024/128 = 8*8 = 64, общее количество рабочих элементов 16 384*64 или 1024*1024 = 1 048 576.
      Процессорный элемент (processing element, PE) Физический вычислительный блок для устройства поддержки, например центральный или графический процессор.
      Вычислительный блок (compute unit)
      		 Набор процессорных элементов в устройстве OpenCL, который может выполнять несколько потоков параллельно.


      Диаграмма классов OpenCL


      Диаграмма классов OpenCL 1.2

      Диаграмма классов OpenCL. Взято из Справочного руководства по OpenCL API 1.2 (Courtesy of OpenCL API 1.2 Reference Guide by Khronos)


      Ниже перечислены наиболее часто используемые классы в OpenCL.
      Далее мы используем их в реальном коде.

      Класс Описание
      cl::Platform Информация о платформе OpenCL, например имя, поставщик, профиль и расширения OpenCL.
      cl::Device Устройство OpenCL, например центральный процессор, графический процессор или другой тип процессора, который реализует стандарт OpenCL.
      cl::Context Логический контейнер для других классов. Пользователи могут начать с контекста, чтобы запросить другую информацию.
      cl::CommandQueue Очередь команд, которые будут выполняться на устройстве OpenCL.
      cl::Program Набор функций ядра, которые могут выполняться на устройстве OpenCL. Методы для создания программы из кода ядра, сборки программы, возможность запрашивать информацию о программе, например, о количестве ядер, имени, размере двоичного файла и т. д.
      cl::Kernel Точка входа имени функции OpenCL для выполнения всего ядра. Всякий раз, когда пользователи создают ядро, ему требуется правильное имя в качестве точки входа для выполнения. Пользователи могут устанавливать аргументы до выполнения.
      cl::Buffer Буфер памяти OpenCL, представляющий собой линейную область памяти, в которой хранятся данные для ввода и вывода при выполнении ядра.
      cl::Event Асинхронное представление события OpenCL для определения состояния команды OpenCL. Можно использовать для синхронизации операций между хостом и устройством.


      Отправной точкой, которая может привести ко всем другим классам, является Платформа (cl::Platform). Скоро мы увидим это в реальном коде. Обычно рабочий процесс выполнения вычислений с OpenCL выглядит следующим образом.

      1. Начинаем с cl::Platform, чтобы получить cl::Device для сортировки желаемого типа устройств
      2. Создаем cl::Context из cl::Device
      3. Создаем cl::CommandQueue из cl::Context
      4. Создаем cl::Program из кода ядра
      5. Создаем cl::Kernel из cl::Program
      6. Создаем несколько cl::Buffer для хранения входных и выходных данных до выполнения ядра (сколько зависит от предметной области)
      7. (Опционально) Создаем cl::Event для синхронизации операций в основном для достижения высокой производительности и согласованности данных между хостом и устройством в случае продолжения асинхронного вызова API
      8. Начинаем выполнение ядра с помощью cl::CommandQueue и ждем, пока результирующие данные не будут переданы обратно с устройства на хост.


      Этап I - Проверка с использованием устройства GPU с OpenCL путем разработки простой программы тестирования OpenCL

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

      Если вы знакомы с кросс-компиляцией, то быстро освоитесь, поскольку мы будем делать то же самое.

      Вы можете загрузить весь проект из прикрепленного zip-файла и просмотреть его раздел за разделом. Я не буду перечислять здесь все шаги, так как это сделает статью слишком длинной.

      Структура файлов проекта следующая:

      standalonetest

      Файл проекта (.zip) приложен к статье.

      • opencltest.cpp
      • Makefile
      • (зависимости .dll в виде файлов символических ссылок на установленные файлы .dll)
        • libgcc_s_seh-1.dll - ссылка на /usr/lib/gcc/x86_64-w64-mingw32/9.3-posix/libgcc_s_seh-1.dll
        • libstdc++-6.dll - ссылка на /usr/lib/gcc/x86_64-w64-mingw32/9.3-posix/libstdc++-6.dll
        • libwinpthread-1.dll - ссылка на /usr/x86_64-w64-mingw32/lib/libwinpthread-1.dll
        • opencl.dll - ссылка на ~/.mt5/drive_c/windows/system32/opencl.dll
      Возможно, вам придется отредактировать файлы символических ссылок, чтобы они указывали на правильное расположение в вашей системе. Расположение по умолчанию - на основе Ubuntu 20.04 и расположения префикса Wine для MetaTrader 5. Его можно редактировать с помощью команды ln -sf <путь-к-новому-файлу> <имя-ссылки>.

      Makefile

      .PHONY: all clean main.exe

COMPILER := x86_64-w64-mingw32-g++-posix
FLAGS := -O2 -fno-rtti -std=c++17 -Wall -Wextra
MORE_FLAGS ?=

all: main.exe

main.exe: opencltest.cpp
        $(COMPILER) $(FLAGS) $(MORE_FLAGS) -I. -o $@ $< -L. -lopencl

clean:
        rm -f main.exe

      В этом случае мы можем игнорировать сборку для нативного Linux и перейти сразу к кросс-компиляции для Windows. Если сомневаетесь, вернитесь к моей предыдущей статье "Разработка экспериментальной DLL с поддержкой многопоточности в C++ для MetaTrader 5 на Linux", чтобы узнать, как выполнить кросс-компиляцию с системой сборки Makefile.


      opencltest.cpp

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

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

      Приступим к работе с кодом.

      #define CL_HPP_TARGET_OPENCL_VERSION 120
#define CL_HPP_MINIMUM_OPENCL_VERSION 120

#include <CL/cl2.hpp>

#include <iostream>
#include <chrono>
#include <vector>
#include <numeric>
#include <cassert>
#include <cmath>

int main() {
        // Get the platform
        std::vector<cl::Platform> platforms;
        cl::Platform::get(&platforms);
        cl::Platform platform = platforms[0];

        // Get the device
        std::vector<cl::Device> devices;
        platform.getDevices(CL_DEVICE_TYPE_GPU, &devices);
        cl::Device device = devices[0];

        // Create the context
        cl::Context context(device);

        // Create the command queue
        cl::CommandQueue queue(context, device);

        // Create the kernel
        std::string kernelCode = "__kernel void add(__global int* a, __global int* b, __global int* c, int size) { "
                "    int i = get_global_id(0);"
                "         if (i < size)"
                "               c[i] = a[i] + b[i];"
                "}";
        cl::Program::Sources sources;
        sources.push_back({kernelCode.c_str(), kernelCode.length()});
        cl::Program program(context, sources);
        if (auto ret_code = program.build({device});
                ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Program::build, code=" << ret_code << std::endl;
                return -1;
        }
        cl::Kernel kernel(program, "add");

        // Create the input and output arrays
        const int SIZE = 10000000;
        std::vector<int> a(SIZE);
        std::vector<int> b(SIZE);
        std::vector<int> c(SIZE, 0);

        // prepare data
        std::iota(a.begin(), a.end(), 1);
        std::iota(b.rbegin(), b.rend(), 1);

        // Create the buffer
        cl::Buffer bufferA(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * a.size(), a.data());
        cl::Buffer bufferB(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * b.size(), b.data());
        cl::Buffer bufferC(context, CL_MEM_WRITE_ONLY, sizeof(int) * c.size());

        {
                cl::Event write_bufferA_event, write_bufferB_event;

                if (auto ret_code = queue.enqueueWriteBuffer(bufferA, CL_FALSE, 0, sizeof(int) * a.size(), a.data(), nullptr, &write_bufferA_event);
                        ret_code != CL_SUCCESS) {
                        std::cerr << "1 Error enqueueWriteBuffer() code=" << ret_code << std::endl;
                        return -1;
                }
                if (auto ret_code = queue.enqueueWriteBuffer(bufferB, CL_FALSE, 0, sizeof(int) * b.size(), b.data(), nullptr, &write_bufferB_event);
                        ret_code != CL_SUCCESS) {
                        std::cerr << "2 Error enqueueWriteBuffer() code=" << ret_code << std::endl;
                        return -1;
                }

                cl::Event::waitForEvents({write_bufferA_event, write_bufferB_event});
        }

        // Set the kernel arguments
        kernel.setArg(0, bufferA);
        kernel.setArg(1, bufferB);
        kernel.setArg(2, bufferC);
        kernel.setArg(3, SIZE);

        auto start = std::chrono::steady_clock::now();

        // Execute the kernel
        if (auto ret_code = queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(SIZE), cl::NullRange);
                ret_code != CL_SUCCESS) {
                std::cerr << "Error enqueueNDRangeKernel() code=" << ret_code << std::endl;
                return -1;
        }

        // Read the result
        if (auto ret_code = queue.enqueueReadBuffer(bufferC, CL_TRUE, 0, sizeof(int) * c.size(), c.data());
                ret_code != CL_SUCCESS) {
                std::cerr << "Error enqueueReadBuffer() code=" << ret_code << std::endl;
                return -1;
        }

        std::chrono::duration<double, std::milli> exec_time = std::chrono::steady_clock::now() - start;
        std::cout << "elapsed time: " << exec_time.count() << "ms" << std::endl;

        // check the result
        for (int i = 0; i < SIZE; i++) {
                assert(c[i] == SIZE + 1);
        }

        return 0;
}

      Прежде всего взглянем на верхнюю часть исходного файла.

      #define CL_HPP_TARGET_OPENCL_VERSION 120
#define CL_HPP_MINIMUM_OPENCL_VERSION 120

#include <CL/cl2.hpp>

      Обратите внимание на первые две строки. Эти определения через препроцессор #define  необходимы, чтобы OpenCL знал, что нам нужен OpenCL 1.2. Эти две строки должны быть там, прежде чем мы включим заголовочный файл CL/cl2.hpp.

      Хотя мы стремимся к OpenCL 1.2, мы решили включить заголовочный файл cl2.hpp, потому что есть некоторые функции поддержки из OpenCL, например, SVM (Shared Virtual Memory, общая виртуальная память) для более эффективного доступа к памяти в приложениях определенного типа, что делает его более удобным в использовании особенно для согласования с последней версией стандарта C++. Хотя мы еще не использовали функции, связанные с OpenCL, но всякий раз, когда MetaTrader 5 будет обновляться для поддержки OpenCL 2.x, работа по миграции нашей базы кодов станет проще.

      Короче говоря, нам необходима версия 1.2, так как это официальная версия OpenCL, поддерживаемая MetaTrader 5.


      Далее создадим cl::Context, включающий cl::Platform, и cl::Device.

              ...
        // Get the platform
        std::vector<cl::Platform> platforms;
        cl::Platform::get(&platforms);
        cl::Platform platform = platforms[0];

        // Get the device
        std::vector<cl::Device> devices;
        platform.getDevices(CL_DEVICE_TYPE_GPU, &devices);
        cl::Device device = devices[0];

        // validate that it is GPU
        assert(device.getInfo<CL_DEVICE_TYPE>() == CL_DEVICE_TYPE_GPU);
        ...

      Платформа состоит из нескольких устройств. Чтобы найти желаемое устройство для работы, нам необходимо перебрать все платформы и все устройства для каждой платформы. Необходимо проверить тип устройства: центральный процессор, графический процессор, ускоритель (более специализированное устройство) или пользовательское устройство. В нашем случае для быстрого тестирования я жестко запрограммировал использование первой платформы и первого порядкового номера, т.е. 0, указывающего на графический процессор, указанный на вкладке "Журнал", как упоминалось в начале статьи. Мы также проверяем, что получаемое нами устройство является графическим процессором.

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

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


      Далее напишем функцию ядра, создадим cl::Program из написанного кода ядра, затем создадим из него cl::Kernel.

              ...
        // Create the command queue
        cl::CommandQueue queue(context, device);

        // Create the kernel
        std::string kernelCode = "__kernel void add(__global int* a, __global int* b, __global int* c, int size) { "
                "    int i = get_global_id(0);"
                "         if (i < size)"
                "               c[i] = a[i] + b[i];"
                "}";
        cl::Program::Sources sources;
        sources.push_back({kernelCode.c_str(), kernelCode.length()});
        cl::Program program(context, sources);
        if (auto ret_code = program.build({device});
                ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Program::build, code=" << ret_code << std::endl;
                return -1;
        }
        cl::Kernel kernel(program, "add");
        ...

       Кроме того, создадим cl::CommandQueue для последующего использования. Код ядра можно записать в отдельный файл (обычно с расширением .cl) или встроенную строку вместе с кодом C++, как мы сделали здесь.

      Ниже перечислены несколько важных аннотаций.

      • __kernel - важная аннотация, которая должна быть помещена перед именем функции, чтобы показать, что такая функция является функцией ядра. Функция ядра будет выполняться на устройстве. Ее можно вызывать с хоста вместе с настройкой аргументов. Функция ядра также может вызывать другие нормальные функции, так что эти нормальные функции не нужно аннотировать с помощью__kernel. Таким образом, мы можем рассматривать функцию ядра как точку входа для выполнения или вычислительную работу, выполняемую OpenCL.
      • __global - используется для объявления того, что данные, переданные в качестве параметра, хранятся в глобальной памяти, которая совместно используется всеми рабочими элементами (вспомните диаграммы вверху).
      • get_global_id(dimindx) - вернуть уникальное значение глобального идентификатора рабочего элемента из указанного значения индекса измерения (dimindx). В нашем случае у нас есть только одно измерение, значение которого равно размеру входного массива. Таким образом, эта функция будет однозначно определять, какой элемент массива будет вычисляться полностью параллельно (можно провести аналогию с индексам массива).

      Используя три приведенные выше аннотации, сделаем следующие пояснения:

      • аргументы a, b и c аннотированы __global. Это означает, что данные, на которые указывают эти указатели, взяты из глобальной памяти, совместно используемой всеми рабочими элементами. Таким образом, все рабочие элементы будут работать вместе, чтобы вычислить суммирование и установить значение результата в массив c.
      • Он имеет безопасную проверку size для ограничения вычисления размером входного/выходного массива, хотя это и не обязательно, так как он уже ограничено значением в глобальном измерении.
      • get_global_id(0) используется для возврата того индекса, над которым будет работать текущий рабочий элемент.

      После этого мы создаем cl::Program::Sources в качестве входного параметра в cl::Kernel. Как я писал выше, имя ядра должно быть правильным и совпадать с функцией ядра.


      Далее мы выделяем и подготавливаем данные.

              ...
        // Create the input and output arrays
        const int SIZE = 10000000;
        std::vector<int> a(SIZE);
        std::vector<int> b(SIZE);
        std::vector<int> c(SIZE, 0);

        // prepare data
        std::iota(a.begin(), a.end(), 1);
        std::iota(b.rbegin(), b.rend(), 1);
        ...

      Используем std::vector здесь, чтобы выделить достаточно памяти, которую предполагается использовать для a, b и c. Затем используем std::iota, чтобы заполнить все элементы следующим образом:

      • массив a - от 1 до SIZE. Заполняется всё - от первого элемента массива до конца
      • массив b - от 1 до SIZE. Заполняется всё - от последнего элемента массива до начала

      Это означает, что всякий раз, когда суммируется каждый элемент массива a и b, он будет равен SIZE+1, как и всегда. Позже мы будем использовать это как подтверждение утверждения.

      Далее создаем буферы и подготавливаем данные для устройства.

              ...
        // Create the buffer
        cl::Buffer bufferA(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * a.size(), a.data());
        cl::Buffer bufferB(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * b.size(), b.data());
        cl::Buffer bufferC(context, CL_MEM_WRITE_ONLY, sizeof(int) * c.size());

        // (This block is OPTIONAL as CL_MEM_COPY_HOST_PTR already took care of copying data from host to device for us. It is for a demonstration of cl::Event usage)
        {
                cl::Event write_bufferA_event, write_bufferB_event;

                if (auto ret_code = queue.enqueueWriteBuffer(bufferA, CL_FALSE, 0, sizeof(int) * a.size(), a.data(), nullptr, &write_bufferA_event);
                        ret_code != CL_SUCCESS) {
                        std::cerr << "1 Error enqueueWriteBuffer() code=" << ret_code << std::endl;
                        return -1;
                }
                if (auto ret_code = queue.enqueueWriteBuffer(bufferB, CL_FALSE, 0, sizeof(int) * b.size(), b.data(), nullptr, &write_bufferB_event);
                        ret_code != CL_SUCCESS) {
                        std::cerr << "2 Error enqueueWriteBuffer() code=" << ret_code << std::endl;
                        return -1;
                }

                cl::Event::waitForEvents({write_bufferA_event, write_bufferB_event});
        }
        ...

      bufferA, bufferB и bufferC связаны с массивами a, b и cсоответственно. Обратите внимание на флаги, используемые для создания каждого буфера, так как они влияют на то, как мы должны подготовить данные, связанные с соответствующим буфером.

      Флаг Значение
      CL_MEM_READ_WRITE Объект памяти разрешен для чтения и записи ядром
      CL_MEM_WRITE_ONLY Объект памяти разрешен для записи только ядром
      CL_MEM_READ_ONLY Объект памяти разрешен для чтения только ядром
      CL_MEM_USE_HOST_PTR Реализация OpenCL должна использовать память, которая уже выделена приложением.
      CL_MEM_ALLOC_HOST_PTR Реализация OpenCL должна выделить память для объекта памяти, а также разрешить доступ к хосту
      CL_MEM_COPY_HOST_PTR  То же, что и CL_MEM_ALLOC_HOST_PTR, но он также автоматически копирует данные на устройство

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

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

      bufferA и bufferB создаются с использованием CL_MEM_READ_ONLY, что означает, что ядро может только читать данные из него. Это имеет смысл, потому что мы уже подготовили данные до запуска ядра. Ядро просто должно прочитать из него, а затем вычислить. bufferC создан с CL_MEM_WRITE_ONLY. Это означает, что ядро помещает в него результат. Это также имеет смысл, поскольку на данный момент у хоста нет причин снова менять результат.

      (опционально) Для кода внутри скобочного блока

      Код внутри блока необязателен. Так как мы используем CL_MEM_COPY_HOST_PTR для bufferA и bufferB, базовая система OpenCL позаботится о копировании данных с хоста на устройство для нас без необходимости делать это снова. Мы добавили туда такой код для демонстрации использования cl::Event для синхронизации операций.

      cl::CommandQueue::enqueueWriteBuffer() поставит в очередь команду записи данных, связанных с указанным буфером, в устройство. Здесь у нас есть два варианта:

      1. CL_FLASE - не ждет завершения операции, возвращается немедленно (асинхронный, или неблокирующий)
      2. CL_TRUE - ждет завершения поставленной в очередь операции, а затем возвращается (синхронный, или блокирующий)

      Как видите, если указать параметр с CL_FALSE для связанного вызова API, такого как cl::CommandQueue::enqueueWriteBuffer(), нам нужно объединить его использование с одним из примитивов синхронизации, например cl::Event.
      Преимущество в том, что оба вызова cl::CommandQueue::enqueueWriteBuffer() вернутся немедленно, и мы подождем, пока они все закончат работу. Вместо того, чтобы ждать, пока они завершатся один за другим, мы экономим время, ставя в очередь другую операцию и ожидая завершения всех операций.


      Далее установим все аргументы функции ядра.

              ...
        // Set the kernel arguments
        kernel.setArg(0, bufferA);
        kernel.setArg(1, bufferB);
        kernel.setArg(2, bufferC);
        kernel.setArg(3, SIZE);
        ...

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

      Далее мы, наконец, приступаем к выполнению ядра и ждем результата.

              ...
        auto start = std::chrono::steady_clock::now();

        // Execute the kernel
        if (auto ret_code = queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(SIZE), cl::NullRange);
                ret_code != CL_SUCCESS) {
                std::cerr << "Error enqueueNDRangeKernel() code=" << ret_code << std::endl;
                return -1;
        }

        // Read the result
        if (auto ret_code = queue.enqueueReadBuffer(bufferC, CL_TRUE, 0, sizeof(int) * c.size(), c.data());
                ret_code != CL_SUCCESS) {
                std::cerr << "Error enqueueReadBuffer() code=" << ret_code << std::endl;
                return -1;
        }

        std::chrono::duration<double, std::milli> exec_time = std::chrono::steady_clock::now() - start;
        std::cout << "elapsed time: " << exec_time.count() << "ms" << std::endl;
        ...

      cl::CommandQueue::enqueueNDRangeKernel() используется для постановки в очередь команды (или операции) для выполнения указанного ядра с указанными размерами.

      Сразу после этого мы ждем, пока результат не будет записан обратно в наш результирующий массив, связанный с bufferC с помощью cl::CommandQueue::enqueueReadBuffer(). Обратите внимание также на параметрCL_TRUE, который будет ждать, пока операция не будет выполнена, как я объяснял ранее. Это означает, что устройство записывает результат обратно в результирующий массив в памяти устройства, а затем копирует такие данные обратно на хост, который является массивом c, связанным с bufferC. В этом смысл вызова enqueueReadBuffer().

      Обертывание кода путем сравнения времени выполнения с использованием обычного std::chrono::steady_clockдля монотонных часов, работающих только в прямом направлении, на которые не влияют внешние факторы, например настройка системных часов.


      Наконец, мы проверяем правильность полученных результатов.

              ...
        // check the result
        for (int i = 0; i < SIZE; i++) {
                assert(c[i] == SIZE + 1);
        }
        ...

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

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

      $ make
x86_64-w64-mingw32-g++-posix -O2 -fno-rtti -std=c++17 -Wall -Wextra  -I. -o main.exe opencltest.cpp -L. -lopencl

$ WINEPREFIX=~./mt5 wine main.exe
elapsed time: 9.1426ms

      Отлично! Если вы не видите сообщения об ошибке или завершения работы программы (результат assert()), то она работает совершенно нормально от начала до конца.

      Итак, мы только что подтвердили, что поиск и использование графического процессора на ПК с OpenCL не вызывает никаких проблем! Мы также подтверждаем наличие проблемы с кодом выбора устройства в самом MetaTrader 5.


      Что дальше?

      Мы изучили OpenCL сверху донизу, поняли его концепцию, архитектуру, модель памяти, а затем прошли практический пример кода от начала до конца. Эти знания будут использованы в реальной правильной реализации DLL или даже в обычном применении встроенного API OpenCL MQL5. Почему? Потому что понимание концепции OpenCL поможет нам обобщить нашу проблемную область, чтобы она соответствовала глобальному/локальному измерениям OpenCL и его концепции рабочего элемента/рабочей группы, не говоря уже об эффективном использовании памяти и полном использовании возможностей устройства при выполнении параллельной задачи.

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

      Перевод с английского произведен MetaQuotes Ltd.
      Оригинальная статья: https://www.mql5.com/en/articles/12108

      Прикрепленные файлы |
      Загрузить ZIP
      OpenCLSimple.zip (2.93 KB)
      Предупреждение: все права на данные материалы принадлежат MetaQuotes Ltd. Полная или частичная перепечатка запрещена.

      Другие статьи автора

      Перейти к обсуждению на форуме трейдеров
      Машинное обучение и Data Science (Часть 11): Наивный байесовский классификатор и теория вероятностей в трейдинге Машинное обучение и Data Science (Часть 11): Наивный байесовский классификатор и теория вероятностей в трейдинге
      Торговлю по вероятностям можно сравнить с ходьбой по канату — она требует точности, баланса и четкого понимания риска. В мире трейдинга вероятность решает все. Именно от нее зависит результат — успех или неудача, прибыль или убыток. Используя возможности вероятности, трейдеры могут принимать более обоснованные решения, эффективнее управлять рисками и достигать своих финансовых целей. Неважно, опытный вы инвестор или начинающий трейдер, понимание вероятности может стать ключом к раскрытию вашего торгового потенциала. В этой статье мы познакомимся с увлекательным миром вероятностного трейдинга и покажем, как вывести игру в торговлю на новый уровень.
      Эксперименты с нейросетями (Часть 5): Нормализация входных параметров для передачи в нейросеть Эксперименты с нейросетями (Часть 5): Нормализация входных параметров для передачи в нейросеть
      Нейросети наше все. Проверяем на практике, так ли это. MetaTrader 5 как самодостаточное средство для использования нейросетей в трейдинге. Простое объяснение.
      Разработка торговой системы на основе Индекса облегчения рынка MFI от Билла Вильямса Разработка торговой системы на основе Индекса облегчения рынка MFI от Билла Вильямса
      Это новая статья из серии, в которой мы учимся создавать торговые системы на основе популярных технических индикаторов. В этой новой статье мы рассмотрим Индекс облегчения рынка (Market Facilitation Index, MFI), разработанный Биллом Вильямсом.
      Мультибот в MetaTrader: запуск множества роботов с одного графика Мультибот в MetaTrader: запуск множества роботов с одного графика
      В этой статье мы рассмотрим простой шаблон для создания универсального робота в MetaTrader, который можно использовать на нескольких графиках, но прицепив его лишь к одному графику, без необходимости настройки каждого экземпляра робота на каждом отдельном графике.