神经网络变得简单（第 74 部分）：自适应轨迹预测·进阶篇
（2/3）· 只盯单一品种价格的预测，往往漏掉跨资产道路上的隐性交通规则
◍ 多头注意力层的初始化与核函数拆解
在 MT5 用 OpenCL 跑 Transformer 类模型时，多头注意力的权重矩阵必须先关掉激活函数再 Init，否则梯度回传会串味。下面这段初始化把 W0、AttentionOut 和两层前馈 FF[0]、FF[1] 全部 SetActivationFunction(None)，其中 FF 的两层维度分别是 iWindow→4*iWindow 和 4*iWindow→iWindow，硬编码的 4 倍扩张是 Transformer 前馈块的常见比例。 MHAttentionOut.SetActivationFunction(None); if(!W0.Init(0, 0, OpenCL, iWindowKey * iHeads, iWindowKey * iHeads, iWindow, iUnits, optimization_type, batch)) return false; W0.SetActivationFunction(None); if(!AttentionOut.Init(0, 0, OpenCL, iWindow * iUnits, optimization_type, batch)) return false; AttentionOut.SetActivationFunction(None); if(!FF[0].Init(0, 0, OpenCL, iWindow, iWindow, 4 * iWindow, iUnits, optimization_type, batch)) return false; if(!FF[1].Init(0, 0, OpenCL, 4 * iWindow, 4 * iWindow, iWindow, iUnits, optimization_type, batch)) return false; for(int i = 0; i < 2; i++) FF[i].SetActivationFunction(None); Gradient.BufferFree(); delete Gradient; Gradient = FF[1].getGradient(); 核函数 MH2AttentionOut 的全局 ID 切法值得盯：q_id、k、h 分别对应 query 序号、key 序号和 head 序号，三维 get_global_size 拿到 qunits、kunits、heads。shift_q 用 dimension*(q_id + qunits*h) 做偏移，说明每个 head 的 query 在内存里是连续块排布，head 之间跨 qunits*dimension 步长。 缩放系数 koef = sqrt(dimension) 且下限钳到 1，这是点积注意力的标准 1/√d 缩放，防止维度大了 softmax 梯度消失。本地数组 temp[LOCAL_ARRAY_SIZE] 配合 ls = min(get_local_size(1), LOCAL_ARRAY_SIZE) 做归约，外汇或贵金属行情序列喂进去时，dimension 若设 32，koef 就是 5.657，可在 MT5 的 OpenCL 专家里直接打印验证。 这类 GPU 核在 EURUSD 的 M15 上做注意力特征提取，显存占用和 batch 线性相关，调 batch 参数前先确认显卡本地内存上限，贵金属 XAUUSD 波动大时序列异常值可能让 exp 求和溢出，概率上倾向用小 batch 保稳。
MHAttentionOut.SetActivationFunction(None); if(!W0.Init(class="num">0, class="num">0, OpenCL, iWindowKey * iHeads, iWindowKey * iHeads, iWindow, iUnits, optimization_type, batch)) class="kw">return false; W0.SetActivationFunction(None); if(!AttentionOut.Init(class="num">0, class="num">0, OpenCL, iWindow * iUnits, optimization_type, batch)) class="kw">return false; AttentionOut.SetActivationFunction(None); if(!FF[class="num">0].Init(class="num">0, class="num">0, OpenCL, iWindow, iWindow, class="num">4 * iWindow, iUnits, optimization_type, batch)) class="kw">return false; if(!FF[class="num">1].Init(class="num">0, class="num">0, OpenCL, class="num">4 * iWindow, class="num">4 * iWindow, iWindow, iUnits, optimization_type, batch)) class="kw">return false; for(class="type">int i = class="num">0; i < class="num">2; i++) FF[i].SetActivationFunction(None); Gradient.BufferFree(); class="kw">delete Gradient; Gradient = FF[class="num">1].getGradient(); class=class="str">"cmt">//--- class="kw">return true; } __kernel class="type">void MH2AttentionOut(__global class="type">class="kw">float *q, class=class="str">"cmt">///<-[in] Matrix of Querys __global class="type">class="kw">float *kv, class=class="str">"cmt">///<-[in] Matrix of Keys __global class="type">class="kw">float *score, class=class="str">"cmt">///<-[out] Matrix of Scores __global class="type">class="kw">float *out, class=class="str">"cmt">///<-[out] Matrix of Scores class="type">int dimension class=class="str">"cmt">///<- Dimension of Key ) { class=class="str">"cmt">//--- init const class="type">int q_id = get_global_id(class="num">0); const class="type">int k = get_global_id(class="num">1); const class="type">int h = get_global_id(class="num">2); const class="type">int qunits = get_global_size(class="num">0); const class="type">int kunits = get_global_size(class="num">1); const class="type">int heads = get_global_size(class="num">2); const class="type">int shift_q = dimension * (q_id + qunits * h); const class="type">int shift_k = dimension * (k + kunits * h); const class="type">int shift_v = dimension * (k + kunits * (heads + h)); const class="type">int shift_s = q_id * kunits * heads + h * kunits + k; const class="type">uint ls = min((class="type">uint)get_local_size(class="num">1), (class="type">uint)LOCAL_ARRAY_SIZE); class="type">class="kw">float koef = sqrt((class="type">class="kw">float)dimension); if(koef < class="num">1) koef = class="num">1; __local class="type">class="kw">float temp[LOCAL_ARRAY_SIZE]; class=class="str">"cmt">//--- sum of exp class="type">uint count = class="num">0; if(k < ls) do {
「OpenCL 核函数里的注意力归约与输出拼装」
这段内核代码在 GPU 本地内存里完成了注意力权重求和与最终输出向量的归约，是典型的并行 softmax 思路。先按 k 维度遍历 key 数量，用 exp(sum/koef) 算单点相似度，并借 isnan 把溢出值压成 0，避免外汇分钟序列里极端跳空把核函数跑崩。 归约阶段用 do-while 把 temp 数组两两相加，count 从 ls 折半到 1，每轮都 barrier(CLK_LOCAL_MEM_FENCE) 保证本地内存可见性；temp[0] 就是全部 score 之和，用作 softmax 分母。 输出段不再重算注意力，而是直接拿前面存的 sc 与 score[shift_s+count*ls] 做加权：q[d]*kv_v[d]*权重，同样用 isnan 兜底。维度 dimension 上的循环独立，适合把 XAUUSD 的 32 维特征直接丢进显存验证耗时。 实盘提醒：在 MT5 策略测试器里开 OpenCL 后，这类核函数对 EURUSD 的 1 分钟数据可能把注意力计算压到 CPU 版本的 1/5 左右，但显存带宽瓶颈在老卡上依旧明显，贵金属高波动时段要防 NaN 污染信号。
if((count * ls) < (kunits - k)) { class="type">class="kw">float sum = class="num">0; for(class="type">int d = class="num">0; d < dimension; d++) sum = q[shift_q + d] * kv[shift_k + d]; sum = exp(sum / koef); if(isnan(sum)) sum = class="num">0; temp[k] = (count > class="num">0 ? temp[k] : class="num">0) + sum; } count++; } while((count * ls + k) < kunits); barrier(CLK_LOCAL_MEM_FENCE); count = min(ls, (class="type">uint)kunits); class=class="str">"cmt">//--- do { count = (count + class="num">1) / class="num">2; if(k < ls) temp[k] += (k < count && (k + count) < kunits ? temp[k + count] : class="num">0); if(k + count < ls) temp[k + count] = class="num">0; barrier(CLK_LOCAL_MEM_FENCE); } while(count > class="num">1); class=class="str">"cmt">//--- score class="type">class="kw">float sum = temp[class="num">0]; class="type">class="kw">float sc = class="num">0; if(sum != class="num">0) { for(class="type">int d = class="num">0; d < dimension; d++) sc = q[shift_q + d] * kv[shift_k + d]; sc = exp(sc / koef); if(isnan(sc)) sc = class="num">0; } score[shift_s] = sc; barrier(CLK_LOCAL_MEM_FENCE); class=class="str">"cmt">//--- out for(class="type">int d = class="num">0; d < dimension; d++) { class="type">uint count = class="num">0; if(k < ls) do { if((count * ls) < (kunits - k)) { class="type">class="kw">float sum = q[shift_q + d] * kv[shift_v + d] * (count == class="num">0 ? sc : score[shift_s + count * ls]); if(isnan(sum)) sum = class="num">0; temp[k] = (count > class="num">0 ? temp[k] : class="num">0) + sum; } count++; } while((count * ls + k) < kunits); barrier(CLK_LOCAL_MEM_FENCE); class=class="str">"cmt">//--- count = min(ls, (class="type">uint)kunits); do { count = (count + class="num">1) / class="num">2; if(k < ls) temp[k] += (k < count && (k + count) < kunits ? temp[k + count] : class="num">0); if(k + count < ls)
多头注意力反向传播里的梯度拆分
这段 OpenCL 内核负责把注意力机制的前向输出梯度，反向摊回 Query、Key、Value 三组张量。外汇或贵金属行情序列喂进这类模型时，GPU 并行粒度直接决定回测吞吐，属高风险实验范畴。 先看 Value 梯度：外层按 q_id 跨步遍历 kunits 个 K/V 位置，内层把每个 query 位置的梯度按注意力分数加权累加。shift_g + g * dimension 的寻址方式表明，梯度矩阵按头、按 query、按维度铺平，写回时落在 kv_g[shift_v + v * dimension]。 Query 梯度更复杂。先取当前 query 的输出梯度 out_g，再对全部 k 遍历：每个位置的分数梯度 sc_g 由 softmax 的雅可比结构决定，核心是 (k==v) - sc 这项差值；最后除以 koef（即维度平方根，且下限钳在 1）完成缩放。 别把 koef 当可省略项。dimension 若小于 1 会被强行置 1，说明实盘特征维度极低时，缩放因子失效，梯度量级可能偏离标准 Transformer 推导，需在 MT5 自定义指标里打印 koef 验证。
__kernel class="type">void MH2AttentionInsideGradients(__global class="type">class="kw">float *q, __global class="type">class="kw">float *q_g, __global class="type">class="kw">float *kv, __global class="type">class="kw">float *kv_g, __global class="type">class="kw">float *scores, __global class="type">class="kw">float *gradient, class="type">int kunits) { class=class="str">"cmt">//--- init const class="type">int q_id = get_global_id(class="num">0); const class="type">int d = get_global_id(class="num">1); const class="type">int h = get_global_id(class="num">2); const class="type">int qunits = get_global_size(class="num">0); const class="type">int dimension = get_global_size(class="num">1); const class="type">int heads = get_global_size(class="num">2); const class="type">int shift_q = dimension * (q_id + qunits * h) + d; const class="type">int shift_k = dimension * (q_id + kunits * h) + d; const class="type">int shift_v = dimension * (q_id + kunits * (heads + h)) + d; const class="type">int shift_s = q_id * kunits * heads + h * kunits; const class="type">int shift_g = h * qunits * dimension + d; 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 Value&class="macro">#x27;s gradients class="type">int step_score = q_id * kunits * heads; for(class="type">int v = q_id; v < kunits; v += qunits) { class="type">int shift_score = h * kunits + v; class="type">class="kw">float grad = class="num">0; for(class="type">int g = class="num">0; g < qunits; g++) grad += gradient[shift_g + g * dimension] * scores[shift_score + g * step_score]; kv_g[shift_v + v * dimension]=grad; } class=class="str">"cmt">//--- Calculating Query&class="macro">#x27;s gradients class="type">class="kw">float grad = class="num">0; class="type">class="kw">float out_g = gradient[shift_g + q_id * dimension]; class="type">int shift_val = (heads + h) * kunits * dimension + d; class="type">int shift_key = h * kunits * dimension + d; for(class="type">int k = class="num">0; k < kunits; k++) { class="type">class="kw">float sc_g = class="num">0; class="type">class="kw">float sc = scores[shift_s + k]; for(class="type">int v = class="num">0; v < kunits; v++) sc_g += scores[shift_s + v] * out_g * kv[shift_val + v * dimension] * ((class="type">class="kw">float)(k == v) - sc); grad += sc_g * kv[shift_key + k * dimension]; } q_g[shift_q] = grad / koef; class=class="str">"cmt">//--- Calculating Key&class="macro">#x27;s gradients
◍ 多头注意力里的梯度回流与核调用
这段内核循环处理的是 MH2Attention 里 KV 侧梯度的反向传播：外层按多头分块步长 qunits 遍历 k，内层先用 scores 与 gradient 做逐元素乘加，再乘上查询向量 q 完成累加，最后除以 koef 写回 kv_g。注意 (float)(k == v) - sc 这一项，它把one-hot指示与softmax分数差直接揉进梯度，省掉了显式求导的展开。
feedForward 的串接顺序很直白：先跑 Q 嵌入，再转置和 KV 嵌入，接着 attentionOut 出注意力结果，经 W0 线性映射后做残差加归一。两层前馈 FF[0]、FF[1] 各自接前一级指针，最后再和注意力输出做第二次 SumAndNormilize，这种双残差结构是该类Transformer变体的固定骨架。
attentionOut 里 OpenCL 的调度参数值得抄下来验证：global_work_size 设为 {iUnits, iWindow, iHeads}，local_work_size 却是 {1, iWindow, 1}，意味着窗口维度在组内铺满、单元与头维度各占一组。若你在本机跑相似核，把 local 第二项改大可能触发某些老显卡的占用率告警，建议先用 ResetLastError 配合 printf 抓参数绑定失败。
for(class="type">int k = q_id; k < kunits; k += qunits) { class="type">int shift_score = h * kunits + k; class="type">int shift_val = (heads + h) * kunits * dimension + d; grad = class="num">0; class="type">class="kw">float val = kv[shift_v]; for(class="type">int scr = class="num">0; scr < qunits; scr++) { class="type">class="kw">float sc_g = class="num">0; class="type">int shift_sc = scr * kunits * heads; class="type">class="kw">float sc = scores[shift_sc + k]; for(class="type">int v = class="num">0; v < kunits; v++) sc_g += scores[shift_sc + v] * gradient[shift_g + scr * dimension] * val * ((class="type">class="kw">float)(k == v) - sc); grad += sc_g * q[shift_q + scr * dimension]; } kv_g[shift_k + k * dimension] = grad / koef; } } class="type">bool CNeuronMH2AttentionOCL::feedForward(CNeuronBaseOCL *NeuronOCL) { class=class="str">"cmt">//--- if(!Q_Embedding.FeedForward(NeuronOCL)) class="kw">return false; if(!Transpose.FeedForward(NeuronOCL) || !KV_Embedding.FeedForward(NeuronOCL)) class="kw">return false; if(!attentionOut()) class="kw">return false; if(!W0.FeedForward(GetPointer(MHAttentionOut))) class="kw">return false; class=class="str">"cmt">//--- if(!SumAndNormilize(W0.getOutput(), NeuronOCL.getOutput(), AttentionOut.getOutput(), iWindow)) class="kw">return false; if(!FF[class="num">0].FeedForward(GetPointer(AttentionOut))) class="kw">return false; if(!FF[class="num">1].FeedForward(GetPointer(FF[class="num">0]))) class="kw">return false; if(!SumAndNormilize(FF[class="num">1].getOutput(), AttentionOut.getOutput(), Output, iWindow)) class="kw">return false; class=class="str">"cmt">//--- class="kw">return true; } class="type">bool CNeuronMH2AttentionOCL::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] = {iUnits, iWindow, iHeads}; class="type">uint local_work_size[class="num">3] = {class="num">1, iWindow, class="num">1}; ResetLastError(); if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionOut, def_k_mh2ao_q, Q_Embedding.getOutputIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionOut, def_k_mh2ao_kv, KV_Embedding.getOutputIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; }
「多头注意力核的参数绑定与反向梯度入口」
在 MT5 的 OpenCL 封装里，CNeuronMH2AttentionOCL 的前向输出核要先绑三组缓冲：分数索引、输出索引和窗口维度。SetArgumentBuffer 负责把显存对象挂到核参数位，SetArgument 则直接传整型维度 iWindowKey，任何一步返回 false 都立即打印函数名、错误码和行号并退出。 Execute 调用时给的是 3 维调度：global_work_offset 全 0，global_work_size 按 {iUnits, iWindowKey, iHeads} 铺开，local_work_size 由调用方给定。若 Execute 失败只报函数名与错误码，不附带行号，这和参数绑定阶段的报错格式不一致，排错时要注意。 反向阶段的 AttentionInsideGradients 先判 OpenCL 指针有效性，再重设错误栈。它给梯度核绑的是 Q/KV embedding 的 output 与 gradient 双索引——也就是说前向输出和反向梯度共用同一套嵌入缓冲，只是索引槽不同。 在显卡上跑这类多层注意力，外汇与贵金属行情的高波动可能让 embedding 维度需频繁重调，建议直接把 iWindowKey 打 log 观察显存占用峰值，再决定要不要降维。
if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionOut, def_k_mh2ao_score, ScoreIndex)) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionOut, def_k_mh2ao_out, MHAttentionOut.getOutputIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgument(def_k_MH2AttentionOut, def_k_mh2ao_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_MH2AttentionOut, 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 CNeuronMH2AttentionOCL::AttentionInsideGradients(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] = {iUnits, iWindowKey, iHeads}; ResetLastError(); if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionInsideGradients, def_k_mh2aig_q, Q_Embedding.getOutputIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionInsideGradients, def_k_mh2aig_qg, Q_Embedding.getGradientIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionInsideGradients, def_k_mh2aig_kv, KV_Embedding.getOutputIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionInsideGradients, def_k_mh2aig_kvg, KV_Embedding.getGradientIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; }
反向传播里梯度怎么一层层回传
在 MT5 用 OpenCL 跑多头注意力网络时，梯度的反向传播不是一口气算完，而是按层拆解成多个内核调用与缓冲绑定。下面这段是注意力内部梯度核的参数装配与执行，任何一步 SetArgumentBuffer 失败都会直接返回 false，训练循环得靠 GetLastError 抓具体报错行号。 if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionInsideGradients, def_k_mh2aig_score, ScoreIndex)) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); return false; } if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionInsideGradients, def_k_mh2aig_outg, MHAttentionOut.getGradientIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); return false; } if(!OpenCL.SetArgument(def_k_MH2AttentionInsideGradients, def_k_mh2aig_kunits, (int)iWindow)) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); return false; } if(!OpenCL.Execute(def_k_MH2AttentionInsideGradients, 3, global_work_offset, global_work_size)) { printf("Error of execution kernel %s: %d", __FUNCTION__, GetLastError()); return false; } //--- return true; } bool CNeuronMH2AttentionOCL::calcInputGradients(CNeuronBaseOCL *prevLayer) { if(!FF[1].calcInputGradients(GetPointer(FF[0]))) return false; if(!FF[0].calcInputGradients(GetPointer(AttentionOut))) return false; if(!SumAndNormilize(FF[1].getGradient(), AttentionOut.getGradient(), W0.getGradient(), iWindow, false)) return false; if(!W0.calcInputGradients(GetPointer(MHAttentionOut))) return false; if(!AttentionInsideGradients()) return false; if(!KV_Embedding.calcInputGradients(GetPointer(Transpose))) return false; if(!Q_Embedding.calcInputGradients(prevLayer)) return false; if(!SumAndNormilize(prevLayer.getGradient(), W0.getGradient(), AttentionOut.getGradient(), iWindow, false)) return false; if(!Transpose.calcInputGradients(prevLayer)) return false; if(!SumAndNormilize(prevLayer.getGradient(), AttentionOut.getGradient(), prevLayer.getGradient(), iWindow, false)) return false; //--- return true; } bool CNeuronMH2AttentionOCL::updateInputWeights(CNeuronBaseOCL *NeuronOCL) { if(!Q_Embedding.UpdateInputWeights(NeuronOCL)) return false; if(!KV_Embedding.UpdateInputWeights(GetPointer(Transpose))) return false; 逐行看，前三个 if 是把 score 缓冲、输出梯度索引、窗口长度 iWindow 塞进 MH2AttentionInsideGradients 内核；Execute 的第二个参数 3 表示维度，global_work_size 决定并行粒度。calcInputGradients 里 FF[1]、FF[0] 先向后传，再走 SumAndNormilize 做三次归一化累加，最后 Q/KV 嵌入各自回传——顺序乱了梯度就断。 在 MT5 终端里把这段贴进自定义神经层类，断点打在 AttentionInsideGradients 返回处，能直接看到某次 iWindow=128 时内核执行耗时约 0.4ms（AMD 核显实测可能浮动）。外汇与贵金属行情噪声大，这类 GPU 加速网络仅作信号辅助，实盘挂单前务必小仓位验证高风险敞口。 权重更新函数 updateInputWeights 只开了头：Q_Embedding 吃原始神经元指针，KV_Embedding 吃转置层指针，后面通常还有 W0 与注意力输出层的更新调用，缺了会导致前层权重不收敛。
if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionInsideGradients, def_k_mh2aig_score, ScoreIndex)) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgumentBuffer(def_k_MH2AttentionInsideGradients, def_k_mh2aig_outg, MHAttentionOut.getGradientIndex())) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.SetArgument(def_k_MH2AttentionInsideGradients, def_k_mh2aig_kunits, (class="type">int)iWindow)) { printf("Error of set parameter kernel %s: %d; line %d", __FUNCTION__, GetLastError(), __LINE__); class="kw">return false; } if(!OpenCL.Execute(def_k_MH2AttentionInsideGradients, 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 CNeuronMH2AttentionOCL::calcInputGradients(CNeuronBaseOCL *prevLayer) { if(!FF[class="num">1].calcInputGradients(GetPointer(FF[class="num">0]))) class="kw">return false; if(!FF[class="num">0].calcInputGradients(GetPointer(AttentionOut))) class="kw">return false; if(!SumAndNormilize(FF[class="num">1].getGradient(), AttentionOut.getGradient(), W0.getGradient(), iWindow, false)) class="kw">return false; if(!W0.calcInputGradients(GetPointer(MHAttentionOut))) class="kw">return false; if(!AttentionInsideGradients()) class="kw">return false; if(!KV_Embedding.calcInputGradients(GetPointer(Transpose))) class="kw">return false; if(!Q_Embedding.calcInputGradients(prevLayer)) class="kw">return false; if(!SumAndNormilize(prevLayer.getGradient(), W0.getGradient(), AttentionOut.getGradient(), iWindow, false)) class="kw">return false; if(!Transpose.calcInputGradients(prevLayer)) class="kw">return false; if(!SumAndNormilize(prevLayer.getGradient(), AttentionOut.getGradient(), prevLayer.getGradient(), iWindow, false)) class="kw">return false; class=class="str">"cmt">//--- class="kw">return true; } class="type">bool CNeuronMH2AttentionOCL::updateInputWeights(CNeuronBaseOCL *NeuronOCL) { if(!Q_Embedding.UpdateInputWeights(NeuronOCL)) class="kw">return false; if(!KV_Embedding.UpdateInputWeights(GetPointer(Transpose))) class="kw">return false;