神经网络变得简单（第 78 部分）：带有变换器的无解码对象检测器（DFFT）·进阶篇
（2/3）· 用计算机视觉的DFFT思路绕开轨迹预测误差，让智能体直接基于历史数据挑动作
「OpenCL 里算局部注意力分数」
这段内核跑在显卡上，负责给时序单元 u 算前后各 1 根的局部注意力分数。窗口由 start 和 stop 卡死：start = max(u-1,0)，stop = min(u+1,units-1)，所以每个单元最多看相邻 3 个位置，计算量被压到 O(3·dimension) 而非全序列。 局部内存 temp[LOCAL_ARRAY_SIZE][3] 存的是 q·k 的累加值，第一维按线程局部分块 ls_d = min(dimension, LOCAL_ARRAY_SIZE)。如果 dimension 比共享内存还大，就用 stride 循环 dim += ls_d 分段搬数据，避免溢出。 分数归约用了二叉树加法：count 从 ls_d 起每次 (count+1)/2 折半，d 和 d+count 两两相加，直到 count=1。这样 256 维向量在多 wave 的 GPU 上大约 8 次 barrier 就能求和，比串行快一个数量级。 最后线程 0 把 temp[0][i] 过 exp 并叠上相对位置偏置 rpb，得到未归一化的注意力权重。外汇和贵金属行情用这种局部窗口推理延迟低，但 GPU 并行结果受浮点精度影响，回测和实盘可能有细微偏差，属高风险品类需自行验证。
class="kw">return true; } __kernel class="type">void DOTFeedForward(__global class="type">float *qkv, __global class="type">float *score, __global class="type">float *rpb, __global class="type">float *out) { const class="type">size_t d = get_local_id(class="num">0); const class="type">size_t dimension = get_local_size(class="num">0); const class="type">size_t u = get_global_id(class="num">1); const class="type">size_t units = get_global_size(class="num">1); const class="type">size_t h = get_global_id(class="num">2); const class="type">size_t heads = get_global_size(class="num">2); class="type">uint step = class="num">3 * dimension * heads; class="type">uint start = max((class="type">int)u - class="num">1, class="num">0); class="type">uint stop = min((class="type">int)u + class="num">1, (class="type">int)units - class="num">1); class="type">uint shift_q = u * step + h * dimension; class="type">uint shift_k = start * step + dimension * (heads + h); class="type">uint shift_score = u * class="num">3 * heads; const class="type">uint ls_d = min((class="type">uint)dimension, (class="type">uint)LOCAL_ARRAY_SIZE); __local class="type">float temp[LOCAL_ARRAY_SIZE][class="num">3]; class=class="str">"cmt">//--- Score if(d < ls_d) { for(class="type">uint pos = start; pos <= stop; pos++) { temp[d][pos - start] = class="num">0; } for(class="type">uint dim = d; dim < dimension; dim += ls_d) { class="type">float q = qkv[shift_q + dim]; for(class="type">uint pos = start; pos <= stop; pos++) { class="type">uint i = pos - start; temp[d][i] = temp[d][i] + q * qkv[shift_k + i * step + dim]; } } barrier(CLK_LOCAL_MEM_FENCE); class="type">int count = ls_d; do { count = (count + class="num">1) / class="num">2; if(d < count && (d + count) < dimension) for(class="type">uint i = class="num">0; i <= (stop - start); i++) { temp[d][i] += temp[d + count][i]; temp[d + count][i] = class="num">0; } barrier(CLK_LOCAL_MEM_FENCE); } while(count > class="num">1); } if(d == class="num">0) { class="type">float sum = class="num">0; for(class="type">uint i = class="num">0; i <= (stop - start); i++) { temp[class="num">0][i] = exp(temp[class="num">0][i] + rpb[shift_score + i]);
◍ 多头注意力在 OpenCL 里的落地与前向链路
上面这段内核收尾做的是 softmax 归一与加权求和：先对 temp[0][i] 按 sum 做除法得到注意力权重，写回 score 缓冲区，再用 qkv 中 value 分支（shift_v 偏移）乘权重累加，输出到 out。注意 shift_v 的偏移公式是 dimension * (heads * (u * 3 + 2) + h)，说明 qkv 在内存里是按 query、key、value 三段连续排布的，value 固定在第 3 段（索引 2）。 CNeuronDOTOCL::DOT 负责把缓冲区绑给 OpenCL kernel 并发动执行。global_work_size 设为 {iDimension, iUnits, iHeads}，local_work_size 是 {iDimension, 1, 1}，意味着每个计算单元只跑一个 head、一个 unit，但横向铺满 dimension 个线程——这种映射在维度较大时可能压满 GPU 占用率，维度小时反而有调度碎片。 feedForward 揭示了完整前向顺序：窗口尺寸不一致时先过 cProjInput 投影，再产 qkv；一致则直接由上游输入产 qkv。之后 DOT 算注意力，cProj 做输出投影，SumAndNormilize 做残差加层归一，最后接 cFF1、cFF2 两层前馈。外汇与贵金属行情高频跳变，这类结构在 MT5 上跑实时推理须警惕显存与核发起伏带来的延迟尖刺，属于高风险验证项。 想确认你本地头数与维度是否匹配，直接把 iDimension、iHeads、iUnits 打印出来，对照 kernel 里 shift_out = dimension * (u * heads + h) + d 的寻址，能立刻看出有没有越界写。
sum += temp[class="num">0][i]; } for(class="type">uint i = class="num">0; i <= (stop - start); i++) { temp[class="num">0][i] = temp[class="num">0][i] / sum; score[shift_score + i] = temp[class="num">0][i]; } } barrier(CLK_LOCAL_MEM_FENCE); class="type">int shift_out = dimension * (u * heads + h) + d; class="type">int shift_v = dimension * (heads * (u * class="num">3 + class="num">2) + h); class="type">float sum = class="num">0; for(class="type">uint i = class="num">0; i <= (stop - start); i++) sum += qkv[shift_v + i] * temp[class="num">0][i]; out[shift_out] = sum; } class="type">bool CNeuronDOTOCL::DOT(class="type">void) { if(!OpenCL) class="kw">return false; class=class="str">"cmt">//--- class="type">uint global_work_offset[class="num">3] = {class="num">0, class="num">0, class="num">0}; class="type">uint global_work_size[class="num">3] = {iDimension, iUnits, iHeads}; class="type">uint local_work_size[class="num">3] = {iDimension, class="num">1, class="num">1}; if(!OpenCL.SetArgumentBuffer(def_k_DOTFeedForward, def_k_dot_qkv, cQKV.getOutputIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_DOTFeedForward, def_k_dot_score, iScoreBuffer)) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_DOTFeedForward, def_k_dot_rpb, cRelativePositionsBias.getOutputIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_DOTFeedForward, def_k_dot_out, MHAttentionOut.getOutputIndex())) class="kw">return false; ResetLastError(); if(!OpenCL.Execute(def_k_DOTFeedForward, class="num">3, global_work_offset, global_work_size, local_work_size)) { printf("Error of execution kernel %s: %d", __FUNCTION__, GetLastError()); class="kw">return false; } class=class="str">"cmt">//--- class="kw">return true; } class="type">bool CNeuronDOTOCL::feedForward(CNeuronBaseOCL *NeuronOCL) { CNeuronBaseOCL* inputs = NeuronOCL; if(iPrevWindowSize != iWindowSize) { if(!cProjInput.FeedForward(inputs) || !cQKV.FeedForward(GetPointer(cProjInput))) class="kw">return false; inputs = GetPointer(cProjInput); } else if(!cQKV.FeedForward(inputs)) class="kw">return false; if(!DOT()) class="kw">return false; if(!cProj.FeedForward(GetPointer(MHAttentionOut))) class="kw">return false; if(!SumAndNormilize(inputs.getOutput(), cProj.getOutput(), AttentionOut.getOutput(), iWindowSize, true)) class="kw">return false; if(!cFF1.FeedForward(GetPointer(AttentionOut))) class="kw">return false; if(!cFF2.FeedForward(GetPointer(cFF1))) class="kw">return false;
注意力核里的反向梯度怎么算
这段 OpenCL 内核负责把前向注意力里的梯度沿着 Q、K、V 三个分支反传回去。核函数 DOTInsideGradients 一上来先用 get_global_id 把 u、d、h 三个维度锁死，分别对应序列位置、特征维度和注意力头，global size 则给出 units、dimension、heads 的总量。 step 被定成 3 * dimension * heads，意味着每个序列位置在 qkv 缓冲里连续排了 Q、K、V 三组向量；shift_q / shift_k / shift_v 就是按这个步长偏移到当前头当前维的指针。start 和 stop 只取 u-1 到 u+1，说明这里跑的是窗口大小为 3 的局部注意力，边界用 max/min 夹住防越界。 Value 的梯度最直接：在 start..stop 内遍历 i，用 forward 的 scores 乘上游 gradient 累加进 qkv_g[shift_v]。Query 的梯度则复杂些，shift_score 基址是 u * heads * 3，再按 k-start 偏移取分系数，里面还嵌了双层循环对 v 和 dim 做 scores * V * K 的链式法则累乘。 在 MT5 里把这段塞进自定义指标算梯度时，先确认你显存里 scores 的布局是不是按 (u, 3, heads) 排的——布局错一位，回测里注意力权重就会静默失真，外汇与贵金属杠杆环境下面这种数值偏差会被放大成异常信号。
if(!SumAndNormilize(AttentionOut.getOutput(), cFF2.getOutput(), SAttenOut.getOutput(), iWindowSize, true)) class="kw">return false; if(!cCAtten.FeedForward(GetPointer(SAttenOut))) class="kw">return false; class=class="str">"cmt">//--- class="kw">return true; } __kernel class="type">void DOTInsideGradients(__global class="type">float *qkv, __global class="type">float *qkv_g, __global class="type">float *scores, __global class="type">float *rpb, __global class="type">float *rpb_g, __global class="type">float *gradient) { class=class="str">"cmt">//--- init const class="type">uint u = get_global_id(class="num">0); const class="type">uint d = get_global_id(class="num">1); const class="type">uint h = get_global_id(class="num">2); const class="type">uint units = get_global_size(class="num">0); const class="type">uint dimension = get_global_size(class="num">1); const class="type">uint heads = get_global_size(class="num">2); class="type">uint step = class="num">3 * dimension * heads; class="type">uint start = max((class="type">int)u - class="num">1, class="num">0); class="type">uint stop = min((class="type">int)u + class="num">1, (class="type">int)units - class="num">1); const class="type">uint shift_q = u * step + dimension * h + d; const class="type">uint shift_k = u * step + dimension * (heads + h) + d; const class="type">uint shift_v = u * step + dimension * (class="num">2 * heads + h) + d; class=class="str">"cmt">//--- Calculating Value&class="macro">#x27;s gradients class="type">float sum = class="num">0; for(class="type">uint i = start; i <= stop; i ++) { class="type">int shift_score = i * class="num">3 * heads; if(u == i) { shift_score += (class="type">uint)(u > class="num">0); } else { if(u > i) shift_score += (class="type">uint)(start > class="num">0) + class="num">1; } class="type">uint shift_g = dimension * (i * heads + h) + d; sum += gradient[shift_g] * scores[shift_score]; } qkv_g[shift_v] = sum; class=class="str">"cmt">//--- Calculating Query&class="macro">#x27;s gradients class="type">float grad = class="num">0; class="type">uint shift_score = u * heads * class="num">3; for(class="type">int k = start; k <= stop; k++) { class="type">float sc_g = class="num">0; class="type">float sc = scores[shift_score + k - start]; for(class="type">int v = start; v <= stop; v++) for(class="type">int dim=class="num">0;dim<dimension;dim++) sc_g += scores[shift_score + v - start] * qkv[v * step + dimension * (class="num">2 * heads + h) + dim] *
「Key 梯度与 OpenCL 内核参数绑定」
在多头注意力反向传播里，Key 的梯度计算沿着 query 与 value 两个轴展开。外层循环遍历 q，内层对 v 做偏移索引 shift_score，当 u 与 q 或 v 相等时加 (uint)(u>0)，否则若 u 大于对方则补 (uint)(start>0)+1，这套偏移规则直接决定了 scores 缓冲区的寻址正确性。 sc_g 的累加式为 scores[shift_score] * qkv[shift_v-d+dim] * gradient[dimension*(v*heads+h)+d] * ((float)(d==v)-sc)，其中 (d==v) 的浮点强转把对角线项单独剥离，非对角则乘负 sc。循环结束后 grad += sc_g * qkv[q*step + dimension*h + d]，最终写回 qkv_g[shift_k]。 梯度算完并不进 CPU，而是交给 OpenCL 内核 DOTInsideGradients 继续跑。global_work_size 设为 {iUnits, iDimension, iHeads} 三维，把 qkv、qkv_g、scores、rpb、rpb_g、gradient 六个缓冲逐一用 SetArgumentBuffer 绑到 def_k_DOTInsideGradients 内核——任何一次绑定返回 false 就直接退出，说明显存指针或索引号对不上。 在 MT5 里把 iUnits/iDimension/iHeads 打印出来，若三维乘积超过几万，内核启动耗时会明显抬头；外汇与贵金属行情序列短，把 iUnits 压到 128 以内可能更稳，但高频重训仍有 GPU 显存溢出风险。
gradient[dimension * (u * heads + h) + dim] * ((class="type">float)(k == v) - sc); grad += sc_g * qkv[k * step + dimension * (heads + h) + d]; if(d == class="num">0) rpb_g[shift_score + k - start] = sc_g; } qkv_g[shift_q] = grad; class=class="str">"cmt">//--- Calculating Key&class="macro">#x27;s gradients grad = class="num">0; for(class="type">int q = start; q <= stop; q++) { class="type">float sc_g = class="num">0; shift_score = q * heads * class="num">3; if(u == q) { shift_score += (class="type">uint)(u > class="num">0); } else { if(u > q) shift_score += (class="type">uint)(start > class="num">0) + class="num">1; } class="type">float sc = scores[shift_score]; for(class="type">int v = start; v <= stop; v++) { shift_score = v * heads * class="num">3; if(u == v) { shift_score += (class="type">uint)(u > class="num">0); } else { if(u > v) shift_score += (class="type">uint)(start > class="num">0) + class="num">1; } for(class="type">int dim=class="num">0;dim<dimension;dim++) sc_g += scores[shift_score] * qkv[shift_v-d+dim] * gradient[dimension * (v * heads + h) + d] * ((class="type">float)(d == v) - sc); } grad += sc_g * qkv[q * step + dimension * h + d]; } qkv_g[shift_k] = grad; } class="type">bool CNeuronDOTOCL::DOTInsideGradients(class="type">void) { if(!OpenCL) class="kw">return false; class=class="str">"cmt">//--- class="type">uint global_work_offset[class="num">3] = {class="num">0, class="num">0, class="num">0}; class="type">uint global_work_size[class="num">3] = {iUnits, iDimension, iHeads}; if(!OpenCL.SetArgumentBuffer(def_k_DOTInsideGradients, def_k_dotg_qkv, cQKV.getOutputIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_DOTInsideGradients, def_k_dotg_qkv_g, cQKV.getGradientIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_DOTInsideGradients, def_k_dotg_scores, iScoreBuffer)) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_DOTInsideGradients, def_k_dotg_rpb, cRelativePositionsBias.getOutputIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_DOTInsideGradients, def_k_dotg_rpb_g, cRelativePositionsBias.getGradientIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_DOTInsideGradients, def_k_dotg_gradient, MHAttentionOut.getGradientIndex()))
◍ 反向传播里的梯度链路与 Adam 更新核
这段 CNeuronDOTOCL::calcInputGradients 把各子模块的梯度串成一条回传链：注意力输出、两层前馈、投影矩阵依次调 calcInputGradients，任何一步返回 false 就直接中断。注意当 iPrevWindowSize 不等于 iWindowSize 时，会多走 cProjInput 这条分支，否则直接把梯度累加到 prevLayer，这是窗口尺寸变化时才有的额外归一化开销。 DOTInsideGradients 内部先 ResetLastError，再调 OpenCL.Execute 跑 def_k_DOTInsideGradients 核，全局工作项维度写死为 3。若执行失败，printf 会打出 __FUNCTION__ 和 GetLastError 码，方便在 MT5 Experts 日志里定位是哪一层 OpenCL 调用崩了。 末尾的 RPBUpdateAdam 是给相对位置偏置做参数更新的 OpenCL 核，入参带 matrix_m 与 matrix_v 分别存一阶、二阶动量，b1、b2 为对应衰减系数。在 MT5 里把这类核接进自己的优化器，能把外汇序列的相对位置项用 Adam 规则迭代；不过 GPU 浮点误差可能让回测与实盘权重漂移，贵金属与外汇杠杆品种属高风险，参数更新效果需用历史 Tick 复跑验证。
class="kw">return false; ResetLastError(); if(!OpenCL.Execute(def_k_DOTInsideGradients, class="num">3, global_work_offset, global_work_size)) { printf("Error of execution kernel %s: %d", __FUNCTION__, GetLastError()); class="kw">return false; } class=class="str">"cmt">//--- class="kw">return true; } class="type">bool CNeuronDOTOCL::calcInputGradients(CNeuronBaseOCL *prevLayer) { if(!prevLayer) class="kw">return false; if(!cCAtten.calcInputGradients(GetPointer(SAttenOut))) class="kw">return false; if(!cFF2.calcInputGradients(GetPointer(cFF1))) class="kw">return false; if(!cFF1.calcInputGradients(GetPointer(AttentionOut))) class="kw">return false; if(!SumAndNormilize(AttentionOut.getGradient(), SAttenOut.getGradient(), cProj.getGradient(), iWindowSize, false)) class="kw">return false; if(!cProj.calcInputGradients(GetPointer(MHAttentionOut))) class="kw">return false; if(!DOTInsideGradients()) class="kw">return false; if(iPrevWindowSize != iWindowSize) { if(!cQKV.calcInputGradients(GetPointer(cProjInput))) class="kw">return false; if(!SumAndNormilize(cProjInput.getGradient(), cProj.getGradient(), cProjInput.getGradient(), iWindowSize, false)) class="kw">return false; if(!cProjInput.calcInputGradients(prevLayer)) class="kw">return false; } else { if(!cQKV.calcInputGradients(prevLayer)) class="kw">return false; if(!SumAndNormilize(prevLayer.getGradient(), cProj.getGradient(), prevLayer.getGradient(), iWindowSize, false)) class="kw">return false; } class=class="str">"cmt">//--- class="kw">return true; } __kernel class="type">void RPBUpdateAdam(__global class="type">float *target, __global const class="type">float *gradient, __global class="type">float *matrix_m, class=class="str">"cmt">///<[in,out] Matrix of first momentum __global class="type">float *matrix_v, class=class="str">"cmt">///<[in,out] Matrix of seconfd momentum const class="type">float b1, class=class="str">"cmt">///< First momentum multiplier const class="type">float b2 class=class="str">"cmt">///< Second momentum multiplier
权重更新与窗口切换的衔接细节
这段 OpenCL 内核把 Adam 的一阶、二阶动量直接搬进 GPU 算：每个神经元先取历史 m、v 与梯度 g，按 b1、b2 做指数滑动平均，再除以 sqrt(v) 得到更新步长 delta，最后用 clamp 把权重锁在 [-MAX_WEIGHT, MAX_WEIGHT] 区间内，避免相对位置偏置在训练里发散。 宿主端 updateRelativePositionsBias 负责把输出缓冲、梯度、两个动量矩阵以及 b1/b2 常量绑给内核，再按神经元总数发起一次全局调度；任一 SetArgument 或 Execute 失败就直接返回 false，并在控制台打出内核名与 GetLastError 代码，方便在 MT5 专家日志里定位是哪一步脱钩。 网络前馈权重的刷新逻辑跟窗口尺寸挂钩：当 iWindowSize 不等于 iPrevWindowSize 时，投影输入层改从 NeuronOCL 取权，QKV 则接投影输入层；尺寸不变则 QKV 直接连 NeuronOCL。后面 cProj、cFF1、cFF2、cCAtten 依次链式 UpdateInputWeights，最后才调 updateRelativePositionsBias，顺序反了会让注意力层读到上一轮的偏置。 在 MT5 里跑这套时，若改了注意力窗口长度，务必确认 iPrevWindowSize 已被正确赋值，否则权重重连分支不触发，模型可能倾向用错位的历史表征做推断，外汇与贵金属行情下这属于高风险的静默偏差。
const class="type">int i = get_global_id(class="num">0); class="type">float m, v, weight; m = matrix_m[i]; v = matrix_v[i]; weight = target[i]; class="type">float g = gradient[i]; m = b1 * m + (class="num">1 - b1) * g; v = b2 * v + (class="num">1 - b2) * pow(g, class="num">2); class="type">float delta = m / (v != class="num">0.0f ? sqrt(v) : class="num">1.0f); target[i] = clamp(weight + delta, -MAX_WEIGHT, MAX_WEIGHT); matrix_m[i] = m; matrix_v[i] = v; } class="type">bool CNeuronDOTOCL::updateRelativePositionsBias(class="type">void) { if(!OpenCL) class="kw">return false; class=class="str">"cmt">//--- class="type">uint global_work_offset[class="num">1] = {class="num">0}; class="type">uint global_work_size[class="num">1] = {cRelativePositionsBias.Neurons()}; if(!OpenCL.SetArgumentBuffer(def_k_RPBUpdateAdam, def_k_rpbw_rpb, cRelativePositionsBias.getOutputIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_RPBUpdateAdam, def_k_rpbw_gradient, cRelativePositionsBias.getGradientIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_RPBUpdateAdam, def_k_rpbw_matrix_m, cRelativePositionsBias.getFirstMomentumIndex())) class="kw">return false; if(!OpenCL.SetArgumentBuffer(def_k_RPBUpdateAdam, def_k_rpbw_matrix_v, cRelativePositionsBias.getSecondMomentumIndex())) class="kw">return false; if(!OpenCL.SetArgument(def_k_RPBUpdateAdam, def_k_rpbw_b1, b1)) class="kw">return false; if(!OpenCL.SetArgument(def_k_RPBUpdateAdam, def_k_rpbw_b2, b2)) class="kw">return false; ResetLastError(); if(!OpenCL.Execute(def_k_RPBUpdateAdam, class="num">1, global_work_offset, global_work_size)) { printf("Error of execution kernel %s: %d", __FUNCTION__, GetLastError()); class="kw">return false; } class=class="str">"cmt">//--- class="kw">return true; } if(iWindowSize != iPrevWindowSize) { if(!cProjInput.UpdateInputWeights(NeuronOCL)) class="kw">return false; if(!cQKV.UpdateInputWeights(GetPointer(cProjInput))) class="kw">return false; } else { if(!cQKV.UpdateInputWeights(NeuronOCL)) class="kw">return false; } if(!cProj.UpdateInputWeights(GetPointer(MHAttentionOut))) class="kw">return false; if(!cFF1.UpdateInputWeights(GetPointer(AttentionOut))) class="kw">return false; if(!cFF2.UpdateInputWeights(GetPointer(cFF1))) class="kw">return false; if(!cCAtten.UpdateInputWeights(GetPointer(SAttenOut))) class="kw">return false; if(!updateRelativePositionsBias()) class="kw">return false; class=class="str">"cmt">//--- class="kw">return true; }