Код OpenCL на MQL5 не получает распределенные задания для каждого ядра GPU
Я создал индикатор на основе графического процессора для платформы терминала MetaTrader, используя OpenCL и MQL5.
Я изо всех сил старался, чтобы моя работа по оптимизации [ MetaTrader Terminal: Strategy Tester ] была перенесена на GPU по максимуму. Большая часть расчетов выполняется по индикатору. Следовательно, я внес изменения в индикатор и полностью перешел на GPU.
Но настоящая проблема возникает, когда я пытаюсь перейти к процессу оптимизации в разделе тестера стратегий.
Процесс, который я вижу, использует мой GPU и CPU, но не влияет на весь процесс.
Я подозреваю, что процесс не распределяется по каждому ядру графического процессора для обработки, вместо этого все ядра графического процессора работают на одном и том же процессе или функции для выполнения.
Пожалуйста, дайте мне знать, что мне нужно сделать, чтобы заставить один GPU работать над выполнением одной функции, чтобы обеспечить более быстрый вывод.
Вот моя ссылка на код прилагается: полный код с экспертом
Ядро моего кода:
__kernel void calSMA(
int limit,
int rates_total,
__global double *price,
__global double *ExtLineBuffer,
int InpMAPeriod
)
{
int count = 0;
int len = get_global_id(2);
for(int i=limit;i<rates_total;i++)
ExtLineBuffer[len+i] = ExtLineBuffer[len+ i-1]+(price[len+i]-price[len+i-InpMAPeriod])/InpMAPeriod;
}
__kernel void calcSMALoop(int begin, int limit, __global double *price, __global double *firstValue, int InpMAPeriod)
{
int i, len = get_global_id(2);
for(i=begin;i<limit;i++)
firstValue[len]+=price[i];
firstValue[len]/=InpMAPeriod;
}
__kernel void calcEMA(int begin, int limit, __global double *price, __global double *ExtLineBuffer, double SmoothFactor)
{
int len = get_global_id(2);
for(int i=begin;i<limit;i++)
ExtLineBuffer[len + i]=price[len + i]*SmoothFactor+ExtLineBuffer[len + i-1]*(1.0-SmoothFactor);
}
__kernel void calcSSMA(int limit, int rates_total, __global double *price, __global double *ExtLineBuffer, int InpMAPeriod)
{
int len = get_global_id(2);
for(int i=limit;i<rates_total;i++)
ExtLineBuffer[len + i]=(ExtLineBuffer[len + i-1]*(InpMAPeriod-1)+price[len + i])/InpMAPeriod;
}
__kernel void calcLWMALoop(int begin, int limit, __global double *price, __global double *firstValue, int weightsum, __global int *weightreturn)
{
weightsum = 0;
int len = get_global_id(2);
for(int i=begin;i<limit;i++)
{
weightsum+=(i-begin+1);
firstValue[len]+=(i-begin+1)*price[i];
}
firstValue[len]/=(double)weightsum;
weightreturn[0] = weightsum;
}
//__global int counter = 0;
double returnCalculation(int InpMAPeriod, double price, int j)
{
return ((InpMAPeriod-j)*price);
}
__kernel void calcLWMA(int limit, int rates_total, __global double *price, __global double *ExtLineBuffer, int InpMAPeriod, int weightsum)
{
int len = get_global_id(2);
for(int i=limit;i<rates_total;i++)
{
double sum = 0;
for(int j=0;j<InpMAPeriod;j++) sum+=returnCalculation(InpMAPeriod,price[len + i-j],j);
ExtLineBuffer[len + i]=sum/weightsum;
}
}
Пожалуйста, предложите мне выход для распределения функции с разными значениями или фреймами в MQL5 с использованием GPU на OpenCL.
РЕДАКТИРОВАНИЕ
Это большой вызов для тех, кто ищет вызов... Даже мне не терпится узнать, можно ли что-нибудь сделать с OpenCL и MQL5 для задачи оптимизации. Я надеюсь, что получу ответы на то, что я ищу.
ИЗДАНО ВНОВЬ MAGPU.mqh файл
#include "CHECKMA.mq5"
#define CUDA_CORE 2
int Execute_SMA(
const double &price[],
int rates_total,
int limit
)
{
int cl_mem = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE),
cl_price = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE);
Check_Memory_Initialization(cl_mem, cl_price, cl_CommonKernel1, "Execute_SMA function error");
if(!CLSetKernelArgMem(cl_CommonKernel1,2,cl_price))
Print("Input Bufer Not Set");
//else Print("Input Buffer Set");
if(!CLSetKernelArgMem(cl_CommonKernel1,3,cl_mem))
Print("Output Bufer Not Set");
//else Print("Output Buffer Set");
if(!CLBufferWrite(cl_price, price))
Print("Could not copy Input buffer");
//else Print("Copied: ",cl_price);
if(!CLBufferWrite(cl_mem, ExtLineBuffer))
Print("Could not copy Input buffer");
//else Print("Copied: ",cl_mem);
//else Print("Input Buffer Copied");
if(!CLSetKernelArg(cl_CommonKernel1,0,limit))
Print("Could Not Set Arg 0");
//else Print("Set Arg 0");
if(!CLSetKernelArg(cl_CommonKernel1,1,rates_total))
Print("Could Not Set Arg 1");
//else Print("Set Arg 1");
//if(!CLSetKernelArg(cl_CommonKernel1,4,previous_value))
//Print("Could Not Set Arg2");
//else Print("Set Arg 2");
if(!CLSetKernelArg(cl_CommonKernel1,4,InpMAPeriod))
Print("Could Not Set Arg3: ",GetLastError());
//Print(CLGetInfoInteger(cl_ctx,CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS));
if(!CLExecute(cl_CommonKernel1,CUDA_CORE,offset,work))
Print("Kernel not executed",GetLastError());
//else Print("Executing Now!");
//if(CLExecutionStatus(cl_krn) == 0) Print("Completed");
//if(CLExecutionStatus(cl_krn) == 1) Print("CL_RUNNING");
//if(CLExecutionStatus(cl_krn) == 2) Print("CL_SUBMITTED");
//if(CLExecutionStatus(cl_krn) == 3) Print("CL_QUEUED");
//if(CLExecutionStatus(cl_krn) == -1)Print("Error Occurred:", GetLastError());
//if(!CLExecutionStatus(cl_krn))
//Print(CLExecutionStatus(cl_krn));
if(!CLBufferRead(cl_mem,ExtLineBuffer))
Print("Buffer Copy Nothing: ", GetLastError());
CLBufferFree(cl_price);
CLBufferFree(cl_mem);
return(1);
}
double ExecuteLoop(
int begin,
int limit,
const double &price[]
)
{
int cl_mem = CLBufferCreate(cl_ctx,sizeof(double),CL_MEM_READ_WRITE),
cl_price = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE);
double temp[];
ArrayResize(temp,1);
temp[0] = 0;
Check_Memory_Initialization(cl_mem, cl_price, cl_CommonKernel2, "ExecuteLoop function error");
if(!CLSetKernelArgMem(cl_CommonKernel2,2,cl_price))
Print("Input Bufer Not Set 2");
if(!CLSetKernelArgMem(cl_CommonKernel2,3,cl_mem))
Print("Output Bufer Not Set 2");
if(!CLBufferWrite(cl_price, price))
Print("Could not copy Input buffer 2");
if(!CLSetKernelArg(cl_CommonKernel2,0,begin))
Print("Could Not Set Arg 0");
if(!CLSetKernelArg(cl_CommonKernel2,1,limit))
Print("Could Not Set Arg 1");
if(!CLSetKernelArg(cl_CommonKernel2,4,InpMAPeriod))
Print("Could Not Set Arg3: ",GetLastError());
if(!CLExecute(cl_CommonKernel2,CUDA_CORE,offset,work))
Print("Kernel not executed",GetLastError());
if(!CLBufferRead(cl_mem,temp))
Print("Buffer Copy Nothing: ", GetLastError());
CLBufferFree(cl_price);
CLBufferFree(cl_mem);
return(temp[0]);
}
int ExecuteEMA(int begin, int limit, const double &price[], double SmoothFactor)
{
int cl_mem = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE),
cl_price = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE);
Check_Memory_Initialization(cl_mem, cl_price, cl_CommonKernel1, "ExecuteEMA function error");
if(!CLSetKernelArgMem(cl_CommonKernel1,2,cl_price))
Print("Input Bufer Not Set");
if(!CLSetKernelArgMem(cl_CommonKernel1,3,cl_mem))
Print("Output Bufer Not Set");
if(!CLBufferWrite(cl_price, price))
Print("Could not copy Input buffer");
if(!CLBufferWrite(cl_mem, ExtLineBuffer))
Print("Could not copy Input buffer");
if(!CLSetKernelArg(cl_CommonKernel1,0,begin))
Print("Could Not Set Arg 0");
if(!CLSetKernelArg(cl_CommonKernel1,1,limit))
Print("Could Not Set Arg 1");
if(!CLSetKernelArg(cl_CommonKernel1,4,SmoothFactor))
Print("Could Not Set Arg3: ",GetLastError());
if(!CLExecute(cl_CommonKernel1,CUDA_CORE,offset,work))
Print("Kernel not executed",GetLastError());
if(!CLBufferRead(cl_mem,ExtLineBuffer))
Print("Buffer Copy Nothing: ", GetLastError());
CLBufferFree(cl_price);
CLBufferFree(cl_mem);
return(1);
}
int Execute_SSMA(
const double &price[],
int rates_total,
int limit
)
{
int cl_mem = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE),
cl_price = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE);
Check_Memory_Initialization(cl_mem, cl_price, cl_CommonKernel1, "Execute_SSMA function error");
if(!CLSetKernelArgMem(cl_CommonKernel1,2,cl_price))
Print("Input Bufer Not Set");
if(!CLSetKernelArgMem(cl_CommonKernel1,3,cl_mem))
Print("Output Bufer Not Set");
if(!CLBufferWrite(cl_price, price))
Print("Could not copy Input buffer");
if(!CLBufferWrite(cl_mem, ExtLineBuffer))
Print("Could not copy Input buffer");
//
//else Print("Input Buffer Copied");
if(!CLSetKernelArg(cl_CommonKernel1,0,limit))
Print("Could Not Set Arg 0");
if(!CLSetKernelArg(cl_CommonKernel1,1,rates_total))
Print("Could Not Set Arg 1");
if(!CLSetKernelArg(cl_CommonKernel1,4,InpMAPeriod))
Print("Could Not Set Arg3: ",GetLastError());
if(!CLExecute(cl_CommonKernel1,CUDA_CORE,offset,work))
Print("Kernel not executed",GetLastError());
if(!CLBufferRead(cl_mem,ExtLineBuffer))
Print("Buffer Copy Nothing: ", GetLastError());
CLBufferFree(cl_price);
CLBufferFree(cl_mem);
return(1);
}
double ExecuteLWMALoop(
int begin,
int limit,
const double &price[],
int weightsumlocal
)
{
int cl_mem = CLBufferCreate(cl_ctx,sizeof(double),CL_MEM_READ_WRITE),
cl_price = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE),
cl_weightsumlocal = CLBufferCreate(cl_ctx,sizeof(int),CL_MEM_READ_WRITE);
double temp[];
int weight[];
ArrayResize(temp,1);
ArrayResize(weight,1);
weight[0] = 0;
temp[0] = 0;
Check_Memory_Initialization(cl_mem, cl_price, cl_CommonKernel2, "ExecuteLWMALoop function error");
if(!CLSetKernelArgMem(cl_CommonKernel2,2,cl_price))
Print("Input Bufer Not Set 2");
if(!CLSetKernelArgMem(cl_CommonKernel2,3,cl_mem))
Print("Output Bufer Not Set 2");
if(!CLSetKernelArgMem(cl_CommonKernel2,5,cl_weightsumlocal))
Print("Output Bufer Not Set 2");
if(!CLBufferWrite(cl_price, price))
Print("Could not copy Input buffer 2");
if(!CLSetKernelArg(cl_CommonKernel2,0,begin))
Print("Could Not Set Arg 0");
if(!CLSetKernelArg(cl_CommonKernel2,1,limit))
Print("Could Not Set Arg 1");
if(!CLSetKernelArg(cl_CommonKernel2,4,weightsumlocal))
Print("Could Not Set Arg3: ",GetLastError());
if(!CLExecute(cl_CommonKernel2,CUDA_CORE,offset,work))
Print("Kernel not executed",GetLastError());
if(!CLBufferRead(cl_mem,temp))
Print("Buffer Copy Nothing: ", GetLastError());
if(!CLBufferRead(cl_weightsumlocal,weight))
Print("Buffer Copy Nothing: ", GetLastError());
weightsum = weight[0];
CLBufferFree(cl_weightsumlocal);
CLBufferFree(cl_price);
CLBufferFree(cl_mem);
return(temp[0]);
}
int Execute_LWMA(const double &price[], int rates_total, int limit, int weightsum1)
{
int cl_mem = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE),
cl_price = CLBufferCreate(cl_ctx,ArraySize(price)*sizeof(double),CL_MEM_READ_WRITE);
Check_Memory_Initialization(cl_mem, cl_price, cl_CommonKernel1, "Execute_SSMA function error");
if(!CLSetKernelArgMem(cl_CommonKernel1,2,cl_price))
Print("Input Bufer Not Set");
if(!CLSetKernelArgMem(cl_CommonKernel1,3,cl_mem))
Print("Output Bufer Not Set");
if(!CLBufferWrite(cl_price, price))
Print("Could not copy Input buffer");
if(!CLBufferWrite(cl_mem, ExtLineBuffer))
Print("Could not copy Input buffer");
//else Print("Input Buffer Copied");
if(!CLSetKernelArg(cl_CommonKernel1,0,limit))
Print("Could Not Set Arg 0");
if(!CLSetKernelArg(cl_CommonKernel1,1,rates_total))
Print("Could Not Set Arg 1");
if(!CLSetKernelArg(cl_CommonKernel1,4,InpMAPeriod))
Print("Could Not Set Arg4: ",GetLastError());
if(!CLSetKernelArg(cl_CommonKernel1,5,weightsum1))
Print("Could Not Set Arg5: ",GetLastError());
if(!CLExecute(cl_CommonKernel1,CUDA_CORE,offset,work))
Print("Kernel not executed",GetLastError());
if(!CLBufferRead(cl_mem,ExtLineBuffer))
Print("Buffer Copy Nothing: ", GetLastError());
CLBufferFree(cl_price);
CLBufferFree(cl_mem);
return(1);
}
void checkKernel(int cl_kernel, string var_name)
{
if(cl_kernel==INVALID_HANDLE )
{
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL kernel create failed: ERR_OPENCL_INVALID_HANDLE ", var_name);
return;
}
if(cl_kernel==ERR_INVALID_PARAMETER )
{
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL kernel create failed: ERR_INVALID_PARAMETER ", var_name);
return;
}
if(cl_kernel==ERR_OPENCL_TOO_LONG_KERNEL_NAME )
{
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL kernel create failed: ERR_OPENCL_TOO_LONG_KERNEL_NAME ", var_name);
return;
}
if(cl_kernel==ERR_OPENCL_KERNEL_CREATE )
{
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL kernel create failed 1: ERR_OPENCL_KERNEL_CREATE ", var_name);
return;
}
}
int Check_Memory_Initialization(int cl_mem, int cl_price, int cl_ker, string name_process_call)
{
if(cl_mem==INVALID_HANDLE)
{
CLKernelFree(cl_ker);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL buffer create failed: cl_mem INVALID_HANDLE: ", name_process_call);
return(0);
}
if(cl_mem==ERR_NOT_ENOUGH_MEMORY )
{
CLKernelFree(cl_ker);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL buffer create failed: cl_mem ERR_NOT_ENOUGH_MEMORY: ", name_process_call);
return(0);
}
if(cl_mem==ERR_OPENCL_BUFFER_CREATE )
{
CLKernelFree(cl_ker);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL buffer create failed: cl_mem ERR_OPENCL_BUFFER_CREATE: ", name_process_call);
return(0);
}
if(cl_price==INVALID_HANDLE)
{
CLKernelFree(cl_ker);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
CLBufferFree(cl_mem);
Print("OpenCL buffer create failed: cl_price: ", name_process_call);
return(0);
}
if(cl_price==ERR_NOT_ENOUGH_MEMORY)
{
CLKernelFree(cl_ker);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
CLBufferFree(cl_mem);
Print("OpenCL buffer create failed: cl_price ERR_NOT_ENOUGH_MEMORY: ", name_process_call);
return(0);
}
if(cl_price==ERR_OPENCL_BUFFER_CREATE)
{
CLKernelFree(cl_ker);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
CLBufferFree(cl_mem);
Print("OpenCL buffer create failed: cl_price ERR_OPENCL_BUFFER_CREATE: ", name_process_call);
return(0);
}
return(1);
}
ФАЙЛ ОСНОВНОГО ИНДИКАТОРА CHECKMA.mq5 файл
#resource "program_MA_GPU.cl" as string cl_program
#include "MAGPU.mqh"
#property indicator_chart_window
#property indicator_buffers 1
#property indicator_plots 1
#property indicator_type1 DRAW_LINE
#property indicator_color1 Red
input int InpMAPeriod=13; // Period
input int InpMAShift=0; // Shift
input ENUM_MA_METHOD InpMAMethod=MODE_SMA; // Method
//--- indicator buffers
double ExtLineBuffer[];
int offset[CUDA_CORE], work[CUDA_CORE];//={0,19,38,57,76,95,114,123};
string str;
int cl_ctx, cl_prg, cl_CommonKernel1, cl_CommonKernel2;
static int weightsum;
void CalculateSimpleMA(int rates_total,int prev_calculated,int begin,const double &price[])
{
int limit;
if(prev_calculated==0)
{
limit=InpMAPeriod+begin;
ArrayFill(ExtLineBuffer,0,limit-1,0.0);
ExtLineBuffer[limit-1]=ExecuteLoop(begin,limit,price);
}
else limit=prev_calculated-ArraySize(price)+InpMAPeriod+17;
Execute_SMA(price,rates_total,limit);
}
void CalculateEMA(int rates_total,int prev_calculated,int begin,const double &price[])
{
int limit;
double SmoothFactor=2.0/(1.0+InpMAPeriod);
if(prev_calculated==0)
{
limit=InpMAPeriod+begin;
ExtLineBuffer[begin]=price[begin];
ExecuteEMA(begin+1,limit,price,SmoothFactor);
}
else limit=prev_calculated;
ExecuteEMA(begin+99900,limit,price,SmoothFactor);
}
void CalculateLWMA(int rates_total,int prev_calculated,int begin,const double &price[])
{
int limit;
if(prev_calculated==0)
{
weightsum=0;
limit=InpMAPeriod+begin;
//--- set empty value for first limit bars
ArrayFill(ExtLineBuffer,0,limit,0.0);
//--- calculate first visible value
ExtLineBuffer[limit-1]=ExecuteLWMALoop(begin,limit,price,weightsum);
}
else limit=prev_calculated-ArraySize(price)+InpMAPeriod+17;
//--- main loop
Execute_LWMA(price,rates_total,limit,weightsum);
}
void CalculateSmoothedMA(int rates_total,int prev_calculated,int begin,const double &price[])
{
int limit;
//--- first calculation or number of bars was changed
if(prev_calculated==0)
{
limit=InpMAPeriod+begin;
//--- set empty value for first limit bars
ArrayFill(ExtLineBuffer,0,limit-1,0.0);
ExtLineBuffer[limit-1]=ExecuteLoop(begin,limit,price);
}
else limit=prev_calculated-ArraySize(price)+InpMAPeriod+17;
Execute_SSMA(price,rates_total,limit);
//---
}
void OnInit()
{
//--- indicator buffers mapping
SetIndexBuffer(0,ExtLineBuffer,INDICATOR_DATA);
//--- set accuracy
IndicatorSetInteger(INDICATOR_DIGITS,_Digits+1);
//--- sets first bar from what index will be drawn
PlotIndexSetInteger(0,PLOT_DRAW_BEGIN,InpMAPeriod);
//---- line shifts when drawing
PlotIndexSetInteger(0,PLOT_SHIFT,InpMAShift);
//--- name for DataWindow
//---- sets drawing line empty value--
PlotIndexSetDouble(0,PLOT_EMPTY_VALUE,0.0);
//---- initialization done
cl_ctx = CLContextCreate(CL_USE_GPU_ONLY);
cl_prg=CLProgramCreate(cl_ctx,cl_program,str);
if(cl_ctx==INVALID_HANDLE)
{
Print("OpenCL not found: ", GetLastError() );
return;
}
if(cl_prg==INVALID_HANDLE)
{
CLContextFree(cl_ctx);
Print("OpenCL program create failed: ", str);
return;
}
if(cl_prg==ERR_INVALID_PARAMETER )
{
CLContextFree(cl_ctx);
Print("OpenCL program create failed: ", str);
return;
}
if(cl_prg==ERR_NOT_ENOUGH_MEMORY )
{
CLContextFree(cl_ctx);
Print("OpenCL program create failed: ", str);
return;
}
if(cl_prg==ERR_OPENCL_PROGRAM_CREATE )
{
CLContextFree(cl_ctx);
Print("OpenCL program create failed: ", str);
return;
}
int c = 1;
ArrayFill(work,0,CUDA_CORE,c);
//ArrayInitialize(offset,0);
int enter = -c;
for (int i =0; i < CUDA_CORE; i++)
{
offset[i] = enter + c;
enter = offset[i];
}
switch(InpMAMethod)
{
case MODE_SMA : cl_CommonKernel1 = CLKernelCreate(cl_prg,"calSMA");
checkKernel(cl_CommonKernel1,"cl_CommonKernel1 SMA");
cl_CommonKernel2 = CLKernelCreate(cl_prg,"calcSMALoop");
checkKernel(cl_CommonKernel2,"cl_CommonKernel2 SMA");
break;
case MODE_EMA : cl_CommonKernel1 = CLKernelCreate(cl_prg,"calcEMA");
checkKernel(cl_CommonKernel1,"cl_CommonKernel1 EMA");
break;
case MODE_LWMA : cl_CommonKernel1 = CLKernelCreate(cl_prg,"calcLWMA");
checkKernel(cl_CommonKernel1,"cl_CommonKernel1 LWMA");
cl_CommonKernel2 = CLKernelCreate(cl_prg,"calcLWMALoop");
checkKernel(cl_CommonKernel2,"cl_CommonKernel2 LWMA");
break;
case MODE_SMMA : cl_CommonKernel1 = CLKernelCreate(cl_prg,"calcSSMA");
checkKernel(cl_CommonKernel1,"cl_CommonKernel1 SSMA");
cl_CommonKernel2 = CLKernelCreate(cl_prg,"calcSMALoop");
checkKernel(cl_CommonKernel2,"cl_CommonKernel2 SSMA");
break;
}
}
int OnCalculate(const int rates_total,
const int prev_calculated,
const int begin,
const double &price[])
{
if(rates_total<InpMAPeriod-1+begin)
return(0);
if(prev_calculated==0)
ArrayInitialize(ExtLineBuffer,0);
PlotIndexSetInteger(0,PLOT_DRAW_BEGIN,InpMAPeriod-1+begin);
switch(InpMAMethod)
{
case MODE_EMA: CalculateEMA(rates_total,prev_calculated,begin,price); break;
case MODE_LWMA: CalculateLWMA(rates_total,prev_calculated,begin,price); break;
case MODE_SMMA: CalculateSmoothedMA(rates_total,prev_calculated,begin,price); break;
case MODE_SMA: CalculateSimpleMA(rates_total,prev_calculated,begin,price); break;
}
//--- return value of prev_calculated for next call
return(rates_total);
}
void OnDeinit(const int reason)
{
CLKernelFree(cl_CommonKernel1);
CLKernelFree(cl_CommonKernel2);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
}
1 ответ
Помогите мне написать этот код надлежащим образом, чтобы ускорить процесс и получить правильный результат с моим графическим процессором.
Сначала факты,
с 2 апреля
ты уже знаешь,
MQL5CustomIndicator не будет работать таким образом...
Факты имеют значение - если архитектура выполнения кода MQL4/5 явно задокументировала, что нет места для каких-либо операций с расширенной задержкой / асинхронностью / блокировкой, которые когда-либо помещались бы внутри любого MQL4/5 CustomIndicator Блок выполнения кода, возможно, есть разумное время, чтобы остановить все эти атаки членов сообщества S/O и начать просто уважать документированный и опубликованный факт.
Документация по MQL5 является довольно явной и предупреждает о неблагоприятных последствиях ее архитектуры с одним общим потоком для производительности и / или предупреждает о риске полной, непреднамеренной блокировки всей системы:
Все индикаторы, рассчитанные по одному символу, даже если они прикреплены к разным графикам, работают в одном потоке. Таким образом, все индикаторы на одном символе совместно используют ресурсы одного потока.
Бесконечный цикл (блокировка / увеличение задержки / неожиданная задержка) в одном индикаторе остановит все остальные индикаторы на этом символе.
Конечно, можно попросить многих добрых профессионалов OpenCL, присутствующих в сообществе S/O, помочь, выразить крайнюю заинтересованность в получении своих спонсируемых знаний и так далее, если они решат потратить свое драгоценное время на спонсирование и расширение таких усилий.
Все это по-прежнему должно соответствовать реальности и лучше всего уважать известные факты, прежде чем любое разумное усилие (при любых усилиях) может по крайней мере начать работать в правильном направлении.
Как внутренне работает топология CPU: GPU?
С использованием асинхронной работы с доставкой вне порядка, асинхронной по определению, с сохранением в очереди, последовательностью вычислительных заданий, нацеленных на устройства с GPU, все с начальным и окончательным преимущественно недетерминированным RTT-Длительность.
Конец стороны доступа к GPU-устройству на стороне процессора. Очередь может быть проинструктирована отправлять задание в GPU (задача ~ программа для выполнения):
В Fiteite Sate Состояние рабочего процесса управления GPU-задачами имеет следующий топологический граф:
< START____________> s = GetMicrosecondCount();
( MQL5 RQSTs )
|
|
|
T0:+---+--> CL_QUEUED : 3 == "queued", i.e. waiting for its turn ( submission )
? | | |
? +----+ |
? v
T0+? +--> CL_SUBMITTED : 2 == "submitted" for an OpenCL-device execution
? | | |
? +----+ |
? v
T0+? +--> CL_RUNNING : 1 == "running" the kernel-code on a mapped OpenCL resource pool
? | | |
? +----+ |
? v
T0+? CL_COMPLETE : 0 == "program complete", processing has finished its remote outputs
? |
T0+?-----------------+
|
( MQL5 FREEs )
< END_____________> e = GetMicrosecondCount();
PrintFormat( "RTT-COST WAS ~ %9d [us] ( CLES==0 ? %d )",
( e - s ),
CLExecutionStatus( _gpuKernelHANDLE )
);
Далее, давайте уважать сферу GPU-архитектур:
Вычислительные устройства на базе графических процессоров отличаются, имея другие кремниевые архитектуры, чем любые универсальные вычислительные устройства на базе CISC/RISC.
Причина, почему здесь очень важна.
В устройствах с графическим процессором используются блоки управления S treaming M ultiprocessor e X (SMX), которые упоминаются в некоторых инструментах проверки оборудования.
В то время как буква M в аббревиатуре S M X подчеркивает, на SMX-модуль загружается несколько исполнений, но все такие случаи действительно выполняются (конечно, только если даны инструкции таким образом, что выходит за рамки этого чтобы охватить / охватить все SM-ядра, представленные в SMX) одни и те же вычислительные инструкции - это единственный способ, которым они могут работать - это называется SIMT/SIMD -типа ограниченного объема параллелизма, достижимого (совместно) по периметру SMX, где s- единственная инструкция- м ультипле- { t hreads | d ata} может выполняться в рамках существующих возможностей планировщика SIMT/SIMD-( WARP-wide | half-WARP-wide | WARP-ignoring-GreedyMode). Важно отметить, что чем меньше ширина запланированного SIMT/SIMD - выполнение выполняется, чем меньше SMX/SM-ядер фактически выполняют какую-либо полезную часть глобального выполнения задания, и тем больше потраченного времени опустошает борьбу за производительность из-за падения числа N-(CPUs) в действительности, как описано ниже.
Перечисление этих 384 ядер, опубликованное выше, означает аппаратное ограничение, за пределы которого этот совместно-организованный параллельный параллелизм SIMT / SIMD-типа не может расти, и все попытки в этом направлении приведут к чисто [SERIAL] внутреннее планирование GPU-заданий (да, т.е. одно за другим).
Понимание этих основ является кардинальным, поскольку без этих особенностей архитектуры можно ожидать поведения, которое на самом деле принципиально невозможно организовать в какой-либо системе GPGPU, имеющей формальную форму [ 1-CPU-host : N-GPU-device(s) ] композиции автономных, асинхронных распределенных систем звезда узлов.
Любое ядро GPU, загруженное с CPU-host на GPU, будет отображаться на непустом наборе SMX-блоков, где применяется указанное количество ядер (другое, более мелкое зерно-геометрия вычислительных ресурсов, снова выходя далеко за рамки этого поста) загружается потоком SIMT / SIMD-инструкций, не нарушая ограничений GPU-устройства:
...
+----------------------------------------------------------------------------------------
Max work items dimensions: 3 // 3D-geometry grids possible
Max work items[0]: 1024 // 1st dimension max.
Max work items[1]: 1024
Max work items[2]: 64 // theoretical max. 1024 x 1024 x 64 BUT...
+----------------------------------------------------------------------------------------
Max work group size: 1024 // actual max. "geometry"-size
+----------------------------------------------------------------------------------------
...
Так,
если ядро 1-SM было внутренне проинструктировано выполнить некоторый блок задачи GPU (задание GPU), только это одно ядро SM будет извлекать одну инструкцию GPU-RISC за другой (игнорируя любую возможную ILP для простоты здесь) и выполняйте его по одному, шагая по потоку SIMD-инструкций указанного GPU-задания. Все остальные ядра SM, присутствующие на одном и том же модуле SMX, обычно ничего не делают в течение этого времени, пока эта задача GPU не будет завершена, и внутренняя система управления процессами GPU не решит отобразить какую-либо другую работу для этого SMX.
если 2-SM-ядрам было дано указание выполнить какое-то GPU-задание, только эта пара SM-ядер будет получать одну (и ту же самую) инструкцию GPU-RISC за другой (игнорируя любую возможную ILP для простоты здесь) и обе выполняйте его по одному, шагая по потоку SIMT / SIMD-инструкций указанного GPU-задания. В этом случае, если одно SM-ядро попадает в состояние, при котором
ifили аналогично разветвленный поток выполнения заставляет одно SM-ядро идти в другой путь потока выполнения кода, чем другой, параллелизм SIMT / SIMD попадает в расходящийся сценарий, где одно SM-ядро получает следующее SIMT / SIMD-инструкция, принадлежащая его пути выполнения кода, тогда как другая ничего не делает (получаетGPU_NOP(s)), пока первый не закончил всю работу (или не был вынужден остановиться на каком-то барьере синхронизации, попавшем в состояние ожидания немаскируемой задержки, при ожидании получения части данных из "далеко" (медленно) расположение нелокальной памяти, опять же, детали выходят далеко за рамки этого поста)- только после того, как что-либо из этого случится, расходящийся путь, пока простоGPU_NOPSM-ядро может получить любую следующую SIMT / SIMD-инструкцию, принадлежащую его (расходящемуся) пути выполнения кода, чтобы двигаться вперед. Все остальные ядра SM, присутствующие на одном и том же модуле SMX, обычно ничего не делают в течение этого времени, пока эта задача GPU не будет завершена, и внутренняя система управления процессами GPU не решит отобразить какую-либо другую работу для этого SMX.если 16-SM-ядрам было дано указание выполнить какое-либо задание GPU с помощью специфической для задачи "геометрии", то только это "стадо" SM-ядер получит одну (и ту же самую) SIMU / SIMD-инструкцию GPU-RISC после другой (игнорируя любой возможный ILP для простоты здесь) и все выполняют его по одному, шагая по потоку SIMT / SIMD-инструкций упомянутого задания GPU. Любое расхождение внутри "стада" уменьшает SIMT / SIMD-эффект и
GPU_NOPблокированные ядра остаются в ожидании, пока основная часть "стада" завершит работу (так же, как было нарисовано прямо над этой точкой).если было доступно больше SIMT / SIMD-потоков-чем-SM-cores для выполнения некоторого GPU-задания в соответствии с "геометрией" задачи, кремний GPU-устройства будет работать так, чтобы поток выполнялся как
[SERIAL]последовательность как много{ WARP-wide | half-WARP-wide }-Пакеты SIMT / SIMD-нитей, пока такая последовательность не завершит все указанное число потоков SIMT / SIMD, сопоставленных с SMX. Следовательно, временная согласованность такой финальной финализации пакета принципиально невозможна, так как они достигают своих соответствующих целей определенным образом в WARP-планировщике, но никогда синхронно (да, ваш код на стороне ЦП здесь должен будет ждать до самого последнего (самый ленивый (по какой-либо причине, будь то причина планирования с ограниченной пропускной способностью, либо причина планирования расхождения кода, либо плохая взаимная (пере) причина синхронизации) поток выполнения кода) в конечном итоге в какое-то неизвестное время в будущем закончи__kernelобработки кода и устройство с OpenCL позволит "удаленно" обнаруживатьCL_COMPLETEперед тем, как вы сможете получить какие-либо значимые результаты (как вы спросите с удивлением в одном из ваших других вопросов).
в любом случае, все остальные SM-ядра, не отображаемые "геометрией" конкретной задачи на SMX-модуле соответствующих GPU-устройств, как правило, вообще не будут делать ничего полезного - так что важность знания аппаратных деталей для правильной задачи специфическая "геометрия" действительно важна, и профилирование может помочь определить пиковую производительность для любого такого созвездия задачи GPU (различия могут варьироваться в пределах нескольких порядков - от лучшего к общему и худшему - среди всех возможных настроек "геометрии" для конкретной задачи).
Во-вторых, когда у меня много ядер, как openCL распределяет задачу, это на каждом ядре один и тот же процесс и одни и те же данные или другое ядро с разными данными?
Как объяснено вкратце выше - SIMT/SIMD кремниевая архитектура устройства такого типа не позволяет ни одному из SMX-ядер SMX выполнять что-либо, кроме той же самой инструкции SIMT / SIMD для всего "стада" SM-ядер, которое было отображено задачей- геометрия "на SMX-блок (не считая GPU_NOP (s) как делать " что-то еще ", так как это просто тратит впустую CPU: время системы GPU).
Так что, да, ".. на каждом ядре одного и того же процесса.." (лучше, если никогда не расходятся во внутренних путях выполнения кода после if или же whileили любой другой вид ветвления пути выполнения кода), поэтому, если алгоритм, основанный на управляемых данными значениях, приводит к различному внутреннему состоянию, каждое ядро может иметь различное локальное состояние потока, в зависимости от которого обработка может отличаться (как показано на примере if -расходные пути выполнения кода выше). Дополнительные сведения о локальных SM-регистрах, локальном SM-кэшировании, ограниченном использовании общей памяти (и затратах на задержку), использовании глобальной памяти устройства GPU (а также затратах на задержку, длине строк кэширования и ассоциативности для лучшего объединения шаблонов доступа для задержки опции маскировки - многие сведения об аппаратной части + программирование экосистемы входят в небольшие тысячи страниц документации по аппаратному обеспечению + программному обеспечению и выходят далеко за рамки этого упрощенного для ясности сообщения)
одни и те же данные или это разные ядра с разными данными?
Это последняя, но не менее важная дилемма: любая хорошо параметризованная активация ядра GPU может также передавать некоторое количество данных внешнего мира в ядро GPU, что может отличать локальные данные потока SMX от ядра SM до ядра SM. ядро. Методы сопоставления и лучшая производительность для этого в основном зависят от устройства ( { SMX | SM-регистры | GPU_GDDR gloMEM: shaMEM: constMEM | GPU-локальная иерархия кэша SMX}-детали и возможности
...
+---------------------------------------------------------
... 901 MHz
Cache type: Read/Write
Cache line size: 128
Cache size: 32768
Global memory size: 4294967296
Constant buffer size: 65536
Max number of constant args: 9
Local memory size: 49152
+---------------------------------------------------------
... 4000 MHz
Cache type: Read/Write
Cache line size: 64
Cache size: 262144
Global memory size: 536838144
Constant buffer size: 131072
Max number of constant args: 480
Local memory size: 32768
+---------------------------------------------------------
... 1300 MHz
Cache type: Read/Write
Cache line size: 64
Cache size: 262144
Global memory size: 1561123226
Constant buffer size: 65536
Max number of constant args: 8
Local memory size: 65536
+---------------------------------------------------------
... 4000 MHz
Cache type: Read/Write
Cache line size: 64
Cache size: 262144
Global memory size: 2147352576
Constant buffer size: 131072
Max number of constant args: 480
Local memory size: 32768
принципиально настолько разные устройства для устройства, что каждый высокопроизводительный кодовый проект принципиально может, но профилировать свою соответствующую задачу задачи устройства GPU - "составление карт геометрии и использования ресурсов для фактического устройства развертывания. Что может работать быстрее на одном устройстве GPU / GPU-привод стека, не нужно работать умнее на другом (или после GPU-драйвера + экзо-программирование обновления / обновления экосистемы), просто скажет только реальный тест (как теория может быть легко напечатана, но вряд ли как легко выполняется, так как многие ограничения, связанные с конкретным устройством и рабочей нагрузкой, будут применяться в реальном развертывании).
Предложите мне выход для распределения функции с разными значениями или фреймами в MQL5 с использованием графического процессора в OpenCL.
Честные и лучшие предложения такие же, как они были представлены вам уже 2 апреля.
Не пытайтесь заблокировать / задержать поток выполнения любого MQL5 CustomIndicator -тип блока выполнения кода с любым кодом задержки / асинхронности / блокировки. Никогда, пока документация платформы Терминала MetaTrader явно не удалит такие предупреждения (присутствующие там еще в 2018/Q2) и не даст явных советов по методам, использующим неблокируемые средства связи для неблокирующих распределенных агентов для скоординированного (почти)-синхронного обмена данными обработки / результаты между стороной MQL5 и стороной устройства GPU (которые скоро будут недоступны из-за SIMT/SIMD характер внепланового планирования GPU-заданий в современных классах доступных GPU-устройств.
Это было задокументировано для естественного потока времени, вызванного потоком внешних событий на валютном рынке (распространяемых через брокера), насчитывающих около нескольких сотен [us] Каденция от события к событию.
Если перейти к синтетическому течению времени, как это организовано в экосистеме симулятора [ Strategy Tester ] Терминала, проблема, описанная выше, будет на много порядков хуже, так как симулятор фактически ускоряет поток времени / частоту и что-либо, не способное идти в ногу, будет (снова) блокировать любое ускорение (что уже было плохо в естественном темпе вышеупомянутого потока времени). Итак, нет, это очень плохое направление, чтобы вкладывать деньги в очередной шаг усилий (опять же, по крайней мере, пока обе платформы не изменят свои архитектурные ограничения).
... так что мой процесс стал быстрее...
Эта часть определения проблемы была решена доктором Джином AMDAHL уже ~ 60 лет назад.
Его (затем упрощенный) Закон убывающей отдачи объясняет, ПОЧЕМУ основной потолок любого ускорения процесса связан с неподвижным [SERIAL] часть, учитывая различие между [SERIAL] часть и потенциально N-(CPU) раз истинно [PARALLEL] части четко обозначены.
Это помогает предварительно оценить влияние затрат / выгод на процесс реинжиниринга.
Итак, здесь ваш код ядра GPU является своего рода (почти)- [PARALLEL] обрабатывающая часть. Все остальное пока чисто [SERIAL] обрабатывающая часть.
Этого достаточно, чтобы угадать пределы эффекта от попыток перейти к перепроектированию процессов в OpenCL.
Но дьявол скрыт в деталях...
Реальные затраты намного выше.
-
[SERIAL]часть никогда не станет быстрее как таковая. -
[SERIAL]-part на самом деле будет "медленнее" и "расширеннее", так как будет еще много шагов для выполнения, прежде чем первая SIMT / SIMD-инструкция полезной нагрузки (ий)… будет "удаленно" доставлена в OpenCL-очередь + OpenCL-Data-Transfer(s) + Ожидание управления задачами OpenCL-Queue... + Отправка управления задачами OpenCL-Queue на устройство... даже начнет выполняться + задача == запланированное OpenCL-Device WARP-запланировано / SIMT/SIMD-выполнение + весь путь назад, из удаленного цирка --- издержки завершения задачи OpenCL-Device + задержки асинхронного обнаружения завершения асинхронного завершения на стороне MQL5 + OpenCL-Data-Transfer (s) -
[PARALLEL]-part будет выполняться только " после " или " при " всех начисленных дополнительных издержках (не показано на рисунке выше, из-за необходимости избегать усложнения и усложнения понимания предела теоретических накладных расходов). игнорирование, ускорение (не-) масштабирование), но еще хуже, поскольку выполнение выполняется только примерно в ~ 4 раза нижеGPU_CLOCK-рейт (не говоря уже о~ 10x ~ 1000xболее медленное время задержки доступа к памяти и кешу), и, как и там, телепортированный алгоритм все еще остается[SERIAL]- только линейно-извилистая обработка данных TimeSeries, таким образом, не может иметь, но отрицательный чистый эффект<< 1.0фактор улучшения теоретического ускорения обработки (достигнутая результирующая производительность становится хуже, чем без такой попытки "улучшить").
Для полного ознакомления с этими чистыми эффектами, пожалуйста, прочитайте раздел " Критика", где более подробное изложение переформулировки закона Амдала в строгих накладных расходах, а также переформулирование в строгих накладных расходах и с учетом ресурсов:
1
S = __________________________; where s, ( 1 - s ), N were defined above
( 1 - s ) pSO:= [PAR]-Setup-Overhead add-on
s + pSO + _________ + pTO pTO:= [PAR]-Terminate-Overhead add-on
N
1 where s, ( 1 - s ), N
S = ______________________________________________ ; pSO, pTO
/ ( 1 - s ) \ were defined above
s + pSO + max| _________ , atomicP | + pTO atomicP:= further indivisible duration of atomic-process-block
\ N /
Граф заголовка, приведенный в верхней части этого поста, предоставляет ссылку на живой графический интерфейс с интерактивными входами и анимированными выходами, где можно проверить влияние значений для p == ( 1 - s )где-то ниже 1,00 (что является просто теоретическим, абсолютно 100% [PARALLEL] график (что технически невозможно в любом реальном сценарии)), а также настройка влияния всех дополнительных издержек в o (выражается как просто скалярная дробь по соображениям простоты) в пределах редактируемого диапазона ~ < 0.0 ~ 0.0001 > ценности, чтобы лучше понять основные ограничения реального поведения многоядерных устройств и стать способными принимать лучшие инженерные решения, прежде чем даже думать о каких-либо этапах кодирования.
И учитывая известный (легко измеримый до одного [us] -решение на стороне MQL5 выполнения кода, используя вызов GetMicrosecondCount()) значения для дополнительных накладных расходов и атомарности обработки - pSO, pTO, atomicP - чистый эффект от попытки продвинуться к простой скользящей средней в оболочке OpenCL, как было показано в коде ядра GPU:
kernel void SMA_executeSMA( float ExtLineBufferi_1,
float price1,
float price2,
int InpMAPeriod,
__global float *output
)
{ // 1: .STO 0x0001, REG
int len = get_global_id( 1 ); // 2: .JMP intrinsic_OpenCL_fun(), ... may get masked by reading a hardwired-const-ID#
// 3: .GET len, REG
output[len] = // 4: .STO MEM[*],
ExtLineBufferi_1 // 5: .ADD const,
+ ( price1 - price2 ) // ( .SUB const, const
/ InpMAPeriod; // .FDIV REG, const )
} // 6: .RET
который имеет только несколько команд с тактовой частотой 900 МГц - т.е. p = ( 1 - s ) -фактор в анимированном графике-визуализация будет идти где-то близко к p == 0 конец, что делает игру в конечном итоге доминирующим [SERIAL] часть ЦП:GPU-состав распределенной вычислительной системы - ( ~ несколько, макс. небольшие десятки [ns] + голый (немаскируемый, поскольку здесь нет повторного использования) задержка доступа к памяти на GPU-устройстве ~ 350 - 700+ [ns]).
Имея такой низкий p плохой знак настройки производительности (если не ANTI-PATTERN) для любых попыток сделать это.
Потому что даже если идти в N-(CPUs) ~ +INF, он все равно никогда не достигнет желаемого ускорения (ссылка: может попытаться изменить такие факторы в интерактивном графике, предложенном выше, и визуально увидеть эффект - насколько низкими будут цифры), хотя то же самое можно было бы вычислить. почти меньше, чем ~ 0.5 [ns] Кроме того, все еще векторизованные, инструкции процессора, здесь также вообще без затрат на добавление).
Это причины "экономии затрат"
(кроме основного MQL5 1) ПОЧЕМУ лучше этого не делать
что никогда не окупит сумму всех [SERIAL] дополнительные расходы, введенные в течение всего цирка OpenCL-re-wrapping-there-send-there-calc-and-after-found-back-send на стороне CPU-code / MQL5 (все от имени делая не больше, чем просто эти действительно очень немногие GPU_INSTR-ы), которые были только кратко упомянуты выше, даже если использовалось бесконечное количество ядер GPU.
Вы просто все еще пытаетесь заплатить больше, чем когда-либо получите.