外汇EA编写教程:OpenCL:通往平行世界的桥梁

简介

本文是关于OpenCL或开放操作语言编程的一系列简短出版物中的第一篇。在提供对OpenCL的支持之前,当前形式的MetaTrader 5平台不允许直接(即本地)享受多核处理器加速的优势。

显然,开发人员总是可以重复“这个终端是多线程的”,并且“每个EA/脚本都在一个单独的线程中运行”,但是编码人员没有机会相对简单地并行执行以下简单循环(计算pi=3.14159265…):

long num_steps = 1000000000;
double step = 1.0 / num_steps;
double x, pi, sum = 0.0;
 
for (long i = 0; i<num_steps; i++)
{
   x = (i + 0.5)*step;
   sum += 4.0/(1.0 + x*x);
}
pi = sum*step;

然而,早在18个月前,一篇题为《元交易者5中的并行计算》的有趣文章出现在"第A部分第6035条中。然而,本文的印象是,尽管这种方法设计得很精细,但有点不自然——在上面的循环中编写的整个程序层次结构(EA事务和两个指标)以加快计算速度可能会适得其反。

我们已经了解到没有计划支持OpenMP,并且添加OMP需要对编译器进行大量的重新编程。唉,难道没有易学、简单和负担得起的代码解决方案吗?

因此,一旦宣布在QL5中对OpenCL最初的支持,它就受到了热烈的欢迎。从同一个新闻稿的第22页,metadriver开始发布脚本,用于评估CPU和GPU之间的实现差异。有一段时间,OpenCL引起了广泛的关注。

起初,作者决定不去注意这些现象。非常低端的计算机配置(Pentium G840/8GB DDR-II1333/无图形卡)似乎无法实现OpenCL的有效使用。

然而,在安装了AMD应用软件开发包(AMD加速并行计算软件开发包)之后,AMD开发的一个专业软件(第一个脚本是由元驱动程序提出的,只有在有独立图形卡的情况下才能工作)已经在作者的计算机上成功运行,其速度的提高也远远不能忽略。与单核处理器上的标准脚本运行时间相比。快25倍!后来,由于在支持团队的帮助下成功安装了Intel OpenCL Runtime,同一脚本的运行时间加快了75倍。

在仔细研究了ixbt.com论坛和网站上的信息后,作者发现英特尔的集成图形处理器(IGP)支持Opencl 1.1,但只支持上面的Ivy桥接处理器和处理器。这样,在具有上述配置的PC机上实现的加速度与IGP无关,在这种特殊情况下,opencl程序代码只能在86核的CPU上执行。

当作者与ixbt.com的专家合作研究加速数据时,他们立即回应说,这实际上是由于源语言(mql5)的优化不够。众所周知,如果用++语言对源代码进行了正确的优化(当然,假设使用了一些多核处理器和SSEX矢量指令),那么在最佳状态下,OpenCL的模拟效果可以提高几十个百分点;在最坏的情况下,甚至数据丢失(如数据传输)也可以得到改善。可能会发生轻微的费用)。时间太长)。

因此还有另一个假设:应充分注意OpenCL在MetaTrader 5中完全模拟的“难以置信的加速数据”,而不是将其归因于OpenCL自身的卓越。为了充分发挥GPU在++程序中合理优化的强大优势,它只能通过一个功能强大的独立图形卡来实现,因为它在某些算法中的计算能力远远超过目前任何CPU的计算能力。

根据终端开发者的说法,它还没有得到合理的优化。他们还建议优化后的加速度将增加几倍。Opencl中的加速数据也将相应地减少几倍“和多次”。但它们仍然比一个大得多。

这是学习Opencl语言的一个很好的理由(即使您的图形卡不支持Opencl 1.1或根本没有图形卡),我们将讨论这个方面。但首先,我想简要介绍一下必要的基础知识,即支持开放式cl的软件和相应的硬件。

1。必要的硬件和软件

AM6033

相应的软件分别由AMD、Intel和Nvidia提供。这三家公司都是Khronos Group的成员,Khronos Group是一家非营利性行业协会成员,负责制定与异构环境中计算相关的不同语言规范。

Khronos集团的官方网站提供了一些有用的信息,例如:

  • Opencl 1.1规范,
  • Opencl 1.1参考。

这些文档通常用于学习opencl,因为终端没有提供任何与opencl相关的帮助信息(只提供了opencl api的摘要)。这三家公司(AMD、Intel和Nvidia)都是视频硬件供应商,每个公司都有自己的OpenCL运行时实现和自己的软件开发工具包SDK。让我们以AMD产品为例来解释图形卡选择的一些特性。

如果您的ADM图形卡不是特别旧(从2009-2010年或更高版本开始),那么很容易升级图形卡,并且它足以立即工作。有关与OpenCL兼容的图形卡的列表,请参阅此处。另一方面,即使是好的图形卡,如RadeonHD4850(4870),在处理OpenCL时也有很多问题。

如果您还没有AMD图形卡,但想要添加它,请先检查它的规格。您可以在这里看到一个相当全面的“最新AMD图形卡规格表”。以下是最重要的考虑事项:

  • 车载存储器&mdash;本地存储器的数量。越大越好。通常1 GB就足够了。
  • 核心时钟(核心频率)核心频率。这也很清楚:GPU多处理器的频率越高,越好。650-700兆赫也不错。
  • [内存]类型([内存]类型)&mdash;内存类型。最好的显示存储器是高速显示存储器,即GDDR5。但是GDDR3也很好,尽管它的显示带宽比GDDR5低两倍。
  • [内存]时钟(eff.)([显示]频率(eff.))-内存的工作(有效)频率。从技术上讲,这个参数与前一个参数密切相关。平均而言,GDDR5的工作频率是GDDR3的两倍。这与显示类型以更高的频率工作这一事实无关,但这是由于显示内存使用的数据传输通道的数量所致。换句话说,它与显示内存的带宽有关。
  • [内存]总线([内存]总线)-总线的数据带宽。比256强。
  • mbw&mdash;显示内存带宽。这个参数实际上是上述三个内存参数的组合。越高越好。
  • 配置核心(spu:tmu(tf:rop)&mdash;配置GPU核心单元。我们的价值(在非图形操作的情况下)是第一个数字。1024:64:32意味着我们需要1024个(统一流处理器或流处理器的数量)。显然,越高越好。
  • 处理能力(处理能力)浮点运算的理论性能(fp32(单精度)/fp64(双精度)。规格表始终包含对应于fp32的值(所有图形卡都可以单精度计算),这与fp64明显不同,因为并非每个图形卡都支持双精度。如果您确定在GPU操作中根本不需要双精度(双类型),则可以忽略第二个参数。但在这两种情况下,参数越高越好。
  • TDP&mdash;热设计功耗。一般来说,这是指图形卡在最复杂的操作中消耗的最大功率。如果EA事务频繁访问GPU,那么图形卡不仅会消耗大量的电源(如果返回良好),而且会产生相当大的噪声。

现在让我们来看第二个场景:没有图形卡,或者现有的图形卡不支持Opencl1.1,但是您有一个AMD处理器。您可以下载AMD应用软件开发工具包,其中包括运行时和软件开发工具包、内核分析器和性能评估器。安装AMD应用软件开发工具包后,处理器将被标识为OpenCL设备。此外,您将能够在CPU模拟模式下开发全功能的OpenCL应用程序。

与AMD不同,SDK的主要功能是它还可以与Intel处理器兼容(尽管在Intel CPU上开发本机SDK时效率会显著提高,因为它支持最近才可用于AMD处理器的SSE 4.1、SSE 4.2和AVX指令集)。

1.2。英特尔

最好在开始使用Intel处理器之前下载Intel Opencl SDK/Runtime。

我们必须指出以下几个方面:

  • 如果计划仅使用CPU(OpenCL模拟模式)开发OpenCL应用程序,请理解Intel CPU图形核心在Sandy Bridge处理器(包括此版本的处理器)之前不支持OpenCL 1.1。这种支持仅限于常春藤桥处理器,但即使是最强大的IntelHD4000集成图形单元,也很难做出任何改变。对于Ivy桥之前的处理器,在MQL5环境中可以实现的加速仅是由于使用了s s(s)ex矢量指令。但它看起来仍然很明显。
  • 安装Intel Opencl SDK后,必须对注册表项hkey_local_machine/software/khronos/opencl/vendors进行以下修改:将“名称”列中的intel opencl64.dll替换为intelocl.dll。然后重新启动并打开MetaTrader 5。CPU现在被识别为Opencl 1.1设备。

老实说,Opencl对Intel的支持还没有完全解决,所以我们希望终端的开发人员在未来给出一些解释。基本上,关键是没有人会为您关注内核代码错误(OpenCL内核是在GPU上执行的程序)-它不是一个MQL5编译器。编译器将只获取一整行内核,然后尝试运行它。例如,如果您不声明内核中使用的内部变量,那么内核仍然会被严格执行,并且还会出现错误。

但是,您将在终端中遇到的所有错误都将减少到大约12个,其中一些错误已在有关clkernel create()和clp rogram create()函数的api opencl相关帮助中进行了描述。语言语法与C非常相似,并通过向量函数和数据类型加以增强(事实上,1999年该语言被采用为ANSI C标准)。

在调试Opencl代码时,作者使用了Intel Opencl SDK脱机编译器;它比在Metaeditor中盲目查找内核错误方便得多。我希望将来情况会好转。

1.3。NVIDIA

不幸的是,作者没有搜索有关此主题的信息。然而,总体建议仍然是相同的。新NVIDIA图形卡的驱动程序将自动支持OpenCL。

基本上,作者对Nvidia图形卡没有异议,但根据在论坛讨论中发现的信息和获得的知识,得出以下结论:AMD图形卡在非图形操作的成本性能上似乎优于Nvidia图形卡。

现在我们来谈谈编程。

2。使用opencl的第一个MQL5程序

为了开发第一个简单的程序,我们需要用这种方式定义任务。在并行编程的过程中,我们必须使用pi值的计算结果(例如,大约3.14159265)。

为此,应采用以下公式(作者以前从未见过这样的特殊公式,但似乎是真的):

Pi 计算公式

我们想把这个值计算到12位小数。基本上,这个精度可以通过大约一百万次迭代来实现,但是这个数字并不能让我们评估OpenCL中操作的优势,因为GPU上的操作持续时间太短。

在GPGPU编程中,建议选择相应的计算量,以便GPU任务至少持续20毫秒。在这种情况下,由于getTickCount()函数与100 ms相比存在显著错误,因此应将限制值设置得更高。

以下MQL5程序用于实现此操作:

//+------------------------------------------------------------------+
//|                                                           pi.mq5 |
//+------------------------------------------------------------------+
#property copyright "Copyright (c) 2012, Mthmt"
#property link      "https://www.mql5.com"

long     _num_steps        = 1000000000;
long     _divisor          = 40000;
double   _step             = 1.0 / _num_steps;
long     _intrnCnt         = _num_steps / _divisor;
//+------------------------------------------------------------------+
//| Script program start function                                    |
//+------------------------------------------------------------------+
int OnStart()
  {
   uint start,stop;
   double x,pi,sum=0.0;

   start=GetTickCount();
//--- first option - direct calculation
   for(long i=0; i<_num_steps; i++)
     {
      x=(i+0.5)*_step;
      sum+=4.0/(1.+x*x);
     }
   pi=sum*_step;
   stop=GetTickCount();

   Print("The value of PI is "+DoubleToString(pi,12));
   Print("The time to calculate PI was "+DoubleToString(( stop-start)/1000.0,3)+" seconds");

//--- calculate using the second option
   start=GetTickCount();
   sum=0.;
   long divisor=40000;
   long internalCnt=_num_steps/divisor;
   double partsum=0.;
   for(long i=0; i<divisor; i++)
     {
      partsum=0.;
      for(long j=i*internalCnt; j<(i+1)*internalCnt; j++)
        {
         x=(j+0.5)*_step;
         partsum+=4.0/(1.+x*x);
        }
      sum+=partsum;
     }
   pi=sum*_step;
   stop=GetTickCount();

   Print("The value of PI is "+DoubleToString(pi,12));
   Print("The time to calculate PI was "+DoubleToString(( stop-start)/1000.0,3)+" seconds");
   Print("_______________________________________________");
   return(0);
  }
//+------------------------------------------------------------------+

编译并运行此脚本后,我们得到以下结果:

2012.05.03 02:02:23    pi (EURUSD,H1)    The time to calculate PI was 8.783 seconds
2012.05.03 02:02:23    pi (EURUSD,H1)    The value of PI is 3.141592653590
2012.05.03 02:02:15    pi (EURUSD,H1)    The time to calculate PI was 7.940 seconds
2012.05.03 02:02:15    pi (EURUSD,H1)    The value of PI is 3.141592653590

pi值~3.14159265的计算方法略有不同。

第一种几乎是证明多线程库(如openmp、intel tpp、intel mkl等)功能的经典方法。

第二种是相同的计算,只以双环的形式。由10亿次迭代组成的整个计算过程被分解成几个大型的外部循环块(包括40000次迭代),每个循环块执行25000次迭代,形成一个内部循环:ldquo;基本&rdquo;迭代。

可以看出,操作慢了10-15%。但是,在转换为OpenCL时,我们打算以这个特定的计算为基础。主要原因在于内核的选择(GPU上的基本计算任务),它在数据从一个存储区域传输到另一个存储区域所需的时间与运行内核所需的计算量之间实现了合理的权衡。因此,对于当前的任务,一般来说,内核将成为第二个计算算法的内部循环。

现在,我们使用opencl来计算这个值。完整的程序代码后面将简要描述绑定到OpenCL的宿主语言(MQL5)函数的特性。但首先,我想强调几个与可能干扰OpenCL编码的典型“屏障”相关的要点:

  1. 内核外没有声明变量。因此,全局变量step和intrncnt必须在内核代码的开头再次声明(见下文)。此外,它们各自的值必须转换为字符串,以适合在内核代码中读取。然而,在OpenCL中编程的这一特性在未来将非常有用,例如在生成不存在于C语言中的向量数据类型时。
  2. 尽量给内核尽可能多的计算量,同时保持其数量的合理性。这对代码并不特别重要,因为在现有硬件上,内核在代码中的速度并不特别快。但是如果你使用一个强大的独立图形卡,这个因素将帮助你加速计算。

下面是OpenCL内核的脚本代码:

//+------------------------------------------------------------------+
//|                                                 OCL_pi_float.mq5 |
//+------------------------------------------------------------------+
#property copyright "Copyright (c) 2012, Mthmt"
#property link      "https://www.mql5.com"
#property version   "1.00"
#property script_show_inputs;

input int _device=0;        /// OpenCL device number (0, I have CPU)

#define  _num_steps        1000000000 
#define  _divisor          40000
#define  _step             1.0 / _num_steps
#define  _intrnCnt         _num_steps / _divisor

string d2s(double arg,int dig) { return DoubleToString(arg,dig); }
string i2s(int arg)            { return IntegerToString(arg); }

const string clSrc=
                   "#define _step "+d2s(_step,12)+"                   /r/n"
                   "#define _intrnCnt "+i2s(_intrnCnt)+"             /r/n"
                   "                                                   /r/n"
                   "__kernel void pi( __global float *out )            /r/n"   // type float
                   "{                                                  /r/n"
                   "  int i = get_global_id( 0 );                      /r/n"
                   "  float partsum = 0.0;                             /r/n"   // type float
                   "  float x = 0.0;                                   /r/n"   // type float
                   "  long from = i * _intrnCnt;                       /r/n"
                   "  long to = from + _intrnCnt;                      /r/n"
                   "  for( long j = from; j < to; j ++ )               /r/n"
                   "  {                                                /r/n"
                   "     x = ( j + 0.5 ) * _step;                      /r/n"
                   "     partsum += 4.0 / ( 1. + x * x );              /r/n"
                   "  }                                                /r/n"
                   "  out[ i ] = partsum;                              /r/n"
                   "}                                                  /r/n";
//+------------------------------------------------------------------+
//| Script program start function                                    |
//+------------------------------------------------------------------+
int OnStart()
  {
   Print("FLOAT: _step = "+d2s(_step,12)+"; _intrnCnt = "+i2s(_intrnCnt));
   int clCtx=CLContextCreate(_device);

   int clPrg = CLProgramCreate( clCtx, clSrc );
   int clKrn = CLKernelCreate( clPrg, "pi" );

   uint st=GetTickCount();

   int clMem=CLBufferCreate(clCtx,_divisor*sizeof(float),CL_MEM_READ_WRITE); // type float
   CLSetKernelArgMem(clKrn,0,clMem);

   const uint offs[ 1 ]  = { 0 };
   const uint works[ 1 ] = { _divisor };
   bool ex=CLExecute(clKrn,1,offs,works);
//--- Print( "CL program executed: " + ex );

   float buf[];                                          // type float
   ArrayResize(buf,_divisor);
   uint read=CLBufferRead(clMem,buf);
   Print("read = "+i2s(read)+" elements");

   float sum=0.0;                                        // type float
   for(int cnt=0; cnt<_divisor; cnt++) sum+=buf[cnt];
   float pi=float(sum*_step);                            // type float

   Print("pi = "+d2s(pi,12));

   CLBufferFree(clMem);
   CLKernelFree(clKrn);
   CLProgramFree(clPrg);
   CLContextFree(clCtx);

   double gone=(GetTickCount()-st)/1000.;
   Print("OpenCl: gone = "+d2s(gone,3)+" sec.");
   Print("________________________");

   return(0);
  }
//+------------------------------------------------------------------+

稍后将更详细地描述脚本代码。

同时编译和启动程序以获得以下信息:

2012.05.03 02:20:20    OCl_pi_float (EURUSD,H1)    ________________________
2012.05.03 02:20:20    OCl_pi_float (EURUSD,H1)    OpenCl: gone = 5.538 sec.
2012.05.03 02:20:20    OCl_pi_float (EURUSD,H1)    pi = 3.141622066498
2012.05.03 02:20:20    OCl_pi_float (EURUSD,H1)    read = 40000 elements
2012.05.03 02:20:15    OCl_pi_float (EURUSD,H1)    FLOAT: _step = 0.000000001000; _intrnCnt = 25000

如您所见,运行时间略有缩短。但我们不能满足这一点:它的π值约为3.14159265,显然只精确到最后三位小数。计算不准确的原因是在实际操作中,内核的精度低于所需的12位十进制浮点数据。

根据mql5文档,浮点数据的精度只有7个有效数字。双精度型数据精度达到15位有效数字。

因此,我们需要使实际数据类型“更加精确”。在上面的代码中,要用双精度类型替换的浮点类型数据行用注释&mdash;&mdash;/浮点类型标记。使用相同的输入数据编译后,我们将得到以下内容(ocl_pi_double.mq5,一个带有源代码的新文件):

2012.05.03 03:25:35    OCL_pi_double (EURUSD,H1)    ________________________
2012.05.03 03:25:35    OCL_pi_double (EURUSD,H1)    OpenCl: gone = 12.480 sec.
2012.05.03 03:25:35    OCL_pi_double (EURUSD,H1)    pi = 3.141592653590
2012.05.03 03:25:35    OCL_pi_double (EURUSD,H1)    read = 40000 elements
2012.05.03 03:25:23    OCL_pi_double (EURUSD,H1)    DOUBLE: _step = 0.000000001000; _intrnCnt = 25000

运行时间急剧增加,甚至超过了在没有OpenCL的情况下运行源代码所需的时间(8.783秒)。

&ldquo;很明显,双精度类型会降低计算速度&rdquo;您可能会这么认为。但是,让我们尝试将输入参数除数从40000显著更改为40000000:

2012.05.03 03:26:55    OCL_pi_double (EURUSD,H1)    ________________________
2012.05.03 03:26:55    OCL_pi_double (EURUSD,H1)    OpenCl: gone = 5.070 sec.
2012.05.03 03:26:55    OCL_pi_double (EURUSD,H1)    pi = 3.141592653590
2012.05.03 03:26:55    OCL_pi_double (EURUSD,H1)    read = 40000000 elements
2012.05.03 03:26:50    OCL_pi_double (EURUSD,H1)    DOUBLE: _step = 0.000000001000; _intrnCnt = 25

这不会影响精度,运行时间比浮点情况下稍短。但是,如果我们只将所有整数从“long integer”更改为“integer”,并恢复先前的值“divisor=40000”,则内核运行时将减少一半:

2012.05.16 00:22:46     OCL_pi_double (EURUSD,H1)       ________________________
2012.05.16 00:22:46     OCL_pi_double (EURUSD,H1)       OpenCl: gone = 2.262 sec.
2012.05.16 00:22:46     OCL_pi_double (EURUSD,H1)       pi = 3.141592653590
2012.05.16 00:22:46     OCL_pi_double (EURUSD,H1)       read = 40000 elements
2012.05.16 00:22:44     OCL_pi_double (EURUSD,H1)       DOUBLE: _step = 0.000000001000; _intrnCnt = 25000

您应该始终记住,如果您有一个相当于“long&rdquo;but&ldquo;light&rdquo;的循环,即包含大量迭代的循环,每个迭代都没有太多的计算,则只需将数据类型从" Heavy "(long-8 bytes)更改为" Light "(integer-4 bytes)To显著缩短内核的运行时间。

现在,让我们暂停一会儿,集中讨论核心代码ldquo;binding&rdquo;的含义,这样我们就可以了解我们将要做什么。现在让我们使用内核代码“binding”来参考Opencl API,它是一种允许内核与主程序(在本例中是MQL5中的程序)通信的指令系统。

三。OpenCL API函数

3.1。创设语境

下面是一个创建上下文的命令,一个用于管理OpenCL对象和资源的环境。

int clCtx = CLContextCreate( _device );

首先,简要介绍了该平台模型。

图1。操作平台的抽象模型

此图显示了操作平台的抽象模型。图形卡的硬件结构描述不是很详细,但与实际情况非常接近,总体思路很好。

主机是控制整个程序执行过程的主CPU。它可以识别几个OpenCL设备(计算设备)。在大多数情况下,如果交易者在系统单元中配备了用于计算的图形卡,图形卡将被视为一个设备(双处理器图形卡将被视为两个设备!)此外,主机本身(CPU)一直被视为OpenCL设备。平台中的每个设备都有一个唯一的编号。

如果设备配备了与86核对应的CPU(包括Intel CPU“虚拟”核,即“由超线程技术创建”核),那么每个设备都将配备多个计算单元;对于图形卡,它将是SIMD引擎(单指令多数据引擎),也就是说,符合GPU计算。AMD/ATI Radeon体系结构特征条款的IMD核心和微处理器。功能强大的图形卡通常有20个SIMD核心。

每个SIMD核心都包含流处理器,如Radeon HD 5870图形卡的每个SIMD引擎中的16个流处理器。最后,每个流处理器包含四到五个处理组件,即同一图形卡中的算术逻辑单元(ALU)。

请注意,所有主要图形硬件供应商使用的术语都非常混乱,特别是对于新来者。"beesa 6035这个词在OpenCL相关论坛和帖子中常用,其含义并不总是很清楚。然而,线程的数量(即,今天在图形卡中计算的同步线程)非常大。例如,RadeonHD5870图形卡中的线程估计数高达5000个。

下图显示了显卡的标准技术规格。

ATI Radeon 5870 GPU 特性

图2。Radeon HD 5870 GPU功能

下面进一步指定的所有内容(opencl resources)都需要与通过clcontextcreate()函数创建的上下文关联:

  • OpenCL设备是操作过程中使用的硬件。
  • 程序对象,即执行内核的程序代码;
  • 内核是在设备上运行的函数。
  • 存储对象,即由设备操作的数据(如缓冲器、二维图像和三维图像);
  • 命令队列(当前实现的终端语言不由各自的API提供)。

创建的上下文可以描述为一个空字段,在该字段下填充设备名称。

图3。OpenCL上下文

执行函数后,上下文本段当前为空。

需要注意的是,MQL5中的opencl上下文仅用于处理一个设备。

3.2。创建程序

int clPrg = CLProgramCreate( clCtx, clSrc );

clprogramcreate()函数创建一个" opencl programm"资源。

" program "对象实际上是Opencl内核的一个选择(稍后讨论),但是在实现metaQuotes的过程中,很明显Opencl程序中只有一个内核。要创建" Program"对象,请确保将源代码(此处为-clsrc)读取到字符串中。

在这种情况下不需要这样做,因为clsrc字符串已声明为全局变量:

clSrc

下图将程序显示为先前创建的上下文的一部分。

图4。程序是上下文
的一部分

如果程序未能编译,开发人员在编译器的输出端独立地启动一个数据请求。完全功能的opencl api具有api函数clgetprogrambuildinfo(),它在调用编译器输出时返回字符串。

当前版本(B.642)不支持此函数,但有必要将其添加到Opencl API中,以向Opencl开发人员提供有关内核代码正确性的更多详细信息。

来自设备(图形卡)的" Tonguesa 6035(语言)是命令队列,显然,MQL5不会在API级别提供支持。

3.3。创建内核

clkernelcreate()函数创建opencl资源" kernel "(内核)。

int clKrn = CLKernelCreate( clPrg, "pi" );

所谓的内核是在OpenCL设备上运行的程序中声明的函数。

在这种情况下,它是名为"PI"的pi()函数。" kernel "对象是内核的函数加上它自己的自变量。函数的第二个独立变量是函数名,它应该与程序中的函数名相匹配。

图5。内核

只要为一个内核设置了不同的自变量,并且将同一个函数声明为内核," kernel"对象就可以尽可能多地使用。

现在,让我们讨论一下clsetkernelarg()和clsetkernelargmem()函数,但让我们简单介绍一下存储在设备内存中的对象。

3.4。内存对象

首先,我们需要明确的是,要在GPU上处理的任何“大型”对象必须在GPU本身的内存中创建,或者从主机内存(RAM)中删除。我们称之为“较大”的对象是指缓冲区(一维阵列)或二维(二维)、三维(三维)图像。

缓冲区是一个包含不同相邻缓冲区元素的大存储区。它可以是简单的数据类型(字符、双精度、浮点、长类型等)或复杂的数据类型(结构、联合等)。可以直接访问、读取和写入独立的缓冲区元素。

现在,我们需要深入研究图像,因为它是一种独特的数据类型。终端开发人员在OpenCL相关文章的主页上提供了代码,表明开发人员没有参与图片的使用。

在引入的代码中,创建缓冲区的功能如下:

int clMem = CLBufferCreate( clCtx,  _divisor * sizeof( double ), CL_MEM_READ_WRITE );

第一个参数是OpenCL缓冲区,作为资源实现关联的上下文句柄;第二个参数是分配给缓冲区的内存;第三个参数显示了可以对对象执行什么操作。如果由于错误而无法创建opencl缓冲区,则返回值是opencl缓冲区的句柄或-1。

在这种情况下,缓冲区直接在GPU内存(即OpenCL设备)中创建。如果它是在RAM中创建的,而不使用此功能,它将移动到OpenCL设备内存(GPU),如下图所示:

图 6. OpenCL 内存对象

图6。OpenCL内存对象

左侧显示了除OpenCL内存对象以外的输入/输出缓冲区(不一定是图像-这里的蒙娜丽莎仅用于说明目的!).empty,未初始化的opencl内存对象在主上下文字段中显示得更多。初始数据“蒙娜丽莎”将被移到OpenCL的上下文中,无论OpenCL程序输出的内容是什么,它都需要移回左边,也就是说,移到RAM中。

opencl中用于在主机/opencl设备之间复制数据的术语如下:

  • 将数据从主机复制到设备内存的过程称为写入(clbufferwrite()函数)。
  • 将数据从设备内存复制到主机内存的过程称为读取(clbufferread()函数,见下文)。

写入命令(主机-&gt;设备)通过数据初始化内存对象,并将该对象放入设备内存中。

请记住,OpenCL规范中没有指定设备中可用内存对象的有效性,因为它取决于设备相应硬件的供应商。所以在创建内存对象时要小心。

当内存对象初始化并写入设备时,图像如下:

图 7. OpenCL 内存对象的初始化结果

图7。opencl内存对象
初始化结果

现在,我们可以继续解释用于设置内核参数的函数。

3.5。设置内核参数

CLSetKernelArgMem( clKrn, 0, clMem );

clsetkernelargmem()函数将先前创建的缓冲区定义为内核的零参数。

如果现在在内核代码中查看相同的参数,可以看到它们如下所示:

__kernel void pi( __global float *out )

在内核中,通过API函数clbufferCreate()创建的相同类型是out[]数组。

类似功能用于设置非缓冲区参数:

bool  CLSetKernelArg( int   kernel,          // handle to the kernel of the OpenCL program
                       uint  arg_index,       // OpenCL function argument number
                      void  arg_value );     // function argument value

例如,如果我们决定将双精度X0设置为内核的第二个参数,那么首先需要在mql5程序中声明和初始化:

double x0 = -2;

接下来,您需要调用函数(也在mql5代码中):

CLSetKernelArg( cl_krn, 1, x0 );

以上操作完成后,如下图所示:

Pi 计算公式0

图8。内核参数设置

3.6。程序执行

bool ex = CLExecute( clKrn, 1, offs, works );

作者在OpenCL规范中没有找到该函数的直接模拟。此函数使用指定的参数执行内核clkrn。最后一个参数“works”设置每次执行的操作数。这个函数证明了spmd的原理:每次调用函数时,它都用自己的参数创建内核实例(实例数等于工作参数的值);通常,这些内核实例是同时执行的,但在不同的流核心上(对于amd)。

opencl的普遍性在于语言不受代码执行所涉及的底层硬件基础设施的约束:编码人员可以在不知道硬件规范的情况下正常执行opencl程序。程序仍将被执行。但是,强烈建议理解这些规范,以提高代码效率(例如速度)。

例如,在作者没有单独图形卡的硬件设备上,代码可以很好地工作。也就是说,作者不知道整个仿真过程中CPU本身的结构。

到目前为止,opencl程序已经成功执行,现在我们可以在主程序中使用结果。

3.7。读取输出数据

以下是主程序从设备读取的数据的一部分:

float buf[  ]; 
ArrayResize( buf, _divisor );                                                   
uint read = CLBufferRead( clMem, buf );

请记住,从OpenCL读取数据意味着将数据从设备复制到主机。上面的三行显示了它们是如何完成的。只需在主程序中声明与读取opencl缓冲区相同类型的buf[]缓冲区,然后调用函数。在主程序中创建的缓冲区类型(本文中的mql5语言)可以不同于内核中的类型,但它们的大小必须完全匹配。

此数据现在已被复制到主机内存中,并且完全可用于主程序(即MQL5中的程序)。

在OpenCL设备上完成所有必要的操作后,应释放所有对象中的内存。

3.8。销毁所有OpenCL对象的

此操作使用以下命令执行:

CLBufferFree( clMem );
CLKernelFree( clKrn );
CLProgramFree( clPrg );
CLContextFree( clCtx );

这些函数系列的主要特点是,对象应该按照与创建相反的顺序进行销毁。

现在,让我们简单看看内核代码本身。

3.9。内核

如您所见,整个内核代码是由多个字符串组成的单个长字符串。

内核头类似于标准函数:

__kernel void pi( __global float *out )

内核头部有几个要求:

  • 返回值的类型始终为空。
  • 描述符内核不需要包含两个下划线字符,它也可以是内核。
  • 如果参数是数组(缓冲区),则只能通过引用传递。内存说明符“全局(或全局)”表示缓冲区存储在设备的全局内存中。
  • 简单数据类型的独立变量按值传递。

内核的主体与C.
中的标准代码没有什么不同。

重要提示:字符串:

int i = get_global_id( 0 ); 

GPU中的i用于确定该单元中的计算结果。此结果将进一步写入输出数组(在本例中为out[]),当数组从GPU内存读取到CPU内存时,该值将添加到主程序中。

注意,opencl代码中可能有多个函数。例如,您可以调用简单的内联函数,而不是&ldquo;master&rdquo;中的pi()函数和内核函数pi()。稍后我们将进一步讨论这个问题。

既然我们已经对metaquotes实现中的opencl api有了大致的了解,那么我们可以继续尝试。在本文中,作者不打算详细介绍允许最大运行时优化的硬件。目前的主要任务也是介绍OpenCL编程。

换句话说,代码相当简单,因为它不考虑硬件规格。同时,它是非常通用的,因此它可以在任何硬件上执行——AMD生产的CPU、IGP(集成在CPU中的GPU)或AMD/NVIDIA生产的独立图形卡。

在进一步研究使用向量数据类型的简单优化之前,我们需要熟悉它们。

4。矢量数据类型

矢量数据类型是OpenCL特有的,它将它们与C99区别开来。这些类型包括任意(u)字符,(u)短,(u)intn,(u)longn,flotan类型,其中n=2 3 4 8 16。

当我们知道(或假定)内置编译器将尝试实现其他并行计算时,应该使用这些类型。这里我们需要注意的是,情况并非总是这样的,即使内核代码只是与n值不同,但在所有其他方面都完全相同(作者将自己判断)。

下表显示了内置数据类型:

Pi 计算公式1

表1。OpenCL中的内置矢量数据类型

中的任何设备都支持上述数据类型。每个数据类型都有一个对应的API类型,用于内核和主程序之间的通信。这在当前的MQL5实现过程中没有提供,但问题并不严重。

其他类型可用,但使用时应明确指定,因为并非所有设备都支持这些类型:

Pi 计算公式2

表2。OpenCL中的其他内置数据类型

此外,OpenCL还支持一些保留的数据类型。它们在语言“规范”中占据了相当大的空间。

要声明向量类型的常量或变量,应遵循简单和直观的原则。

以下列举了几个例子:

float4 f = ( float4 ) ( 1.0f, 2.0f, 3.0f, 4.0f);
uint4  u = ( uint4  ) ( 1 );             /// u is converted to a vector (1, 1, 1, 1).
float4 f = ( float4 ) ( ( float2 )( 1.0f, 2.0f ), ( float2 )( 3.0f, 4.0f ) );
float4 f = ( float4 ) ( 1.0f, ( float2 )( 2.0f, 3.0f ), 4.0f );
float4 f = ( float4 ) ( 1.0f, 2.0f );   /// error

如您所见,您只需将右侧数据类型与左侧声明的变量“宽度”匹配(此处为4)匹配即可。唯一的例外是当组件等于标量时发生的标量到矢量转换(第2行)。

所有矢量数据类型都采用简单的矢量分量处理机制。它们一方面是向量(数组),另一方面是结构。例如,如果向量的第一个分量宽度为2(例如float2 u),则可以将其视为U.X,而第二个分量则视为U.Y.

long3 u型矢量的三个分量是u.x、u.y和u.z。

对于float4 U型向量,顺序为。X Y Z W,即U.X,U.Y,U.Z,U.W。

float2 pos;
pos.x = 1.0f; // valid
pos.z = 1.0f; // invalid because pos.z does not exist

float3 pos;
pos.z = 1.0f; // valid
pos.w = 1.0f; // invalid because pos.w does not exist

您可以一次选择多个组件,甚至可以重新排列它们(组标记):

float4 c;
c.xyzw = ( float4 ) ( 1.0f, 2.0f, 3.0f, 4.0f );
c.z = 1.0f;
c.xy =  ( float2 ) ( 3.0f, 4.0f );
c.xyz = ( float3 ) ( 3.0f, 4.0f, 5.0f );

float4 pos = ( float4 ) ( 1.0f, 2.0f, 3.0f, 4.0f );
float4 swiz= pos.wzyx;                          // swiz = ( 4.0f, 3.0f, 2.0f, 1.0f )
float4 dup = pos.xxyy;                          // dup = ( 1.0f, 1.0f, 2.0f, 2.0f )

组标记分量,即多个分量的规格,可能会出现在赋值语句的左侧(即左值):

float4 pos = ( float4 ) ( 1.0f, 2.0f, 3.0f, 4.0f );
pos.xw     = ( float2 ) ( 5.0f, 6.0f );            // pos = ( 5.0f, 2.0f, 3.0f, 6.0f )
pos.wx     = ( float2 ) ( 7.0f, 8.0f );            // pos = ( 8.0f, 2.0f, 3.0f, 7.0f )
pos.xyz    = ( float3 ) ( 3.0f, 5.0f, 9.0f );      // pos = ( 3.0f, 5.0f, 9.0f, 4.0f )
pos.xx     = ( float2 ) ( 3.0f, 4.0f );           // invalid as 'x' is used twice
pos.xy     = ( float4 ) (1.0f, 2.0f, 3.0f, 4.0f ); // mismatch between float2 and float4

float4 a, b, c, d;

float16 x;
x = ( float16 ) ( a, b, c, d );
x = ( float16 ) ( a.xxxx, b.xyz, c.xyz, d.xyz, a.yzw );
x = ( float16 ) ( a.xxxxxxx, b.xyz, c.xyz, d.xyz );  // invalid as the component a.xxxxxxx is not a valid vector type

可以使用另一个标记访问每个组件,该标记包括在十六进制数字之前插入的字符s(或s),或在组标记中插入多个数字:

Pi 计算公式3

表3。索引
用于访问矢量数类型的每个组件

如果您声明一个向量变量f

float8 f;

则 f.s0 是向量的第 1 个分量,f.s7 是第 8 个分量。

同样,如果我们声明一个 16 维向量 x,

float16 x;

则 x.sa (或 x.sA)是向量 x 的第 11 个分量;而 x.sf (或 x.sF)则指向量 x 的第 16 个分量。

数字索引(。S0123456789 abcdef)和字母标记(.XYZW)不应与带有组件标记的相同标识符混淆。

float4 f, a;
a = f.x12w;             // invalid as numeric indices are intermixed with the letter notations .xyzw
a.xyzw = f.s0123;       // valid

最后,还有另一种方法可以使用来操纵向量类型的组件。瞧,你好,。甚至。奇怪的。

后缀如下:

  • Lo是指指定向量的下半部分。
  • Hi是指指定向量的上半部分。
  • eve是指向量的所有偶数分量。
  • 奇数是指向量的所有奇数分量。

例如:

float4 vf;
float2 low  = vf.lo;           // vf.xy
float2 high = vf.hi;           // vf.zw
float2 even = vf.even;         // vf.xz
float2 odd  = vf.odd;          // vf.yw

此标记可以重复使用,直到出现标量(非矢量数据类型)。

float8 u = (float8) ( 1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f );
float2 s = u.lo.lo;                              //  ( 1.0f, 2.0f )
float2 t = u.hi.lo;                              //  ( 5.0f, 6.0f )
float2 q = u.even.lo;                            //  ( 1.0f, 3.0f )
float  r = u.odd.lo.hi;                           //  4.0f

对于三分量矢量类型,情况稍微复杂一点:严格来说,它是四分量矢量类型,四分量的值未定义。

float3 vf   = (float3) (1.0f, 2.0f, 3.0f);
float2 low  = vf.lo;                             // ( 1.0f, 2.0f );
float2 high = vf.hi;                             // ( 3.0f, undefined );

简单操作规则(+、-、*、/)。

为同一维度的向量定义所有指定的算术运算,并根据其分量逐个完成。

float4  d   = (float4) ( 1.0f, 2.0f, 3.0f, 4.0f );
float4  w   = (float4) ( 5.0f, 8.0f, 10.0f, -1.0f );
float4  _sum = d + w;                             // ( 6.0f, 10.0f, 13.0f, 3.0f )
float4  _mul = d * w;                             // ( 5.0f, 16.0f, 30.0f, -4.0f )
float4  _div = w / d;                             // ( 5.0f,  4.0f, 3.333333f, -0.25f )

唯一的例外是一个操作数是标量,另一个是向量。在这种情况下,标量类型将作为在向量中声明的数据类型进行计算,并将标量本身转换为具有相同维度和向量操作数的向量。接下来,执行算术运算。关系运算符(&lt;、&gt;、&lt;=、&gt;=)相同。

OpenCL语言还支持C99本地数据类型(如结构、联合、数组等),这些数据类型可以从本章第一个图表中列出的内置数据类型派生。

最后一件事:如果你想使用GPU进行精确的计算,毫无疑问你必须使用双精度数据类型和得到的双精度数据类型。

为此,只需要以下内容:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable 

在内核代码的开头插入它。

这些信息足以让您了解以下大部分内容。如果您有任何问题,请参考Opencl 1.1规范。

5。使用
矢量数据类型实现核心

老实说,作者没有立即编写一个向量数据类型的工作代码。

一开始,作者不太重视阅读语言规范,但认为只要在内核中声明一个向量数据类型(如double8),所有的问题都会自然地得到解决。此外,作者试图将一个输出数组声明为一个双向量数组失败。

我花了一段时间才意识到,内核的矢量化和实际加速度是绝对不够的。将结果输出到向量数组是不可能解决这个问题的,因为它不仅需要快速的输入和输出数据,而且需要快速的计算。如果顺利实现,可以加速流程并提高效率,从而最终能够开发一组更快的代码。

但这并不简单。虽然上述内核代码几乎可以实现盲调优,但由于使用矢量数据,很难发现错误。我们可以用这个标准信息

ERR_OPENCL_INVALID_HANDLE - invalid handle to the OpenCL program 

或以下信息中的

ERR_OPENCL_KERNEL_CREATE - internal error while creating an OpenCL object

你得到了什么建设性的信息?

因此,作者不得不求助于SDK。在这种情况下,考虑到作者的硬件配置,恰好是由Intel Opencl SDK提供的Intel Opencl SDK脱机编译器(32位)(适用于Intel之外的CPU/GPU,SDK还应包括相关的脱机编译器)。因为它允许您在不绑定到主机API的情况下调试内核代码,所以非常方便。

您只需要将内核代码插入编译器窗口(虽然不是以mql5代码的形式),用没有外部引用字符和"/R/N"(回车字符)替换它,然后用齿轮图标按构建按钮。

这样,“生成日志”窗口将显示有关生成过程及其进度的信息:

Pi 计算公式4

图9。Intel Opencl SDK脱机编译器中的程序编译

为了获得不带引用字符的内核代码,可以使用主机语言(mql5)编写一个简单的程序writeclprogram(),它可以将内核代码输出到一个文件中。它现在包含在主程序代码中。

编译器的消息并不总是清晰的,但它比当前的MQL5更可用。错误可以在编译器中立即修复,一旦确定没有错误,就可以使用metaeditor将修复发送到内核代码。

最后一件事。作者最初的想法是开发一个矢量化代码,通过设置通道A 6035的单个全局变量",可以处理Double4、Double8和Double16。几天来,作者一直在费力地使用拼接操作符(出于某种原因,该符号不能在内核代码中被操纵),最终它完成了。

在此期间,作者成功地开发了一个具有三个内核代码的脚本的工作代码,每个内核代码都适合其大小-4、8或16。尽管本文不提供这个中间代码,但值得一提的是,因为您可能希望轻松地编写内核代码。本文的结尾附加了由这个脚本实现的代码(ocl_Pi_Double_Several_Simple_Kernels.mq5)。

下面是矢量化内核的代码:

"/// enable extensions with doubles                                                          /r/n"
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable                                               /r/n"
"#define _ITERATIONS "        + i2s( _intrnCnt ) + "                                     /r/n"
"#define _STEP "             + d2s( _step, 12 )  + "                                     /r/n"     
"#define _CH "               + i2s( _ch )       + "                                     /r/n"
"#define _DOUBLETYPE double"   + i2s( _ch )       + "                                     /r/n"   
"                                                                                            /r/n"
"/// extensions for 4-, 8- and 16- scalar products                                           /r/n"     
"#define dot4( a, b )         dot( a, b )                                                    /r/n"
"                                                                                            /r/n"
"inline double dot8( double8 a, double8 b )                                                  /r/n"
"{                                                                                           /r/n"     
"     return dot4( a.lo, b.lo ) + dot4( a.hi, b.hi );                                        /r/n"
"}                                                                                           /r/n"
"                                                                                            /r/n"      
"inline double dot16( double16 a, double16 b )                                               /r/n"     
"{                                                                                           /r/n"     
"     double16 c  = a * b;                                                                   /r/n" 
"     double4 _1  = ( double4 ) ( 1., 1., 1., 1. );                                          /r/n" 
"     return dot4( c.lo.lo + c.lo.hi + c.hi.lo  + c.hi.hi, _1 );                             /r/n" 
"}                                                                                           /r/n"     
"                                                                                            /r/n"
"__kernel void pi( __global double *out )                                                    /r/n"
"{                                                                                           /r/n"
"  int i = get_global_id( 0 );                                                               /r/n"
"                                                                                            /r/n"
"  /// define vector constants                                                               /r/n"     
"  double16 v16  = ( double16 ) ( 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 );    /r/n"
"  double8  v8   = v16.lo;                                                                   /r/n"
"  double4  v4   = v16.lo.lo;                                                                /r/n"
"  double2  v2   = v16.lo.lo.lo;                                                             /r/n"          
"                                                                                            /r/n"    
"  /// all vector-related with the calculated type                                           /r/n"
"  _DOUBLETYPE in;                                                                           /r/n"    
"  _DOUBLETYPE xVect;                                                                        /r/n"
"  _DOUBLETYPE sumVect          = ( _DOUBLETYPE ) ( 0.0 );                                   /r/n"
"  _DOUBLETYPE doubleOneVect    = ( _DOUBLETYPE ) ( 1.0 );                                   /r/n"
"  _DOUBLETYPE doubleCHVect     = ( _DOUBLETYPE ) ( _CH + 0. );                              /r/n"
"  _DOUBLETYPE doubleSTEPVect   = ( _DOUBLETYPE ) ( _STEP );                                 /r/n"
"                                                                                            /r/n"    
"  for( long j = 0; j < _ITERATIONS; j ++ )                                                  /r/n"
"  {                                                                                         /r/n"
"     in =  v" + i2s( _ch ) + " + doubleCHVect * ( i * _ITERATIONS + j );                  /r/n"           
"     xVect = ( in + 0.5 ) * doubleSTEPVect;                                                 /r/n"
"     sumVect += 4.0 / ( xVect * xVect + 1. );                                               /r/n"
"  }                                                                                         /r/n"
"  out[ i ] = dot" + i2s( _ch ) + "(  sumVect, doubleOneVect );                            /r/n"
"}                                                                                           /r/n";   

除了设置“矢量通道”的全局常数ch、新的全局常数ch的个数和全局常数intrncnt已被多次减少外,外部主程序几乎没有变化。因此,作者决定不在这里显示主程序。程序(ocl_pi_double_parallel_direct.mq5)可以在本文末尾附加的脚本中找到。

可以看出,除了核心pi()的“主”函数外,我们现在还有两个用于确定向量dotn(a,b)的无向积和一个宏置换的内联函数。因为opencl中的dot()函数被定义为维数小于4的向量,所以它们被包括在内。

使用dot4()宏(它重新定义了dot()函数)只是为了方便调用计算名称为
的dotn()函数。

"  out[ i ] = dot" + i2s( _ch ) + "(  sumVect, doubleOneVect );                              /r/n"

如果我们使用一个通常不采用指数4形式的点()函数,那么当_ch=4(矢量化通道数为4)时,我们就不能这么容易地调用它。

这一行说明了一个特定内核形式的另一个有用函数,它的特点是在主程序中作为字符串进行处理:我们可以在内核中使用计算出的标识符,不仅包括函数标识符,还包括数据类型标识符!

包含内核(ocl_pi_double_parallel_direct.mq5)的完整主程序代码附在下面。

使用" width"的16(_ch=16)矢量运行脚本后,可以获得以下结果:

2012.05.15 00:15:47     OCL_pi_double2_parallel_straight (EURUSD,H1)    ==================================================
2012.05.15 00:15:47     OCL_pi_double2_parallel_straight (EURUSD,H1)    CPUtime / GPUtime = 4.130
2012.05.15 00:15:47     OCL_pi_double2_parallel_straight (EURUSD,H1)    SMARTER: The time to calculate PI was 8.830 seconds
2012.05.15 00:15:47     OCL_pi_double2_parallel_straight (EURUSD,H1)    SMARTER: The value of PI is 3.141592653590
2012.05.15 00:15:38     OCL_pi_double2_parallel_straight (EURUSD,H1)    DULL: The time to calculate PI was 8.002 seconds
2012.05.15 00:15:38     OCL_pi_double2_parallel_straight (EURUSD,H1)    DULL: The value of PI is 3.141592653590
2012.05.15 00:15:30     OCL_pi_double2_parallel_straight (EURUSD,H1)    OPENCL: gone = 2.138 sec.
2012.05.15 00:15:30     OCL_pi_double2_parallel_straight (EURUSD,H1)    OPENCL: pi = 3.141592653590
2012.05.15 00:15:30     OCL_pi_double2_parallel_straight (EURUSD,H1)    read = 20000 elements
2012.05.15 00:15:28     OCL_pi_double2_parallel_straight (EURUSD,H1)    CLProgramCreate: unknown error.
2012.05.15 00:15:28     OCL_pi_double2_parallel_straight (EURUSD,H1)    DOUBLE2: _step = 0.000000001000; _intrnCnt = 3125
2012.05.15 00:15:28     OCL_pi_double2_parallel_straight (EURUSD,H1)    ==================================================

正如您所看到的,即使使用向量数据类型进行优化,也不能使内核更快。

但是,如果您在GPU上运行相同的代码,那么速度增加就更明显了。

根据metadriver提供的信息(HIS Radeon HD 6930用于图形卡,AMD Phenom II X6 1100T用于CPU),相同的代码产生以下结果:

2012.05.14 11:36:07     OCL_pi_double2_parallel_straight (AUDNZD,M5)    ==================================================
2012.05.14 11:36:07     OCL_pi_double2_parallel_straight (AUDNZD,M5)    CPUtime / GPUtime = 84.983
2012.05.14 11:36:07     OCL_pi_double2_parallel_straight (AUDNZD,M5)    SMARTER: The time to calculate PI was 14.617 seconds
2012.05.14 11:36:07     OCL_pi_double2_parallel_straight (AUDNZD,M5)    SMARTER: The value of PI is 3.141592653590
2012.05.14 11:35:52     OCL_pi_double2_parallel_straight (AUDNZD,M5)    DULL: The time to calculate PI was 14.040 seconds
2012.05.14 11:35:52     OCL_pi_double2_parallel_straight (AUDNZD,M5)    DULL: The value of PI is 3.141592653590
2012.05.14 11:35:38     OCL_pi_double2_parallel_straight (AUDNZD,M5)    OPENCL: gone = 0.172 sec.
2012.05.14 11:35:38     OCL_pi_double2_parallel_straight (AUDNZD,M5)    OPENCL: pi = 3.141592653590
2012.05.14 11:35:38     OCL_pi_double2_parallel_straight (AUDNZD,M5)    read = 20000 elements
2012.05.14 11:35:38     OCL_pi_double2_parallel_straight (AUDNZD,M5)    CLProgramCreate: unknown error.
2012.05.14 11:35:38     OCL_pi_double2_parallel_straight (AUDNZD,M5)    DOUBLE2: _step = 0.000000001000; _intrnCnt = 3125
2012.05.14 11:35:38     OCL_pi_double2_parallel_straight (AUDNZD,M5)    ==================================================

6。点睛之笔

还有一个内核(可以在下面附加的ocl_pi_double_several_simple_kernels.mq5文件中找到,但这里没有演示)。

该脚本实现了作者之前的想法,当他暂时放弃编写“单内核模块”时,计划为不同的向量维(4、8、16、32)编写四个简单的内核模块:

"#pragma OPENCL EXTENSION cl_khr_fp64 : enable                                               /r/n"
"#define _ITERATIONS " + i2s( _itInKern ) + "                                             /r/n"
"#define _STEP "      + d2s( _step, 12 ) + "                                              /r/n"
"                                                                                            /r/n"
"typedef struct                                                                              /r/n"
"{                                                                                           /r/n"
"  double16    lo;                                                                           /r/n"
"  double16    hi;                                                                           /r/n"
"} double32;                                                                                 /r/n"
"                                                                                            /r/n"
"inline double32 convert2double32( double a )                                                /r/n"        
"{                                                                                           /r/n"
"  double32 b;                                                                               /r/n"    
"    b.lo = ( double16 )( a );                                                               /r/n"
"    b.hi = ( double16 )( a );                                                               /r/n"
"    return b;                                                                               /r/n"
"}                                                                                           /r/n"
"                                                                                            /r/n"
"inline double dot32( double32 a, double32 b )                                               /r/n"     
"{                                                                                           /r/n"     
"     double32 c;                                                                            /r/n" 
"     c.lo = a.lo * b.lo;                                                                    /r/n" 
"     c.hi = a.hi * b.hi;                                                                    /r/n" 
"     double4 _1  = ( double4 ) ( 1., 1., 1., 1. );                                          /r/n" 
"     return dot( c.lo.lo.lo + c.lo.lo.hi + c.lo.hi.lo  + c.lo.hi.hi +                       /r/n" 
"                 c.hi.lo.lo + c.hi.lo.hi + c.hi.hi.lo  + c.hi.hi.hi, _1 );                  /r/n" 
"}                                                                                           /r/n"     
"                                                                                            /r/n"
"__kernel void pi( __global double *out )                                                    /r/n"
"{                                                                                           /r/n"
"  int i = get_global_id( 0 );                                                               /r/n"
"                                                                                            /r/n"
"  /// define vector constants                                                               /r/n"
"  double32 _v32;                                                                            /r/n"
"  _v32.lo = ( double16 ) (  0.,  1.,  2.,  3.,  4.,  5.,  6.,  7.,                          /r/n"
"                            8.,  9., 10., 11., 12., 13., 14., 15. );                        /r/n"
"  _v32.hi = ( double16 ) ( 16., 17., 18., 19., 20., 21., 22., 23.,                          /r/n"
"                           24., 25., 26., 27., 28., 29., 30., 31. );                        /r/n"
"                                                                                            /r/n"
"  /// all vector-related with undefined type                                                /r/n"
"  double32 xVect;                                                                           /r/n"
"  double32 sumVect           = convert2double32( 0.0 );                                     /r/n"
"  double32 double1Vect       = convert2double32( 1.0 );                                     /r/n"
"                                                                                            /r/n"
"  double32 in;                                                                              /r/n"
" /// work only with 16-vectors in the loop!                                                 /r/n"
"  for( long j = 0; j < _ITERATIONS; j ++ )                                                  /r/n"
"  {                                                                                         /r/n"
"     in.lo = _v32.lo + 32. * ( i * _ITERATIONS + j );                                       /r/n"
"     in.hi = _v32.hi + 32. * ( i * _ITERATIONS + j );                                       /r/n"
"     xVect.lo = ( in.lo  + 0.5 ) * _STEP;                                                   /r/n"
"     xVect.hi = ( in.hi  + 0.5 ) * _STEP;                                                   /r/n"
"     sumVect.lo += 4. / ( xVect.lo * xVect.lo + 1. );                                       /r/n"
"     sumVect.hi += 4. / ( xVect.hi * xVect.hi + 1. );                                       /r/n"                             
"  }                                                                                         /r/n"
"  out[ i ] = dot32( sumVect, double1Vect );                                                 /r/n"
"}                                                                                           /r/n";    

内核模块实现向量维32。新的向量类型和几个必要的内联函数是在内核的主函数之外定义的。除此之外(这很重要!)主循环中的所有计算都是有意使用标准向量数据类型进行的;非标准类型在循环之外进行处理。这大大缩短了代码的执行时间。

在我们的计算中,内核似乎并不比用于16个宽度的向量慢,但速度不快。

根据metadriver提供的信息,包含内核(_ch=32)的脚本产生以下结果:

2012.05.14 12:05:33 OCL_pi_double32-01 (AUDNZD,M5) OPENCL: gone = 0.156 sec.
2012.05.14 12:05:33 OCL_pi_double32-01 (AUDNZD,M5) OPENCL: pi = 3.141592653590
2012.05.14 12:05:33 OCL_pi_double32-01 (AUDNZD,M5) read = 10000 elements
2012.05.14 12:05:32 OCL_pi_double32-01 (AUDNZD,M5) CLProgramCreate: unknown error or no error.
2012.05.14 12:05:32 OCL_pi_double32-01 (AUDNZD,M5) GetLastError returned .. 0
2012.05.14 12:05:32 OCL_pi_double32-01 (AUDNZD,M5) DOUBLE2: _step = 0.000000001000; _itInKern = 3125; vectorization channels - 32
2012.05.14 12:05:32 OCL_pi_double32-01 (AUDNZD,M5) =================================================================

总结与结论

作者非常清楚解释OpenCL资源的任务对于这种语言来说并不是特别典型的。

打开一本教科书,复制一个大型矩阵乘法的标准例子会简单得多。显然,这样的例子肯定令人印象深刻。但是,是否有大量的mql5.com论坛用户参与需要大量矩阵乘法的金融操作?这值得怀疑。作者要选择自己的榜样,克服所有的困难,努力与他人分享自己的经历。当然,你是评委,亲爱的论坛用户。

事实证明,与使用元驱动程序脚本获得的大量效率提高相比,OpenCL模拟(仅通过CPU)获得的效率提高微不足道。但是在一个合适的GPU上,即使我们忽略了它在AMD CPU上运行时间更长的事实,它至少比模拟实现的效率高出一个数量级。尽管OpenCL的速度提升非常有限,但它仍然值得学习!

作者的下一篇文章将讨论与真实硬件上的OpenCL抽象模式显示功能相关的问题。对这些方面的知识有时可以在相当大的程度上进一步加速操作。

作者要特别感谢metadriver提供了宝贵的编程和性能优化建议,并感谢支持团队提供了使用IntelOpencl SDK的可能性。

随附文件目录如下:

  • 圆周率。mq5—完全使用mql5编写的脚本,其中包含两种计算"pi"值的方法;
  • ocl ou pi ou float.mq5——第一次实现了一个带有opencl内核的脚本,涉及浮点操作。
  • ocl_pi_double.mq5-同上,仅涉及双精度类型的实际操作;
  • ocl ou pi ou double ou several ou simple ou kernels.mq5-包含多个特定内核的脚本,适用于多个向量&ldquo;width&rdquo;(4、8、16、32);
  • ocl ou pi ou double_parallel ou straight.mq5-一个具有单个内核的脚本,适用于某些向量“width”(4、8、16)。

本文由MetaQuotes Software Corp.翻译自俄语原文
,网址为https://www.mql5.com/ru/articles/405。

附加文件下载zip pi.mq5(1.84 kb)ocl_pi_float.mq5(3.59 kb)ocl_pi_double.mq5(3.85 kb)ocl_u pi_double_几个简单内核.mq5(22.42 kb)ocl_pi_double_parallel_straight.mq5(11.3 kb)

 

 


MyFxtop迈投(www.myfxtop.com)-靠谱的外汇跟单社区,免费跟随高手做交易!

 

免责声明:本文系转载自网络,如有侵犯,请联系我们立即删除,另:本文仅代表作者个人观点,与迈投财经(www.myfxtop.cn)无关。其原创性以及文中陈述文字和内容未经本站证实,对本文以及其中全部或者部分内容、文字的真实性、完整性、及时性本站不作任何保证或承诺,请读者仅作参考,并请自行核实相关内容。

著作权归作者所有。
商业转载请联系作者获得授权,非商业转载请注明出处。

風險提示

MyFxtops邁投所列信息僅供參考,不構成投資建議,也不代表任何形式的推薦或者誘導行為。MyFxtops邁投非外匯經紀商,不接觸妳的任何資金。 MYFXTOPS不保證客戶盈利,不承擔任何責任。從事外彙和差價合約等金融產品的槓桿交易具有高風險,損失有可能超過本金,請量力而行,入市前需充分了解潛在的風險。過去的交易成績並不代表以後的交易成績。依據各地區法律法規,MyFxtops邁投不向中國大陸、美國、加拿大、朝鮮居民提供服務。

邁投公眾號

聯繫我們

客服QQ:981617007
Email: service@myfxtop.com

MyFxtops 邁投