OpenCL：从朴素到更具深度的编程·进阶篇
（2/3）· 多数人只做了朴素优化，忽略硬件规格后矩阵乘法还能再快一个量级
接上篇，我们继续深挖 OpenCL 在实际硬件上的优化空间。大多数交易者写并行内核时只管把循环丢给 GPU，却从不去看存储模型——全局、常量、本地内存的延迟差可能直接决定你的策略回测是秒级还是分钟级。
◍ 千阶矩阵乘法在 CPU 上的耗时实测
在 MT5 脚本里做二维矩阵乘法，核心就是三重循环把 first 和 second 乘进 third。下面这段是内层累加的逻辑，乘积累加的写法直接决定了缓存命中率。 for(int c=0; c<COLS2; c++) third[r][c]+=first[r][cr]*second[cr][c]; return;
- 年 5 月 19 日跑在 EURUSD H1 上的日志很有参考价值：1000×1000 乘 1000×1000 的矩阵，CPU 耗时 10.530 秒；规模翻到 2000 阶之后，耗时跳到 83.663 秒，粗略看计算量翻 8 倍、时间也接近 8 倍。
外汇与贵金属交易里用这类运算做多资产协方差矩阵时，要清楚 CPU 路径在两千阶以上就会明显卡顿，属于高计算负载场景，实盘前务必在策略测试器里先压一遍。
for(class="type">int c=class="num">0; c<COLS2; c++) third[r][c]+=first[r][cr]*second[cr][c]; class="kw">return;
「用 OpenCL 把矩阵乘法丢给显卡」
在 MT5 里做重计算，CPU 单核跑大矩阵会卡死策略线程。把乘法核写进 OpenCL，让 GPU 并行算，是绕开瓶颈的直接办法。 下面这段内核用 float 精度定义 REALTYPE，__kernel 函数 matricesMul 接收三个全局缓冲区：两个输入矩阵和一个输出矩阵。get_global_id(0) 和 (1) 分别取行、列索引，循环 COLSROWS 次做乘加累积。 [CODE] #define REALTYPE float __kernel void matricesMul( __global REALTYPE *in1, __global REALTYPE *in2, __global REALTYPE *out ) { int r = get_global_id( 0 ); int c = get_global_id( 1 ); for( int cr = 0; cr < COLSROWS; cr ++ ) out[ r * COLS2 + c ] += in1[ r * COLSROWS + cr ] * in2[ cr * COLS2 + c ]; } [/CODE] 逐行拆一下：第1行把 REALTYPE 绑死为 float，省掉 double 的显存带宽；第3-5行是内核签名，三个指针都带 __global 说明数据在显存；第7-8行用两个维度全局 ID 定位到当前线程负责的单元格；第9行循环累加，COLSROWS 是左矩阵列数也是右矩阵行数；第10-11行完成 out[r][c] += in1[r][cr]*in2[cr][c] 的归约。 对照参考实现，同样规模矩阵在 CPU 单核 c-r-cr 顺序跑完花了 11.544 秒，GPU 内核理论上能把这个数压到秒级以下。外汇与贵金属杠杆高，回测若依赖此类密集计算，用 OpenCL offload 可能显著缩短优化周期，但实盘算力抖动仍是高风险点。
class="macro">#define REALTYPE class="type">float __kernel class="type">void matricesMul( __global REALTYPE *in1, __global REALTYPE *in2, __global REALTYPE *out ) { class="type">int r = get_global_id( class="num">0 ); class="type">int c = get_global_id( class="num">1 ); for( class="type">int cr = class="num">0; cr < COLSROWS; cr ++ ) out[ r * COLS2 + c ] += in1[ r * COLSROWS + cr ] * in2[ cr * COLS2 + c ]; }
单核与GPU矩阵乘法的耗时对照
这段脚本把同一套矩阵乘法分别丢给 CPU 单核和 OpenCL 设备跑，用 GetTickCount() 抓前后时间戳，最后除以 1000 转成秒。CPU 部分走 mulCPUOneCore()，三层 for 循环做 r*COLS2+c 的累加；GPU 部分走 initAllDataGPU() 加 executeGPU()，核函数跑完再把输出缓冲读回 buf[]。 计时打印出来是两行关键数据：CPUTime 和 OpenCL total time。脚本末尾还打了 CPUTime / GPUTotalTime 的比值，若矩阵规模够大（如 ROWS1、COLS2 在千级），该比值往往明显大于 1，说明 GPU 总耗时倾向更短，但已包含数据回传 RAM 的代价。 校验没走全量比对，只随机抽 10 组用 checkRandom() 验精度。外汇与贵金属相关的算法回测若套用此类并行计算，需注意 GPU 浮点误差可能让信号边界偏移，属高风险操作，结果仅作概率参考。 想复现就开 MT5 建脚本，把下文代码贴入，调大 COLSROWS 看比值变化。
class="type">void OnStart() { initAllDataCPU(); class=class="str">"cmt">//--- start working with non-parallel version("bare" CPU, single core) class=class="str">"cmt">//--- calculate the output matrix on a single core CPU class="type">uint st=GetTickCount(); mulCPUOneCore(); class=class="str">"cmt">//--- output total calculation time class="type">class="kw">double timeCPU=(GetTickCount()-st)/class="num">1000.; Print("CPUTime = "+d2s(timeCPU,class="num">3)); class=class="str">"cmt">//--- start working with OCL class="type">int clCtx; class=class="str">"cmt">// context handle class="type">int clPrg; class=class="str">"cmt">// handle to the program on the device class="type">int clKrn; class=class="str">"cmt">// kernel handle class="type">int clMemIn1; class=class="str">"cmt">// first(input) buffer handle class="type">int clMemIn2; class=class="str">"cmt">// second(input) buffer handle class="type">int clMemOut; class=class="str">"cmt">// third(output) buffer handle class=class="str">"cmt">//--- start calculating the program runtime on GPU initAllDataGPU(clCtx,clPrg,clKrn,clMemIn1,clMemIn2,clMemOut); class=class="str">"cmt">//--- start calculating total OCL code runtime st=GetTickCount(); executeGPU(clKrn); class=class="str">"cmt">//--- create a buffer for reading and read the result; we will need it later REALTYPE buf[]; readOutBuf(clMemOut,buf); class=class="str">"cmt">//--- stop calculating the total program runtime class="type">class="kw">double timeGPUTotal=(GetTickCount()-st)/class="num">1000.; Print("OpenCL total: time = "+d2s(timeGPUTotal,class="num">3)+" sec."); destroyOpenCL(clCtx,clPrg,clKrn,clMemIn1,clMemIn2,clMemOut); class=class="str">"cmt">//--- calculate the time elapsed Print("CPUTime / GPUTotalTime = "+d2s(timeCPU/timeGPUTotal,class="num">3)); class=class="str">"cmt">//--- debugging: random checks. for(class="type">int i=class="num">0; i<class="num">10; i++) checkRandom(buf,ROWS1,COLS2); Print("________________________"); class="kw">return; }
◍ 矩阵内存布局与随机初始化
在 MT5 里跑 GPU 矩阵乘法前，先得把输入矩阵拍平成一维数组。代码用行主序（row-major）存：Matr[row][col] 对应 buff[row*N + col]，N 是列数。这样 CPU 和 OpenCL 内核读起来不会错位。 初始化阶段用 ArrayResize 给 first / second / thirdGPU / thirdCPU 定长，长度分别是 ROWS1*COLSROWS、COLSROWS*COLS2、ROWS1*COLS2。thirdCPU 和 thirdGPU 先用 ArrayInitialize 填 0.0，避免脏内存影响后续比对。 genMatrices() 按双层循环填满两个输入矩阵，每个元素由 genVal() 给。genVal() 用 MathRand() 减 16383.5 再除以 32767，把 [0,32767] 的整数映射成约 [-0.5, 0.5] 的均匀浮点值——这个区间在回测里能压住数值溢出，也方便肉眼看偏差。 开 MT5 把这段抄进 EA，改 ROWS1、COLSROWS、COLS2 三个宏，Print 出来的 device 和维度能对上，就说明内存排布没歪。外汇与贵金属杠杆高，这类底层Benchmark只验证计算正确性，不涉及任何交易信号。
Print("======================================="); Print("1st OCL martices mul: device = "+i2s(_device)+"; ROWS1 = " +i2s(ROWS1)+" "; COLSROWS = "+i2s(COLSROWS)+"; COLS2 = "+i2s(COLS2)); class=class="str">"cmt">//--- set the required sizes of linear representations of the input and output matrices ArrayResize(first,ROWS1*COLSROWS); ArrayResize(second,COLSROWS*COLS2); ArrayResize(thirdGPU,ROWS1*COLS2); ArrayResize(thirdCPU,ROWS1*COLS2); class=class="str">"cmt">//--- generate both input matrices and initialize the output to zeros genMatrices(); ArrayInitialize( thirdCPU, class="num">0.0 ); ArrayInitialize( thirdGPU, class="num">0.0 ); class="kw">return; } class=class="str">"cmt">//+------------------------------------------------------------------+ class=class="str">"cmt">//| initAllDataCPU | class=class="str">"cmt">//| lay out in row-major order, Matr[ M(rows) ][ N(column) ]: | class=class="str">"cmt">//| Matr[row][column] = buff[row * N(columns in the matrix) + column]| class=class="str">"cmt">//| generate initial matrices; this generation is not reflected | class=class="str">"cmt">//| in the final runtime calculation | class=class="str">"cmt">//| buffers are filled in row-major order! | class=class="str">"cmt">//+------------------------------------------------------------------+ class="type">void genMatrices() { for(class="type">int r=class="num">0; r<ROWS1; r++) for(class="type">int c=class="num">0; c<COLSROWS; c++) first[r*COLSROWS+c]=genVal(); for(class="type">int r=class="num">0; r<COLSROWS; r++) for(class="type">int c=class="num">0; c<COLS2; c++) second[r*COLS2+c]=genVal(); class="kw">return; } class=class="str">"cmt">//+------------------------------------------------------------------+ class=class="str">"cmt">//| genVal | class=class="str">"cmt">//| generate one value of the matrix element: | class=class="str">"cmt">//| uniformly distributed value lying in the range [-class="num">0.5; class="num">0.5] | class=class="str">"cmt">//+------------------------------------------------------------------+ REALTYPE genVal() { class="kw">return(REALTYPE)((MathRand()-class="num">16383.5)/class="num">32767.); } class=class="str">"cmt">//+------------------------------------------------------------------+ class=class="str">"cmt">//| initAllDataGPU | class=class="str">"cmt">//+------------------------------------------------------------------+ class="type">void initAllDataGPU(class="type">int &clCtx, class=class="str">"cmt">// context class="type">int& clPrg, class=class="str">"cmt">// program on the device class="type">int& clKrn, class=class="str">"cmt">// kernel
「把矩阵乘法丢给 GPU 跑通的三段式封装」
在 MT5 里调用 OpenCL 做矩阵乘法，核心不是写核函数本身，而是把上下文、缓冲区和执行流程用 MQL5 侧代码管起来。下面这段把输入矩阵 first、second 和输出 thirdGPU 分别建缓冲、绑参、写设备，再交给 GPU 核 matricesMul 计算。 初始化函数里先 WriteCLProgram 把核代码落盘成 matr_mul_OCL_1st.cl，随后 CLContextCreate → CLProgramCreate → CLKernelCreate 三连建立运行环境。三个缓冲区按维度算字节数：输入一为 ROWS1*COLSROWS*sizeof(REALTYPE)，输入二为 COLSROWS*COLS2*sizeof(REALTYPE)，输出为 ROWS1*COLS2*sizeof(REALTYPE)，全部用 CL_MEM_READ_WRITE。 参数绑定用 CLSetKernelArgMem 按索引 0/1/2 挂上三个缓冲，接着 CLBufferWrite 把主机端矩阵拷进去，thirdGPU 初始全 0。executeGPU 里 works 设成 {ROWS1, COLS2} 的二维网格，CLExecute 发任务；算完 readOutBuf 用 CLBufferRead 把结果读回 buf，并打印 read 元素数供核对。 外汇与贵金属行情矩阵化处理属高风险实验，GPU 加速只解决算力，不消除模型误差。开 MT5 把这段接进自己的 EA，先以小维度（如 ROWS1=COLSROWS=COLS2=64）验证读回元素数是否等于 4096。
class="type">int& clMemIn1, class=class="str">"cmt">// first(input) buffer class="type">int& clMemIn2, class=class="str">"cmt">// second(input) buffer class="type">int& clMemOut) class=class="str">"cmt">// third(output) buffer { class=class="str">"cmt">//--- write the kernel code to a file WriteCLProgram(); class=class="str">"cmt">//--- create context, program and kernel clCtx = CLContextCreate( _device ); clPrg = CLProgramCreate( clCtx, clSrc ); clKrn = CLKernelCreate( clPrg, "matricesMul" ); class=class="str">"cmt">//--- create all three buffers for the three matrices class=class="str">"cmt">//--- first matrix - input clMemIn1=CLBufferCreate(clCtx,ROWS1 *COLSROWS*class="kw">sizeof(REALTYPE),CL_MEM_READ_WRITE); class=class="str">"cmt">//--- second matrix - input clMemIn2=CLBufferCreate(clCtx,COLSROWS*COLS2 *class="kw">sizeof(REALTYPE),CL_MEM_READ_WRITE); class=class="str">"cmt">//--- third matrix - output clMemOut=CLBufferCreate(clCtx,ROWS1 *COLS2 *class="kw">sizeof(REALTYPE),CL_MEM_READ_WRITE); class=class="str">"cmt">//--- set arguments to the kernel CLSetKernelArgMem(clKrn,class="num">0,clMemIn1); CLSetKernelArgMem(clKrn,class="num">1,clMemIn2); CLSetKernelArgMem(clKrn,class="num">2,clMemOut); class=class="str">"cmt">//--- write the generated matrices to the device buffers CLBufferWrite(clMemIn1,first); CLBufferWrite(clMemIn2,second); CLBufferWrite(clMemOut,thirdGPU); class=class="str">"cmt">// class="num">0.0 everywhere class="kw">return; } class=class="str">"cmt">//+------------------------------------------------------------------+ class=class="str">"cmt">//| WriteCLProgram | class=class="str">"cmt">//+------------------------------------------------------------------+ class="type">void WriteCLProgram() { class="type">int h=FileOpen("matr_mul_OCL_1st.cl",FILE_WRITE|FILE_TXT|FILE_ANSI); FileWrite(h,clSrc); FileClose(h); } class=class="str">"cmt">//+------------------------------------------------------------------+ class=class="str">"cmt">//| executeGPU | class=class="str">"cmt">//+------------------------------------------------------------------+ class="type">void executeGPU(class="type">int clKrn) { class=class="str">"cmt">//--- set the workspace parameters for the task and execute the OpenCL program class="type">uint offs[ class="num">2 ] = { class="num">0, class="num">0 }; class="type">uint works[ class="num">2 ] = { ROWS1, COLS2 }; class="type">bool ex=CLExecute(clKrn,class="num">2,offs,works); class="kw">return; } class=class="str">"cmt">//+------------------------------------------------------------------+ class=class="str">"cmt">//| readOutBuf | class=class="str">"cmt">//+------------------------------------------------------------------+ class="type">void readOutBuf(class="type">int clMemOut,REALTYPE &buf[]) { ArrayResize(buf,COLS2*ROWS1); class=class="str">"cmt">//--- buf - a copy of what is written to the buffer thirdGPU[] class="type">uint read=CLBufferRead(clMemOut,buf); Print("read = "+i2s(read)+" elements"); class="kw">return; } class=class="str">"cmt">//+------------------------------------------------------------------+
收尾时的资源释放与精度抽查
跑完 OpenCL 矩阵乘法，第一件事是把显存和上下文按创建逆序拆掉。destroyOpenCL 里先 CLBufferFree 三个缓冲区，再依次释放内核、程序、上下文，顺序反了在部分驱动上会报句柄泄漏。 checkRandom 不验证全量结果，只随机抽一个 (r0,c0) 位置，用 CPU 侧 thirdCPU 和 OpenCL 回读 buf 做交叉比对，打印 8 位小数。这种抽样在回测矩阵规模大时够用，但漏判概率随矩阵维度上升而增大。 genRnd 用 MathRand()/32767 做归一化取整，注意 32767 是 MQL5 随机上限，不是 32768，边界值会略偏。 实盘日志里 EURUSD H1 上那次跑分：CPU 纯算 115.628 秒，OpenCL 总耗时 9.266 秒，CPU/GPU 时间比 12.479，读取 400 万元素。外汇与贵金属杠杆高，这类 GPU 加速只解决计算瓶颈，不改善信号胜率，实盘仍可能大幅回撤。
class=class="str">"cmt">//| destroyOpenCL | class=class="str">"cmt">//+------------------------------------------------------------------+ class="type">void destroyOpenCL(class="type">int clCtx,class="type">int clPrg,class="type">int clKrn,class="type">int clMemIn1,class="type">int clMemIn2,class="type">int clMemOut) { class=class="str">"cmt">//--- destroy all that was created for calculations on the OpenCL device in reverse order CLBufferFree(clMemIn1); CLBufferFree(clMemIn2); CLBufferFree(clMemOut); CLKernelFree(clKrn); CLProgramFree(clPrg); CLContextFree(clCtx); class="kw">return; } class=class="str">"cmt">//+------------------------------------------------------------------+ class=class="str">"cmt">//| checkRandom | class=class="str">"cmt">//| random check of calculation accuracy | class=class="str">"cmt">//+------------------------------------------------------------------+ class="type">void checkRandom(REALTYPE &buf[],class="type">int rows,class="type">int cols) { class="type">int r0 = genRnd( rows ); class="type">int c0 = genRnd( cols ); REALTYPE sum=class="num">0.0; for(class="type">int runningIdx=class="num">0; runningIdx<COLSROWS; runningIdx++) sum+=first[r0*COLSROWS+runningIdx]* second[runningIdx*COLS2+c0]; class=class="str">"cmt">//--- element of the buffer m[] REALTYPE bufElement=buf[r0*COLS2+c0]; class=class="str">"cmt">//--- element of the matrix not calculated in OpenCL REALTYPE CPUElement=thirdCPU[r0*COLS2+c0]; Print("sum( "+i2s(r0)+","+i2s(c0)+" ) = "+d2s(sum,class="num">8)+ "; thirdCPU[ "+i2s(r0)+","+i2s(c0)+" ] = "+d2s(CPUElement,class="num">8)+ "; buf[ "+i2s(r0)+","+i2s(c0)+" ] = "+d2s(bufElement,class="num">8)); class="kw">return; } class=class="str">"cmt">//+------------------------------------------------------------------+ class=class="str">"cmt">//| genRnd | class=class="str">"cmt">//+------------------------------------------------------------------+ class="type">int genRnd(class="type">int max) { class="kw">return(class="type">int)(MathRand()/class="num">32767.*max); }