神经网络变得轻松(第三十七部分):分散关注度·进阶篇
📘

神经网络变得轻松(第三十七部分):分散关注度·进阶篇

第 2/3 篇

在 GPU 上反向传播稀疏注意力的梯度

这段 OpenCL 内核处理多头稀疏注意力(MHSparseAttention)的反向传播,内核索引在宏定义里固定为 44(def_k_MHSparseAttentionScore 44),说明它挂在某套自定义神经算子体系的第 44 号位置。 内核先取全局 ID 算出当前单元 u 与头 h,再用 get_global_size 拿总单元数 units 和头数 heads。缩放系数 koef 取 dimension 的平方根,若小于 1 则钳为 1,避免除数量级爆炸。 分数梯度循环里,对 v 从 0 到 units 遍历:只处理 scores 为正的位置,负分直接 continue。每个有效分数用对应 Value 与上游 gradient 做点积得到 sg,再乘 sigmoid 近似门控 (s<1 ? s*(1-s) : 1) 并除以 koef,写回 scores_g。 下半段算 Q/K/V 三者的梯度。对每个维度 d,遍历 l 累加:vg 来自 gradient 乘分数,kg/qg 则来自分数梯度乘对端 Q/K 向量。最终按 3*u 偏移写回 qkv_g 对应槽位。barrier(CLK_GLOBAL_MEM_FENCE) 保证分数梯度先全员落盘再读。 在 MT5 里若你要改这套算子的头数或维度,直接调 dimension 与 heads 的传入值即可,但内核 44 的索引不要动,否则神经网络描述符会找不到反向入口。外汇与贵金属策略接这类 GPU 算子时波动放大,回测通过不代表实盘稳健,属于高风险用法。

MQL5 / C++
 __global class="type">class="kw">float *scores, __global class="type">class="kw">float *scores_g,
 __global class="type">class="kw">float *gradient, class="type">int dimension)
 {
  class="type">int u = get_global_id(class="num">0);
  class="type">int h = get_global_id(class="num">1);
  class="type">int units = get_global_size(class="num">0);
  class="type">int heads = get_global_size(class="num">1);
  class="type">class="kw">float koef = sqrt((class="type">class="kw">float)dimension);
  if(koef < class="num">1)
     koef = class="num">1;
class=class="str">"cmt">//--- Calculating score&class="macro">#x27;s gradients
  class="type">uint shift_s = units * (h + u * heads);
  for(class="type">int v = class="num">0; v < units; v++)
    {
      class="type">class="kw">float s = scores[shift_s + v];
      if(s <= class="num">0)
        class="kw">continue;
      class="type">class="kw">float sg = class="num">0;
      class="type">int shift_v = dimension * (h + heads * (class="num">3 * v + class="num">2));
      class="type">int shift_g = dimension * (h + heads * v);
      for(class="type">int d = class="num">0; d < dimension; d++)
        sg += qkv[shift_v + d] * gradient[shift_g + d];
      scores_g[shift_s + v] = sg * (s < class="num">1 ? s * (class="num">1 - s) : class="num">1) / koef;
    }
  barrier(CLK_GLOBAL_MEM_FENCE);
class=class="str">"cmt">//--- Calculating gradients for Query, Key and Value
  class="type">uint shift_qg = dimension * (h + class="num">3 * u * heads);
  class="type">uint shift_kg = dimension * (h + (class="num">3 * u + class="num">1) * heads);
  class="type">uint shift_vg = dimension * (h + (class="num">3 * u + class="num">2) * heads);
  for(class="type">int d = class="num">0; d < dimension; d++)
    {
      class="type">class="kw">float vg = class="num">0;
      class="type">class="kw">float qg = class="num">0;
      class="type">class="kw">float kg = class="num">0;
      for(class="type">int l = class="num">0; l < units; l++)
       {
        class="type">class="kw">float sg = scores[shift_s + l];
        if(sg <= class="num">0)
          class="kw">continue;
        class="type">uint shift_q = dimension * (h + class="num">3 * l * heads) + d;
        class="type">uint shift_k = dimension * (h + (class="num">3 * l + class="num">1) * heads) + d;
        class="type">uint shift_g = dimension * (h + heads * l) + d;
        class=class="str">"cmt">//---
        vg += gradient[shift_g] * sg;
        sg = scores_g[shift_s + l];
        kg += sg * qkv[shift_q];
        qg += sg * qkv[shift_k];
        }
      qkv_g[shift_qg + d] = qg;
      qkv_g[shift_kg + d] = kg;
      qkv_g[shift_vg + d] = vg;
    }
 }
class="macro">#define def_k_MHSparseAttentionScore    class="num">44 class=class="str">"cmt">///< Index of the kernel of the multi-heads sparse attention neuron

◍ 稀疏多头注意力在 OpenCL 侧的落地参数

把多头稀疏注意力搬进 MT5 的 OpenCL 管线,先得把核函数数量与索引钉死。示例代码里 opencl.SetKernelsCount(46) 说明整套网络至少挂了 46 个核,其中稀疏注意力用了两个:评分核 MHSparseAttentionScore 索引 45、输出核 MHSparseAttentionOut 索引 45 之前的 3 号(宏 def_k_mhas_sparse 值为 3,注释标明稀疏系数低于 1.0)。 类 CNeuronMLMHSparseAttention 继承自 CNeuronMLMHAttentionOCL,构造函数把成员 m_dSparse 默认置为 0.3f,也就是默认丢弃约七成非关键注意力权重。通过 Sparse(float) 方法可在运行时改写该系数,值越小越稀疏,显存与算力占用倾向更低,但信息损失概率上升。 核创建失败必须拦在初始化阶段:KernelCreate 返回 false 时直接打印错误码与行号并 return false,避免在后续前向传播里出现空核崩溃。外汇与贵金属行情高频跳变,这类 GPU 加速模型若稀疏度调错,回测与实盘信号可能偏离,属高风险验证项,建议先在策略测试器用历史 tick 跑通再上模拟盘。

MQL5 / C++
class="macro">#define def_k_mhas_sparse        class="num">3   class=class="str">"cmt">///< less than class="num">1.0 coefficient of sparse
class="macro">#define def_k_MHSparseAttentionOut  class="num">45 class=class="str">"cmt">///< Index of the kernel of the multi-heads sparse attention neuron to calculate multi-heads out matrix(class="macro">#MHSparseAttentionOut)
   opencl.SetKernelsCount(class="num">46);
   if(!opencl.KernelCreate(def_k_MHSparseAttentionScore, "MHSparseAttentionScore"))
     {
       PrintFormat("Error of create kernell: %d line %d", GetLastError(), __LINE__);
       class="kw">return false;
     }
   if(!opencl.KernelCreate(def_k_MHSparseAttentionOut, "MHSparseAttentionOut"))
     {
       PrintFormat("Error of create kernell: %d line %d", GetLastError(), __LINE__);
       class="kw">return false;
     }
   class="type">bool      Create(CArrayObj *Description);
   class="type">bool      Load(class="type">class="kw">string file_name, class="type">class="kw">float &error, class="type">class="kw">float &undefine, class="type">class="kw">float &forecast, class="type">class="kw">datetime &time,
                          class="type">bool common = true);
   class="kw">virtual class="type">bool      Load(const class="type">int file_handle);
class CNeuronMLMHSparseAttention  : class="kw">public CNeuronMLMHAttentionOCL
  {
class="kw">protected:
   class="type">class="kw">float      m_dSparse;
   class="kw">virtual class="type">bool      AttentionScore(CBufferFloat *qkv, CBufferFloat *scores, class="type">bool mask = true);
   class="kw">virtual class="type">bool      AttentionOut(CBufferFloat *qkv, CBufferFloat *scores, CBufferFloat *out);
class="kw">public:
                     CNeuronMLMHSparseAttention(class="type">void)  :  m_dSparse(class="num">0.3f) {};
                    ~CNeuronMLMHSparseAttention(class="type">void) {};
   class="type">void      Sparse(class="type">class="kw">float value)  { m_dSparse = value;}
   class="type">class="kw">float     Sparse(class="type">void)         { class="kw">return m_dSparse; }
   class="kw">virtual class="type">int       Type(class="type">void) const  {  class="kw">return defNeuronMLMHSparseAttentionOCL;  }

「稀疏注意力类的存档与前向传播骨架」

在 MT5 的神经网络扩展里,CNeuronMLMHSparseAttention 继承自 CNeuronMLMHAttentionOCL,它额外多了一个 m_dSparse 浮点成员。Save 方法先调父类存档,再用 FileWriteFloat 把 m_dSparse 写进文件句柄,写入字节数不足 sizeof(float)(即 4 字节)就返回 false,这说明稀疏系数必须随模型权重一起落盘,否则重载会错位。 feedForward 是多层注意力堆叠的核心循环。代码里 iLayers 控制层数,每层先拿输入(首层用 NeuronOCL.getOutput(),其余用 FF_Tensors 偏移 6*i-4 的缓冲),过一次卷积生成 QKV;卷积核步长在 SGD 优化下权重索引跳 2,其他优化跳 3。 随后 AttentionScore 算注意力分数,AttentionOut 做多头拼接,再用 ConvolutionForward 把维度从 iWindowKey*iHeads 压回 iWindow(SGD 下权重跳 6,否则跳 9)。SumAndNormilize 做残差归一,最后接一个 LReLU 激活的 4*iWindow 扩展卷积。外汇与贵金属行情噪声大,直接拿这套结构跑实盘前,建议在策略测试器里先验证各层缓冲偏移是否与你改过的张量布局一致。

MQL5 / C++
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 CNeuronMLMHSparseAttention::Save(const class="type">int file_handle)
  {
   if(!CNeuronMLMHAttentionOCL::Save(file_handle))
      class="kw">return false;
   if(FileWriteFloat(file_handle, m_dSparse) < class="kw">sizeof(class="type">class="kw">float))
      class="kw">return false;
class=class="str">"cmt">//---
   class="kw">return true;
  }
class="type">bool CNeuronMLMHAttentionOCL::feedForward(CNeuronBaseOCL *NeuronOCL)
  {
   if(CheckPointer(NeuronOCL) == POINTER_INVALID)
      class="kw">return false;
class=class="str">"cmt">//---
   for(class="type">uint i = class="num">0; (i < iLayers && !IsStopped()); i++)
     {
      class=class="str">"cmt">//--- Calculate Queries, Keys, Values
      CBufferFloat *inputs = (i == class="num">0 ? NeuronOCL.getOutput() : FF_Tensors.At(class="num">6 * i - class="num">4));
      CBufferFloat *qkv = QKV_Tensors.At(i * class="num">2);
      if(IsStopped() || !ConvolutionForward(QKV_Weights.At(i * (optimization == SGD ? class="num">2 : class="num">3)),
inputs, qkv, iWindow, class="num">3 * iWindowKey * iHeads, None))
         class="kw">return false;
      class=class="str">"cmt">//--- Score calculation
      CBufferFloat *temp = S_Tensors.At(i * class="num">2);
      if(IsStopped() || !AttentionScore(qkv, temp, true))
         class="kw">return false;
      class=class="str">"cmt">//--- Multi-heads attention calculation
      CBufferFloat *out = AO_Tensors.At(i * class="num">2);
      if(IsStopped() || !AttentionOut(qkv, temp, out))
         class="kw">return false;
      class=class="str">"cmt">//--- Attention out calculation
      temp = FF_Tensors.At(i * class="num">6);
      if(IsStopped() || !ConvolutionForward(FF_Weights.At(i * (optimization == SGD ? class="num">6 : class="num">9)),
out, temp, iWindowKey * iHeads, iWindow, None))
         class="kw">return false;
      class=class="str">"cmt">//--- Sum and normilize attention
      if(IsStopped() || !SumAndNormilize(temp, inputs, temp))
         class="kw">return false;
      class=class="str">"cmt">//--- Feed Forward
      inputs = temp;
      temp = FF_Tensors.At(i * class="num">6 + class="num">1);
      if(IsStopped() || !ConvolutionForward(FF_Weights.At(i * (optimization == SGD ? class="num">6 : class="num">9) + class="num">1),
inputs, temp, iWindow, class="num">4 * iWindow, LReLU))
         class="kw">return false;
      out = FF_Tensors.At(i * class="num">6 + class="num">2);

稀疏注意力的 OpenCL 核调度细节

在 MT5 的自定义神经网络类里,多头稀疏注意力靠两个独立内核完成:先算得分再算输出。AttentionScore 里先校验 OpenCL 上下文与三个浮点缓冲指针有效性,任一为 POINTER_INVALID 直接返回 false,避免空指针在显存侧炸核。 内核启动采用二维全局工组尺寸,global_work_size[0] 取 iUnits、[1] 取 iHeads,相当于每个头、每个单元各跑一个线程。m_dSparse 以 float 形式塞进 def_k_mhas_sparse 参数,控制注意力矩阵的稀疏度,这个值调大可能让远端窗口权重衰减更狠。 若 OpenCL.Execute 返回失败,代码用 CLGetInfoString 抓错误描述并打印 __FUNCSIG__,方便在专家日志里定位是哪一层内核挂了。AttentionOut 同理,但多校验了 out 缓冲,且只 SetArgument 不执行返回判断的片段截在这里——实际完整版仍需 Execute 调用。 直接把这段逻辑抄进你的 EA 神经元模块,开 MT5 用 Print 看 iUnits*iHeads 的线程规模,外汇与贵金属杠杆品种下跑 GPU 核有爆显存风险,建议先用小窗口回测。

MQL5 / C++
if(IsStopped() || !ConvolutionForward(FF_Weights.At(i * (optimization == SGD ? class="num">6 : class="num">9) + class="num">2),
temp, out, class="num">4 * iWindow, iWindow, activation))
      class="kw">return false;
      class=class="str">"cmt">//--- Sum and normilize out
      if(IsStopped() || !SumAndNormilize(out, inputs, out))
         class="kw">return false;
   }
class=class="str">"cmt">//---
   class="kw">return true;
   }
class="type">bool CNeuronMLMHSparseAttention::AttentionScore(CBufferFloat *qkv, CBufferFloat *scores, class="type">bool mask = true)
   {
   if(CheckPointer(OpenCL) == POINTER_INVALID || CheckPointer(qkv) == POINTER_INVALID ||
CheckPointer(scores) == POINTER_INVALID)
      class="kw">return false;
class=class="str">"cmt">//---
   if(qkv.GetIndex() < class="num">0)
      class="kw">return false;
   if(scores.GetIndex() < class="num">0)
      class="kw">return false;
class=class="str">"cmt">//---
   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] = iHeads;
   OpenCL.SetArgumentBuffer(def_k_MHSparseAttentionScore, def_k_mhas_qkv, qkv.GetIndex());
   OpenCL.SetArgumentBuffer(def_k_MHSparseAttentionScore, def_k_mhas_score, scores.GetIndex());
   OpenCL.SetArgument(def_k_MHSparseAttentionScore, def_k_mhas_dimension, (class="type">int)iWindowKey);
   OpenCL.SetArgument(def_k_MHSparseAttentionScore, def_k_mhas_sparse, (class="type">class="kw">float)m_dSparse);
   if(!OpenCL.Execute(def_k_MHSparseAttentionScore, class="num">2, global_work_offset, global_work_size))
      {
         class="type">class="kw">string error;
         CLGetInfoString(OpenCL.GetContext(), CL_ERROR_DESCRIPTION, error);
         printf("Error of execution kernel %s: %s", __FUNCSIG__, error);
         class="kw">return false;
      }
class=class="str">"cmt">//---
   class="kw">return true;
   }
class="type">bool CNeuronMLMHSparseAttention::AttentionOut(CBufferFloat *qkv, CBufferFloat *scores, CBufferFloat *out)
   {
   if(CheckPointer(OpenCL) == POINTER_INVALID || CheckPointer(qkv) == POINTER_INVALID ||
      CheckPointer(scores) == POINTER_INVALID || CheckPointer(out) == POINTER_INVALID)
      class="kw">return false;
   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] = iHeads;
   if(qkv.GetIndex() < class="num">0)
      class="kw">return false;
   if(scores.GetIndex() < class="num">0)
      class="kw">return false;
   if(out.GetIndex() < class="num">0)
      class="kw">return false;
class=class="str">"cmt">//---
   OpenCL.SetArgumentBuffer(def_k_MHSparseAttentionOut, def_k_mhao_qkv, qkv.GetIndex());
   OpenCL.SetArgumentBuffer(def_k_MHSparseAttentionOut, def_k_mhao_score, scores.GetIndex());
   OpenCL.SetArgumentBuffer(def_k_MHSparseAttentionOut, def_k_mhao_out, out.GetIndex());
   OpenCL.SetArgument(def_k_MHSparseAttentionOut, def_k_mhao_dimension, (class="type">int)iWindowKey);

◍ 稀疏注意力层的 OpenCL 执行与实例化

在 MT5 的 OpenCL 路径里,内核真正跑起来之前要先确认 Execute 返回成功。下面这段把 def_k_MHSparseAttentionOut 内核以 2 维全局偏移和工作组尺寸下发,一旦失败就通过 CLGetInfoString 抓上下文错误描述并用 printf 打出函数签名和错误串,随后返回 false 终止。 if(!OpenCL.Execute(def_k_MHSparseAttentionOut, 2, global_work_offset, global_work_size)) { string error; CLGetInfoString(OpenCL.GetContext(), CL_ERROR_DESCRIPTION, error); printf("Error of execution kernel %s: %s", __FUNCSIG__, error); return false; } //--- return true; } 神经元工厂分支里,defNeuronMLMHSparseAttentionOCL 负责 new 一个 CNeuronMLMHSparseAttention。若指针校验为 POINTER_INVALID 就删掉临时对象 temp 并返回 false;Init 调用传入 outputs、窗口 window/window_out、步长 step、头数 count、层数 layers 等描述符字段,任一环节失败都先 delete 再退出。 case defNeuronMLMHSparseAttentionOCL: neuron_sparseattention = new CNeuronMLMHSparseAttention(); if(CheckPointer(neuron_sparseattention) == POINTER_INVALID) { delete temp; return false; } if(!neuron_sparseattention.Init(outputs, 0, opencl, desc.window, desc.window_out, desc.step, desc.count, desc.layers, desc.optimization, desc.batch)) { delete neuron_sparseattention; delete temp; return false; } neuron_sparseattention.SetActivationFunction(desc.activation); neuron_sparseattention.Sparse(desc.probability); 稀疏化由 Sparse(desc.probability) 控制,probability 即随机置零比例,调大它网络容量下降但过拟合概率降低。Add 进容器失败也要清理已分配的 neuron_mlattention_ocl 与 temp,避免 MT5 终端内存泄漏。 另一个同型 case 在 OpenCL 指针无效时直接返回 false,new 出 temp_mlat_ocl 后若指针无效把 result 置 false;Init 用写死的 1,1,1,1,0 和 ADAM 优化器、batch=1,适合单样本前向验证而非批量训练。

MQL5 / C++
if(!OpenCL.Execute(def_k_MHSparseAttentionOut, class="num">2, global_work_offset, global_work_size))
  {
    class="type">class="kw">string error;
    CLGetInfoString(OpenCL.GetContext(), CL_ERROR_DESCRIPTION, error);
    printf("Error of execution kernel %s: %s", __FUNCSIG__, error);
    class="kw">return false;
  }
class=class="str">"cmt">//---
  class="kw">return true;
}

case defNeuronMLMHSparseAttentionOCL:
    neuron_sparseattention = new CNeuronMLMHSparseAttention();
    if(CheckPointer(neuron_sparseattention) == POINTER_INVALID)
      {
        class="kw">delete temp;
        class="kw">return false;
      }
    if(!neuron_sparseattention.Init(outputs, class="num">0, opencl, desc.window, desc.window_out, desc.step,
                                   desc.count, desc.layers, desc.optimization, desc.batch))
      {
        class="kw">delete neuron_sparseattention;
        class="kw">delete temp;
        class="kw">return false;
      }
    neuron_sparseattention.SetActivationFunction(desc.activation);
    neuron_sparseattention.Sparse(desc.probability);
    if(!temp.Add(neuron_sparseattention))
      {
        class="kw">delete neuron_mlattention_ocl;
        class="kw">delete temp;
        class="kw">return false;
      }
    neuron_sparseattention = NULL;
    break;

case  defNeuronMLMHSparseAttentionOCL:
    if(CheckPointer(OpenCL) == POINTER_INVALID)
      class="kw">return false;
    temp_mlat_ocl = new CNeuronMLMHSparseAttention();
    if(CheckPointer(temp_mlat_ocl) == POINTER_INVALID)
      result = false;
    if(temp_mlat_ocl.Init(iOutputs, index, OpenCL, class="num">1, class="num">1, class="num">1, class="num">1, class="num">0, ADAM, class="num">1))

「前向传播里的类型分发与梯度回传」

CNeuronBaseOCL::FeedForward 先用 CheckPointer 判空,遇到 POINTER_INVALID 直接返回 false,避免后续对悬空对象调用虚函数导致 MT5 终端崩溃。 紧接着用 SourceObject.Type() 做 switch 分发,覆盖 defNeuronBaseOCL 到 defNeuronSoftMaxOCL 共 12 种神经元类型;命中后把 SourceObject 赋给 temp 指针并调用 feedForward(temp),由具体子类实现计算。 未命中任何 case 时函数落空返回 false,说明喂给该层的上游对象类型不被支持,调试时可在 return false 前打印 Type() 值快速定位。 梯度侧片段里,defNeuronMLMHAttentionOCL 与 defNeuronMLMHSparseAttentionOCL 共用 mlat 指针;当 bTrain 为 false 且 mlat.TrainMode() 也为 false 时直接返回 true,跳过梯度计算以省显存。 否则取 this 指针交给 mlat.calcInputGradients 做反向传播,这意味着多头注意力层持有上游梯度入口,训练模式下必走此分支。外汇与贵金属模型训练属高风险实验,回测结论仅具概率意义,实盘前务必在 MT5 策略测试器核验。

MQL5 / C++
          {
            m_data[index] = temp_mlat_ocl;
            class="kw">return true;
          }
          break;
class="type">bool CNeuronBaseOCL::FeedForward(CObject *SourceObject)
  {
   if(CheckPointer(SourceObject) == POINTER_INVALID)
      class="kw">return false;
class=class="str">"cmt">//---
   CNeuronBaseOCL *temp = NULL;
   class="kw">switch(SourceObject.Type())
     {
      case defNeuronBaseOCL:
      case defNeuronProofOCL:
      case defNeuronConvOCL:
      case defNeuronAttentionOCL:
      case defNeuronMHAttentionOCL:
      case defNeuronMLMHAttentionOCL:
      case defNeuronMLMHSparseAttentionOCL:
      case defNeuronDropoutOCL:
      case defNeuronBatchNormOCL:
      case defNeuronVAEOCL:
      case defNeuronLSTMOCL:
      case defNeuronSoftMaxOCL:
          temp = SourceObject;
          class="kw">return feedForward(temp);
          break;
     }
class=class="str">"cmt">//---
   class="kw">return false;
  }
        case defNeuronMLMHAttentionOCL:
        case defNeuronMLMHSparseAttentionOCL:
            mlat = TargetObject;
            if(!bTrain && !mlat.TrainMode())
              class="kw">return true;
            temp = GetPointer(this);
            class="kw">return mlat.calcInputGradients(temp);

常见问题

在OpenCL核里按反向拓扑依次调度梯度核,先算输出梯度再逐层乘稀疏掩码回传,注意掩码要在显存中以紧凑格式存储以减少带宽。
重点设头数、每头维度、稀疏度阈值与序列长度上限,按实际显卡显存留20%余量,避免实例化时直接分配最大缓冲。
可以,小布能读取你的执行日志并标出超显存层级与核调度耗时,直接给压缩建议,你照着调参数即可。
必须存掩码索引、头配置、缩放因子与随机种子,否则重载后前向数值会漂移,无法接续梯度回传。
核对各核函数的入参类型声明是否一致,尤其是半精度与单精度分支,建议先跑单头单序列验证数值再扩规模。