//+------------------------------------------------------------------+
//|                                                       CLFast.mqh |
//|                                             Copyright 2025, Leo. |
//|                     https://www.mql5.com/es/users/nique_372/news |
//+------------------------------------------------------------------+
#property copyright "Copyright 2025, Leo."
#property link      "https://www.mql5.com/es/users/nique_372/news"
#property strict

#ifndef CL_FAST_BY_LEO_MQH
#define CL_FAST_BY_LEO_MQH

//--- Incluimos el repo MQLArticles
#include "..\\MQLArticles\\Utils\\Funciones Array.mqh"

//+------------------------------------------------------------------+
//| Class for working with OpenCL                                    |
//+------------------------------------------------------------------+
class COpenCL : public CLoggerBase
 {
protected:
  //--- General
  int                m_context;
  int                m_program;

  //--- Kernel
  string             m_kernel_names[];
  int                m_kernels[];
  int                m_kernels_total;

  //--- Buffers
  int                m_buffers[];
  int                m_buffers_total;
  string             m_device_extensions;

  //---
  bool               m_support_cl_khr_fp64;
  bool               m_support_cl_khr_int64; //Soporta int64, long, ulong?

  //--- Custom
  void               CLLastError(const string &type_err, const string &funcion_Act);

  //--- Contexto \ Programa
  bool               m_context_create;
  bool               m_program_create;

  //--- Chekear validez [Contexto \ Programa]
  inline bool        CheckProgram(string funcion) const;
  inline bool        CheckContext(string funcion) const;

  //--- Chekear validez, y rango  [Kernels \ Buffers]
  inline bool        CheckKernel(const string &funcion, int kernel_index, bool check_is_not_invalid_handle) const;
  inline bool        CheckBuffer(const string &funcion, int buffer_index, bool check_is_not_invalid_handle) const;


public:
                     COpenCL(void);
                    ~COpenCL(void);

  //--- Get handles
  inline int         GetContext(void) const { return(m_context); } // Obtener contextos
  inline int         GetProgram(void) const { return(m_program); } // Obtener el programa
  int                GetKernel(const int kernel_index) const; // Obtener un handle
  string             GetKernelName(const int kernel_index) const; // Obtener el nombre de un kernel

  //--- Global and local memory size
  bool               GetGlobalMemorySize(long &global_memory_size);
  bool               GetLocalMemorySize(long &local_memory_size);

  //--- Maximal workgroup size
  bool               GetMaxWorkgroupSize(long &max_workgroup_size);

  //--- Check support
  inline bool        SupportDouble(void) const { return(m_support_cl_khr_fp64); }
  inline bool        SupportInt64(void) const { return(m_support_cl_khr_int64); }

  //--- Initialization and shutdown
  bool               Initialize(const string program, const bool show_log = true);
  void               Shutdown(void);

  bool               ContextCreate(const int device = CL_USE_ANY);
  void               ContextClean(void);

  bool               ProgramCreate(const string program, const bool show_log = true);
  void               ProgramDelete(void);

  //--- Set buffers/kernels count
  bool               SetBuffersCount(const int total_buffers);
  bool               SetKernelsCount(const int total_kernels);

  //--- Kernel operations
  bool               KernelCreate(const int kernel_index, const string kernel_name);
  bool               KernelFree(const int kernel_index);

  //--- Device and kernel info
  long               GetDeviceInfo(const int prop);
  long               GetDeviceInfoInteger(ENUM_OPENCL_PROPERTY_INTEGER prop);
  long               GetKernelInfoInteger(const int kernel_index, ENUM_OPENCL_PROPERTY_INTEGER prop);

  //--- Buffers
  //- Limpieza
  bool               BufferFree(const int buffer_index);

  //- Creacion
  bool               BufferCreate(const int buffer_index, const uint size_in_bytes, const uint flags);

  //- Creacion / Estritura
  template<typename T>
  bool               BufferFromArray(const int buffer_index, T &data[], const uint data_array_offset, const uint data_array_count, const uint flags);

  template<typename T>
  bool               BufferFromMatrix(const int buffer_index, matrix<T> &data, const uint flags);

  template<typename T>
  bool               BufferFromVector(const int buffer_index, vector<T> &data, const uint flags);

  //- Escritura
  template<typename T>
  inline bool        BufferWrite(const int buffer_index, T &data[], const uint cl_buffer_offset, const uint data_array_offset, const uint data_array_count);

  //- Lectura
  template<typename T>
  inline bool        BufferToMatrix(const int buffer_index, matrix<T> &data, const ulong rows = -1, const ulong cols = -1);

  template<typename T>
  inline bool        BufferToVector(const int buffer_index, vector<T> &data, const ulong size = -1);

  template<typename T>
  inline bool        BufferRead(const int buffer_index, T &data[], const uint cl_buffer_offset, const uint data_array_offset, const uint data_array_count);

  //--- Set kernel arguments
  template<typename T>
  inline bool        SetArgument(const int kernel_index, const int arg_index, T value);
  inline bool        SetArgumentBuffer(const int kernel_index, const int arg_index, const int buffer_index);
  bool               SetArgumentLocalMemory(const int kernel_index, const int arg_index, const int local_memory_size);

  //--- Kernel execution
  inline bool        Execute(const int kernel_index, const int work_dim, const uint &work_offset[], const uint &work_size[]);
  inline bool        Execute(const int kernel_index, const int work_dim, const uint &work_offset[], const uint &work_size[], const uint &local_work_size[]);
 };



//+------------------------------------------------------------------+
//| COpenCL class constructor                                        |
//+------------------------------------------------------------------+
COpenCL::COpenCL(void)
  : m_program_create(false), m_context_create(false)
 {
  m_context = INVALID_HANDLE;
  m_program = INVALID_HANDLE;
  m_buffers_total = 0;
  m_kernels_total = 0;
  m_device_extensions = "";
  m_support_cl_khr_fp64 = false;
 }
//+------------------------------------------------------------------+
//| COpenCL class destructor                                         |
//+------------------------------------------------------------------+
COpenCL::~COpenCL(void)
 {
  Shutdown();
 }

//+------------------------------------------------------------------+
//|                                                                  |
//+------------------------------------------------------------------+
void COpenCL::CLLastError(const string &type_err, const string &funcion_Act)
 {
  static const string separator = " | ";

  int error_code = (int)CLGetInfoInteger(m_context, CL_LAST_ERROR);
  string desc = "";

  if(!CLGetInfoString(m_context, CL_ERROR_DESCRIPTION, desc))
    desc = "No se pudo obtener la descripción del error OpenCL, código: " + (string)error_code;

  Print("[", type_err, "] ", funcion_Act, separator, desc, " (", error_code, ")");
 }

//+------------------------------------------------------------------+
//|                                                                  |
//+------------------------------------------------------------------+
inline bool COpenCL::CheckContext(string funcion) const
 {
  if(m_context_create)
    return true;

  LogFatalError("El contexto de OpenCL es invalido", funcion);
  Remover();
  return false;
 }

//+------------------------------------------------------------------+
//|                                                                  |
//+------------------------------------------------------------------+
inline bool COpenCL::CheckProgram(string funcion) const
 {
  if(m_program_create)
    return true;

  LogFatalError("El programa de OpenCL es invalido", funcion);
  Remover();
  return false;
 }

//+------------------------------------------------------------------+
//|                                                                  |
//+------------------------------------------------------------------+
inline bool COpenCL::CheckKernel(const string &funcion, int kernel_index, bool check_is_not_invalid_handle) const
 {
  if(kernel_index < 0 || kernel_index >= m_kernels_total)
   {
    LogError(StringFormat("Indice de kernel %d fuera de rango. Rango disponible: [0 - %d]", kernel_index, m_kernels_total), funcion);
    return false;
   }

  if(!check_is_not_invalid_handle)
    return true;

  if(m_kernels[kernel_index] == INVALID_HANDLE)
   {
    LogError(StringFormat("El handle del kernel[%d] - %s, es invalido", kernel_index, m_kernel_names[kernel_index]), funcion);
    return false;
   }

  return true;
 }


//+------------------------------------------------------------------+
//|                                                                  |
//+------------------------------------------------------------------+
inline bool COpenCL::CheckBuffer(const string &funcion, int buffer_index, bool check_is_not_invalid_handle) const
 {
  if(buffer_index < 0 || buffer_index >= m_buffers_total)
   {
    LogError(StringFormat("Indice de buffer %d fuera de rango. Rango disponible: [0 - %d]", buffer_index, m_buffers_total - 1), funcion);
    return false;
   }

  if(!check_is_not_invalid_handle)
    return true;

  if(m_buffers[buffer_index] == INVALID_HANDLE)
   {
    LogError(StringFormat("El handle del buffer %d, es invalido", buffer_index), funcion);
    return false;
   }

  return true;
 }

//+------------------------------------------------------------------+
//| GetKernel                                                        |
//+------------------------------------------------------------------+
int COpenCL::GetKernel(const int kernel_index) const
 {
//--- check parameters
  if(!CheckKernel(FUNCION_ACTUAL, kernel_index, true))
    return INVALID_HANDLE;
//---
  return m_kernels[kernel_index];
 }

//+------------------------------------------------------------------+
//| GetKernelName                                                    |
//+------------------------------------------------------------------+
string COpenCL::GetKernelName(const int kernel_index) const
 {
//--- check parameters
  if(!CheckKernel(FUNCION_ACTUAL, kernel_index, true))
    return ("");
//---
  return m_kernel_names[kernel_index];
 }

//+------------------------------------------------------------------+
//| GetGlobalMemorySize                                              |
//+------------------------------------------------------------------+
bool COpenCL::GetGlobalMemorySize(long &global_memory_size)
 {
//--- check parameters
  if(CheckContext(FUNCION_ACTUAL) == false)
    return false;

//--- get global memory size
  ResetLastError();
  global_memory_size = CLGetInfoInteger(m_context, CL_DEVICE_GLOBAL_MEM_SIZE);

  if(global_memory_size == -1)
   {
    CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
    return (false);
   }

//---
  return(true);
 }

//+------------------------------------------------------------------+
//| GetLocalMemorySize                                               |
//+------------------------------------------------------------------+
bool COpenCL::GetLocalMemorySize(long &local_memory_size)
 {
//--- check parameters
  if(CheckContext(FUNCION_ACTUAL) == false)
    return false;

//--- get local memory size
  ResetLastError();
  local_memory_size = CLGetInfoInteger(m_context, CL_DEVICE_LOCAL_MEM_SIZE);

  if(local_memory_size == -1)
   {
    CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
    return (false);
   }
//---
  return(true);
 }

//+------------------------------------------------------------------+
//| GetMaxWorkgroupSize                                              |
//+------------------------------------------------------------------+
bool COpenCL::GetMaxWorkgroupSize(long &max_workgroup_size)
 {
//--- check parameters
  if(CheckContext(FUNCION_ACTUAL) == false)
    return false;

//--- get maximal workgroup size
  max_workgroup_size = CLGetInfoInteger(m_context, CL_DEVICE_MAX_WORK_GROUP_SIZE);

  if(max_workgroup_size == -1)
   {
    CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
    return(false);
   }
//---
  return(true);
 }

//+------------------------------------------------------------------+
//| Initialize                                                       |
//+------------------------------------------------------------------+
bool COpenCL::Initialize(const string program, const bool show_log)
 {
//--- create context
  if(!ContextCreate(CL_USE_ANY))
    return(false);
//---
  return(ProgramCreate(program, show_log));
 }

//+------------------------------------------------------------------+
//| ContextCreate                                                    |
//+------------------------------------------------------------------+
bool COpenCL::ContextCreate(const int device)
 {
//--- remove context
  if(m_context != INVALID_HANDLE)
   {
    this.m_context_create = false;
    CLContextFree(m_context);
    m_context = INVALID_HANDLE;
    LogWarning("Se esta limpiando la instancia anterior del contexto CL, creando un nuevo contexto..", FUNCION_ACTUAL);
   }

//--- create context
  ResetLastError();
  if((m_context = CLContextCreate(device)) == INVALID_HANDLE)
   {
    LogFatalError(StringFormat("No se pudo crear un contexto CL, ultimo error: [%d], descripcion: ", GetLastError()), FUNCION_ACTUAL);
    CLLastError(FATAL_ERROR_TEXT, FUNCION_ACTUAL);
    return false;
   }

//--- check support working with doubles (cl_khr_fp64)
  m_support_cl_khr_fp64 = false;
  m_support_cl_khr_int64 = false;

  ResetLastError();
  if(CLGetInfoString(m_context, CL_DEVICE_EXTENSIONS, m_device_extensions))
   {
    string res[];
    StringSplit(m_device_extensions, ' ', res);

    for(int i = 0; i < ArraySize(res); i++)
     {
      if(!m_support_cl_khr_fp64 && res[i] == "cl_khr_fp64")
       {
        m_support_cl_khr_fp64 = true;
        LogWarning("Dispositivo soporta double", FUNCION_ACTUAL);
       }
      else
        if(!m_support_cl_khr_int64 && res[i] == "cl_khr_int64")
         {
          m_support_cl_khr_int64 = true;
          LogWarning("Dispositivo soporta ulong/long", FUNCION_ACTUAL);
         }


      if(m_support_cl_khr_fp64 || m_support_cl_khr_int64)
        break;
     }
   }
  else
   {
    LogError(StringFormat("No se pudieron obtener los divice para el contexto CL, ultimo erorr = %d, descripcion del error: ", GetLastError()), FUNCION_ACTUAL);
    CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
   }

  LogCaution("Contexto CL creado exitosamente", FUNCION_ACTUAL);
  this.m_context_create = true;
//---
  return(true);
 }

//+------------------------------------------------------------------+
//| ProgramCreate                                                    |
//+------------------------------------------------------------------+
bool COpenCL::ProgramCreate(const string program, const bool show_log)
 {
//--- check parameters
  if(CheckContext(FUNCION_ACTUAL) == false)
    return(false);

//--- remove program
  if(m_program != INVALID_HANDLE)
   {
    this.m_program_create = false;
    CLProgramFree(m_program);
    m_program = INVALID_HANDLE;
    LogWarning("Se esta liberando el programa anterior CL, creando otro programa CL...", FUNCION_ACTUAL);
   }

//--- compile the program
  string build_error_log;
  ResetLastError();

  if((m_program = CLProgramCreate(m_context, program, build_error_log)) == INVALID_HANDLE)
   {
    //--- show details
    if(show_log)
     {
      int    lines_count;
      string lines[];
      //---
      StringSplit(build_error_log, '\n', lines);

      lines_count = ArraySize(lines);
      for(int i = 0; i < lines_count; i++)
        Print(lines[i]);
     }

    //---
    CLContextFree(m_context);
    LogFatalError(StringFormat("No se pudo crear el programa cl, ultimo error = [%d], descripcion: ", GetLastError()), FUNCION_ACTUAL);
    CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
    return(false);
   }

  LogCaution("Programa CL creado exitosamente", FUNCION_ACTUAL);
  this.m_program_create = true;
//---
  return(true);
 }

//+------------------------------------------------------------------+
//| ProgramDelete                                                    |
//+------------------------------------------------------------------+
void COpenCL::ProgramDelete(void)
 {
//--- remove program
  if(m_program_create)
   {
    CLProgramFree(m_program);
    m_program = INVALID_HANDLE;
    m_program_create = false;
    LogWarning("El programa CL se ha eliminado con exito", FUNCION_ACTUAL);
   }
 }

//+------------------------------------------------------------------+
//| ContextClean                                                     |
//+------------------------------------------------------------------+
void COpenCL::ContextClean(void)
 {
//--- remove buffers
  if(m_buffers_total > 0)
   {
    for(int i = 0; i < m_buffers_total; i++)
      BufferFree(i);

    m_buffers_total = 0;
   }

//--- remove buffers
  if(m_kernels_total > 0)
   {
    for(int i = 0; i < m_kernels_total; i++)
      KernelFree(i);

    m_kernels_total = 0;
   }

//---
  ArrayFree(m_kernel_names);
  ArrayFree(m_kernels);
  ArrayFree(m_buffers);

//--- remove program
  if(m_program_create)
   {
    CLProgramFree(m_program);
    m_program = INVALID_HANDLE;
    m_program_create = false;
   }

  LogWarning("Programa CL limpiado exitosamente, numero de buffer, y kernels = 0", FUNCION_ACTUAL);
 }

//+------------------------------------------------------------------+
//| Shutdown                                                         |
//+------------------------------------------------------------------+
void COpenCL::Shutdown(void)
 {
  ContextClean();

//--- remove context
  if(m_context_create)
   {
    CLContextFree(m_context);
    m_context = INVALID_HANDLE;
    m_context_create = false;
   }

  LogWarning("Contexto CL limpiado exitosamente", FUNCION_ACTUAL);
 }

//+------------------------------------------------------------------+
//| SetBuffersCount                                                  |
//+------------------------------------------------------------------+
bool COpenCL::SetBuffersCount(const int total_buffers)
 {
//--- check parameters
  if(total_buffers <= 0)
   {
    LogError(StringFormat("El numero de buffers %d a establecer es invalido", total_buffers), FUNCION_ACTUAL);
    return(false);
   }

//---
  LogCaution(StringFormat("Numero de buffers %d, establecido correctamente", total_buffers), FUNCION_ACTUAL);

//---
  m_buffers_total = total_buffers;

  if(ArraySize(m_buffers) < m_buffers_total)
    ArrayResize(m_buffers, m_buffers_total);

  for(int i = 0; i < m_buffers_total; i++)
    m_buffers[i] = INVALID_HANDLE;

//---
  return(true);
 }
//+------------------------------------------------------------------+
//| SetKernelsCount                                                  |
//+------------------------------------------------------------------+
bool COpenCL::SetKernelsCount(const int total_kernels)
 {
//--- check parameters
  if(total_kernels <= 0)
   {
    LogError(StringFormat("El numero de kernels %d a establcer es invalido", total_kernels), FUNCION_ACTUAL);
    return(false);
   }

//---
  LogCaution(StringFormat("Numero de kernels %d establecido correctamente", total_kernels), FUNCION_ACTUAL);

//---
  m_kernels_total = total_kernels;

  if(ArraySize(m_kernels) < m_kernels_total)
    ArrayResize(m_kernels, m_kernels_total);

  if(ArraySize(m_kernel_names) < m_kernels_total)
    ArrayResize(m_kernel_names, m_kernels_total);

//---
  for(int i = 0; i < m_kernels_total; i++)
   {
    m_kernel_names[i] = "";
    m_kernels[i] = INVALID_HANDLE;
   }

//---
  return(true);
 }
//+------------------------------------------------------------------+
//| KernelCreate                                                     |
//+------------------------------------------------------------------+
bool COpenCL::KernelCreate(const int kernel_index, const string kernel_name)
 {
//--- check parameters
  if(!CheckContext(FUNCION_ACTUAL) || !CheckProgram(FUNCION_ACTUAL))
    return(false);

  if(!CheckKernel(FUNCION_ACTUAL, kernel_index, false)) //Aqui no es necesario verificar el kernel, dado que recien se crea
    return(false);

//---
  int kernel_handle = m_kernels[kernel_index];
  bool is_invalid = kernel_handle == INVALID_HANDLE;
  bool change_name = m_kernel_names[kernel_index] != kernel_name;

  if(is_invalid || change_name)
   {
    if(change_name && IsWarningLogEnabled() && is_invalid == false)
     {
      FastLog(FUNCION_ACTUAL, WARNING_TEXT, StringFormat("El kernel[%d], name = %s, se esta cambiando", kernel_index, m_kernel_names[kernel_index]));
      FastLog(FUNCION_ACTUAL, WARNING_TEXT, StringFormat("Nuevo kernel_name = %s", kernel_name));
     }

    ResetLastError();

    //--- Create kernel
    if((kernel_handle = CLKernelCreate(m_program, kernel_name)) == INVALID_HANDLE)
     {
      //--- cleanup
      CLProgramFree(m_program);
      m_program = INVALID_HANDLE;

      CLContextFree(m_context);
      m_context = INVALID_HANDLE;

      LogFatalError(StringFormat("No se pudo crear el kernel[%d] >> %s\nUltimo error =%d, descripion del error: ", kernel_index, kernel_name, GetLastError()), FUNCION_ACTUAL);
      CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
      return(false);
     }

    //---
    m_kernels[kernel_index] = kernel_handle;
    m_kernel_names[kernel_index] = kernel_name;
   }



  LogCaution(StringFormat("kernel[%d] - %s, creado existosamente", kernel_index, kernel_name), FUNCION_ACTUAL);
//---
  return(true);
 }
//+------------------------------------------------------------------+
//| KernelFree                                                       |
//+------------------------------------------------------------------+
bool COpenCL::KernelFree(const int kernel_index)
 {
//--- check kernel index
  if(CheckKernel(FUNCION_ACTUAL, kernel_index, false) == false)
    return(false);

  if(m_kernels[kernel_index] == INVALID_HANDLE)
    return (true);

//--- free kernel handle
  CLKernelFree(m_kernels[kernel_index]);
  m_kernels[kernel_index] = INVALID_HANDLE;
  m_kernel_names[kernel_index] = "";

//---
  return(true);
 }
//+------------------------------------------------------------------+
//| GetDeviceInfo                                                    |
//+------------------------------------------------------------------+
long COpenCL::GetDeviceInfo(const int prop)
 {
//--- check parameters
  if(CheckContext(FUNCION_ACTUAL) == false)
    return(-1);

//---
  uchar data[];
  uint  size = 0;

  if(!CLGetDeviceInfo(m_context, prop, data, size))
    return(-1);

  if(size < 4)
    return(-1);

//---
  union ___res_data
   {
    uchar cdata[8];
    long  ldata;
   } res;

  if(size <= 8)
   {
    ZeroMemory(res);
    ArrayCopy(res.cdata, data);
   }
  else
    ArrayCopy(res.cdata, data, 0, 8);

//---
  return(res.ldata);
 }
//+------------------------------------------------------------------+
//| GetDeviceInfoInteger                                             |
//+------------------------------------------------------------------+
long COpenCL::GetDeviceInfoInteger(ENUM_OPENCL_PROPERTY_INTEGER prop)
 {
//---
  return(CLGetInfoInteger(m_context, prop));
 }
//+------------------------------------------------------------------+
//| GetKernelInfoInteger                                             |
//+------------------------------------------------------------------+
long COpenCL::GetKernelInfoInteger(const int kernel_index, ENUM_OPENCL_PROPERTY_INTEGER prop)
 {
//---
  return(CLGetInfoInteger(m_kernels[kernel_index], prop));
 }
//+------------------------------------------------------------------+
//| BufferCreate                                                     |
//+------------------------------------------------------------------+
bool COpenCL::BufferCreate(const int buffer_index, const uint size_in_bytes, const uint flags)
 {
//---
  int buffer_handle = CLBufferCreate(m_context, size_in_bytes, flags);

  if(buffer_handle != INVALID_HANDLE)
   {
    m_buffers[buffer_index] = buffer_handle;
    return(true);
   }
//---
  return(false);
 }
//+------------------------------------------------------------------+
//| BufferFree                                                       |
//+------------------------------------------------------------------+
bool COpenCL::BufferFree(const int buffer_index)
 {
//--- check buffer index - handle
  if(!CheckBuffer(FUNCION_ACTUAL, buffer_index, false))
    return (false);

  if(m_buffers[buffer_index] == INVALID_HANDLE)
    return true;

//--- free buffer handle
  CLBufferFree(m_buffers[buffer_index]);
  m_buffers[buffer_index] = INVALID_HANDLE;

//---
  return(true);
 }
//+------------------------------------------------------------------+
//| BufferFromArray                                                  |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::BufferFromArray(const int buffer_index, T &data[], const uint data_array_offset, const uint data_array_count, const uint flags)
 {
//---
  if(m_buffers[buffer_index] != INVALID_HANDLE)
    return(CLBufferWrite(m_buffers[buffer_index], data, 0, data_array_offset, data_array_count) == data_array_count);

//---
  uint size_in_bytes = data_array_count * sizeof(T);
  int buffer_handle = CLBufferCreate(m_context, size_in_bytes, flags);

//---
  if(buffer_handle == INVALID_HANDLE)
    return(false);

  m_buffers[buffer_index] = buffer_handle;

//---
  return (CLBufferWrite(m_buffers[buffer_index], data, 0, data_array_offset, data_array_count) == data_array_count);
 }
//+------------------------------------------------------------------+
//| BufferWriteFromMatrix                                            |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::BufferFromMatrix(const int buffer_index, matrix<T> &data, const uint flags)
 {
//---
  if(m_buffers[buffer_index] != INVALID_HANDLE)
    return(CLBufferWrite(m_buffers[buffer_index], 0, data));

//---
  uint matrix_size  = uint(data.Rows() * data.Cols());
  uint size_in_bytes = matrix_size * sizeof(T);
  int buffer_handle = CLBufferCreate(m_context, size_in_bytes, flags);

//---
  if(buffer_handle == INVALID_HANDLE)
    return(false);

  m_buffers[buffer_index] = buffer_handle;
//---
  return (CLBufferWrite(m_buffers[buffer_index], 0, data));
 }
//+------------------------------------------------------------------+
//| BufferWriteFromVector                                            |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::BufferFromVector(const int buffer_index, vector<T> &data, const uint flags)
 {
//---
  if(m_buffers[buffer_index] != INVALID_HANDLE)
    return(CLBufferWrite(m_buffers[buffer_index], 0, data));

//---
  uint size_in_bytes = (uint)data.Size() * sizeof(T);
  int buffer_handle = CLBufferCreate(m_context, size_in_bytes, flags);

//---
  if(buffer_handle == INVALID_HANDLE)
    return(false);

  m_buffers[buffer_index] = buffer_handle;
//---
  return (CLBufferWrite(m_buffers[buffer_index], 0, data));
 }

//+------------------------------------------------------------------+
//| BufferReadToMatrix                                               |
//+------------------------------------------------------------------+
template<typename T>
inline bool COpenCL::BufferToMatrix(const int buffer_index, matrix<T> &data, const ulong rows, const ulong cols)
 {
//--- read data from OpenCL buffer
  return(CLBufferRead(m_buffers[buffer_index], 0, data, rows, cols));
 }

//+------------------------------------------------------------------+
//| BufferReadToVector                                               |
//+------------------------------------------------------------------+
template<typename T>
inline bool COpenCL::BufferToVector(const int buffer_index, vector<T> &data, const ulong size)
 {
  return(CLBufferRead(m_buffers[buffer_index], 0, data, size));
 }

//+------------------------------------------------------------------+
//| BufferRead                                                       |
//+------------------------------------------------------------------+
template<typename T>
inline bool COpenCL::BufferRead(const int buffer_index, T &data[], const uint cl_buffer_offset, const uint data_array_offset, const uint data_array_count)
 {
  return(CLBufferRead(m_buffers[buffer_index], data, cl_buffer_offset, data_array_offset, data_array_count) == data_array_count);
 }

//+------------------------------------------------------------------+
//| BufferWrite                                                      |
//+------------------------------------------------------------------+
template<typename T>
inline bool COpenCL::BufferWrite(const int buffer_index, T &data[], const uint cl_buffer_offset, const uint data_array_offset, const uint data_array_count)
 {
  return(CLBufferWrite(m_buffers[buffer_index], data, cl_buffer_offset, data_array_offset, data_array_count) == data_array_count);
 }

//+------------------------------------------------------------------+
//| SetArgument                                                      |
//+------------------------------------------------------------------+
template<typename T>
inline bool COpenCL::SetArgument(const int kernel_index, const int arg_index, T value)
 {
  return CLSetKernelArg(m_kernels[kernel_index], arg_index, value);
 }

//+------------------------------------------------------------------+
//| SetArgumentBuffer                                                |
//+------------------------------------------------------------------+
inline bool COpenCL::SetArgumentBuffer(const int kernel_index, const int arg_index, const int buffer_index)
 {
  return CLSetKernelArgMem(m_kernels[kernel_index], arg_index, m_buffers[buffer_index]);
 }

//+------------------------------------------------------------------+
//| SetArgumentLocalMemory                                           |
//+------------------------------------------------------------------+
bool COpenCL::SetArgumentLocalMemory(const int kernel_index, const int arg_index, const int local_memory_size)
 {
//--- check device local memory size
  long device_local_memory_size = CLGetInfoInteger(m_context, CL_DEVICE_LOCAL_MEM_SIZE);

  if(local_memory_size > device_local_memory_size)
    return(false);
//---
  return CLSetKernelArgMemLocal(m_kernels[kernel_index], arg_index, local_memory_size);
 }

//+------------------------------------------------------------------+
//| Execute                                                          |
//+------------------------------------------------------------------+
inline bool COpenCL::Execute(const int kernel_index, const int work_dim, const uint &work_offset[], const uint &work_size[])
 {
//---
  return CLExecute(m_kernels[kernel_index], work_dim, work_offset, work_size);
 }

//+------------------------------------------------------------------+
//| Execute                                                          |
//+------------------------------------------------------------------+
inline bool COpenCL::Execute(const int kernel_index, const int work_dim, const uint &work_offset[], const uint &work_size[], const uint &local_work_size[])
 {
//---
  return CLExecute(m_kernels[kernel_index], work_dim, work_offset, work_size, local_work_size);
 }
//+------------------------------------------------------------------+
#endif