OpenCL:并行世界的桥梁(基础篇)
📘

OpenCL:并行世界的桥梁(基础篇)

第 1/3 篇

用 OpenCL 给 MT5 装上并行引擎

MT5 从架构层面接入了 OpenCL,让 CPU 之外闲置的 GPU 算力能直接参与指标与策略计算。对跑大规模样本回测或多品种扫描的交易者来说,这意味着同样的逻辑可能从分钟级等待压缩到秒级。 实际收益取决于显存带宽与内核写法。官方示例里一个典型矩阵运算在中等独显上相对纯 CPU 实现有数倍吞吐提升,但是杂乱的内存访问 pattern 会直接吃掉并行红利。 开 MT5 按 F4 进 MetaEditor,新建 OpenCL 程序并绑定上下文,是验证这套机制的最低成本动作。先跑通官方 sample,再替换成你自己的价格数组,比空谈架构有用。

「MT5 原生多核为何要等 OpenCL」

在 MT5 提供 OpenCL 支持前,终端虽然号称多线程,每个 EA 或脚本跑在独立线程里,但写代码的人没法简单地把一段纯计算循环并行丢到多核上跑。比如下面这段算 π 的循环,以前只能在单核里硬算。 早在一年半前就有文章用 EA 加两个指标搭出并行结构算东西,但那种做法层次太重,反而容易拖慢。官方也明确过不会上 OpenCL 之外的 OpenMP,因为改编译器成本太高。 作者自己用奔腾 G840、8G 内存、无独显的机器试过,装了 AMD APP SDK 后,MetaDriver 发的第一个脚本居然能跑,比单核标准脚本快了约 25 倍;后来装了 Intel OpenCL Runtime,同一脚本快到 75 倍。 查了 ixbt.com 的资料才发现,这台机器的 Ivy Bridge 以下 CPU 核显只支持 OpenCL 1.1 且用不上,所以加速纯粹来自 x86 核心的 CPU 在跑 OpenCL 仿真。论坛专家点破:这其实是 MQL5 源语言优化不足导致的“虚高”加速,真用 C++ 配好多核加 SSEx 向量指令,OpenCL 仿真提升可能只有几十个百分点,差的时候还会因传输耗时丢数据。 所以别把 MT5 里那些夸张的加速比当成 OpenCL 本身多神。开发人员也承认还没做合理优化,优化后加速会降几倍,但仍大于一。即便你显卡不支持 OpenCL 1.1 或没显卡,学这套也值得,起码能在 CPU 上验证并行雏形。 下面这段是原文里算 π 的串行参考,逐行看清楚循环本质: long num_steps = 1000000000; // 定义循环总步数十亿级,决定精度 double step = 1.0 / num_steps; // 单步宽度,把 0~1 区间切薄 double x, pi, sum = 0.0; // 中间变量 x、结果 pi、累加器 sum 清零下 for (long i = 0; i<num_steps; i++) // 串行跑十亿次,单核吃满 { x = (i + 0.5)*step; // 取每个小区间中点 sum += 4.0/(1.0 + x*x); // 累加矩形面积近似积分 } pi = sum*step; // 乘步宽得出 π 近似值,外汇贵金属测试请自担高风险的算力开销

MQL5 / C++
class="type">long num_steps = class="num">1000000000;
class="type">class="kw">double step = class="num">1.0 / num_steps;
class="type">class="kw">double x, pi, sum = class="num">0.0;

for (class="type">long i = class="num">0; i<num_steps; i++)
{
  x = (i + class="num">0.5)*step;
  sum += class="num">4.0/(class="num">1.0 + x*x);
}
pi = sum*step;

◍ 挑显卡跑 OpenCL:AMD 与 Intel 的硬门槛

在 MT5 里做 GPU 加速,先得确认硬件和运行时到位。Khronos 集团只管 OpenCL 规范,终端本身几乎不提供 OpenCL 帮助,只给了一段 API 摘要,所以选卡和装 SDK 是第一步。 AMD 这边,2009–2010 年后首发的卡基本升个驱动就能用;但像 Radeon HD 4850/4870 这种老卡跑 OpenCL 麻烦很多。没卡想加卡,看板载内存至少 1 GB、核心频率 650–700 MHz 还行、显存优先 GDDR5(GDDR3 带宽大约少一半)、总线 256 位以上、流处理器数(如 1024:64:32 里的 1024)越高越好。TDP 要注意:EA 频繁调 GPU 时噪声明摆着会变大。 没有支持 OpenCL 1.1 的卡、但有 AMD CPU,装 AMD APP SDK 后处理器会被识别成 OpenCL 设备,还能在 CPU 仿真模式开发。Intel 处理器也能用这个 SDK,不过原生 Intel SDK 因支持 SSE4.1/4.2/AVX 会更高效。 Intel 用户装完 OpenCL SDK 必须改注册表:把 HKEY_LOCAL_MACHINE\SOFTWARE\Khronos\OpenCL\Vendors 下 Name 列的 IntelOpenCL64.dll 换成 intelocl.dll,重启 MT5 后 CPU 才被认作 OpenCL 1.1 设备。Sandy Bridge 及更早的核显不支持 OpenCL 1.1,只能靠 SS(S)Ex 向量指令加速,Ivy Bridge 才开始有核显支持。 内核报错不会像 MQL5 编译器那样定位到行,作者用 Intel 离线编译器调 OpenCL 比在 MetaEditor 里盲找省事太多。外汇和贵金属交易杠杆高、风险大,GPU 加速只是工具,不代表策略胜率提升。

用圆周率算例摸透 OpenCL 内核边界

想在 MT5 里写第一个吃 OpenCL 螃蟹的程序,拿 π 值估算当练手最合适。目标精度锁在 12 位小数(约 3.14159265),纯 CPU 侧用 10 亿次迭代拆成外层 40000 块、内层 25000 次的双循环,实测比单循环经典写法慢 10–15%,但这段双循环的内层正好能原样塞进 GPU 内核。 GPGPU 有个潜规则:内核任务最好跑满 20 毫秒以上,否则数据搬移开销会吃掉算力收益。本例里 GetTickCount() 本身有百毫秒级误差,所以迭代量选得比理论下限更高。初版内核用 float 算,结果只精确到小数点后 3 位——MQL5 文档写明 float 仅 7 位有效数,离 12 位差得远。 把注释标了「///浮点型」的行全换成 double,精度达标但耗时飙到 8.783 秒,比不用 OpenCL 还慢。怀疑是双精度拖后腿?把 _divisor 从 40000 改 40000000,精度不变且比 float 版稍快;再把所有长整型退回整型、_divisor 复原,内核时间直接砍半。结论很硬:长而轻的循环,把 8 字节「重」类型换成 4 字节「轻」类型,能大幅压缩内核耗时。 内核里读不到外部全局变量,_step 和 _intrnCnt 得重新声明并把值转成字符串写进内核代码头部。这种隔离以后做向量类型时反而是利器。下面这段宿主脚本是双循环 CPU 版的完整骨架,逐行拆完你就能照着改 GPU 版。

MQL5 / C++
class=class="str">"cmt">//+------------------------------------------------------------------+
class=class="str">"cmt">//|                                                 pi.mq5 |
class=class="str">"cmt">//+------------------------------------------------------------------+
class="macro">#class="kw">property copyright "Copyright(c) class="num">2012, Mthmt"
class="macro">#class="kw">property link      "[MQL5官方文档]
class="type">long    _num_steps        = class="num">1000000000;
class="type">long    _divisor          = class="num">40000;
class="type">class="kw">double  _step            = class="num">1.0 / _num_steps;
class="type">long    _intrnCnt         = _num_steps / _divisor;
class=class="str">"cmt">//+------------------------------------------------------------------+
class=class="str">"cmt">//| Script program start function                                      |
class=class="str">"cmt">//+------------------------------------------------------------------+
class="type">int OnStart()
  {
   class="type">uint start,stop;
   class="type">class="kw">double x,pi,sum=class="num">0.0;
   start=GetTickCount();
class=class="str">"cmt">//--- first option - direct calculation
   for(class="type">long i=class="num">0; i<_num_steps; i++)
     {
      x=(i+class="num">0.5)*_step;
      sum+=class="num">4.0/(class="num">1.+x*x);
     }
   pi=sum*_step;
   stop=GetTickCount();
   Print("The value of PI is "+DoubleToString(pi,class="num">12));
   Print("The time to calculate PI was "+DoubleToString(( stop-start)/class="num">1000.0,class="num">3)+" seconds");
class=class="str">"cmt">//--- calculate class="kw">using the second option
   start=GetTickCount();
   sum=class="num">0.;
   class="type">long divisor=class="num">40000;
   class="type">long internalCnt=_num_steps/divisor;
   class="type">class="kw">double partsum=class="num">0.;
   for(class="type">long i=class="num">0; i<divisor; i++)
     {
      partsum=class="num">0.;
      for(class="type">long j=i*internalCnt; j<(i+class="num">1)*internalCnt; j++)
        {

「用 OpenCL 在 MT5 里算 π 看设备算力」

把莱布尼茨型积分拆成十亿步(_num_steps=1000000000),再用 _divisor=40000 切成小块丢给 OpenCL 设备并行跑,是测 MT5 端异构计算延迟的直接办法。下面这段宿主侧收尾代码先把局部累加合总,再打点计时。 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); 日志里两次实跑都给出 π=3.141592653590,耗时分别是 8.783 秒与 7.940 秒,误差在 1e-12 级,说明 CPU 做这个规模浮点循环的稳定度够用。外汇与贵金属交易用这类脚本做算力基准时,仍属高性能计算实验,账户资金高风险属性不变。 脚本头部写死 _device=0 指向本机 CPU,想换显卡算就把 input int _device 改对应序号,重跑对比秒数即可。

MQL5 / C++
x=(j+class="num">0.5)*_step;
partsum+=class="num">4.0/(class="num">1.+x*x);
}
sum+=partsum;
}
pi=sum*_step;
stop=GetTickCount();
Print("The value of PI is "+DoubleToString(pi,class="num">12));
Print("The time to calculate PI was "+DoubleToString((stop-start)/class="num">1000.0,class="num">3)+" seconds");
Print("_______________________________________________");
class="kw">return(class="num">0);

◍ 用 OpenCL 把 π 算例跑在显卡上

这段内核把 π 的积分拆给 GPU 并行算:每个工作项负责一段区间,用梯形法近似 4/(1+x²) 的积分和。_step 是单步长,_intrnCnt 是单个线程要扫的点数,两者共同决定精度与耗时。 内核里 x 用 float 存,这意味着单精度浮点误差会随 _intrnCnt 增大而累积;若把 float 换成 double,结果可能更接近 3.141592653589,但部分老显卡不支持双精度而直接报错。 在 MT5 里把这段内核交给 CLContextCreate 后的上下文执行,Print 会先打出 _step=7.957747e-7(当总量 4e6 时)与 _intrnCnt 的值,方便你确认并行粒度。打开终端的 OpenCL 日志,能直接看到每张卡的分组执行时间。

MQL5 / C++
__kernel class="type">void pi( __global class="type">float *out )
{
  class="type">int i = get_global_id( class="num">0 );
  class="type">float partsum = class="num">0.0;
  class="type">float x = class="num">0.0;
  class="type">long from = i * _intrnCnt;
  class="type">long to = from + _intrnCnt;
  for( class="type">long j = from; j < to; j ++ )
  {
    x = ( j + class="num">0.5 ) * _step;
    partsum += class="num">4.0 / ( class="num">1. + x * x );
  }
  out[ i ] = partsum;
}

class="type">int OnStart()
{
  Print("FLOAT: _step = "+d2s(_step,class="num">12)+"; _intrnCnt = "+i2s(_intrnCnt));
  class="type">int clCtx=CLContextCreate(_device);

常见问题

MT5 原生多核调度有主线程瓶颈,密集型数值计算丢给 OpenCL 走显卡并行,延迟可能明显下降。
部分老款 AMD 核显驱动不全,Intel 某些型号不支持特定浮点扩展,选卡前先查厂商 OpenCL 支持列表。
小布可读取你的硬件信息并比对算力门槛,直接告诉你当前机器跑并行内核是勉强还是富余。
能看到单精度/双精度吞吐差、工作组大小上限,以及显存拷贝成为瓶颈的临界点。
能直观摸清自己设备的并行上限,后续把指标或优化任务搬上显卡时少踩坑。