performance - MQL5 中的 OpenCL 代码不会将分布式作业分配给每个 GPU 内核

Question

我使用 OpenCL 和 MQL5 为 MetaTrader 终端平台创建了一个基于 GPU 的指标。

我已经努力让我的 [ MetaTrader Terminal: Strategy Tester ] 优化工作必须在 GPU 上传输到最大。大多数计算都是由指标完成的。因此，我对指标进行了更改，并已完全转移到 GPU 上。

但是当我尝试在策略测试器部分进行优化时，真正的问题就出现了。
我看到的过程同时使用了我的 GPU 和 CPU，但对整个过程没有影响。

我怀疑该进程没有分配给每个 GPU 核心进行处理，而是所有 GPU 核心都在同一个进程或函数上工作以执行。

请让我知道我需要做什么才能让单个 GPU 工作以执行单个函数以提供更快的输出。

这是我的代码链接：Expert 的完整代码

我的代码的内核是：

__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;
            }
          }

请建议我使用 OpenCL 上的 GPU 在 MQL5 中分配具有不同值或帧的函数的出路。

已编辑

对于寻求挑战的人来说，这是一个巨大的挑战……甚至我也很想知道 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);
      }

Answer 1

帮助我以正确的方式编写此代码，以便我的过程变得更快，并使用我的 GPU 给出正确的结果。

首先，事实是，
自 4 月 2日^以来
，您已经知道，
MQL5 CustomIndicator不会以这种方式工作...

事实很重要- 如果 MQL4/5 代码执行架构已明确记录，任何扩展延迟/异步/阻塞操作都不能放置在任何 MQL4/5CustomIndicator代码执行单元中，那么可能有一个合理的是时候停止所有这些 S/O 社区成员的攻击并开始简单地尊重记录和公布的事实了。

MQL5文档非常明确并警告其单共享线程架构对性能的不利影响和/或警告整个系统完全、无意死锁的风险：

在一个交易品种上计算的所有指标，即使它们附加到不同的图表，也在同一个线程中工作。因此，一个品种上的所有指标共享一个线程的资源。

---

一个指标中的无限循环（阻塞/增加延迟/意外延迟）将停止该符号上的所有其他指标。

当然，如果他们决定进一步花费宝贵的时间来赞助和扩展此类工作，人们可能会要求 S/O 社区中的许多善良的 OpenCL 专业人士提供帮助，并表达对获得他们赞助的知识等的极大兴趣。

在任何合理的努力（在任何表达的力量下）至少可以开始朝着正确的方向工作之前，所有这一切仍然必须符合现实并最好地尊重已知事实。

---

CPU：GPU 拓扑如何在内部工作？

使用异步、无序交付操作、定义为异步、队列存储、零对多GPU 设备目标计算作业的序列，所有这些都具有从头到尾主要是不确定的 RTT -期间。

GPU 设备访问队列的 CPU 端可以被指示向 GPU 发送作业（一个任务 ~ 一个要执行的程序）：

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 架构的领域：

GPU 计算设备与任何通用 CPU CISC/RISC 计算设备不同，具有其他硅硬连线架构。

WHY的原因在这里非常重要。

GPU 设备使用流式多处理器 e X执行单元（SMX单元），在某些硬件检查工具中引用了这些单元。

虽然S M X 缩写中的字母M强调，有多个执行可加载到 SMX 单元上，然而，所有这些情况实际上都会执行（当然，只有以这种方式指示，这超出了本文的范围主题，涵盖/跨越每个 SMX 存在的 SM 核心）完全相同的计算指令 - 这是它们可以操作的唯一方式 - 它被称为可实现的有限并行范围类型（​​共本地）仅在 SMX 的外围，其中单指令-多- {线程| dSIMT/SIMData } 可以在当前 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) ]由自治、异步分布式系统星形组成的正式形式。节点。

任何从 CPU 主机加载到 GPU 上的 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）和一次执行一个，单步执行所述 GPU 作业的 SIMD 指令流。同一 SMX 单元上的所有其他 SM 核心通常在此期间什么都不做，直到这个 GPU 作业完成并且内部 GPU 进程管理系统决定为这个 SMX 映射一些其他工作。

• 如果2-SM 核心被指示执行一些 GPU 作业，那么这对 SM 核心将一个接一个地获取一个（并且非常相同的）GPU-RISC 指令（为了简单起见，忽略任何可能的 ILP）并且两者一次执行一个，单步执行所述 GPU 作业的 SIMT/SIMD 指令流。在这种情况下，如果一个 SM-core 进入一个条件，其中一个if-ed 或类似的分支，执行流程使一个 SM-core 进入另一个代码执行流程路径而不是另一个，SIMT/SIMD -并行性进入不同的场景，其中一个 SM 核心得到下一个 SIMT/SIMD 指令，属于它的代码执行路径，而另一个什么都不做（得到一个GPU_NOP（s）），直到第一个完成整个工作（或在等待从“远”（慢）获取一条数据时，强制停止在某个同步障碍处陷入不可屏蔽的延迟等待状态非本地内存位置，同样，细节超出了这篇文章的范围）——只有在任何一个发生之后，发散路径，到目前为止，刚刚GPU_NOP-ed SM-core 可以接收任何下一个 SIMT/SIMD 指令，属于其（发散的）代码执行路径以向前移动。同一 SMX 单元上的所有其他 SM 核心通常在此期间什么都不做，直到这个 GPU 作业完成并且内部 GPU 进程管理系统决定为这个 SMX 映射一些其他工作。

• 如果任务特定的“几何”指示16 个 SM 内核执行一些 GPU 作业，那么这组 SM 内核将在之后获取一个（并且非常相同的）GPU-RISC SIMT/SIMD 指令另一个（为了简单起见，忽略任何可能的 ILP）并一次执行一个，逐步执行所述 GPU 作业的 SIMT/SIMD 指令流。“牛群”内部的任何分歧都会降低 SIMT/SIMD 效应，并且GPU_NOP阻塞的核心仍在等待“牛群”的主要部分完成工作（与此点上方的草图相同）。

• 如果任务特定的“几何”指示更多的 SIMT/SIMD-threads-than-SM-cores-available执行某些 GPU 作业，则 GPU 设备芯片将运行它以作为[SERIAL]-sequence of as many { WARP-wide | half-WARP-wide }- SIMT/SIMD-thread 打包，直到这样的序列完成映射到 SMX 上的所有指令数量的 SIMT/SIMD-threads。因此，这种包统一完成的时间一致性原则上是不可能的，因为它们以 WARP 调度程序特定的方式到达各自的终点，但从不同步（是的，您的 CPU 端代码必须等到最后一个（最懒惰的（由于任何原因，无论是容量受限的调度原因，还是代码发散调度原因或糟糕的相互（重新）同步原因）代码执行流程）最终将在未来某个未知时间完成__kernel代码处理，并且 OpenCL 操作的设备将允许“远程”检测CL_COMPLETE状态，然后才能获取任何有意义的结果（如您所问对您的其他问题之一感到惊讶）。

无论如何，所有其他 SM 核心，没有被相应 GPU 设备的 SMX 单元上的特定任务“几何”映射，通常会毫无用处 - 所以了解正确任务的硬件细节的重要性- 特定的“几何”确实很重要，分析可能有助于确定任何此类 GPU 任务星座的峰值性能（差异可能在几个数量级 - 从最佳到常见到更差 - 在所有可能的特定任务“几何”设置中）。

---

其次，当我有很多核心时，openCL如何分配任务，是在每个核心上相同的进程和相同的数据还是不同的核心和不同的数据？

如上所述 -SIMT/SIMD类型设备硅架构不允许任何 SMX SM 内核在整个“群”SM 内核上执行除完全相同的 SIMT/SIMD 指令之外的任何内容，即由任务“几何”映射到 SMX 单元（不将GPU_NOP(s) 算作“其他事情”，因为它只是在浪费 CPU:GPU 系统时间）。

所以，是的， “..在每个核心相同的进程上..”（最好，如果在其内部代码执行路径在或任何其他类型的代码执行路径分支之后永远不会分歧if）while，那么如果算法基于数据驱动值会导致不同的内部状态，每个内核可能具有不同的线程本地状态，基于此处理可能会有所不同（例如if上面驱动的不同代码执行路径）。有关 SM 本地寄存器、SM 本地缓存、受限共享内存使用（和延迟成本）、GPU 设备全局内存使用（以及延迟成本和缓存线长度以及用于延迟的最佳合并访问模式的关联性）的更多详细信息屏蔽选项 - 许多与硬件相关的 + 编程生态系统细节都包含在数千页的硬件 + 软件特定文档中，并且远远超出了为清楚起见而简化后的范围）

相同的数据还是具有不同数据的不同核心？

这是最后但并非最不重要的两难境地——任何参数化良好的 GPU 内核激活也可能将一些外部世界数据向下传递到 GPU 内核，这可能使 SMX 线程本地数据不同于 SM 内核和 SM-核。执行此操作的映射实践和最佳性能主要是特定于设备的（ { SMX | SM-registers | GPU_GDDR gloMEM : shaMEM : constMEM | GPU SMX-local cache-hierarchy }-详细信息和容量

  ...
 +---------------------------------------------------------
  ...                                               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 驱动程序 + 外部编程生态系统更新/升级之后），只是只有现实生活中的基准测试会告诉你（因为理论可以很容易地打印出来，但几乎没有易于执行，因为许多特定于设备和工作负载注入的限制将适用于实际部署）。

---

建议我使用 OpenCL 上的 GPU 在 MQL5 中分配具有不同值或帧的函数的出路。

诚实和最佳的建议与 4 月 2 日已经提供给您的建议完全相同。

不要尝试CustomIndicator使用任何扩展延迟/异步/阻塞代码来阻塞/延迟任何 MQL5 类型的代码执行单元的执行流程。永远不要，直到 MetaTrader 终端平台文档明确删除此类警告（仍然存在于 2018/Q2 那里），并将明确建议使用避免延迟的非阻塞分布式代理通信工具来协调（几乎）同步交换处理数据的技术/ MQL5 端和 GPU 设备端之间的结果（由于SIMT/SIMD当前可用的 GPU 设备类别中 GPU 作业的无序调度的性质，这将不会很快可用。

这被记录为自然时间流，由外部外汇市场（经纪人广播传播）事件流选通，具有大约数百个[us]事件到事件的节奏。

如果进入合成时间流，就像在终端的[策略测试器]模拟器生态系统中精心安排的那样，上面记录的问题会变得更糟，因为模拟器实际上加速了时间流/节奏任何无法跟上步伐的东西都会（再次）阻碍任何加速（这在上述时间流的自然节奏中已经很糟糕了）。所以，不，这是一个非常糟糕的投资方向（再次，至少在两个平台都改变了它们的架构限制之前）。

---

...让我的过程变得更快...

这部分问题定义已经在 60 年前由 Gene AMDAHL 博士决定。

他的（然后简化的）收益递减定律解释了为什么任何进程加速的主要上限都与静止[SERIAL]部分相关联，因为明确识别了纯[SERIAL]部分和潜在的 N-(CPU) 次真实部分[PARALLEL]之间的区别。

这有助于预先估计流程再造的成本/收益效果。

因此，在这里，您的 GPU 内核代码是（几乎）[PARALLEL]处理部分。其余的仍然是纯[SERIAL]处理部分。

这足以猜测尝试进入 OpenCL 包装的流程重新设计的效果限制。

---

但是，魔鬼隐藏在细节之中……

实际成本要高得多。

• -part[SERIAL]本身永远不会变得更快。
• -part[SERIAL]实际上会变得“更慢”和“扩展”，因为在有效载荷的第一个 SIMT/SIMD 指令......被“远程”传递到 OpenCL 之前，还有更多的步骤需要执行队列 + OpenCL-Data-Transfer(s) + OpenCL-Queue Task Management waiting... + OpenCL-Queue TaskManagement 提交到设备... 甚至会开始执行 + 任务 == 预期的 OpenCL-Device WARP-scheduled / SIMT/SIMD-执行 + 从远程马戏团一路返回 --- OpenCL-Device 任务完成开销 + MQL5 端异步完成检测异步附加延迟 + OpenCL-Data-Transfer(s)
• -part[PARALLEL]将仅在“之后”或“在”所有附加成本累积后执行（上图中未显示，因为需要避免使其过于复杂和难以理解理论的限制，忽略开销，加速（不）缩放），更糟糕的是，执行速度仅降低了约 4 倍GPU_CLOCK（更不用说对内存和缓存的访问延迟时间~ 10x ~ 1000x更慢），并且因为那里的传送算法仍然是[SERIAL]-只有线性卷积的时间序列数据处理，因此只能产生不利的净效应<< 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               /

本文顶部引用的标题图提供了指向具有交互式输入和动画输出的实时 GUI 的链接，人们可以在其中测试值p == ( 1 - s )在 1.00 以下的任何地方的影响（这只是理论上的，绝对 100% 的[PARALLEL]时间表（这在任何现实世界的场景中在技术上都是不可能的））并且还在可编辑的值范围内调整所有附加开销的影响o（为简单起见仅表示为标量分数）~ < 0.0 ~ 0.0001 >，以便更好地感知主要限制多核设备的真实世界行为，并能够在考虑任何编码步骤之前做出更好的工程决策。

---

并且考虑到附加开销和处理原子性的已知值（[us]在代码执行的 MQL5 端可轻松测量到单一分辨率，使用调用） ——尝试的净效应继续朝着GPU-kernel-code 中描绘的OpenCL 包装的简单移动平均线：GetMicrosecondCount()pSO, pTO, atomicP

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 MHz 时钟指令 - 即p = ( 1 - s )动画图形可视化中的因素将接近p == 0尾声，使游戏最终由 CPU 的纯部分主导：分布式[SERIAL]的 GPU 组成-计算系统--（〜几个，最多几十个[ns]+裸（不可屏蔽，因为这里有零重用）GPU设备内存访问延迟~ 350 - 700+ [ns]）。

对于任何这样做的尝试，拥有如此低p的性能是一个调整性能的坏信号（如果不是ANTI-PATTERN的话）。

因为即使进入N-(CPUs) ~ +INF，它仍然永远不会实现期望的加速（参考：可能会尝试在上面提供的交互式图表中修改这些因素并直观地看到效果——那里的数字会有多低）——而相同的计算可以在几乎小于~ 0.5 [ns]、更可向量化的 CPU 指令中计算，这里也完全具有零附加成本）。

这些是“成本经济”的原因
_{（除了主要的MQL5）}为什么最好不这样做

这永远不会偿还所有[SERIAL]附加成本的总和，在 CPU 代码上的整个 OpenCL-re-wrapping-there-sending-there-calc'd-and-after-detected-back-sending circus 中引入/MQL5 方面（所有这些都是为了让这些确实很少的 GPU_INSTR-s 发生），即使使用了无限数量的 GPU 核心，上面也只是简单地提到过。

您仍然只是尝试支付超过一个人将收到的回报的方式。