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

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

MetaTrader 5 — Примеры | 11 мая 2023, 15:59
553 0
Wasin Thonkaew
Wasin Thonkaew

Содержание

Введение

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

Это подготовит нас к разработке полноценной поддержки OpenCL в виде DLL в следующей части.


Ключевые моменты

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

Ниже приведены ключевые положения статьи.

  1. Правильная передача строки из DLL в MQL5-программу. Нам нужно убедиться, что используется кодировка UTF-16, так как MetaTrader 5 использует ее для вывода через Print().
  2. Создание DLL, которую может использовать программа MQL5 в MetaTrader 5.
  3. Использование ключевых API-интерфейсов OpenCL C++ API в основном для получения информации о платформах/устройствах и выполнения функций ядра от инициализации до получения результата.


Реализация Simple DLL

Нам необходимо превратить код, написанный в предыдущей статье в простую библиотеку в виде DLL, которую мы можем использовать позже с MQL5.

Структура файла проекта выглядит следующим образом

Файл Определение
Makefile Кросс-компиляция как для Windows, так и для Linux с помощью Mingw64. Полученный файл openclsimple.dll копируется в папку Libraries/, которая используется для пути поиска DLL в MetaTrader 5.
Makefile-g++ Нативная компиляция Linux для тестирования.
openclsimple.h и openclsimple.cpp Основной заголовок и реализация библиотеки openclsimple DLL.
util.h и util.cpp Часть библиотеки openclsimple. Помимо прочего преобразовывает строки для библиотеки.
main.cpp Кроссплатформенная основная программа тестирования.
mql5/OpenCLSimple.mqh и mql5/TestCLSimple.mq5 Программа тестирования заголовков и скриптов MQL5 на MetaTrader 5


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

  1. clsimple_listall(char* out, int len, bool utf16=true)
    Перечисление всех платформ и устройств для важной информации

  2. clsimple_compute(const int arr_1[], const int arr_2[], int arr_3[], int num_elem)
    Вычисление суммы двух входных массивов и их вывод в выходной массив


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

openclsimple.h

#pragma once

#ifdef WINDOWS
        #ifdef CLSIMPLE_API_EXPORT
                #define CLSIMPLE_API __declspec(dllexport)
        #else
                #define CLSIMPLE_API __declspec(dllimport)
        #endif
#else
        #define CLSIMPLE_API
#endif

/**
 * We didn't define CL_HPP_ENABLE_EXCEPTIONS thus there would be no exceptions thrown
 * from any OpenCL related API.
 */
extern "C" {
        /**
         * List all platforms, and devices available.
         * If there any error occurs during the operation of this function, it will
         * print error onto standard error. The resultant text output is still maintained
         * separately.
         *
         * # Arguments
         * - out - output c-string to be filled
         * - len - length of output c-string to be filled
         * - utf16 - whether or not to convert string to UTF-16 encoding. По умолчанию параметр установлен в значение true.
         *                       If used on MetaTrader 5, this flag should be set to true.
         */
        CLSIMPLE_API void clsimple_listall(char* out, int len, bool utf16=true) noexcept;

        /**
         * Compute a summation of two input arrays then output into 3rd array limiting
         * by the number of elements specified.
         *
         * # Arguments
         * - arr_1 - first read-only array input holding integers
         * - arr_2 - second read-only array input holding integers
         * - arr_3 - output integer array to be filled with result of summation of both arr_1 and arr_2
         * - num_elem - number of element to be processed for both arr_1 and arr_2
         *
         * # Return
         * Returned code for result of operation. 0 means success, otherwise means failure.
         */
        CLSIMPLE_API [[nodiscard]] int clsimple_compute(const int arr_1[], const int arr_2[], int arr_3[], int num_elem) noexcept;
};


Думаю, раздел #ifdef знаком читателям, так как для Windows требуется экспортировать функции из DLL. Мы видим определения WINDOWS и CLSIMPLE_API_EXPORT, которые играют ключевую роль в явном экспорте каждой функции.

Раздел extern "C" обертывает общедоступный API для функций, которые могут быть вызваны программой.


openclsimple.cpp

#include "openclsimple.h"
#include "util.h"

#define CL_HPP_TARGET_OPENCL_VERSION 120
#define CL_HPP_MINIMUM_OPENCL_VERSION 120

#include <CL/cl2.hpp>

#include <iostream>
#include <vector>
#include <sstream>

#ifdef ENABLE_DEBUG
#include <cstdarg>
#endif

#ifdef ENABLE_DEBUG
const int LOG_BUFFER_SIZE = 2048;
char log_buffer[LOG_BUFFER_SIZE];

inline void DLOG(const char* ctx, const char* format, ...) {
        va_list args;
        va_start(args, format);
        std::vsnprintf(log_buffer, LOG_BUFFER_SIZE-1, format, args);
        va_end(args);

        std::cerr << "[DEBUG] [" << ctx << "] " << log_buffer << std::endl;
}
#else
        #define DLOG(...)
#endif

CLSIMPLE_API void clsimple_listall(char* out, int len, bool utf16) noexcept {
        // Get the platform
        std::vector<cl::Platform> platforms;
        int ret_code = cl::Platform::get(&platforms);
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Platform::get(), code=" << ret_code << std::endl;
                return;
        }

        std::stringstream output_str;

        for (size_t i=0; i<platforms.size(); ++i) {
                auto& p = platforms[i];

                std::string tmp_str;

                ret_code = p.getInfo(CL_PLATFORM_NAME, &tmp_str);
                if (ret_code != CL_SUCCESS)
                        std::cerr << "Error cl::Platform::getInfo(), code=" << ret_code << std::endl;
                else
                        output_str << "[" << i << "] Platform: " << tmp_str << std::endl;
                
                ret_code = p.getInfo(CL_PLATFORM_VENDOR, &tmp_str);
                if (ret_code != CL_SUCCESS)
                        std::cerr << "Error cl::Platform::getInfo(), code=" << ret_code << std::endl;
                else
                        output_str << "Vendor: " << tmp_str << std::endl;

                std::vector<cl::Device> devices;
                ret_code = p.getDevices(CL_DEVICE_TYPE_ALL, &devices);
                if (ret_code != CL_SUCCESS) {
                        std::cerr << "Error cl::Platform::getDevices(), code=" << ret_code << std::endl;
                        continue;
                }

                for (size_t j=0; j<devices.size(); ++j) {
                        const auto& d = devices[j];
                        cl_device_type tmp_device_type;

                        ret_code = d.getInfo(CL_DEVICE_NAME, &tmp_str);
                        if (ret_code != CL_SUCCESS)
                                std::cerr << "Error cl::Device::getInfo(), code=" << ret_code << std::endl;
                        else
                                output_str << " -[" << j << "] Device name: " << tmp_str << std::endl;

                        ret_code = d.getInfo(CL_DEVICE_TYPE, &tmp_device_type);
                        if (ret_code != CL_SUCCESS)
                                std::cerr << "Error cl::Device::getInfo(), code=" << ret_code << std::endl;
                        else {
                                if (tmp_device_type & CL_DEVICE_TYPE_GPU)
                                        output_str << " -Type: GPU" << std::endl;
                                else if (tmp_device_type & CL_DEVICE_TYPE_CPU)
                                        output_str << " -Type: CPU" << std::endl;
                                else if (tmp_device_type & CL_DEVICE_TYPE_ACCELERATOR)
                                        output_str << " -Type: Accelerator" << std::endl;
                                else
                                        output_str << " -Type: Unknown" << std::endl;
                        }
                }
        }

        // keep a copy of the string from stringstream
        std::string copy_str = output_str.str();
        if (utf16)
                util::str_to_cstr_u16(copy_str, out, len);
        else
                util::str_to_cstr(copy_str, out, len);
}

CLSIMPLE_API int clsimple_compute(const int arr_1[], const int arr_2[], int arr_3[], int num_elem) noexcept {
        cl_int ret_code = CL_SUCCESS;

        // Get the platform
        std::vector<cl::Platform> platforms;
        ret_code = cl::Platform::get(&platforms);
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Platform::get(), code=" << ret_code << std::endl;
                return ret_code;
        }
        DLOG(__FUNCTION__, "found %d platform(s)", platforms.size());

        if (platforms.empty()) {
                std::cerr << "Error found 0 platform." << std::endl;
                return CL_DEVICE_NOT_FOUND;             // reuse this error value
        }

        cl::Platform platform = platforms[0];
        DLOG(__FUNCTION__, "%s", "passed getting platforms");

        // Get the device
        std::vector<cl::Device> devices;
        ret_code = platform.getDevices(CL_DEVICE_TYPE_GPU, &devices);
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Platform::getDevices(), code=" << ret_code << std::endl;
                return ret_code;
        }
        DLOG(__FUNCTION__, "found %d GPU device(s)", devices.size());

        if (devices.empty()) {
                std::cerr << "Error found 0 device." << std::endl;
                return CL_DEVICE_NOT_FOUND;
        }
        cl::Device device = devices[0];
        DLOG(__FUNCTION__, "%s", "passed getting a GPU device");

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

        DLOG(__FUNCTION__, "%s", "passed creating a context");

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

        DLOG(__FUNCTION__, "%s", "passed creating command queue");

        // 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);
        ret_code = program.build({device});
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Program::build(), code=" << ret_code << std::endl;
                return ret_code;
        }

        DLOG(__FUNCTION__, "%s", "passed building a kernel program");

        cl::Kernel kernel(program, "add");

        DLOG(__FUNCTION__, "%s", "passed adding kernel function");

        // Create buffers
        cl::Buffer buffer_a(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(int) * num_elem, const_cast<int*>(arr_1));
        cl::Buffer buffer_b(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(int) * num_elem, const_cast<int*>(arr_2));
        cl::Buffer buffer_c(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(int) * num_elem, arr_3);

        kernel.setArg(0, buffer_a);
        kernel.setArg(1, buffer_b);
        kernel.setArg(2, buffer_c);
        kernel.setArg(3, num_elem);

        DLOG(__FUNCTION__, "%s", "passed setting all arguments");

        // execute the kernel function
        // NOTE: this is a blocking call although enqueuing is async call but the current thread
        // will be blocked until he work is done. Work is done doesn't mean that the result buffer
        // will be written back at the same time.
        //
        ret_code = queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(num_elem), cl::NullRange);
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::CommandQueue::enqueueNDRangeKernel(), code=" << ret_code << std::endl;
                return ret_code;
        }

        // CL_TRUE to make it blocking call
        // it requires for moving data from device back to host
        // NOTE: Important to call this function to make sure the result is sent back to host.
        ret_code = queue.enqueueReadBuffer(buffer_c, CL_TRUE, 0, sizeof(int) * num_elem, arr_3);
        if (ret_code != CL_SUCCESS)
                std::cerr << "Error cl::CommandQueue::enqueueReadBuffer(), code=" << ret_code << std::endl;

        return ret_code;
}


В коде выше есть три основные части, на которые следует обратить внимание.

  1. Утилита DLOG для ведения журнала отладки
  2. clsimple_listall()
    • util.h и util.cpp - утилита преобразования строк для передачи строки из DLL в MQL5
    1. clsimple_compute()

    DLOG

    #ifdef ENABLE_DEBUG
#include <cstdarg>
#endif

#ifdef ENABLE_DEBUG
const int LOG_BUFFER_SIZE = 2048;
char log_buffer[LOG_BUFFER_SIZE];

inline void DLOG(const char* ctx, const char* format, ...) {
        va_list args;
        va_start(args, format);
        std::vsnprintf(log_buffer, LOG_BUFFER_SIZE-1, format, args);
        va_end(args);

        std::cerr << "[DEBUG] [" << ctx << "] " << log_buffer << std::endl;
}
#else
        #define DLOG(...)
#endif

    Это утилита ведения журнала, которая обеспечивает стандартный вывод ошибок. Имеется охранное условие #ifdef для проверки наличия ENABLE_DEBUG при создании проекта. При наличии мы включаем требуемый заголовок, делая DLOG() пригодным к использованию.

    Буфер журнала устанавливается с фиксированным размером 2048 байт на вызов. Мы не ожидаем такого длинного отладочного сообщения, поэтому такого размера нам вполне достаточно.


    clsimple_listall()

    Функция для вывода списка всех устройств на всех платформах. Эти устройства доступны для использования с OpenCL.

    Всё начинается с cl::Platform для получения другой информации.

            ...
        // Get the platform
        std::vector<cl::Platform> platforms;
        int ret_code = cl::Platform::get(&platforms);
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Platform::get(), code=" << ret_code << std::endl;
                return;
        }
        ...

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

    Код возврата успеха для работы с OpenCL API - CL_SUCCESS. Обратите внимание, что в случае ошибки мы всегда будем выводить сообщение об ошибке в стандартную ошибку.

    Перебираем все платформы и устройства для получения необходимой нам информации

            ...
        std::stringstream output_str;

        for (size_t i=0; i<platforms.size(); ++i) {
                ...

                for (size_t j=0; j<devices.size(); ++j) {
                        ...
                }
        }
        ...

    Эта функция основана на записи вывода строки в обозначенный указатель c-строки. Другими словами, мы будем использовать std::stringstream, чтобы избежать создания временного std::string и копирования операции каждый раз, когда нам нужно добавить строку к текущему результату.

    cl::Platform является отправной точкой для получения другой информации. Каждая платформа содержит один или несколько cl::Device. Итак, у нас есть двойной цикл for для выполнения нашей работы.

    Внутри цикла

            ...
        for (size_t i=0; i<platforms.size(); ++i) {
                auto& p = platforms[i];

                // temporary variables to hold temporary platform/device informatin
                std::string tmp_str;

                ret_code = p.getInfo(CL_PLATFORM_NAME, &tmp_str);
                if (ret_code != CL_SUCCESS)
                        std::cerr << "Error cl::Platform::getInfo(), code=" << ret_code << std::endl;
                else
                        output_str << "[" << i << "] Platform: " << tmp_str << std::endl;
                
                ret_code = p.getInfo(CL_PLATFORM_VENDOR, &tmp_str);
                if (ret_code != CL_SUCCESS)
                        std::cerr << "Error cl::Platform::getInfo(), code=" << ret_code << std::endl;
                else
                        output_str << "Vendor: " << tmp_str << std::endl;

                std::vector<cl::Device> devices;
                ret_code = p.getDevices(CL_DEVICE_TYPE_ALL, &devices);
                if (ret_code != CL_SUCCESS) {
                        std::cerr << "Error cl::Platform::getDevices(), code=" << ret_code << std::endl;
                        continue;
                }

                for (size_t j=0; j<devices.size(); ++j) {
			cl_device_type tmp_device_type;
                        const auto& d = devices[j];

                        ret_code = d.getInfo(CL_DEVICE_NAME, &tmp_str);
                        if (ret_code != CL_SUCCESS)
                                std::cerr << "Error cl::Device::getInfo(), code=" << ret_code << std::endl;
                        else
                                output_str << " -[" << j << "] Device name: " << tmp_str << std::endl;

                        ret_code = d.getInfo(CL_DEVICE_TYPE, &tmp_device_type);
                        if (ret_code != CL_SUCCESS)
                                std::cerr << "Error cl::Device::getInfo(), code=" << ret_code << std::endl;
                        else {
                                if (tmp_device_type & CL_DEVICE_TYPE_GPU)
                                        output_str << " -Type: GPU" << std::endl;
                                else if (tmp_device_type & CL_DEVICE_TYPE_CPU)
                                        output_str << " -Type: CPU" << std::endl;
                                else if (tmp_device_type & CL_DEVICE_TYPE_ACCELERATOR)
                                        output_str << " -Type: Accelerator" << std::endl;
                                else
                                        output_str << " -Type: Unknown" << std::endl;
                        }
                }
        }
        ...

    Внутри цикла у нас есть временные переменные, а именно tmp_str и tmp_device_type для хранения временной информации, получаемой с платформы или устройства.
    Попутно, если что-то пойдет не так, мы выводим ошибку в стандартную ошибку, в противном случае добавляем результирующую строку в output_str.

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

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

    • CL_PLATFORM_NAME - наименование платформы
    • CL_PLATFORM_VENDOR - наименование поставщика, например AMD, Nvidia, pocl и т.д.
    • CL_DEVICE_NAME - устройство, например, кодовое имя графического процессора или наименование процессора
    • CL_DEVICE_TYPE - тип устройства, например графический процессор или CPU

    Есть довольно много информации, касающейся платформы и устройства. Разработчики могут заглянуть в заголовочный файл, а именно CL/CL2.h из пути включения системы /usr/include/. Ниже приведен соответствующий отрывок


    Информация о платформе

    ...
/* cl_platform_info */
#define CL_PLATFORM_PROFILE                         0x0900
#define CL_PLATFORM_VERSION                         0x0901
#define CL_PLATFORM_NAME                            0x0902
#define CL_PLATFORM_VENDOR                          0x0903
#define CL_PLATFORM_EXTENSIONS                      0x0904
#ifdef CL_VERSION_2_1
#define CL_PLATFORM_HOST_TIMER_RESOLUTION           0x0905
#endif
...

     

    Информация об устройстве

    ...
/* cl_device_info */
#define CL_DEVICE_TYPE                                   0x1000
#define CL_DEVICE_VENDOR_ID                              0x1001
#define CL_DEVICE_MAX_COMPUTE_UNITS                      0x1002
#define CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS               0x1003
#define CL_DEVICE_MAX_WORK_GROUP_SIZE                    0x1004
#define CL_DEVICE_MAX_WORK_ITEM_SIZES                    0x1005
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR            0x1006
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT           0x1007
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT             0x1008
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG            0x1009
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT           0x100A
#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE          0x100B
#define CL_DEVICE_MAX_CLOCK_FREQUENCY                    0x100C
#define CL_DEVICE_ADDRESS_BITS                           0x100D
#define CL_DEVICE_MAX_READ_IMAGE_ARGS                    0x100E
#define CL_DEVICE_MAX_WRITE_IMAGE_ARGS                   0x100F
#define CL_DEVICE_MAX_MEM_ALLOC_SIZE                     0x1010
#define CL_DEVICE_IMAGE2D_MAX_WIDTH                      0x1011
#define CL_DEVICE_IMAGE2D_MAX_HEIGHT                     0x1012
#define CL_DEVICE_IMAGE3D_MAX_WIDTH                      0x1013
#define CL_DEVICE_IMAGE3D_MAX_HEIGHT                     0x1014
#define CL_DEVICE_IMAGE3D_MAX_DEPTH                      0x1015
#define CL_DEVICE_IMAGE_SUPPORT                          0x1016
#define CL_DEVICE_MAX_PARAMETER_SIZE                     0x1017
#define CL_DEVICE_MAX_SAMPLERS                           0x1018
#define CL_DEVICE_MEM_BASE_ADDR_ALIGN                    0x1019
#define CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE               0x101A
#define CL_DEVICE_SINGLE_FP_CONFIG                       0x101B
#define CL_DEVICE_GLOBAL_MEM_CACHE_TYPE                  0x101C
#define CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE              0x101D
#define CL_DEVICE_GLOBAL_MEM_CACHE_SIZE                  0x101E
#define CL_DEVICE_GLOBAL_MEM_SIZE                        0x101F
#define CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE               0x1020
#define CL_DEVICE_MAX_CONSTANT_ARGS                      0x1021
#define CL_DEVICE_LOCAL_MEM_TYPE                         0x1022
#define CL_DEVICE_LOCAL_MEM_SIZE                         0x1023
#define CL_DEVICE_ERROR_CORRECTION_SUPPORT               0x1024
#define CL_DEVICE_PROFILING_TIMER_RESOLUTION             0x1025
#define CL_DEVICE_ENDIAN_LITTLE                          0x1026
#define CL_DEVICE_AVAILABLE                              0x1027
#define CL_DEVICE_COMPILER_AVAILABLE                     0x1028
#define CL_DEVICE_EXECUTION_CAPABILITIES                 0x1029
#define CL_DEVICE_QUEUE_PROPERTIES                       0x102A    /* deprecated */
#ifdef CL_VERSION_2_0
#define CL_DEVICE_QUEUE_ON_HOST_PROPERTIES               0x102A
#endif
#define CL_DEVICE_NAME                                   0x102B
#define CL_DEVICE_VENDOR                                 0x102C
#define CL_DRIVER_VERSION                                0x102D
#define CL_DEVICE_PROFILE                                0x102E
#define CL_DEVICE_VERSION                                0x102F
#define CL_DEVICE_EXTENSIONS                             0x1030
#define CL_DEVICE_PLATFORM                               0x1031
#ifdef CL_VERSION_1_2
#define CL_DEVICE_DOUBLE_FP_CONFIG                       0x1032
#endif
...

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


    Преобразование строки в UTF-16 для использования в MetaTrader 5

            ...
        // keep a copy of the string from stringstream
        std::string copy_str = output_str.str();
        if (utf16)
                util::str_to_cstr_u16(copy_str, out, len);
        else
                util::str_to_cstr(copy_str, out, len);

    Наконец, вывод строки из output_str необходимо преобразовать в кодировку UTF-16, поскольку MetaTrader 5 использует ее для отображения текста на вкладке "Эксперты".


    Теперь самое время посмотреть, как реализованы util::str_to_cstr и util::str_to_cstr_u16.
    util.h и util.cpp не предназначены для применения пользователями DLL. Файлы используются только внутри библиотеки. Таким образом, нет необходимости соответствовать C, то есть export "C", как и в случае с MetaTrader 5, когда пользователи применяют экспортированные функции из DLL.


    util.h

    #pragma once

#include <string>

namespace util {
        /**
         * Convert via copying from std::string to C-string.
         *
         * # Arguments
         * - str - input string
         * - out - destination c-string pointer to copy the content of string to
         * - len - length of string to copy from
         */
        void str_to_cstr(const std::string& str, char* out, unsigned len);

        /**
         * Convert via copying from std::string to UTF-16 string.
         *
         * # Arguments
         * - str - input string
         * - out - destination c-string pointer to copy the content of converted string
         *                 of UTF-16 to
         * - len - length of string to copy from
         */
        void str_to_cstr_u16(const std::string& str, char* out, unsigned len);
};

    Эти функции будут выполнять преобразование, если это необходимо, а затем копировать в целевую c-строку указатель с указанной длиной буфера.


    util.cpp

    #include "util.h"

#include <cuchar>
#include <locale>
#include <codecvt>
#include <cstring>

namespace util {
        /* converter of byte character to UTF-16 (2 bytes) */
        std::wstring_convert<std::codecvt_utf8<char16_t>, char16_t> ch_converter;

        void str_to_cstr(const std::string& str, char* out, unsigned len) {
                const char* str_cstr = str.c_str();
                const size_t capped_len = strlen(str_cstr) <= (len-1) ? strlen(str_cstr) : (len-1);
                std::memcpy(out, str_cstr, capped_len+1);
        }

        void str_to_cstr_u16(const std::string& str, char* out, unsigned len) {
                const char* str_cstr = str.c_str();

                std::u16string u16_str = ch_converter.from_bytes(str);
                const char16_t* u16_str_cstr = u16_str.c_str();
                
                const size_t capped_len = strlen(str_cstr) <= (len-1) ? strlen(str_cstr) : (len-1);
                std::memcpy(out, u16_str_cstr, capped_len*2+1);
        }
};

    Строки ограничивает длину строки, с которой нужно работать. Если длина строки меньше len, просто используйте длину строки. В противном случае используйте len-1.
    Мы вычитаем на 1, чтобы оставить место для символа с нулем в конце, который будет добавлен в следующей строке.

    Строка умножает на 2, поскольку UTF-16 имеет двойной размер обычной кодировки символов, используемой программой C++ (UTF-8).

    Строка создает конвертер, преобразовывающий

    • UTF-8 в UTF-16 с помощью функции std::wstring_convert::from_bytes()
    • UTF-16 в UTF-8 с помощью функции std::wstring_convert::to_bytes()

    Два аргумента шаблона для std::wstring_convert<_Codecvt, _Elem> можно описать следующим образом

    • _Codecvt - исходная кодировка символов для преобразования
    • _Elem - целевая кодировка символов

    Наконец, мы используем std::memcpy() для копирования потока байтов из преобразованной исходной строки в целевой указатель c-строки.


    Ниже приведены результаты, полученные на моей машине.
    Мы вернемся к этому в разделе тестирования.

    Пример clsimple_listall() на вкладке "Журнал"

    Пример вызова clsimple_listall() в MetaTrader 5


    clsimple_compute()

    Во-первых, посмотрим на сигнатуру функции.

    CLSIMPLE_API int clsimple_compute(const int arr_1[], const int arr_2[], int arr_3[], int num_elem) noexcept {
          ...
}


    Цель состоит в том, чтобы преобразовать код из предыдущей части серии в функцию. К счастью, конкретно в этом случае большую часть кода можно просто переместить в одну функцию.
    Полноценное преобразование оставим на потом. Пока мы проверяем правильность работы DLL в полном цикле с MQL5 в MetaTrader 5.

    Таким образом, код будет почти таким же.

    Функция clsimple_compute принимает следующие аргументы

    • arr_1 - массив целых чисел только для чтения
    • arr_2 - массив целых чисел только для чтения
    • arr_3 - выходной массив для суммирования arr_1 и arr_2
    • num_elem - количество элементов для обработки из обоих входных массивов

    Функция возвращает код возврата, а не результат суммирования обоих массивов.

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


    Получение платформы

            ...
        cl_int ret_code = CL_SUCCESS;

        // Get the platform
        std::vector<cl::Platform> platforms;
        ret_code = cl::Platform::get(&platforms);
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Platform::get(), code=" << ret_code << std::endl;
                return ret_code;
        }
        DLOG(__FUNCTION__, "found %d platform(s)", platforms.size());

        if (platforms.empty()) {
                std::cerr << "Error found 0 platform." << std::endl;
                return CL_DEVICE_NOT_FOUND;             // reuse this error value
        }

        cl::Platform platform = platforms[0];
        DLOG(__FUNCTION__, "%s", "passed getting platforms");
        ...


    По сравнению с предыдущей частью серии была улучшена обширная обработка ошибок. Сообщение об ошибке распечатывается вместе с возвращаемым кодом ошибки.
    Строки с DLOG() можно игнорировать, но они полезны, когда мы строим ENABLE_DEBUG в целях отладки.

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

    Получение устройства

            ...
        // Get the device
        std::vector<cl::Device> devices;
        ret_code = platform.getDevices(CL_DEVICE_TYPE_GPU, &devices);
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Platform::getDevices(), code=" << ret_code << std::endl;
                return ret_code;
        }
        DLOG(__FUNCTION__, "found %d GPU device(s)", devices.size());

        if (devices.empty()) {
                std::cerr << "Error found 0 device." << std::endl;
                return CL_DEVICE_NOT_FOUND;
        }
        cl::Device device = devices[0];
        DLOG(__FUNCTION__, "%s", "passed getting a GPU device");
        ...


    В этом случае мы ищем графический процессор только с такой платформы и используем первый найденный.


    Создание контекста

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

        DLOG(__FUNCTION__, "%s", "passed creating a context");
        ...


    Создание очереди команд

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

        DLOG(__FUNCTION__, "%s", "passed creating command queue");
        ...


    Создание ядра

            ...
        // 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);
        ret_code = program.build({device});
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::Program::build(), code=" << ret_code << std::endl;
                return ret_code;
        }

        DLOG(__FUNCTION__, "%s", "passed building a kernel program");

        cl::Kernel kernel(program, "add");

        DLOG(__FUNCTION__, "%s", "passed adding kernel function");
        ...
    Для создания функции ядра нам нужно создать cl::Program из исходной строки, в которой нам нужно построить cl::Program::Sources, а затем передать его как часть параметров конструктора cl::Kernel.


    Создание буферов

            ...
        // Create buffers
        cl::Buffer buffer_a(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(int) * num_elem, const_cast<int*>(arr_1));
        cl::Buffer buffer_b(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(int) * num_elem, const_cast<int*>(arr_2));
        cl::Buffer buffer_c(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(int) * num_elem, arr_3);
        ...

    Имеются три буфера.

    1. buffer_a - первый входной массив. Он доступен только для чтения и выделен на хосте, который также разрешает доступ с устройства.
    2. buffer_b - второй входной массив. Аналогичен первому.
    3. buffer_c - результирующий массив. Он доступен только для записи и выделен на хосте, который также разрешает доступ с устройства.

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

    Обратите внимание, что для arr_1, и arr_2 мы создаем const_cast<int*> для удаления const из переменной. Это нормально, так как мы получаем переменные const в функцию. Это гарантирует пользователям, что мы ничего не изменим для них.
    Но конструктор cl::Buffer требует передачи указателя определенного типа. При этом конструктор ничего не должен изменять.


    Установка аргумента в функцию ядра

            ...
        kernel.setArg(0, buffer_a);
        kernel.setArg(1, buffer_b);
        kernel.setArg(2, buffer_c);
        kernel.setArg(3, num_elem);

        DLOG(__FUNCTION__, "%s", "passed setting all arguments");
        ...

    Правильно задаем аргументы в соответствии с сигнатурой функции ядра, как показано в коде ядра OpenCL выше.


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

            ...
        // execute the kernel function
        // NOTE: this is a blocking call although enqueuing is async call but the current thread
        // will be blocked until he work is done. Work is done doesn't mean that the result buffer
        // will be written back at the same time.
        //
        ret_code = queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(num_elem), cl::NullRange);
        if (ret_code != CL_SUCCESS) {
                std::cerr << "Error cl::CommandQueue::enqueueNDRangeKernel(), code=" << ret_code << std::endl;
                return ret_code;
        }

        // CL_TRUE to make it blocking call
        // it requires for moving data from device back to host
        // NOTE: Important to call this function to make sure the result is sent back to host.
        ret_code = queue.enqueueReadBuffer(buffer_c, CL_TRUE, 0, sizeof(int) * num_elem, arr_3);
        if (ret_code != CL_SUCCESS)
                std::cerr << "Error cl::CommandQueue::enqueueReadBuffer(), code=" << ret_code << std::endl;

        return ret_code;

    Текст указывает глобальное измерение, которое будет использоваться для выполнения такого ядра. В данном случае это количество элементов входного массива. Укажем cl::NullRange для локального измерения, чтобы OpenCL автоматически определял значение для нас.

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


    Makefile

    .PHONY: all clean openclsimple.dll main.exe

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

all: openclsimple.dll main.exe
        cp -afv $< ~/.mt5/drive_c/Program\ Files/MetaTrader\ 5/MQL5/Libraries/

openclsimple.dll: util.o openclsimple.o
        @# check if symbolic link file to wine's opencl.dll exists, if not then create one
        test -h opencl.dll && echo "opencl.dll exists, no need to create symbolic link again" || ln -s ~/.mt5/drive_c/windows/system32/opencl.dll ./opencl.dll
        $(COMPILER) -shared $(FLAGS) $(MORE_FLAGS) -fPIC -o $@ $^ -L. -lopencl

openclsimple.o: openclsimple.cpp openclsimple.h
        $(COMPILER) $(FLAGS) $(MORE_FLAGS) -DCLSIMPLE_API_EXPORT -DWINDOWS -I. -fPIC -o $@ -c $<

util.o: util.cpp util.h
        $(COMPILER) $(FLAGS) $(MORE_FLAGS) -I. -fPIC -o $@ -c $<

main.exe: main.cpp openclsimple.dll
        $(COMPILER) $(FLAGS) $(MORE_FLAGS) -I. -DWINDOWS -o $@ $< -L. -lopenclsimple

clean:
        rm -f openclsimple.dll main.exe opencl.dll util.o openclsimple.o

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

    Строка указывает на улучшение по сравнению с Makefile из предыдущей части статьи. Теперь вместо создания файла символической ссылки, указывающего на opencl.dll в виде префикса wine (место установки MetaTrader 5 по адресу ~/.mt5 с другим именем пользователя как части пути к домашнему каталогу), мы динамически и заново создаем файл символической ссылки. Таким образом файл символической ссылки будет указывать на правильный путь в соответствии с их именем пользователя и домашним каталогом без необходимости перезаписывать путь, указанный файлом символической ссылки, который мы упаковали и доставили пользователю.

    Строка указывает, что мы копируем полученный файл DLL, а именно openclsimple.dll в Libraries/, который будет использоваться MetaTrader 5 для поиска DLL во время выполнения. Это экономит нам массу времени во время разработки.


    Makefile-g++

    .PHONY: all clean libopenclsimple.so main.exe

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

all: libopenclsimple.so main.out

libopenclsimple.so: util.o openclsimple.o
        @# NOTE: notice capital letters in -lOpenCL
        $(COMPILER) -shared $(FLAGS) $(MORE_FLAGS) -I. -fPIC -o $@ $^ -lOpenCL

openclsimple.o: openclsimple.cpp openclsimple.h util.h
        $(COMPILER) $(FLAGS) $(MORE_FLAGS) -I. -fPIC -o $@ -c $<

util.o: util.cpp util.h
        $(COMPILER) $(FLAGS) $(MORE_FLAGS) -I. -fPIC -o $@ -c $<

main.out: main.cpp libopenclsimple.so
        $(COMPILER) $(FLAGS) $(MORE_FLAGS) -I. -o $@ $< -L. -lopenclsimple

clean:
        rm -f libopenclsimple.so main.out util.o openclsimple.o

    Аналогично для Makefile-g++, который предназначен для использования изначально в системе Linux для быстрого тестирования. Содержание аналогичное, заметная разница заключается в том, что мы будем связываться с библиотекой OpenCL, установленной в системе. Ее имя отлично от библиотеки Windows.


    Тестирование в Linux и Windows (через Wine)

    У нас есть готовая система сборки. Все готово как минимум для тестирования нативно на Linux и на Windows (через Wine).


    Linux

    Выполните следующую команду

    make -f Makefile-g++

    У нас будут следующие выходные файлы

    • libopenclsimple.so
    • main.out


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

    ./main.out


    Мы увидим результат, похожий на следующий

    Результат тестирования в Linux тестовой программы, созданной с помощью Makefile-g++

    Результат тестирования main.out в Linux, собранного из Makefile-g++

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


    Windows (через Wine)

    Выполните следующую команду

    make

    У нас будут следующие выходные файлы

    • openclsimple.dll
    • main.exe


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

    WINEPREFIX=~/.mt5 wine ./main.exe


    Мы увидим результат, похожий на следующий

    Результат тестирования программы в Windows (через Wine)

    Результат тестирования main.exe для Windows (через Wine)

    Мы всегда используем WINEPREFIX=~/.mt5, так как это префикс в случае, когда MetaTrader 5 установлен по умолчанию. Мы тестируем в той же среде, в которой будет работать MetaTrader 5.
    Тот же вывод, что и ранее протестированный в Linux.


    Вы можете дополнительно брать выходные файлы, созданные с помощью Makefile, для нативного тестирования в Windows. Результаты будут аналогичными.


    Тестирование с помощью MetaTrader 5

    Теперь всё готов для тестов в MQL5.


    mql5/OpenCLSimple.mqh

    //+------------------------------------------------------------------+
//|                                                      OpenCLX.mqh |
//|                                          Copyright 2022, haxpor. |
//|                                                 https://wasin.io |
//+------------------------------------------------------------------+
#property copyright "Copyright 2022, haxpor."
#property link      "https://wasin.io"

#import "openclsimple.dll"
void clsimple_listall(string& out, int len);
int clsimple_compute(const int& arr_1[], const int& arr_2[], int& arr_3[], int num_elem);
#import


    Обратите внимание на подсвеченный текст из сигнатуры функции clsimple_listall(). Как видно из DLL, сама функция имеет три аргумента

    CLSIMPLE_API void clsimple_listall(char* out, int len, bool utf16=true) noexcept;

    Нам не нужно включать аргумент utf16 в файле .mqh, поскольку мы всегда устанавливаем для этого аргумента значение true, так как нам нужно преобразовать строку в UTF-16, чтобы ее можно было напечатать на вкладке "Советники" в MetaTrader 5.
    Достаточно определить только первые два параметра.


    mql5/TestCLSimple.mq5

    //+------------------------------------------------------------------+
//|                                                  TestOpenCLX.mq5 |
//|                                          Copyright 2022, haxpor. |
//|                                                 https://wasin.io |
//+------------------------------------------------------------------+
#property copyright "Copyright 2022, haxpor."
#property link      "https://wasin.io"
#property version   "1.00"

#include "OpenCLSimple.mqh"

#define STR_BUFFER_LEN 2048

//+------------------------------------------------------------------+
//| Script program start function                                    |
//+------------------------------------------------------------------+
void OnStart()
{
        // 1: test clsimple_listall()
        // construct a string to hold resultant of platforms/devices listing
        string listall_str;
        StringInit(listall_str, STR_BUFFER_LEN, '\0');

        // get platforms/devices and print the result
        clsimple_listall(listall_str, STR_BUFFER_LEN);
        Print(listall_str);

        // 2: test clsimple_compute()
        int arr_1[];
        int arr_2[];
        int arr_3[];

        ArrayResize(arr_1, 10000000);
        ArrayFill(arr_1, 0, ArraySize(arr_1), 1);

        ArrayResize(arr_2, 10000000);
        ArrayFill(arr_2, 0, ArraySize(arr_2), 1);

        ArrayResize(arr_3, 10000000);

        uint start_time = GetTickCount();

        int ret_code = clsimple_compute(arr_1, arr_2, arr_3, ArraySize(arr_1));
        if (ret_code != 0) {
                Print("Error occurs, code=", ret_code);
                return;
        }
        Print("Elapsed time: " + (string)(GetTickCount() - start_time) + " ms");

        bool is_valid = true;
        for (int i=0; i<ArraySize(arr_3); ++i) {
                if (arr_3[i] != 2) {
                        Print("Something is wrong at index=" + (string)i);
                        is_valid = false;
                }
        }

        if (is_valid) {
                Print("Passed test");
        }
}

    Для получения строкового вывода из DLL (возвращаемого в виде указателя c-строки путем копирования его буфера) нам нужно определить переменную string и инициализировать ее потенциал для максимальной длины, которую мы будем поддерживать.

    Подготовка к вызову clsimple_compute() потребует немного больше усилий. Нам нужно объявить массивы ввода целых чисел, заполнить их правильными значениями и объявить массив целых чисел, используемый для вывода. В любом случае, на самом деле мы будем считывать такие входные данные тик за тиком из цены актива. К тому же, нам нужно очистить или подготовить данные, прежде чем предоставлять их как часть аргументов всякий раз, когда мы вызываем clsimple_compute().

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

    Passed test

    Поместите .mqh туда же, где находится .mq5 или в директорию Includes/ пути установки MetaTrader 5. Затем скомпилируйте .mq5 и наконец перетащите созданную программу на график в MetaTrader 5.
    Мы увидим следующий результат на вкладке "Эксперты".

    Результаты тестирования MQL5-программы в MetaTrader 5 на вкладке "Эксперты"

    Результаты тестирования MQL5-программы (скрипта) в MetaTrader 5.
    Текст отображается правильно благодаря нашей рабочей утилите преобразования строк



    Загрузка исходного кода

    Вы можете загрузить исходный код из zip-файла в самом низу этой статьи или в репозитории Github по адресуgithub.com/haxpor/opencl-simple(откройте каталог simple/, каталог standalonetest/ относится к предыдущей части серии).



    Что дальше?

    В этой статье мы преобразовали код из предыдущей части в DLL-библиотеку, которая может использоваться как обычной программой C++ в Linux и Windows (через Wine или нативно), так и MQL5 в MetaTrader 5.

    Также мы рассмотрели, как правильно передать строку из DLL в программу MQL5, поскольку нам нужно преобразовать ее в кодировку UTF-16, используемую самим MetaTrader 5 для отображения, по крайней мере, на вкладке "Эксперты". Если строка, отображаемая в MetaTrader 5, верна, значит, мы все сделали правильно.

    В следующей части серии мы углубимся в OpenCL C++ API, чтобы разработать полноценную функцию поддержки OpenCL в виде DLL для использования с MQL5.

    Также мы рассмотрим требования для максимально эффективной работы с OpenCL API и используем эти знания для разработки высокопроизводительной программы MQL5 на основе OpenCL.

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

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

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

    Перейти к обсуждению на форуме трейдеров
    Оборачиваем ONNX-модели в классы Оборачиваем ONNX-модели в классы
    Объектно-ориентированное программирование позволяет создавать более компактный код, который легко читать и модифицировать. Представляем пример для трёх ONNX-моделей.
    Машинное обучение и Data Science (Часть 13): Анализируем финансовый рынок с помощью метода главных компонент (PCA) Машинное обучение и Data Science (Часть 13): Анализируем финансовый рынок с помощью метода главных компонент (PCA)
    Попробуем качественно улучшить анализ финансовых рынков с помощью метода главных компонент (Principal Component Analysis, PCA). Узнаем, как этот метод может помочь выявлять скрытые закономерности в данных, определять скрытые рыночные тенденции и оптимизировать инвестиционные стратегии. В этой статье мы посмотрим, как метод PCA дает новую перспективу для анализа сложных финансовых данных, помогая увидеть идеи, которые мы упустили при использовании традиционных подходов. Дает ли применение метода PCA на данных финансовых рынков конкурентное преимущество и поможет ли быть на шаг впереди?
    Нейросети — это просто (Часть 41): Иерархические модели Нейросети — это просто (Часть 41): Иерархические модели
    Статья описывает иерархические модели обучения, которые предлагают эффективный подход к решению сложных задач машинного обучения. Иерархические модели состоят из нескольких уровней, каждый из которых отвечает за различные аспекты задачи.
    Машинное обучение и Data Science (Часть 12): Можно ли выигрывать на рынке с помощью самообучающихся нейронных сетей? Машинное обучение и Data Science (Часть 12): Можно ли выигрывать на рынке с помощью самообучающихся нейронных сетей?
    Наверняка многим надоели постоянные попытки предсказать фондовый рынок. Хотели бы вы иметь хрустальный шар, который бы помогал принимать более обоснованные инвестиционные решения? Самообучающиеся нейронные сети могут стать таким решением. В этой статье мы посмотрим, могут ли такие мощные алгоритмы помочь «оседлать волну» и перехитрить фондовый рынок. Анализируя огромные объемы данных и выявляя закономерности, самообучающиеся нейронные сети могут делать прогнозы, которые зачастую более точны, чем прогнозы от трейдеров. Давайте посмотрим, можно ли использовать эти передовые технологии, чтобы принимать разумные инвестиционные решения и зарабатывать больше.