1114 lines
No EOL
78 KiB
MQL5
1114 lines
No EOL
78 KiB
MQL5
//+------------------------------------------------------------------+
|
|
//| CL.mqh |
|
|
//| Copyright 2025, Leo. |
|
|
//| https://www.mql5.com |
|
|
//+------------------------------------------------------------------+
|
|
#property copyright "Copyright 2025, Leo."
|
|
#property link "https://www.mql5.com"
|
|
#property strict
|
|
|
|
#ifndef CL_BY_LEO_MQH
|
|
#define CL_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>
|
|
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>
|
|
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);
|
|
|
|
//--- 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_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);
|
|
Remover();
|
|
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);
|
|
Remover();
|
|
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);
|
|
Remover();
|
|
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);
|
|
Remover();
|
|
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)
|
|
{
|
|
//--- check parameters
|
|
if(CheckContext(FUNCION_ACTUAL) == false)
|
|
return(-1);
|
|
|
|
//---
|
|
return(CLGetInfoInteger(m_context, prop));
|
|
}
|
|
//+------------------------------------------------------------------+
|
|
//| GetKernelInfoInteger |
|
|
//+------------------------------------------------------------------+
|
|
long COpenCL::GetKernelInfoInteger(const int kernel_index, ENUM_OPENCL_PROPERTY_INTEGER prop)
|
|
{
|
|
//--- check parameters
|
|
if(!CheckKernel(FUNCION_ACTUAL, kernel_index, true)) //Aqui se chekeamos necesitamos que el handle sea true
|
|
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(!CheckBuffer(FUNCION_ACTUAL, buffer_index, false)) //Aqui no checkeamos dado qeu pueda ser qeu el buffer recien se cree pro lo tnato su handle sea invalido
|
|
return (false);
|
|
|
|
if(CheckContext(FUNCION_ACTUAL) == false)
|
|
return(false);
|
|
|
|
//---
|
|
ResetLastError();
|
|
int buffer_handle = CLBufferCreate(m_context, size_in_bytes, flags);
|
|
|
|
if(buffer_handle != INVALID_HANDLE)
|
|
{
|
|
m_buffers[buffer_index] = buffer_handle;
|
|
return(true);
|
|
}
|
|
|
|
LogError(StringFormat("No se pudo crear un handle para el buffer %d, ultimo error = %d, descripcion del error: ", buffer_index, GetLastError()), FUNCION_ACTUAL);
|
|
CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
|
|
|
|
//---
|
|
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)
|
|
{
|
|
//--- check parameters
|
|
if(CheckContext(FUNCION_ACTUAL) == false)
|
|
return(false);
|
|
|
|
if(!CheckBuffer(FUNCION_ACTUAL, buffer_index, false)) //Aqui tampoco chekemos al buffer dado que mas abajo tenemos la opcion de crearlo
|
|
return (false);
|
|
|
|
if(data_array_count <= 0)
|
|
{
|
|
LogError(StringFormat("La data %u establecida en el buffer %d, es invalida", data_array_count, buffer_index), FUNCION_ACTUAL);
|
|
return(false);
|
|
}
|
|
|
|
//--- buffer does not exists, create it
|
|
if(m_buffers[buffer_index] == INVALID_HANDLE)
|
|
{
|
|
uint size_in_bytes = data_array_count * sizeof(T);
|
|
ResetLastError();
|
|
int buffer_handle = CLBufferCreate(m_context, size_in_bytes, flags);
|
|
|
|
//---
|
|
if(buffer_handle == INVALID_HANDLE)
|
|
{
|
|
LogError(StringFormat("No se pudo crear el buffer[%d], ultimo error =%d, descripcion del erorr: ", buffer_index, GetLastError()), FUNCION_ACTUAL);
|
|
CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
|
|
return(false);
|
|
}
|
|
|
|
m_buffers[buffer_index] = buffer_handle;
|
|
}
|
|
|
|
//--- write data to OpenCL buffer
|
|
ResetLastError();
|
|
uint data_written = CLBufferWrite(m_buffers[buffer_index], data, 0, data_array_offset, data_array_count);
|
|
|
|
if(data_written != data_array_count)
|
|
{
|
|
LogError(StringFormat("No se pudo escribir en el buffer[%d] el tamaño total de data %u [array data = %u - data escrita = %u]\nUltimo error = %d", buffer_index, data_array_count, data.Size(), data_written, GetLastError()), FUNCION_ACTUAL);
|
|
return(false);
|
|
}
|
|
|
|
//---
|
|
return(true);
|
|
}
|
|
//+------------------------------------------------------------------+
|
|
//| BufferWriteFromMatrix |
|
|
//+------------------------------------------------------------------+
|
|
template<typename T>
|
|
bool COpenCL::BufferFromMatrix(const int buffer_index, matrix<T> &data, const uint flags)
|
|
{
|
|
//--- check parameters
|
|
if(CheckContext(FUNCION_ACTUAL) == false)
|
|
return(false);
|
|
|
|
if(!CheckBuffer(FUNCION_ACTUAL, buffer_index, false)) //Aqui tampoco chekemos al buffer dado que mas abajo tenemos la opcion de crearlo
|
|
return (false);
|
|
|
|
if(data.Rows() == 0 || data.Cols() == 0)
|
|
{
|
|
LogError(StringFormat("El numero de columnas %I64u, o filas %I64u de la matrix a crear para el buffer %d es invalido", data.Cols(), data.Rows(), buffer_index), FUNCION_ACTUAL);
|
|
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);
|
|
|
|
//---
|
|
ResetLastError();
|
|
int buffer_handle = CLBufferCreate(m_context, size_in_bytes, flags);
|
|
|
|
//---
|
|
if(buffer_handle == INVALID_HANDLE)
|
|
{
|
|
LogError(StringFormat("No se pudo crear un buffer[%d] para la matrix de %I64u Cols y %I64u Filas\nUltimo error = %d, descripcion del error: ", buffer_index, data.Cols(), data.Rows(), GetLastError()), FUNCION_ACTUAL);
|
|
CLLastError(FATAL_ERROR_TEXT, FUNCION_ACTUAL);
|
|
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(CheckContext(FUNCION_ACTUAL) == false)
|
|
return(false);
|
|
|
|
if(!CheckBuffer(FUNCION_ACTUAL, buffer_index, false)) //Aqui tampoco chekemos al buffer dado que mas abajo tenemos la opcion de crearlo
|
|
return (false);
|
|
|
|
if(data.Size() == 0)
|
|
{
|
|
LogError(StringFormat("El tamaño del vector a crear %I64u para el buffer %d es invalido", data.Size(), buffer_index), FUNCION_ACTUAL);
|
|
return(false);
|
|
}
|
|
|
|
//--- buffer does not exists, create it
|
|
if(m_buffers[buffer_index] == INVALID_HANDLE)
|
|
{
|
|
uint size_in_bytes = (uint)data.Size() * sizeof(T);
|
|
|
|
//---
|
|
ResetLastError();
|
|
int buffer_handle = CLBufferCreate(m_context, size_in_bytes, flags);
|
|
//---
|
|
if(buffer_handle == INVALID_HANDLE)
|
|
{
|
|
LogError(StringFormat("No se pudo crear el buffer[%d], tamaño del vector = %I64u\nUltimo error = %d, descripcion del error: ", buffer_index, data.Size(), GetLastError()), FUNCION_ACTUAL);
|
|
CLLastError(FATAL_ERROR_TEXT, FUNCION_ACTUAL);
|
|
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(!CheckBuffer(FUNCION_ACTUAL, buffer_index, true)) //No comprobamos el contexto dado que para agregar un buffer el contexto debe de ser true, si comprobamos 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(!CheckBuffer(FUNCION_ACTUAL, buffer_index, true)) //No comprobamos el contexto dado que para agregar un buffer el contexto debe de ser true
|
|
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(!CheckBuffer(FUNCION_ACTUAL, buffer_index, true)) //No comprobamos el contexto dado que para agregar un buffer el contexto debe de ser true
|
|
return (false);
|
|
|
|
if(data_array_count == 0)
|
|
{
|
|
LogError(StringFormat("El tamaño de data a leer %u, del buffer[%d] es invalida", buffer_index, data_array_count), FUNCION_ACTUAL);
|
|
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)
|
|
{
|
|
LogError(StringFormat("El numero de elementos leidos del buffer[%d] = %u, no es igual a los esperados %u", buffer_index, data_read, data_array_count), FUNCION_ACTUAL);
|
|
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(!CheckBuffer(FUNCION_ACTUAL, buffer_index, true)) //No comprobamos el contexto dado que para agregar un buffer el contexto debe de ser true
|
|
return (false);
|
|
|
|
if(data_array_count == 0)
|
|
{
|
|
LogError(StringFormat("El tamaño de data a leer %u, del buffer[%d] es invalida", buffer_index, data_array_count), FUNCION_ACTUAL);
|
|
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)
|
|
{
|
|
LogError(StringFormat("El numero de elemenos escritos en el buffer[%d] = %u, no es igual a los esperados %u", buffer_index, data_written, data_array_count), FUNCION_ACTUAL);
|
|
return(false);
|
|
}
|
|
|
|
//---
|
|
return(true);
|
|
}
|
|
|
|
//+------------------------------------------------------------------+
|
|
//| SetArgument |
|
|
//+------------------------------------------------------------------+
|
|
template<typename T>
|
|
bool COpenCL::SetArgument(const int kernel_index, const int arg_index, T value)
|
|
{
|
|
//--- check parameters
|
|
if(!CheckKernel(FUNCION_ACTUAL, kernel_index, true)) //Aqui no comprobamos el contexto dado que setear un kernel se debe de perdir el contexto de lo contario sera INVALID_HANDLE
|
|
return false;
|
|
|
|
bool res = CLSetKernelArg(m_kernels[kernel_index], arg_index, value);
|
|
|
|
if(!res)
|
|
{
|
|
LogError(StringFormat("No se pudo establecer el argumento[%d] en el kernel[%d]\n=%s, Ultimo error = %d, descripcion del error: ", arg_index, kernel_index, m_kernel_names[kernel_index], GetLastError()), FUNCION_ACTUAL);
|
|
CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
|
|
}
|
|
|
|
//---
|
|
return res;
|
|
}
|
|
|
|
//+------------------------------------------------------------------+
|
|
//| SetArgumentBuffer |
|
|
//+------------------------------------------------------------------+
|
|
bool COpenCL::SetArgumentBuffer(const int kernel_index, const int arg_index, const int buffer_index)
|
|
{
|
|
//--- check parameters
|
|
if(!CheckKernel(FUNCION_ACTUAL, kernel_index, true)) //Aqui no comprobamos el contexto dado que setear un kernel se debe de perdir el contexto de lo contario sera INVALID_HANDLE
|
|
return false;
|
|
|
|
if(!CheckBuffer(FUNCION_ACTUAL, buffer_index, true))
|
|
return false;
|
|
|
|
bool res = CLSetKernelArgMem(m_kernels[kernel_index], arg_index, m_buffers[buffer_index]);
|
|
|
|
if(!res)
|
|
{
|
|
LogError(StringFormat("No se pudo establecer el arg[%d]-buffer[%d] en el kernel[%d]\n=%s, Ultimo error = %d, descripcion del error: ", arg_index, buffer_index, kernel_index, m_kernel_names[kernel_index], GetLastError()), FUNCION_ACTUAL);
|
|
CLLastError(ERROR_TEXT, FUNCION_ACTUAL);
|
|
}
|
|
|
|
//---
|
|
return res;
|
|
}
|
|
|
|
//+------------------------------------------------------------------+
|
|
//| SetArgumentLocalMemory |
|
|
//+------------------------------------------------------------------+
|
|
bool COpenCL::SetArgumentLocalMemory(const int kernel_index, const int arg_index, const int local_memory_size)
|
|
{
|
|
//--- check parameters
|
|
if(!CheckKernel(FUNCION_ACTUAL, kernel_index, true))
|
|
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(!CheckKernel(FUNCION_ACTUAL, kernel_index, true))
|
|
return false;
|
|
|
|
bool res = CLExecute(m_kernels[kernel_index], work_dim, work_offset, work_size);
|
|
if(!res)
|
|
{
|
|
LogFatalError(StringFormat("Error al ejecutar el kernel[%d] - [ %s ]\nUltimo error = %d, error description: ", kernel_index, m_kernel_names[kernel_index], GetLastError()), FUNCION_ACTUAL);
|
|
CLLastError(FATAL_ERROR_TEXT, FUNCION_ACTUAL);
|
|
}
|
|
|
|
//---
|
|
return res;
|
|
}
|
|
|
|
//+------------------------------------------------------------------+
|
|
//| 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(!CheckKernel(FUNCION_ACTUAL, kernel_index, true))
|
|
return false;
|
|
|
|
bool res = CLExecute(m_kernels[kernel_index], work_dim, work_offset, work_size, local_work_size);
|
|
if(!res)
|
|
{
|
|
LogFatalError(StringFormat("Error al ejecutar el kernel[%d] - [ %s ]\nUltimo error = %d, error description: ", kernel_index, m_kernel_names[kernel_index], GetLastError()), FUNCION_ACTUAL);
|
|
CLLastError(FATAL_ERROR_TEXT, FUNCION_ACTUAL);
|
|
}
|
|
|
|
//---
|
|
return res;
|
|
}
|
|
//+------------------------------------------------------------------+
|
|
#endif |