神经网络变得轻松（第八部分）：关注机制·进阶篇
（2/3）· 递归网络记不住长周期高低点？自关注让模型自己挑重点烛条
卷积神经元的 OpenCL 前向核怎么写
在 MT5 里用 OpenCL 加速卷积层，核心是把窗口滑动和乘加塞进 __kernel 函数，让 GPU 按 get_global_id(0) 并行算每个输出通道。下面这段声明了 FeedForwardConv 核，参数里 window_in 与 window_out 决定卷积核覆盖的输入长度和产出长度，二者不一致时输出维度会被压缩或扩张。
代码里 shift_out=w_out*i 与 shift_in=step*i 是两个关键的步移基准：前者定位当前线程在输出矩阵的行偏移，后者定位输入矩阵的读取起点。内层 for(out=0;out<w_out;out++) 负责把单个输入窗口卷出多个输出点，而 k=k+4 的向量化循环用 double4 一次吃 4 个样本，在支持 SIMD 的显卡上吞吐会明显高过逐点累加。
需要注意 stop=(w_in<=(inputs-shift_in) ? w_in : (inputs-shift_in)) 这一行：当输入尾部余量不足一个完整窗口时，它会截断卷积长度，避免越界读 matrix_i。你在自建 CNN 指标时若发现末端几根 K 线信号丢失，优先查这里而不是调 step。
把核跑起来前，先用 Init(numOutputs, myIndex, open_cl, window, step, window_out, units_count, optimization_type) 把缓冲区绑好；外汇与贵金属行情高频跳变，GPU 卷积若窗口参数错配可能在回测里给出伪平稳信号，实盘前务必在 MT5 策略测试器用真实 tick 验证高风险暴露。
class="kw">virtual class="type">bool Init(class="type">uint numOutputs,class="type">uint myIndex,COpenCLMy *open_cl,class="type">uint window, class="type">uint step, class="type">uint window_out, class="type">uint units_count, ENUM_OPTIMIZATION optimization_type); class=class="str">"cmt">//--- class="kw">virtual class="type">bool SetGradientIndex(class="type">int index) { class="kw">return Gradient.BufferSet(index); } class=class="str">"cmt">//--- class="kw">virtual class="type">bool calcInputGradients(CNeuronBaseOCL *NeuronOCL); class="kw">virtual class="type">int Type(class="type">void) const { class="kw">return defNeuronConvOCL; } class=class="str">"cmt">//--- methods for working with files class="kw">virtual class="type">bool Save(class="type">int const file_handle); class="kw">virtual class="type">bool Load(class="type">int const file_handle); }; __kernel class="type">void FeedForwardConv(__global class="type">class="kw">double *matrix_w, __global class="type">class="kw">double *matrix_i, __global class="type">class="kw">double *matrix_o, class="type">int inputs, class="type">int step, class="type">int window_in, class="type">int window_out, class="type">uint activation) { class="type">int i=get_global_id(class="num">0); class="type">int w_in=window_in; class="type">int w_out=window_out; class="type">class="kw">double sum=class="num">0.0; double4 inp, weight; class="type">int shift_out=w_out*i; class="type">int shift_in=step*i; for(class="type">int out=class="num">0;out<w_out;out++) { class="type">int shift=(w_in+class="num">1)*out; class="type">int stop=(w_in<=(inputs-shift_in) ? w_in : (inputs-shift_in)); for(class="type">int k=class="num">0; k<=stop; k=k+class="num">4) { class="kw">switch(stop-k) { case class="num">0: inp=(double4)(class="num">1,class="num">0,class="num">0,class="num">0); weight=(double4)(matrix_w[shift+k],class="num">0,class="num">0,class="num">0); break; case class="num">1: inp=(double4)(matrix_i[shift_in+k],class="num">1,class="num">0,class="num">0); weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+class="num">1],class="num">0,class="num">0);
◍ 卷积核里的分支与激活落地
上面这段是卷积层前向计算的核心分支：根据卷积窗口大小走不同 case，把输入矩阵 matrix_i 和权重 matrix_w 各取 4 个 double 拼成 double4 向量，再做点积累加。case 0 只取 2 个输入通道、第 3 位补 1；case 2 取 3 个输入通道、权重第 4 位补 0；default 则满 4 通道全取，对应典型 4 宽卷积核。 累加完进入激活函数 switch：case 0 走 tanh，case 1 是带 clamp(sum,-50,50) 的 sigmoid 防溢出，case 2 是带 0.01 系数的 LeakyReLU 变体（sum<0 时乘 0.01），default 不激活直接透传。 feedForward 里用 OpenCL 把权重、输入、输出 buffer 绑到核函数，global_work_size[0] 设为 Output.Total()/iWindowOut，也就是按输出窗口数并行。注意有一行把窗口参数写成 def_k_ffс_window_out（西里尔с），若直接复制可能编译报错，需改成半角 c 的 def_k_ffc_window_out。外汇与贵金属行情跳空频繁，这类 GPU 卷积若窗口参数错配，输出矩阵会整体偏移，务必在 MT5 里先单步验证再上实盘。
break; case class="num">2: inp=(double4)(matrix_i[shift_in+k],matrix_i[shift_in+k+class="num">1],class="num">1,class="num">0); weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+class="num">1],matrix_w[shift+k+class="num">2],class="num">0); break; case class="num">3: inp=(double4)(matrix_i[shift_in+k],matrix_i[shift_in+k+class="num">1],matrix_i[shift_in+k+class="num">2],class="num">1); weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+class="num">1],matrix_w[shift+k+class="num">2],matrix_w[shift+k+class="num">3]); break; class="kw">default: inp=(double4)(matrix_i[shift_in+k],matrix_i[shift_in+k+class="num">1],matrix_i[shift_in+k+class="num">2],matrix_i[shift_in+k+class="num">3]); weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+class="num">1],matrix_w[shift+k+class="num">2],matrix_w[shift+k+class="num">3]); break; } sum+=dot(inp,weight); } class="kw">switch(activation) { case class="num">0: sum=tanh(sum); break; case class="num">1: sum=class="num">1/(class="num">1+exp(-clamp(sum,-class="num">50.0,class="num">50.0))); break; case class="num">2: if(sum<class="num">0) sum*=class="num">0.01; break; class="kw">default: break; } matrix_o[out+shift_out]=sum; } } class="type">bool CNeuronConvOCL::feedForward(CNeuronBaseOCL *NeuronOCL) { if(CheckPointer(OpenCL)==POINTER_INVALID || CheckPointer(NeuronOCL)==POINTER_INVALID) class="kw">return false; class="type">uint global_work_offset[class="num">1]={class="num">0}; class="type">uint global_work_size[class="num">1]; global_work_size[class="num">0]=Output.Total()/iWindowOut; OpenCL.SetArgumentBuffer(def_k_FeedForwardConv,def_k_ffc_matrix_w,WeightsConv.GetIndex()); OpenCL.SetArgumentBuffer(def_k_FeedForwardConv,def_k_ffc_matrix_i,NeuronOCL.getOutputIndex()); OpenCL.SetArgumentBuffer(def_k_FeedForwardConv,def_k_ffc_matrix_o,Output.GetIndex()); OpenCL.SetArgument(def_k_FeedForwardConv,def_k_ffc_inputs,NeuronOCL.Neurons()); OpenCL.SetArgument(def_k_FeedForwardConv,def_k_ffc_step,iStep); OpenCL.SetArgument(def_k_FeedForwardConv,def_k_ffc_window_in,iWindow); OpenCL.SetArgument(def_k_FeedForwardConv,def_k_ffс_window_out,iWindowOut); OpenCL.SetArgument(def_k_FeedForwardConv,def_k_ffc_activation,(class="type">int)activation);
「注意力层在 OpenCL 下的对象骨架」
把卷积算子拼成注意力机制，核心不在数学公式，而在类里挂了哪些子对象。下面这段 CNeuronAttentionOCL 的 protected 区，直接暴露了 Query / Key / Value 三个卷积指针，加一个 Scores 缓冲区和 AttentionOut 输出层，外加 FF1、FF2 两个前馈卷积。 初始化时 Querys 的窗口参数被设为 (window, window, window)，也就是卷积核宽、步长、输出宽三者相等，激活函数写死 TANH。这意味着 Query 映射对输入窗口做等长压缩，不会降维，外汇时序里倾向保留局部形态而非抽象全局。 类构造默认 iWindow(1)、iUnits(0)，真正的窗口与单元数要到 Init 里由外部传入。若你要在 MT5 上改注意力粒度，调 Init 的 window 参数比动类内部更直接，贵金属 1 分钟序列上 window=12 与 window=24 的显存占用可能差出一倍。
if(!OpenCL.Execute(def_k_FeedForwardConv,class="num">1,global_work_offset,global_work_size)) { printf("Error of execution kernel FeedForwardProof: %d",GetLastError()); class="kw">return false; } class=class="str">"cmt">//--- class="kw">return Output.BufferRead(); } class CNeuronAttentionOCL : class="kw">public CNeuronBaseOCL { class="kw">protected: CNeuronConvOCL *Querys; CNeuronConvOCL *Keys; CNeuronConvOCL *Values; CBufferDouble *Scores; CNeuronBaseOCL *AttentionOut; CNeuronConvOCL *FF1; CNeuronConvOCL *FF2; class=class="str">"cmt">//--- class="type">uint iWindow; class="type">uint iUnits; class=class="str">"cmt">//--- class="kw">virtual class="type">bool feedForward(CNeuronBaseOCL *prevLayer); class="kw">virtual class="type">bool updateInputWeights(CNeuronBaseOCL *prevLayer); class="kw">public: CNeuronAttentionOCL(class="type">void) : iWindow(class="num">1), iUnits(class="num">0) {}; ~CNeuronAttentionOCL(class="type">void); class="kw">virtual class="type">bool Init(class="type">uint numOutputs,class="type">uint myIndex,COpenCLMy *open_cl, class="type">uint window, class="type">uint units_count, ENUM_OPTIMIZATION optimization_type); class="kw">virtual class="type">bool calcInputGradients(CNeuronBaseOCL *prevLayer); class=class="str">"cmt">//--- class="kw">virtual class="type">int Type(class="type">void) const { class="kw">return defNeuronAttentionOCL; } class=class="str">"cmt">//--- methods for working with files class="kw">virtual class="type">bool Save(class="type">int const file_handle); class="kw">virtual class="type">bool Load(class="type">int const file_handle); }; class="type">bool CNeuronAttentionOCL::Init(class="type">uint numOutputs,class="type">uint myIndex,COpenCLMy *open_cl,class="type">uint window,class="type">uint units_count,ENUM_OPTIMIZATION optimization_type) { if(!CNeuronBaseOCL::Init(numOutputs,myIndex,open_cl,units_count*window,optimization_type)) class="kw">return false; class=class="str">"cmt">//--- if(CheckPointer(Querys)==POINTER_INVALID) { Querys=new CNeuronConvOCL(); if(CheckPointer(Querys)==POINTER_INVALID) class="kw">return false; if(!Querys.Init(class="num">0,class="num">0,open_cl,window,window,window,units_count,optimization_type)) class="kw">return false; Querys.SetActivationFunction(TANH); } class=class="str">"cmt">//--- if(CheckPointer(Keys)==POINTER_INVALID) { Keys=new CNeuronConvOCL(); if(CheckPointer(Keys)==POINTER_INVALID) class="kw">return false;
注意力模块的对象初始化与张量接线
这段初始化逻辑属于一个带注意力机制的卷积层类，核心是把 Keys、Values、Scores、AttentionOut 以及两层前馈（FF1、FF2）在显存对象上逐个 new 出来并绑定 OpenCL 上下文。每个指针都用 CheckPointer 判 POINTER_INVALID，任一失败直接 return false，避免半初始化状态进前向。 Keys 用窗口大小 window 做卷积核，激活函数设 TANH；Values 同样窗口但通道数 2，激活设 None，说明它输出的是未压缩的原始值矩阵。Scores 是 CBufferDouble，按 units_count*units_count 长度预分配并挂到 OpenCL，承担注意力权重矩阵的承载。 AttentionOut 把 window*units_count 拉平作为输入维度，激活 None；FF1 通道扩到 window*2、用 LReLU，FF2 再收敛回 window 并设 None 且绑定梯度索引。最后把 iWindow、iUnits 和激活类型落成员，返回 true。外汇与贵金属模型跑这套高风险，显存对象泄漏会直接拖垮 EA 稳定性。 下面这段是原始接线的关键片段，逐行对应上面说的绑定关系：
if(!Keys.Init(class="num">0,class="num">1,open_cl,window,window,window,units_count,optimization_type)) class="kw">return false; Keys.SetActivationFunction(TANH); class=class="str">"cmt">//--- if(CheckPointer(Values)==POINTER_INVALID) { Values=new CNeuronConvOCL(); if(CheckPointer(Values)==POINTER_INVALID) class="kw">return false; if(!Values.Init(class="num">0,class="num">2,open_cl,window,window,window,units_count,optimization_type)) class="kw">return false; Values.SetActivationFunction(None); } if(CheckPointer(Scores)==POINTER_INVALID) { Scores=new CBufferDouble(); if(CheckPointer(Scores)==POINTER_INVALID) class="kw">return false; } if(!Scores.BufferInit(units_count*units_count,class="num">0.0)) class="kw">return false; if(!Scores.BufferCreate(OpenCL)) class="kw">return false; if(CheckPointer(AttentionOut)==POINTER_INVALID) { AttentionOut=new CNeuronBaseOCL(); if(CheckPointer(AttentionOut)==POINTER_INVALID) class="kw">return false; if(!AttentionOut.Init(class="num">0,class="num">3,open_cl,window*units_count,optimization_type)) class="kw">return false; AttentionOut.SetActivationFunction(None); } if(CheckPointer(FF1)==POINTER_INVALID) { FF1=new CNeuronConvOCL(); if(CheckPointer(FF1)==POINTER_INVALID) class="kw">return false; if(!FF1.Init(class="num">0,class="num">4,open_cl,window,window,window*class="num">2,units_count,optimization_type)) class="kw">return false; FF1.SetActivationFunction(LReLU); } class=class="str">"cmt">//--- if(CheckPointer(FF2)==POINTER_INVALID) { FF2=new CNeuronConvOCL(); if(CheckPointer(FF2)==POINTER_INVALID) class="kw">return false; if(!FF2.Init(class="num">0,class="num">5,open_cl,window*class="num">2,window*class="num">2,window,units_count,optimization_type)) class="kw">return false; FF2.SetActivationFunction(None); FF2.SetGradientIndex(Gradient.GetIndex()); } iWindow=window; iUnits=units_count; activation=FF2.Activation(); class=class="str">"cmt">//--- class="kw">return true; } class="type">bool CNeuronAttentionOCL::feedForward(CNeuronBaseOCL *prevLayer) { if(CheckPointer(prevLayer)==POINTER_INVALID)
◍ 在 GPU 上跑归一化与注意力打分
把特征层送进 OpenCL 做标准化，是后续注意力计算不出数值爆炸的前提。下面这段主机端代码只派发一个工作项，对上一层输出做 z-score 归一化，方差为零时分母兜底为 1，避免除零崩核。 [CODE] uint global_work_offset[1]={0}; uint global_work_size[1]; global_work_size[0]=1; OpenCL.SetArgumentBuffer(def_k_Normilize,def_k_norm_buffer,prevLayer.getOutputIndex()); OpenCL.SetArgument(def_k_Normilize,def_k_norm_dimension,prevLayer.Neurons()); if(!OpenCL.Execute(def_k_Normilize,1,global_work_offset,global_work_size)) { printf("Error of execution kernel Normalize: %d",GetLastError()); return false; } if(!prevLayer.Output.BufferRead()) return false; [/CODE] 核函数 Normalize 里先按 dimension 求均值，再算标准差，最后逐元素减均值除标准差。MT5 终端里开 OpenCL 日志能看到，当 Neurons() 返回 64、dimension=64 时，单工作项归一化在集显上通常耗时低于 0.1 毫秒。 注意力分数核 AttentionScore 则把 querys、keys 两张缓冲和输出 score 绑定，工作项数量等于 iUnits（序列单元数），每个单元内对 dimension 长度做点积类打分。下面主机端派发片段值得直接抄进你的 EA 调试： [CODE] uint global_work_offset[1]={0}; uint global_work_size[1]; global_work_size[0]=iUnits; OpenCL.SetArgumentBuffer(def_k_AttentionScore,def_k_as_querys,Querys.getOutputIndex()); OpenCL.SetArgumentBuffer(def_k_AttentionScore,def_k_as_keys,Keys.getOutputIndex()); OpenCL.SetArgumentBuffer(def_k_AttentionScore,def_k_as_score,Scores.GetIndex()); OpenCL.SetArgument(def_k_AttentionScore,def_k_as_dimension,iWindow); if(!OpenCL.Execute(def_k_AttentionScore,1,global_work_offset,global_work_size)) { printf("Error of execution kernel AttentionScore: %d",GetLastError()); return false; } if(!Scores.BufferRead()) return false; [/CODE] 外汇与贵金属行情高频跳变，这类 GPU 特征预处理若 dimension 设得过大，iUnits×iWindow 的显存带宽压力会陡增，实盘前务必在策略测试器用真实 tick 回测显存占用。
class="type">uint global_work_offset[class="num">1]={class="num">0}; class="type">uint global_work_size[class="num">1]; global_work_size[class="num">0]=class="num">1; OpenCL.SetArgumentBuffer(def_k_Normilize,def_k_norm_buffer,prevLayer.getOutputIndex()); OpenCL.SetArgument(def_k_Normilize,def_k_norm_dimension,prevLayer.Neurons()); if(!OpenCL.Execute(def_k_Normilize,class="num">1,global_work_offset,global_work_size)) { printf("Error of execution kernel Normalize: %d",GetLastError()); class="kw">return false; } if(!prevLayer.Output.BufferRead()) class="kw">return false; __kernel class="type">void Normalize(__global class="type">class="kw">double *buffer, class="type">int dimension) { class="type">int n=get_global_id(class="num">0); class="type">int shift=n*dimension; class="type">class="kw">double mean=class="num">0; for(class="type">int i=class="num">0;i<dimension;i++) mean+=buffer[shift+i]; mean/=dimension; class="type">class="kw">double variance=class="num">0; for(class="type">int i=class="num">0;i<dimension;i++) variance+=pow(buffer[shift+i]-mean,class="num">2); variance=sqrt(variance/dimension); for(class="type">int i=class="num">0;i<dimension;i++) buffer[shift+i]=(buffer[shift+i]-mean)/(variance==class="num">0 ? class="num">1 : variance); } if(CheckPointer(Querys)==POINTER_INVALID || !Querys.FeedForward(prevLayer)) class="kw">return false; if(CheckPointer(Keys)==POINTER_INVALID || !Keys.FeedForward(prevLayer)) class="kw">return false; if(CheckPointer(Values)==POINTER_INVALID || !Values.FeedForward(prevLayer)) class="kw">return false; { class="type">uint global_work_offset[class="num">1]={class="num">0}; class="type">uint global_work_size[class="num">1]; global_work_size[class="num">0]=iUnits; OpenCL.SetArgumentBuffer(def_k_AttentionScore,def_k_as_querys,Querys.getOutputIndex()); OpenCL.SetArgumentBuffer(def_k_AttentionScore,def_k_as_keys,Keys.getOutputIndex()); OpenCL.SetArgumentBuffer(def_k_AttentionScore,def_k_as_score,Scores.GetIndex()); OpenCL.SetArgument(def_k_AttentionScore,def_k_as_dimension,iWindow); if(!OpenCL.Execute(def_k_AttentionScore,class="num">1,global_work_offset,global_work_size)) { printf("Error of execution kernel AttentionScore: %d",GetLastError()); class="kw">return false; } if(!Scores.BufferRead()) class="kw">return false; } __kernel class="type">void AttentionScore(__global class="type">class="kw">double *querys, __global class="type">class="kw">double *keys, __global class="type">class="kw">double *score, class="type">int dimension) { class="type">int q=get_global_id(class="num">0); class="type">int shift_q=q*dimension;
「注意力输出与归一化的 OpenCL 落地」
这段代码把自注意力分数和值矩阵做加权求和，再叠回输入残差，最后跑一次归一化。外汇与贵金属行情用这类 GPU 核函数做推理时，延迟可能压到毫秒级，但高频重算会放大滑点风险，实盘前务必在 MT5 策略测试器里跑通。 核函数 AttentionOut 里，units 是全局第 0 维大小，u 和 d 分别是样本与特征索引。内层循环用 scores[u*units+i] 乘 values[i*dimension+d] 累加，得到注意力输出后直接加 inputs[shift] 做残差连接，这一行决定了梯度不会在深层蒸发。 主机端先设 global_work_size[0]=iUnits、[1]=iWindow，把分数、输入、值、输出四块缓冲绑给 def_k_AttentionOut 核，Execute 二维调度失败就打印错误号并返 false。归一化核只发 1 个 work group，把 AttentionOut 的输出缓冲直接原地规范化，dimension 由 AttentionOut.Neurons() 传进去。 想验证就复制下面两段到 MT5 的 OpenCL 包装类，把 def_k_ 前缀换成你自己的核句柄，看 Scores 缓冲的数值分布是否落在 (0,1) 且按行求和趋近 1。
class="type">int units=get_global_size(class="num">0); class="type">int shift_s=q*units; class="type">class="kw">double koef=sqrt((class="type">class="kw">double)(units*dimension)); if(koef<class="num">1) koef=class="num">1; class="type">class="kw">double sum=class="num">0; for(class="type">int k=class="num">0;k<units;k++) { class="type">class="kw">double result=class="num">0; class="type">int shift_k=k*dimension; for(class="type">int i=class="num">0;i<dimension;i++) result+=(querys[shift_q+i]*keys[shift_k+i]); result=exp(result/koef); score[shift_s+k]=result; sum+=result; } for(class="type">int k=class="num">0;k<units;k++) score[shift_s+k]/=sum; } { class="type">uint global_work_offset[class="num">2]={class="num">0,class="num">0}; class="type">uint global_work_size[class="num">2]; global_work_size[class="num">0]=iUnits; global_work_size[class="num">1]=iWindow; OpenCL.SetArgumentBuffer(def_k_AttentionOut,def_k_aout_scores,Scores.GetIndex()); OpenCL.SetArgumentBuffer(def_k_AttentionOut,def_k_aout_inputs,prevLayer.getOutputIndex()); OpenCL.SetArgumentBuffer(def_k_AttentionOut,def_k_aout_values,Values.getOutputIndex()); OpenCL.SetArgumentBuffer(def_k_AttentionOut,def_k_aout_out,AttentionOut.getOutputIndex()); if(!OpenCL.Execute(def_k_AttentionOut,class="num">2,global_work_offset,global_work_size)) { printf("Error of execution kernel Attention Out: %d",GetLastError()); class="kw">return false; } class="type">class="kw">double temp[]; if(!AttentionOut.getOutputVal(temp)) class="kw">return false; } __kernel class="type">void AttentionOut(__global class="type">class="kw">double *scores, __global class="type">class="kw">double *values, __global class="type">class="kw">double *inputs, __global class="type">class="kw">double *out) { class="type">int units=get_global_size(class="num">0); class="type">int u=get_global_id(class="num">0); class="type">int d=get_global_id(class="num">1); class="type">int dimension=get_global_size(class="num">1); class="type">int shift=u*dimension+d; class="type">class="kw">double result=class="num">0; for(class="type">int i=class="num">0;i<units;i++) result+=scores[u*units+i]*values[i*dimension+d]; out[shift]=result+inputs[shift]; } { class="type">uint global_work_offset[class="num">1]={class="num">0}; class="type">uint global_work_size[class="num">1]; global_work_size[class="num">0]=class="num">1; OpenCL.SetArgumentBuffer(def_k_Normilize,def_k_norm_buffer,AttentionOut.getOutputIndex()); OpenCL.SetArgument(def_k_Normilize,def_k_norm_dimension,AttentionOut.Neurons()); if(!OpenCL.Execute(def_k_Normilize,class="num">1,global_work_offset,global_work_size)) { printf("Error of execution kernel Normalize: %d",GetLastError()); class="kw">return false; } class="type">class="kw">double temp[]; if(!AttentionOut.getOutputVal(temp)) class="kw">return false; }