神经网络变得简单（第 80 部分）：图形变换器生成式对抗模型（GTGAN）·进阶篇
GPU 内核里的注意力分数怎么算
在 MT5 用 OpenCL 跑 Transformer 类模型时，注意力分数的并行计算被拆进了 GTEFeedForward 这个内核。每个全局线程拿到的 cur_q、cur_k、h 分别对应 query 序号、key 序号和注意力头，通过 get_global_id / get_local_id 直接映射到显存维度。 分数本身是一条点积后 softmax 的近似：scr 先累加 qkv 里 shift_q 与 shift_k 对应维度的乘积，再除以 dimension 的平方根做缩放，最后用 exp(min(..., 30.0f)) 截断防溢出。这里 30.0f 的上限很关键，AUC 类回测里不截断可能在维度 >64 时让单头数值爆掉。 连接关系（connected）和非连接（not connected）两套 score 缓冲分开写：当 cur_q == cur_k 时两者都存 scr；当两者相邻（abs 差为 1）时只给 con 路径存 scr，notcon 置 0。局部内存 local_score 仅缓存前 ls_score 个 key，ls_score = min(units_k, LOCAL_ARRAY_SIZE)，这是为了避免小核显存撑爆。 外汇与贵金属行情用这类 GPU 推理做特征提取时，仍需警惕过拟合与滑点风险，显存并行结果只是概率层面的信号，不是方向保证。
__kernel class="type">void GTEFeedForward(__global class="type">float *qkv, __global class="type">float *score, __global class="type">float *out, class="type">int dimension) { const class="type">size_t cur_q = get_global_id(class="num">0); const class="type">size_t units_q = get_global_size(class="num">0); const class="type">size_t cur_k = get_local_id(class="num">1); const class="type">size_t units_k = get_local_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">int shift_q = dimension * (cur_q + h * units_q); class="type">int shift_k = (cur_k + h * units_k + heads * units_q); class="type">int shift_v = dimension * (h * units_k + heads * (units_q + units_k)); class="type">int shift_score_con = units_k * (cur_q * class="num">2 * heads + h) + cur_k; class="type">int shift_score_notcon = units_k * (cur_q * class="num">2 * heads + heads + h) + cur_k; class="type">int shift_out_con = dimension * (cur_q + h * units_q); class="type">int shift_out_notcon = dimension * (cur_q + units_q * (h + heads)); const class="type">uint ls_score = min((class="type">uint)units_k, (class="type">uint)LOCAL_ARRAY_SIZE); __local class="type">float local_score[LOCAL_ARRAY_SIZE][class="num">2]; class=class="str">"cmt">//--- Score class="type">float scr = class="num">0; for(class="type">int d = class="num">0; d < dimension; d ++) scr += qkv[shift_q + d] * qkv[shift_k + d]; scr = exp(min(scr / sqrt((class="type">float)dimension), class="num">30.0f)); if(cur_q == cur_k) { score[shift_score_con] = scr; score[shift_score_notcon] = scr; if(cur_k < ls_score) { local_score[cur_k][class="num">0] = scr; local_score[cur_k][class="num">1] = scr; } } else { if(abs(cur_q - cur_k) == class="num">1) { score[shift_score_con] = scr; score[shift_score_notcon] = class="num">0; if(cur_k < ls_score) { local_score[cur_k][class="num">0] = scr; local_score[cur_k][class="num">1] = class="num">0; } }
「GPU 上的注意力归约与输出拼装」
这段 OpenCL 内核收尾做的是局部内存里的并行归约：先把每个线程负责的 score 段写进 local_score 的两列（连续/非连续），再用 barrier(CLK_LOCAL_MEM_FENCE) 卡住线程同步，避免读写竞争。 归约用 do-while 折半累加，count 从 ls_score 起每次 (count+1)/2 收敛，直到 count<=1；此时 local_score[0][0] 和 [0][1] 就是两类分数的全域和，随后 score 两基位各自除以该和，得到归一化权重。 最后的输出循环按 dimension 切分，每个线程把 qkv 对应切片用归一权重做加权求和，分别写进 out 的 con / notcon 偏移区；这里 cur_k+d 越界会被 if 拦掉，不会写爆缓冲区。 CNeuronGTE::AttentionOut 只是调度层：检测 OpenCL 上下文存在性，设 global_work_size 为 {iUnits, iUnits, iHeads}、local 为 {1, iUnits, 1}，再 ResetLastError 清状态。外汇/贵金属行情下跑这套核，显存带宽瓶颈可能让 iUnits 超过 256 时延迟非线性抬升，建议用 MT5 的 OpenCL 探针实测吞吐再调 local_work_size。
else { score[shift_score_con] = class="num">0; score[shift_score_notcon] = scr; if(cur_k < ls_score) { local_score[cur_k][class="num">0] = class="num">0; local_score[cur_k][class="num">1] = scr; } } } barrier(CLK_LOCAL_MEM_FENCE); for(class="type">int k = ls_score; k < units_k; k += ls_score) { if((cur_k + k) < units_k) { local_score[cur_k][class="num">0] += score[shift_score_con + k]; local_score[cur_k][class="num">1] += score[shift_score_notcon + k]; } } barrier(CLK_LOCAL_MEM_FENCE); class=class="str">"cmt">//--- class="type">int count = ls_score; do { count = (count + class="num">1) / class="num">2; if(cur_k < count) { if((cur_k + count) < units_k) { local_score[cur_k][class="num">0] += local_score[cur_k + count][class="num">0]; local_score[cur_k][class="num">1] += local_score[cur_k + count][class="num">1]; local_score[cur_k + count][class="num">0] = class="num">0; local_score[cur_k + count][class="num">1] = class="num">0; } } barrier(CLK_LOCAL_MEM_FENCE); } while(count > class="num">1); barrier(CLK_LOCAL_MEM_FENCE); score[shift_score_con] /= local_score[class="num">0][class="num">0]; score[shift_score_notcon] /= local_score[class="num">0][class="num">1]; barrier(CLK_LOCAL_MEM_FENCE); shift_score_con -= cur_k; shift_score_notcon -= cur_k; for(class="type">int d = class="num">0; d < dimension; d += ls_score) { if((cur_k + d) < dimension) { class="type">float sum_con = class="num">0; class="type">float sum_notcon = class="num">0; for(class="type">int v = class="num">0; v < units_k; v++) { sum_con += qkv[shift_v + v * dimension + cur_k + d] * score[shift_score_con + v]; sum_notcon += qkv[shift_v + v * dimension + cur_k + d] * score[shift_score_notcon + v]; } out[shift_out_con + cur_k + d] = sum_con; out[shift_out_notcon + cur_k + d] = sum_notcon; } } } class="type">bool CNeuronGTE::AttentionOut(class="type">void) { if(!OpenCL) class="kw">return false; class="type">uint global_work_offset[class="num">3] = {class="num">0}; class="type">uint global_work_size[class="num">3] = {iUnitsclass=class="str">"cmt">/*Q units*/, iUnitsclass=class="str">"cmt">/*K units*/, iHeads}; class="type">uint local_work_size[class="num">3] = {class="num">1, iUnits, class="num">1}; ResetLastError();
◍ GTE 神经元的 OpenCL 内核参数绑定与前向链路
在 MT5 的 OpenCL 环境里跑 GTE 类 Transformer 编码器，第一步是把各个缓冲区句柄塞进计算内核。下面这段把 QKV 输出、注意力分数、注意力输出分别用 SetArgumentBuffer 绑到 def_k_GTEFeedForward 内核，再把窗口维度 iWindowKey 以 int 形式用 SetArgument 传入；任何一步返回 false 就打印错误码和行号并退出，避免半初始化状态污染后续张量。 内核真正触发靠 OpenCL.Execute，这里传了 3 维的 global_work_offset / global_work_size / local_work_size。若 Execute 失败只报内核名与 GetLastError，不带回退逻辑——实盘加载自定义 GPU 模型时，这一行报错往往意味着显存碎片或 work_size 算错，需要你手动对齐 buffer 长度。 feedForward 成员函数串起了完整前向：cQKV 先算查询键值，cSoftMax 对其做归一，AttentionOut 出注意力结果，cW0 做输出投影，再用 SumAndNormilize 做残差加层归一。之后接两级 cGraphConv 图卷积与两级 cFF 前馈，最后再残差归一写进 Output。整条链任一层 FeedForward 返回 false 就直接中断，没有部分前向结果可用。 外层调用者若拿到 false，应当假定该帧特征提取完全无效；在外汇或贵金属这样的高波动高风险市场，用损坏的 GPU 特征去驱动信号可能放大滑点亏损，建议加一层返回检查再进策略判断。
if(!OpenCL.SetArgumentBuffer(def_k_GTEFeedForward, def_k_gteff_qkv, cQKV.getOutputIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgumentBuffer(def_k_GTEFeedForward, def_k_gteff_score, ScoreIndex)) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgumentBuffer(def_k_GTEFeedForward, def_k_gteff_out, cAttentionOut.getOutputIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgument(def_k_GTEFeedForward, def_k_gteff_dimension, (class="type">int)iWindowKey)) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.Execute(def_k_GTEFeedForward, 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 CNeuronGTE::feedForward(CNeuronBaseOCL *NeuronOCL) { if(!cQKV.FeedForward(NeuronOCL)) class="kw">return false; if(!cSoftMax.FeedForward(GetPointer(cQKV))) class="kw">return false; if(!AttentionOut()) class="kw">return false; if(!cW0.FeedForward(GetPointer(cMHAttentionOut))) class="kw">return false; if(!SumAndNormilize(NeuronOCL.getOutput(), cW0.getOutput(), cAttentionOut.getOutput(), iWindow, true)) class="kw">return false; if(!cGraphConv[class="num">0].FeedForward(GetPointer(cAttentionOut))) class="kw">return false; if(!cGraphConv[class="num">1].FeedForward(GetPointer(cGraphConv[class="num">0]))) class="kw">return false; if(!cFF[class="num">0].FeedForward(GetPointer(cGraphConv[class="num">1]))) class="kw">return false; if(!cFF[class="num">1].FeedForward(GetPointer(cFF[class="num">0]))) class="kw">return false; if(!SumAndNormilize(cAttentionOut.getOutput(), cFF[class="num">1].getOutput(), Output, iWindow, true)) class="kw">return false; class=class="str">"cmt">//--- class="kw">return true; } __kernel class="type">void GTEInsideGradients(__global class="type">float *qkv, __global class="type">float *qkv_g,
反传里 Value 与 Query 梯度的并行拆解
在 MT5 的 OpenCL 内核里做多头注意力反传，最容易被忽略的是连通（con）与非连通（notcon）两套 score 要分开累加。上面这段内核用 get_global_id 把 u、d、h 三维铺到 GPU 线程，units、dimension、heads 由 get_global_size 取回，相当于把样本数、特征维、头数直接映射成并行粒度。
Value 梯度那一段，shift_v 把写回位置算到 dimension*(h*units + 2*heads*units + u) + d，循环上界写成 i <= units 而非 < units，实际会比单元数多跑一次，复制代码到本地编译时建议先改成严格小于，否则越界读 scores 的概率偏高。外汇与贵金属行情序列短、units 常设在 128 以内，这种 off-by-one 在显存窄带宽下更容易触发静默 NaN。
Query 梯度嵌套了 k、v、dim 三层循环，sc_g 的更新把 (k==v) 强转成 float 再减 sc_con，就是 softmax Jacobian 的标准形式。想验证梯度对不对，可以把 heads 设为 1、units 设为 8 跑小规模内核，用 CPU 端双重循环对拍，误差大于 1e-4 就说明 shift 算错了。
__global class="type">float *scores, __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=class="str">"cmt">//--- Calculating Value&class="macro">#x27;s gradients { class="type">int shift_out_con = dimension * h * units + d; class="type">int shift_out_notcon = dimension * units * (h + heads) + d; class="type">int shift_score_con = units * h + u; class="type">int shift_score_notcon = units * (heads + h) + u; class="type">int step_score = units * class="num">2 * heads; class="type">int shift_v = dimension * (h * units + class="num">2 * heads * units + u) + d; class="type">float sum = class="num">0; for(class="type">uint i = class="num">0; i <= units; i ++) { sum += gradient[shift_out_con + i * dimension] * scores[shift_score_con + i * step_score]; sum += gradient[shift_out_notcon + i * dimension] * scores[shift_score_notcon + i * step_score]; } qkv_g[shift_v] = sum; } class=class="str">"cmt">//--- Calculating Query&class="macro">#x27;s gradients { class="type">int shift_q = dimension * (u + h * units) + d; class="type">int shift_out_con = dimension * (h * units + u) + d; class="type">int shift_out_notcon = dimension * (u + units * (h + heads)) + d; class="type">int shift_score_con = units * h; class="type">int shift_score_notcon = units * (heads + h); class="type">int shift_v = dimension * (h * units + class="num">2 * heads * units); class="type">float grad = class="num">0; for(class="type">int k = class="num">0; k < units; k++) { class="type">int shift_k = (k + h * units + heads * units) + d; class="type">float sc_g = class="num">0; class="type">float sc_con = scores[shift_score_con + k]; class="type">float sc_notcon = scores[shift_score_notcon + k]; for(class="type">int v = class="num">0; v < units; v++) for(class="type">int dim = class="num">0; dim < dimension; dim++) { sc_g += scores[shift_score_con + v] * qkv[shift_v + v * dimension + dim] * gradient[shift_out_con + dim] * ((class="type">float)(k == v) - sc_con); sc_g += scores[shift_score_notcon + v] * qkv[shift_v + v * dimension + dim] *
「Key 梯度与层间回传的索引套路」
上面那段是 GTE 结构里 Key 侧梯度的核心循环。先按 (u + (h + heads) * units) + d 算出 shift_k，再把 connected / non-connected 两套输出偏移、score 偏移分别用 dimension * h * units、units * h 等乘积定位，step_score 固定为 units * 2 * heads，说明正负两套注意力分数被连续排布在内存里。 内层三重循环先对 q 遍历、再对 g 和 dim 展开：sc_g 累加时同时吃了 scores[shift_score_con + g] 与 scores[shift_score_notcon + g]，并各自乘上 gradient 里对应 g*dim 位置的误差，权重修正项用 ((float)(u == g) - sc_con) 这种 one-hot 减 softmax 概率的形式，和 Query 侧完全对称。 算完 grad 直接写回 qkv_g[shift_k]，一个 head 的一个维度就处理完了。把这套偏移公式原样抄进 MT5 的 OpenCL 核里，跑小批量序列能直观看到梯度量级随 heads 增大而摊薄——外汇或贵金属特征上训这类结构仍属高风险实验。 calcInputGradients 负责把误差往前一层吐：cFF[1]→cFF[0]→cGraphConv[1]→cGraphConv[0] 和 cAttentionOut，任意一环返回 false 就整体中止。改网络拓扑时若插了新层，这里漏掉一行 GetPointer 调用，反向传播会静默断链。
gradient[shift_out_notcon + dim] *
((class="type">float)(k == v) - sc_notcon);
}
grad += sc_g * qkv[shift_k];
}
qkv_g[shift_q] = grad;
}
class=class="str">"cmt">//--- Calculating Key&class="macro">#x27;s gradients
{
class="type">int shift_k = (u + (h + heads) * units) + d;
class="type">int shift_out_con = dimension * h * units + d;
class="type">int shift_out_notcon = dimension * units * (h + heads) + d;
class="type">int shift_score_con = units * h + u;
class="type">int shift_score_notcon = units * (heads + h) + u;
class="type">int step_score = units * class="num">2 * heads;
class="type">int shift_v = dimension * (h * units + class="num">2 * heads * units);
class="type">float grad = class="num">0;
for(class="type">int q = class="num">0; q < units; q++)
{
class="type">int shift_q = dimension * (q + h * units) + d;
class="type">float sc_g = class="num">0;
class="type">float sc_con = scores[shift_score_con + u + q * step_score];
class="type">float sc_notcon = scores[shift_score_notcon + u + q * step_score];
for(class="type">int g = class="num">0; g < units; g++)
{
for(class="type">int dim = class="num">0; dim < dimension; dim++)
{
sc_g += scores[shift_score_con + g] *
qkv[shift_v + u * dimension + dim] *
gradient[shift_out_con + g * dimension + dim] *
((class="type">float)(u == g) - sc_con);
sc_g += scores[shift_score_notcon + g] *
qkv[shift_v + u * dimension + dim] *
gradient[shift_out_notcon + g * dimension+ dim] *
((class="type">float)(u == g) - sc_notcon);
}
}
grad += sc_g * qkv[shift_q];
}
qkv_g[shift_k] = grad;
}
}
class="type">bool CNeuronGTE::calcInputGradients(CNeuronBaseOCL *prevLayer)
{
if(!cFF[class="num">1].calcInputGradients(GetPointer(cFF[class="num">0])))
class="kw">return false;
if(!cFF[class="num">0].calcInputGradients(GetPointer(cGraphConv[class="num">1])))
class="kw">return false;
if(!cGraphConv[class="num">1].calcInputGradients(GetPointer(cGraphConv[class="num">0])))
class="kw">return false;
if(!cGraphConv[class="num">1].calcInputGradients(GetPointer(cAttentionOut)))
class="kw">return false;◍ 反向传播与编码器骨架的落地细节
GTE 神经元类的梯度回传顺序很固定：先对注意力输出做求和归一，再沿 W0、内部注意力、SoftMax、QKV 逐层反推输入梯度，任一步返回 false 就整体中断。这种链式判断在外汇样本上若某层梯度爆炸，会直接让整轮训练作废，概率随窗口长度 iWindow 上升而加大。 权重更新函数 updateInputWeights 则按 QKV → W0 → 两级图卷积 → 两级前馈的顺序刷参，cGraphConv[0] 吃注意力输出，cGraphConv[1] 吃上一级卷积，cFF 同理。想验证结构，可在 MT5 策略测试器里把这两段塞进自定义神经网络 EA，单步跟一遍返回值。 CreateEncoderDescriptions 负责搭编码器描述数组：输入层节点数 = HistoryBars * BarDescr，激活函数为 None，优化器用 ADAM；第一隐藏层是 BatchNorm，batch 尺寸取 MathMax(1000, GPTBars)。贵金属与外汇行情高频非平稳，用 ADAM 虽收敛快，但过拟合风险高，实盘前务必用历史分段回测。
if(!SumAndNormilize(cAttentionOut.getGradient(), Gradient, cW0.getGradient(), iWindow, false)) class="kw">return false; if(!cW0.calcInputGradients(GetPointer(cMHAttentionOut))) class="kw">return false; if(!AttentionInsideGradients()) class="kw">return false; if(!cSoftMax.calcInputGradients(GetPointer(cQKV))) class="kw">return false; if(!cQKV.calcInputGradients(prevLayer)) class="kw">return false; if(!SumAndNormilize(cW0.getGradient(), prevLayer.getGradient(), prevLayer.getGradient(), iWindow, false)) class="kw">return false; class=class="str">"cmt">//--- class="kw">return true; } class="type">bool CNeuronGTE::updateInputWeights(CNeuronBaseOCL *NeuronOCL) { if(!cQKV.UpdateInputWeights(NeuronOCL)) class="kw">return false; if(!cW0.UpdateInputWeights(GetPointer(cMHAttentionOut))) class="kw">return false; if(!cGraphConv[class="num">0].UpdateInputWeights(GetPointer(cAttentionOut))) class="kw">return false; if(!cGraphConv[class="num">1].UpdateInputWeights(GetPointer(cGraphConv[class="num">0]))) class="kw">return false; if(!cFF[class="num">0].UpdateInputWeights(GetPointer(cGraphConv[class="num">1]))) class="kw">return false; if(!cFF[class="num">1].UpdateInputWeights(GetPointer(cFF[class="num">0]))) class="kw">return false; class=class="str">"cmt">//--- class="kw">return true; } class="type">bool CreateEncoderDescriptions(CArrayObj *encoder, CArrayObj *decoder) { class=class="str">"cmt">//--- CLayerDescription *descr; class=class="str">"cmt">//--- if(!encoder) { encoder = new CArrayObj(); if(!encoder) class="kw">return false; } if(!decoder) { decoder = new CArrayObj(); if(!decoder) class="kw">return false; } class=class="str">"cmt">//--- Encoder encoder.Clear(); class=class="str">"cmt">//--- Input layer if(!(descr = new CLayerDescription())) class="kw">return false; descr.type = defNeuronBaseOCL; class="type">int prev_count = descr.count = (HistoryBars * BarDescr); descr.activation = None; descr.optimization = ADAM; if(!encoder.Add(descr)) { class="kw">delete descr; class="kw">return false; } class=class="str">"cmt">//--- layer class="num">1 if(!(descr = new CLayerDescription())) class="kw">return false; descr.type = defNeuronBatchNormOCL; descr.count = prev_count; descr.batch = MathMax(class="num">1000, GPTBars); descr.activation = None; descr.optimization = ADAM; if(!encoder.Add(descr)) { class="kw">delete descr; class="kw">return false; }