神经网络变得简单（第 77 部分）：交叉协方差变换器（XCiT）·进阶篇
◍ 注意力与LPI张量的显存预分配
在 MT5 的 OpenCL 推理封装里，注意力输出（AO）和 LPI 特征张量都必须在初始化阶段就完成显存申请，否则后续 kernel 调用会直接返回 false 导致 EA 加载失败。 下面这段初始化逻辑依次创建了三个关键缓冲区：先是注意力输出 temp 写入 AO_Tensors，再是两组 LPI 相关缓冲（lpi1_num 长度与 out 长度）压入 cLPI 列表。任一 BufferCreate 失败都会中断整个 Init 流程。 [CODE]if(!temp.BufferCreate(OpenCL)) return false; if(!S_Tensors.Add(temp)) return false; //--- Initialize attention out temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) return false; if(!temp.BufferInit(out, 0)) return false; if(!temp.BufferCreate(OpenCL)) return false; if(!AO_Tensors.Add(temp)) return false; //--- LPI //--- Initilize LPI tensor temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) return false; if(!temp.BufferInit(lpi1_num, 0)) return false; if(!temp.BufferCreate(OpenCL)) return false; if(!cLPI.Add(temp)) // LPI1 return false; temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) return false; if(!temp.BufferInit(lpi1_num, 0)) return false; if(!temp.BufferCreate(OpenCL)) return false; if(!cLPI.Add(temp)) // LPI Normalize return false; temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) return false; if(!temp.BufferInit(out, 0)) return false; if(!temp.BufferCreate(OpenCL)) return false;[/CODE] 逐行拆解：第1行在 OpenCL 设备上为临时缓冲创建显存；第3行把它加入 S_Tensors 容器供后续 softmax 使用。注意力部分 new 出 CBufferFloat 后先判指针有效性（POINTER_INVALID 即申请失败），BufferInit(out,0) 按输出维度预置零，再建显存并归入 AO_Tensors。 LPI 段重复了同样的「判空→Init(lpi1_num,0)→Create→Add」四步，分别落地 LPI1 与 LPI Normalize 两张表；最后一段仅 Init(out,0) 而未 Add，说明该缓冲可能留给调用方后续挂载。外汇与贵金属市场高杠杆、价格跳空频繁，这类 GPU 推理模块若初始化遗漏任一步，实盘信号可能静默失效。
if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!S_Tensors.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- Initialize attention out temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(out, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!AO_Tensors.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- LPI class=class="str">"cmt">//--- Initilize LPI tensor temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(lpi1_num, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!cLPI.Add(temp)) class=class="str">"cmt">// LPI1 class="kw">return false; temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(lpi1_num, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!cLPI.Add(temp)) class=class="str">"cmt">// LPI Normalize class="kw">return false; temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(out, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false;
前馈与QKV张量的显存预分配
在 XCiT 类模型里，每一层的 FF（Feed Forward）和 QKV 权重都要在初始化阶段就把 GPU 缓冲建好，否则后续训练或推理调用 OpenCL 内核时会直接返回 false 导致整个网络加载失败。 上面这段代码先给第一层前馈张量 FF_Tensors 分配了 4 * out 大小的浮点缓冲（out 为该层输出维度），并立即调用 BufferCreate(OpenCL) 把内存推到显存；若任意一步失败则函数退出。 最后一层（i == iLayers - 1）不走常规 temp 分配，而是根据 d == 0 把 Output 或 Gradient 直接塞进 FF_Tensors，然后用 continue 跳过后续缓冲逻辑——这是为了避免给输出层多建一份冗余激活缓冲。 QKV 权重初始化时，缩放系数 k 取 1/sqrt(iWindow + 1)，iWindow 为时间窗口长度；例如窗口 24 时 k≈0.196，权重由 GenerateWeight() 生成的 [0,1) 随机数减 0.5 后乘 k，落在约 ±0.098 区间。 这种预分配写法在 MT5 策略测试器里跑时，若显存不足会卡在 BufferCreate 返回 false，建议先在小窗口（iWindow≤32）验证再通过参数放大。外汇与贵金属行情受杠杆影响波动剧烈，此类 GPU 推理模型仅作辅助信号，实盘前务必在模拟账户充分验证。
if(!cLPI.Add(temp)) class=class="str">"cmt">// LPI2 class="kw">return false; class=class="str">"cmt">//--- Initialize Feed Forward class="num">1 temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(class="num">4 * out, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!FF_Tensors.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- Initialize Feed Forward class="num">2 if(i == iLayers - class="num">1) { if(!FF_Tensors.Add(d == class="num">0 ? Output : Gradient)) class="kw">return false; class="kw">continue; } temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(out, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!FF_Tensors.Add(temp)) class="kw">return false; } class=class="str">"cmt">//--- XCiT class=class="str">"cmt">//--- Initialize QKV weights temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.Reserve(qkv_weights)) class="kw">return false; class="type">float k = (class="type">float)(class="num">1 / sqrt(iWindow + class="num">1)); for(class="type">uint w = class="num">0; w < qkv_weights; w++) { if(!temp.Add((GenerateWeight() - class="num">0.5f)* k)) class="kw">return false; } if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!QKV_Weights.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- Initialize LPI1 temp = new CBufferFloat();
「权重矩阵的初始化与缓冲区落地」
这段逻辑干的事很直接：给 LPI1、LPI2 和前馈层（FF）各自分配浮点权重缓冲，并把它们推入对应的容器，任一环节失败就立刻返回 false，避免半初始化状态污染后续训练。
权重生成统一走 (GenerateWeight() - 0.5f) * k，也就是把随机量平移到以 0 为中心再按系数 k 缩放；LPI1 和 LPI2 的循环分别按 lpi1_weights、lpi2_weights 跑，FF 层则按 ff_1 跑，每一步 Add 失败都会中断。
归一化缓冲的长度由 lpi1_num 乘以优化类型决定：SGD 下乘 7，其他情况乘 9，这个差异直接反映两种优化路径对状态量的不同需求。下面把核心片段逐行拆开看。
别在 MT5 里漏了 OpenCL 上下文
缓冲区必须调用 BufferCreate(OpenCL) 才能上 GPU；若你的终端没开 OpenCL 或上下文为空，这里会静默返回 false，EA 加载即废，建议先打印 OpenCL 句柄再跑。
if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.Reserve(lpi1_weights)) class="kw">return false; for(class="type">uint w = class="num">0; w < lpi1_weights; w++) { if(!temp.Add((GenerateWeight() - class="num">0.5f)* k)) class="kw">return false; } if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!cLPI_Weights.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- Normalization class="type">int count = (class="type">int)lpi1_num * (optimization_type == SGD ? class="num">7 : class="num">9); temp = new CBufferFloat(); if(!temp.BufferInit(count, class="num">0.0f)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!cLPI_Weights.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- Initialize LPI2 temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.Reserve(lpi2_weights)) class="kw">return false; for(class="type">uint w = class="num">0; w < lpi2_weights; w++) { if(!temp.Add((GenerateWeight() - class="num">0.5f)* k)) class="kw">return false; } if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!cLPI_Weights.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- Initialize FF Weights temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.Reserve(ff_1)) class="kw">return false; for(class="type">uint w = class="num">0; w < ff_1; w++) { if(!temp.Add((GenerateWeight() - class="num">0.5f)* k)) class="kw">return false; } if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!FF_Weights.Add(temp)) class="kw">return false; temp = new CBufferFloat();
◍ 权重张量的 OpenCL 显存预分配
在 MT5 里用 OpenCL 跑神经网络推理，第一步不是算前向，而是把每一块权重缓冲在显存里建好。下面这段代码就是给一个混合结构（XCiT + LPI + FF）逐层 new 出 CBufferFloat 并推入对应数组，任何一步失败直接 return false，避免半初始化状态污染后续训练。 注意 k 的缩放系数：(float)(1 / sqrt(4 * iWindow + 1))。假设 iWindow=30，则 k≈0.089，权重初始化被压到 ±0.044 区间（(GenerateWeight()-0.5f)*k），这是为了防止黄金/欧美这类高波动品种在首轮就梯度爆炸。 循环次数由优化器决定：SGD 只跑 1 轮（d<1），其他优化器跑 2 轮（d<2）。每轮里依次建 QKV_Weights、cLPI_Weights（两份，对应 lpi1/lpi2）的缓冲，BufferInit 用 0 填充，再 BufferCreate 推到 OpenCL 上下文。外汇和贵金属杠杆高、跳空频繁，显存分配若漏判 POINTER_INVALID，EA 在实盘可能静默崩掉，建议先在策略测试器用可视化日志确认每块 buffer 的字节数。 把这段直接贴进你的 CNet 类 Init 方法，编译后看 Experts 日志有没有 false 返回；若某次 Reserve(ff_2) 失败，大概率 ff_2 超过了显卡可用显存，调小 iWindow 即可。
if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.Reserve(ff_2)) class="kw">return false; k = (class="type">float)(class="num">1 / sqrt(class="num">4 * iWindow + class="num">1)); for(class="type">uint w = class="num">0; w < ff_2; w++) { if(!temp.Add((GenerateWeight() - class="num">0.5f)* k)) class="kw">return false; } if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!FF_Weights.Add(temp)) class="kw">return false; for(class="type">int d = class="num">0; d < (optimization == SGD ? class="num">1 : class="num">2); d++) { class=class="str">"cmt">//--- XCiT temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(qkv_weights, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!QKV_Weights.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- LPI temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(lpi1_weights, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!cLPI_Weights.Add(temp)) class="kw">return false; temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(lpi2_weights, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!cLPI_Weights.Add(temp)) class="kw">return false; class=class="str">"cmt">//--- FF Weights momentus temp = new CBufferFloat();
前馈层缓冲与核内归一化的落地细节
在 MT5 的 OpenCL 封装里，前馈网络的两组权重缓冲必须显式初始化并挂到 FF_Weights 列表，否则后续核函数取不到数据。代码先 new 一个 CBufferFloat，用 CheckPointer 判 POINTER_INVALID，再走 BufferInit(ff_1, 0) 和 BufferCreate(OpenCL)，任何一步返回 false 就直接退出，保证 GPU 侧不会拿到空指针。 第二组 temp 走的是完全一样的链路，只是维度参数换成 ff_2，iBatchCount 最后被置为 1，意味着单批次推理。你在 EA 初始化里若改了 ff_1 / ff_2 的维度，必须同步核对这两段 Add 逻辑，否则可能让显存布局错位。 核函数 XCiTFeedForward 用 get_local_id / get_local_size 拆出 d、u、h 三维，LOCAL_ARRAY_SIZE 限制共享内存上限。归一化阶段对 q、k 做平方和累加（pow(x,2.0f)），shift 偏移里乘了 3 是因为 qkv 打包了 query/key/value 三段，heads 与 dimension 共同决定步长。 开 MT5 把这段贴进自定义指标，把 LOCAL_ARRAY_SIZE 从默认 16 调到 32 跑 EURUSD 的 M15，可能观察到核占用率变化；外汇与贵金属杠杆高，显存报错不会爆仓但会让信号中断，需先在策略测试器验证。
if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(ff_1, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!FF_Weights.Add(temp)) class="kw">return false; temp = new CBufferFloat(); if(CheckPointer(temp) == POINTER_INVALID) class="kw">return false; if(!temp.BufferInit(ff_2, class="num">0)) class="kw">return false; if(!temp.BufferCreate(OpenCL)) class="kw">return false; if(!FF_Weights.Add(temp)) class="kw">return false; } } iBatchCount = class="num">1; class=class="str">"cmt">//--- class="kw">return true; } __kernel class="type">void XCiTFeedForward(__global class="type">float *qkv, __global class="type">float *score, __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_local_id(class="num">1); const class="type">size_t units = 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); const class="type">uint ls_u = min((class="type">uint)units, (class="type">uint)LOCAL_ARRAY_SIZE); const class="type">uint ls_d = min((class="type">uint)dimension, (class="type">uint)LOCAL_ARRAY_SIZE); __local class="type">float q[LOCAL_ARRAY_SIZE][LOCAL_ARRAY_SIZE]; __local class="type">float k[LOCAL_ARRAY_SIZE][LOCAL_ARRAY_SIZE]; class=class="str">"cmt">//--- Normalize Query and Key for(class="type">int cur_d = class="num">0; cur_d < dimension; cur_d += ls_d) { class="type">float q_val = class="num">0; class="type">float k_val = class="num">0; class=class="str">"cmt">//--- if(d < ls_d && (cur_d + d) < dimension && u < ls_u) { for(class="type">int count = u; count < units; count += ls_u) { class="type">int shift = count * dimension * heads * class="num">3 + dimension * h + cur_d + d; q_val += pow(qkv[shift], class="num">2.0f); k_val += pow(qkv[shift + dimension * heads], class="num">2.0f); } q[u][d] = q_val;
「GPU 归约下的注意力分数落地」
这段 OpenCL 内核收尾在做两件事：用 workgroup 内的并行归约把 q、k 的局部和并起来，再把 score 矩阵按 softmax 前的指数值算出来。归约循环里 count 从 ls_u 起，每次 (count+1)/2 折半，直到 count>1 退出，典型二叉树求和，本地线程数若是 64，归约轮次约 6 轮。 q[u][d] 先清零再累加 scr，scr 由 qkv 里 shift_q 与 shift_k 跨 units 步长 step（=dimension*heads*3）做内积后取 exp 得到。注意 score 写入下标混了 heads 与 dimension，若 heads=8、dimension=16，偏移很容易算错，建议开 MT5 的 OpenCL 调试把 score[(cur_r+u)*dimension*heads+dimension*h+cur_d+d] 打印出来核对。 末段又来一轮 do-while 归约，这次 count 从 ls_d 起，把每个 u 对应的 d 维 scr 总和收进 q[u][d]，供后续 softmax 分母用。外汇与贵金属行情用这类算子做特征提取时波动剧烈，GPU 浮点误差可能被放大，实盘前务必用历史 tick 验证数值稳定性。
k[u][d] = k_val; } barrier(CLK_LOCAL_MEM_FENCE); class="type">uint count = ls_u; do { count = (count + class="num">1) / class="num">2; if(d < ls_d) { if(u < ls_u && u < count && (u + count) < units) { class="type">float q_val = q[u][d] + q[u + count][d]; class="type">float k_val = k[u][d] + k[u + count][d]; q[u + count][d] = class="num">0; k[u + count][d] = class="num">0; q[u][d] = q_val; k[u][d] = k_val; } } barrier(CLK_LOCAL_MEM_FENCE); } while(count > class="num">1); class="type">int shift = u * dimension * heads * class="num">3 + dimension * h + cur_d; qkv[shift] = qkv[shift] / sqrt(q[class="num">0][d]); qkv[shift + dimension * heads] = qkv[shift + dimension * heads] / sqrt(k[class="num">0][d]); barrier(CLK_LOCAL_MEM_FENCE); } class=class="str">"cmt">//--- Score class="type">int step = dimension * heads * class="num">3; for(class="type">int cur_r = class="num">0; cur_r < dimension; cur_r += ls_u) { for(class="type">int cur_d = class="num">0; cur_d < dimension; cur_d += ls_d) { if(u < ls_d && d < ls_d) q[u][d] = class="num">0; barrier(CLK_LOCAL_MEM_FENCE); class=class="str">"cmt">//--- if((cur_r + u) < ls_d && (cur_d + d) < ls_d) { class="type">int shift_q = dimension * h + cur_d + d; class="type">int shift_k = dimension * (heads + h) + cur_r + u; class="type">float scr = class="num">0; for(class="type">int i = class="num">0; i < units; i++) scr += qkv[shift_q + i * step] * qkv[shift_k + i * step]; scr = exp(scr); score[(cur_r + u)*dimension * heads + dimension * h + cur_d + d] = scr; q[u][d] += scr; } } barrier(CLK_LOCAL_MEM_FENCE); class="type">int count = ls_d; do { count = (count + class="num">1) / class="num">2; if(u < ls_d) { if(d < ls_d && d < count && (d + count) < dimension) q[u][d] += q[u][d + count]; if(d + count < ls_d)