//+------------------------------------------------------------------+
//|                                                       OpenCL.mqh |
//|                             Copyright 2000-2025, MetaQuotes Ltd. |
//|                                             https://www.mql5.com |
//+------------------------------------------------------------------+

//+------------------------------------------------------------------+
//| Class for working with OpenCL                                    |
//+------------------------------------------------------------------+
class COpenCL
  {
protected:
   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;

public:
                     COpenCL(void);
                    ~COpenCL(void);
   //--- get handles
   int               GetContext(void) const { return(m_context); }
   int               GetProgram(void) const { return(m_program); }
   int               GetKernel(const int kernel_index) const;
   string            GetKernelName(const int kernel_index) const;
   //--- 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 working with double
   bool              SupportDouble(void) const { return(m_support_cl_khr_fp64); }
   //--- 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);
   //--- buffer operations
   bool              BufferCreate(const int buffer_index,const uint size_in_bytes,const uint flags);
   bool              BufferFree(const int buffer_index);
   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);
   template<typename T>
   bool              BufferToMatrix(const int buffer_index,matrix<T> &data,const ulong rows=-1,const ulong cols=-1);
   template<typename T>
   bool              BufferToVector(const int buffer_index,vector<T> &data,const ulong size=-1);
   template<typename T>
   bool              BufferRead(const int buffer_index,T &data[],const uint cl_buffer_offset,const uint data_array_offset,const uint data_array_count);
   template<typename T>
   bool              BufferWrite(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>
   bool              SetArgument(const int kernel_index,const int arg_index,T value);
   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
   bool              Execute(const int kernel_index,const int work_dim,const uint &work_offset[],const uint &work_size[]);
   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_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();
  }
//+------------------------------------------------------------------+
//| GetKernel                                                        |
//+------------------------------------------------------------------+
int COpenCL::GetKernel(const int kernel_index) const
  {
//--- check parameters
   if(m_kernels_total<=0 || kernel_index<0 || kernel_index>=m_kernels_total)
      return(INVALID_HANDLE);
//---
   return m_kernels[kernel_index];
  }
//+------------------------------------------------------------------+
//| GetKernelName                                                    |
//+------------------------------------------------------------------+
string COpenCL::GetKernelName(const int kernel_index) const
  {
//--- check parameters
   if(m_kernels_total<=0 || kernel_index<0 || kernel_index>=m_kernels_total)
      return("");
//---
   return m_kernel_names[kernel_index];
  }
//+------------------------------------------------------------------+
//| GetGlobalMemorySize                                              |
//+------------------------------------------------------------------+
bool COpenCL::GetGlobalMemorySize(long &global_memory_size)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      return(false);
//--- get global memory size
   global_memory_size=CLGetInfoInteger(m_context,CL_DEVICE_GLOBAL_MEM_SIZE);

   if(global_memory_size==-1)
      return(false);
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| GetLocalMemorySize                                               |
//+------------------------------------------------------------------+
bool COpenCL::GetLocalMemorySize(long &local_memory_size)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      return(false);
//--- get local memory size
   local_memory_size=CLGetInfoInteger(m_context,CL_DEVICE_LOCAL_MEM_SIZE);

   if(local_memory_size==-1)
      return(false);
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| GetMaxWorkgroupSize                                              |
//+------------------------------------------------------------------+
bool COpenCL::GetMaxWorkgroupSize(long &max_workgroup_size)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      return(false);
//--- get maximal workgroup size
   max_workgroup_size=CLGetInfoInteger(m_context,CL_DEVICE_MAX_WORK_GROUP_SIZE);
   if(max_workgroup_size==-1)
      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)
     {
      CLContextFree(m_context);
      m_context=INVALID_HANDLE;
     }
//--- create context
   if((m_context=CLContextCreate(device))==INVALID_HANDLE)
     {
      Print("OpenCL not found, error code=",GetLastError());
      return(false);
     }
//--- check support working with doubles (cl_khr_fp64)
   m_support_cl_khr_fp64=false;
   if(CLGetInfoString(m_context,CL_DEVICE_EXTENSIONS,m_device_extensions))
     {
      int    size;
      string extenstions[];
      //---
      StringSplit(m_device_extensions,' ',extenstions);

      size=ArraySize(extenstions);
      for(int i=0; i<size; i++)
         if(extenstions[i]=="cl_khr_fp64")
           {
            m_support_cl_khr_fp64=true;
            break;
           }
     }
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| ProgramCreate                                                    |
//+------------------------------------------------------------------+
bool COpenCL::ProgramCreate(const string program,const bool show_log)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      return(false);
//--- remove program
   if(m_program!=INVALID_HANDLE)
     {
      CLProgramFree(m_program);
      m_program=INVALID_HANDLE;
     }
//--- compile the program
   string build_error_log;

   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);
      Print("OpenCL program create failed, error code=",GetLastError());
      return(false);
     }
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| ProgramDelete                                                    |
//+------------------------------------------------------------------+
void COpenCL::ProgramDelete(void)
  {
//--- remove program
   if(m_program!=INVALID_HANDLE)
     {
      CLProgramFree(m_program);
      m_program=INVALID_HANDLE;
     }
  }
//+------------------------------------------------------------------+
//| 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;
     }
//--- remove program
   if(m_program!=INVALID_HANDLE)
     {
      CLProgramFree(m_program);
      m_program=INVALID_HANDLE;
     }
  }
//+------------------------------------------------------------------+
//| Shutdown                                                         |
//+------------------------------------------------------------------+
void COpenCL::Shutdown(void)
  {
   ContextClean();
//--- remove context
   if(m_context!=INVALID_HANDLE)
     {
      CLContextFree(m_context);
      m_context=INVALID_HANDLE;
     }
  }
//+------------------------------------------------------------------+
//| SetBuffersCount                                                  |
//+------------------------------------------------------------------+
bool COpenCL::SetBuffersCount(const int total_buffers)
  {
//--- check parameters
   if(total_buffers<=0)
      return(false);
//---
   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)
      return(false);
//---
   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(m_context==INVALID_HANDLE || m_program==INVALID_HANDLE)
      return(false);

   if(kernel_index<0 || kernel_index>=m_kernels_total)
      return(false);
//---
   int kernel_handle=m_kernels[kernel_index];

   if(kernel_handle==INVALID_HANDLE || m_kernel_names[kernel_index]!=kernel_name)
     {
      //--- 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;

         Print("OpenCL kernel create failed, error code=",GetLastError());
         return(false);
        }
      //---
      m_kernels[kernel_index]=kernel_handle;
      m_kernel_names[kernel_index]=kernel_name;
     }
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| KernelFree                                                       |
//+------------------------------------------------------------------+
bool COpenCL::KernelFree(const int kernel_index)
  {
//--- check kernel index
   if(kernel_index<0 || kernel_index>=m_kernels_total)
      return(false);

   if(m_kernels[kernel_index]==INVALID_HANDLE)
      return(false);
//--- free kernel handle
   CLKernelFree(m_kernels[kernel_index]);
   m_kernels[kernel_index]=INVALID_HANDLE;
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| GetDeviceInfo                                                    |
//+------------------------------------------------------------------+
long COpenCL::GetDeviceInfo(const int prop)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      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)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      return(-1);
//---
   return(CLGetInfoInteger(m_context,prop));
  }
//+------------------------------------------------------------------+
//| GetKernelInfoInteger                                             |
//+------------------------------------------------------------------+
long COpenCL::GetKernelInfoInteger(const int kernel_index,ENUM_OPENCL_PROPERTY_INTEGER prop)
  {
//--- check parameters
   if(kernel_index<0 || kernel_index>=m_kernels_total)
      return(-1);
//---
   return(CLGetInfoInteger(m_kernels[kernel_index],prop));
  }
//+------------------------------------------------------------------+
//| BufferCreate                                                     |
//+------------------------------------------------------------------+
bool COpenCL::BufferCreate(const int buffer_index,const uint size_in_bytes,const uint flags)
  {
//--- check parameters
   if(buffer_index<0 || buffer_index>=m_buffers_total)
      return(false);

   if(m_context==INVALID_HANDLE)
      return(false);
//---
   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
   if(buffer_index<0 || buffer_index>=m_buffers_total)
      return(false);

   if(m_buffers[buffer_index]==INVALID_HANDLE)
      return(false);
//--- 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)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      return(false);

   if(buffer_index<0 || buffer_index>=m_buffers_total || data_array_count<=0)
      return(false);
//--- buffer does not exists, create it
   if(m_buffers[buffer_index]==INVALID_HANDLE)
     {
      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;
     }
//--- write data to OpenCL buffer
   uint data_written=CLBufferWrite(m_buffers[buffer_index],data,0,data_array_offset,data_array_count);

   if(data_written!=data_array_count)
      return(false);
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| BufferWriteFromMatrix                                            |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::BufferFromMatrix(const int buffer_index,matrix<T> &data,const uint flags)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      return(false);

   if(buffer_index<0 || buffer_index>=m_buffers_total || data.Rows()==0 || data.Cols()==0)
      return(false);
//--- buffer does not exists, create it
   if(m_buffers[buffer_index]==INVALID_HANDLE)
     {
      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;
     }
//--- write data to OpenCL buffer
   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)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE)
      return(false);

   if(buffer_index<0 || buffer_index>=m_buffers_total || data.Size()==0)
      return(false);
//--- buffer does not exists, create it
   if(m_buffers[buffer_index]==INVALID_HANDLE)
     {
      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;
     }
//--- write data to OpenCL buffer
   return(CLBufferWrite(m_buffers[buffer_index],0,data));
  }
//+------------------------------------------------------------------+
//| BufferReadToMatrix                                               |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::BufferToMatrix(const int buffer_index,matrix<T> &data,const ulong rows,const ulong cols)
  {
//--- check parameters
   if(buffer_index<0 || buffer_index>=m_buffers_total)
      return(false);

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

   if(m_context==INVALID_HANDLE || m_program==INVALID_HANDLE)
      return(false);
//--- read data from OpenCL buffer
   return(CLBufferRead(m_buffers[buffer_index],0,data,rows,cols));
  }
//+------------------------------------------------------------------+
//| BufferReadToVector                                               |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::BufferToVector(const int buffer_index,vector<T> &data,const ulong size)
  {
//--- check parameters
   if(buffer_index<0 || buffer_index>=m_buffers_total)
      return(false);

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

   if(m_context==INVALID_HANDLE || m_program==INVALID_HANDLE)
      return(false);
//--- read data from OpenCL buffer
   return(CLBufferRead(m_buffers[buffer_index],0,data,size));
  }
//+------------------------------------------------------------------+
//| BufferRead                                                       |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::BufferRead(const int buffer_index,T &data[],const uint cl_buffer_offset,const uint data_array_offset,const uint data_array_count)
  {
//--- check parameters
   if(buffer_index<0 || buffer_index>=m_buffers_total || data_array_count<=0)
      return(false);

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

   if(m_context==INVALID_HANDLE || m_program==INVALID_HANDLE)
      return(false);
//--- read data from OpenCL buffer
   uint data_read=CLBufferRead(m_buffers[buffer_index],data,cl_buffer_offset,data_array_offset,data_array_count);

   if(data_read!=data_array_count)
      return(false);
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| BufferWrite                                                      |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::BufferWrite(const int buffer_index,T &data[],const uint cl_buffer_offset,const uint data_array_offset,const uint data_array_count)
  {
//--- check parameters
   if(buffer_index<0 || buffer_index>=m_buffers_total || data_array_count<=0)
      return(false);

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

   if(m_context==INVALID_HANDLE)
      return(false);
//--- write data to OpenCL buffer
   uint data_written=CLBufferWrite(m_buffers[buffer_index],data,cl_buffer_offset,data_array_offset,data_array_count);

   if(data_written!=data_array_count)
      return(false);
//---
   return(true);
  }
//+------------------------------------------------------------------+
//| SetArgument                                                      |
//+------------------------------------------------------------------+
template<typename T>
bool COpenCL::SetArgument(const int kernel_index,const int arg_index,T value)
  {
//--- check parameters
   if(kernel_index<0 || kernel_index>=m_kernels_total)
      return(false);
//---
   int kernel_handle=m_kernels[kernel_index];

   if(kernel_handle==INVALID_HANDLE)
      return(false);
//---
   return CLSetKernelArg(kernel_handle,arg_index,value);
  }
//+------------------------------------------------------------------+
//| SetArgumentBuffer                                                |
//+------------------------------------------------------------------+
bool COpenCL::SetArgumentBuffer(const int kernel_index,const int arg_index,const int buffer_index)
  {
//--- check parameters
   if(m_context==INVALID_HANDLE || m_program==INVALID_HANDLE)
      return(false);

   if(kernel_index<0 || kernel_index>=m_kernels_total)
      return(false);

   if(buffer_index<0 || buffer_index>=m_buffers_total)
      return(false);

   if(m_buffers[buffer_index]==INVALID_HANDLE)
      return(false);
//---
   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 parameters
   if(m_context==INVALID_HANDLE || m_program==INVALID_HANDLE)
      return(false);

   if(kernel_index<0 || kernel_index>=m_kernels_total)
      return(false);
//--- 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                                                          |
//+------------------------------------------------------------------+
bool COpenCL::Execute(const int kernel_index,const int work_dim,const uint &work_offset[],const uint &work_size[])
  {
//--- check parameters
   if(kernel_index<0 || kernel_index>=m_kernels_total)
      return(false);
//---
   int kernel_handle=m_kernels[kernel_index];

   if(kernel_handle==INVALID_HANDLE)
      return(false);
//---
   return CLExecute(kernel_handle,work_dim,work_offset,work_size);
  }
//+------------------------------------------------------------------+
//| Execute                                                          |
//+------------------------------------------------------------------+
bool COpenCL::Execute(const int kernel_index,const int work_dim,const uint &work_offset[],const uint &work_size[],const uint &local_work_size[])
  {
//--- check parameters
   if(kernel_index<0 || kernel_index>=m_kernels_total)
      return(false);
//---
   return CLExecute(m_kernels[kernel_index],work_dim,work_offset,work_size,local_work_size);
  }
//+------------------------------------------------------------------+