**作者**昇腾实战派# 一、背景介绍2026年千问发布了Qwen3.5系列模型模型结构与Qwen3-Next类似出现了Gated Attention模块该系列模型可通过vllm-ascend框架在昇腾平台上部署。本文基于transform和vllm-ascend框架中的代码分析,提取了Qwen3.5-27B的Profiling将对Qwen3.5-27B模型代码展开讲解详细介绍每个API的内容和作用。vllm-ascend相关代码链接[vllm-ascend/vllm\_ascend/ops/gdn.py at e14b89cf3021639ae2f4093a3052ad96c13e8b52 · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/e14b89cf3021639ae2f4093a3052ad96c13e8b52/vllm_ascend/ops/gdn.py#L253)# 二、Qwen3.5-27B模型结构## 2.1 Qwen3.5-27B模型结构介绍Qwen3.5-27B模型结构如图1所示主要包含几个模块**13层Linear Attention主要包含Gated DeltaNet和FNN模块****21层Full Attention主要包含Gated Attention和FNN模型其中Gated Attention与Qwen3系列模型Attention基本一致仍然沿用GQA的结构****3MTP**图1 Qwen3.5-27B模型结构示意图## 2.2 GDN简介### 2.2.1 GDN背景介绍通过注意力核线性化linear attention改造传统自注意力将复杂度从O(n2)降至O(n)采用业界顶尖的**GDNGated Delta Rule线性注意力架构继承 Delta Net 原生动态 Delta 离散化规则并吸收 Mamba 选择性 SSM 的全局遗忘与门控设计在实现跨序列全局平滑衰减的同时以单键值隐状态累积替换**机制完成序列建模无需全局 token 两两交互。Linear Attention主要优势有如下3点* Linear Attention 将标准 Attention 的 softmax 替换为近似核函数并改变 Q/K/V 的计算顺序先算再乘 Q将复杂度从O(L^2d)降为O(Ld^2)* 在自回归场景下Linear Attention 可以转化为 RNN 递推形式有状态 h、时序递归每一步仅需维护一个**固定尺寸的隐状态矩阵 h**无需保存历史全部 KV。* 递推形式每步串行且无法利用 tensor core因此采用** chunkwise 分块并行算法**兼顾效率。随着序列长度增加Flash Attention 的O(L^2d)计算量导致 kernel 逐渐进入 compute-bound而 Linear Attention 的O(Ld^2)复杂度优势开始显现。### 2.2.2 GDN算法简介Linear attention: 每个时间步 t 通过线性投影产生三个列向量: qt, vt, kt则当前的状态St可以表示为其中通过交换律可以推导到原来的公式原始 Linear Attention 在模型效果上远弱于标准 Transformer根本原因是**固定大小的隐状态容量有限导致”记忆碰撞”且缺乏遗忘机制**。为解决这个问题有以下两种改进方向* **Mamba2/GLA**引入标量门控衰减  可以快速擦除全局记忆但不能选择性更新单个键值对。* **DeltaNet**使用 delta 规则精确替换特定键值对但缺乏全局记忆的快速清除机制。**Gated Delta Networks (GDN)** GDN结构如图2所示将两者结合在线性递推的基础上同时引入门控衰减和 delta 更新强度 兼具全局遗忘和定向更新能力但需要用chunkwise算法实现。图2 Gated DeltaNet网络结构示意图### 2.2.3 GDN关键步骤总览**主要流程参考如下**1. 输入投影in_proj_qkv / in_proj_z / in_proj_ba→ 得到 QKV、z、b、a2. 因果卷积causal_conv1d→ 局部特征提取3. 生成 g(alpha)、beta→ 门控参数4. Gated Delta Rule 核心计算循环状态机→ 长序列建模5. 门控归一化 输出投影→ norm(z * out) → out_proj# 三、Qwen3.5-27B模型结构代码走读## 3.1 Qwen3.5-27B模型代码结构总览## 3.2 GatedDeltaNetAttention分析### 3.2.1 GatedDeltaNetAttention初始化self.hidden_size 模型维度 (如 2048 / 4096)self.num_k_heads Q/K 头数self.num_v_heads V 头数 (通常是 K 的整数倍)self.head_k_dim Q/K 每头维度self.head_v_dim V 每头维度self.key_dim num_k_heads * head_k_dimself.value_dim num_v_heads * head_v_dimself.conv_kernel_size 因果卷积核大小### 3.2.2 QKV、z、b、a输入映射if hasattr(self, in_proj_qkv):# LoRA path (Qwen3.5 only): separate in_proj_qkv and in_proj_zmixed_qkv, _ self.in_proj_qkv(hidden_states)ba, _ self.in_proj_ba(hidden_states)z, _ self.in_proj_z(hidden_states)z z.reshape(z.size(0), -1, self.head_v_dim)b, a ba.chunk(2, dim-1)b b.contiguous()a a.contiguous()else:mixed_qkvz, _ self.in_proj_qkvz(hidden_states)ba, _ self.in_proj_ba(hidden_states)if self.gqa_interleaved_layout:# Qwen3-Next: unpack the interleaved GQA layoutquery, key, value, z, b, a self.fix_query_key_value_ordering(mixed_qkvz, ba)query, key, value map(lambda x: rearrange(x, l p d - l (p d)), (query, key, value))mixed_qkv torch.cat((query, key, value), dim-1)else:# Qwen3.5: weights are already in [q, k, v, z] and [b, a] orderqkv_size (self.key_dim * 2 self.value_dim) // self.tp_sizez_size self.value_dim // self.tp_sizemixed_qkv, z mixed_qkvz.split([qkv_size, z_size], dim-1)z z.reshape(z.size(0), -1, self.head_v_dim)b, a ba.chunk(2, dim-1)b b.contiguous()a a.contiguous()* **in_proj_qkv**: 把输入映射成 QKV shape: hidden_size → key_dim key_dim value_dim输入shape[batch_size,seq_len, hidden_size//tp]输出shape[batch_size,seq_len, (key_dimkey_dimvalue_dim)//tp]* **in_proj_z**: 门控信号 shape: hidden_size → value_dim 输入shape[batch_size,seq_len, hidden_size//tp]输出shape[batch_size,seq_len, value_dim//tp]--- before reshape输出shape[batch_size,seq_len, num_v_heads//tp, head_v_dim]--- after reshape* **in_proj_ba**: 循环门控参数输入shape[batch_size,seq_len, hidden_size//tp]输出shape[batch_size,seq_len, 2*num_v_heads//tp]后续每个输入会根据上面输出进行拆解q shape[batch_size,seq_len, key_dim//tp]k shape[batch_size,seq_len, key_dim//tp]v shape[batch_size,seq_len, value_dim//tp]b shape: [batch_size,seq_len, num_v_heads//tp]a shape: [batch_size,seq_len, num_v_heads//tp]### 3.2.3 gdn_attention_core计算整个GatedDeltaNetAttention网络最核心的部分是调用了 gdn_attention_core core_attn_out torch.zeros((num_tokens, self.num_v_heads // self.tp_size, self.head_v_dim),dtypehidden_states.dtype,devicehidden_states.device,)torch.ops.vllm.gdn_attention_core(mixed_qkv,b,a,core_attn_out,_encode_layer_name(self.prefix),)**关键步骤**1. 申请一块全是0的输入内存 vLLM 算子是**in-place 写入**模式必须提前给空间。2. 跑完整的GatedDeltaNet计算结果写入上面空内存**【gdn_attention_core计算核心逻辑】**#### 3.2.3.1 causal_conv1d_update_npu**1API功能介绍****2API入参介绍**decode阶段对应的causal_conv1d备注prefill阶段是调用到了Ascend C固在此不展开描述* KERNEL_WIDTH w 卷积核长度qwen3.5固定width4卷积核大小较小属于small kernel* d通道维度 dim* t时间步 token 位置* xt[d]第 t 个 token、第 d 通道输入* w[d, k]第 d 通道、卷积核第 k 个权重代码 weight: (dim, width)* b[d]偏置 bias* yt[d]卷积原始输出* ot[d]最终输出代码里的 o* 因果规则**只能用 t, t-1, t-2,... 历史不能用未来 t1****3API核心公式**拆开逐段对应代码1. 卷积求和内核循环 j2. SiLU 激活源码 SILU_ACTIVATION最终* **xₜ**当前 token你有* **xₜ₋₁, xₜ₋₂ ...**过去的 token**必须存在 cache 里**decode推理的时候是**一个 token 一个 token 生成**的不可能每次都把整个历史序列重新传一遍。所以**cache_t 保存最近的w个历史 x供下一次卷积直接用**输出是mixed_qkv后续再拆分成q,k,v作为gated delta rule的输入公式如下query_spec, key_spec, value_spec self.rearrange_mixed_qkv(mixed_qkv_spec)query_non_spec, key_non_spec, value_non_spec self.rearrange_mixed_qkv(mixed_qkv_non_spec)#### 3.2.3.2 fused_gdn_gating_patch1API 功能介绍什么时候调用在prefill阶段和decode开启mtp阶段会用到2API 核心公式这里输出是g和这里的就是后面的关于g的操作* Softplus 保证输出**恒正**softplus(x)≥0β是温度系数缩放时间步threshold是溢出保护当βx过大时直接线性截断避免  爆炸。* 负号 - → 保证 exp(g) 永远 ** 1**对于后续GDN作用是让历史状态不断遗忘;* A_log → 每个头固定衰减率因为 beta_output 是**输入门控**必须满足* 范围 **(0,1)*** 对于后续GDN作用**控制当前 token 对状态 h 的贡献强度**也就是**控制**新信息注入强弱****#### **3.2.2.3 chunk_gated_delta_rule****1为什么在prefill阶段要用chunk**把长序列切成大小固定的块 BT64**长序列尤其 Prefill 阶段在 NPU/GPU 上**直接算会极慢、显存爆炸、无法并行且NPU/GPU 不擅长逐 token 循环对于Prefill 长序列无法一次性载入算子**串行依赖**对于每个时间步骤St依赖St-1序列方向完全串行无法利用序列维度的并行性。1. **块内**64 个 token 一次性并行算完内部递归2. **块间**只传递块首尾的状态 S* 块开始时载入上一块的最终状态 S_prev* 块内跑完所有 token 后输出本块最终状态 S_curr* 传给下一个块当初始状态**计算复杂度的下降**对于chunk计算可以写成如下公式这里N代表新项o代表旧项的门控衰减系数均可看成常量进行移项可以得到如下公式得到的因果下三角核如下也就是对应的公式S计算可以通过计算得到。由于这里的A是对应一个chunk大小的矩阵通过等比求和公式可以得到如下公式假设chunk大小是K对于下三角矩阵A^K0,固对于k做截断只需要计算IA...A^k-1即可复杂度降到了O(K^2)总的计算复杂度从如果序列长度是L)O(L^2)降到了O(LK)**为了保持全文的连贯性与可读性后续隐状态Ht均用St来表示****2输入shape介绍**q shape[batch_size,seq_len, num_k_heads//tp, head_k_dim]k shape[batch_size,seq_len, num_k_heads//tp, head_k_dim]v shape[batch_size,seq_len, num_v_heads//tp, head_v_dim]beta shape: [batch_size,seq_len, num_v_heads//tp]g shape: [batch_size,seq_len, num_v_heads//tp]**3调用逻辑**chunk_gated_delta_rule() [入口函数参数校验 格式转换]↓ChunkGatedDeltaRuleFunction.apply()↓forward() [自动求导的前向传播]↓ (可选)l2norm_fwd(q) L2归一化l2norm_fwd(k)↓chunk_gated_delta_rule_fwd() [真正的核心前向计算]↓┌─────────────────────────────────────────┐│ 核心计算流水线分块并行计算 ││ 1. chunk_local_cumsum(g) 门控累积和 ││ 2. chunk_scaled_dot_kkt_fwd() 计算A矩阵 ││ 3. solve_tril() 下三角求解 ││ 4. recompute_w_u_fwd() 计算w, u ││ 5. chunk_gated_delta_rule_fwd_h() 计算h ││ 6. chunk_fwd_o() 计算最终输出o │└─────────────────────────────────────────┘↓返回 o, final_state##### 3.2.2.3.1 l2norm_fwd**1功能介绍**这段代码中主要是通过triton实现L2 归一化实现的公式可以写成如下y x / sqrt( sum(x²) eps )目标**让变成单位向量**作用对于fwd_h里的公式后续会提到有大量的指数衰减项exp)* 如果k很大每次更新的S就会很大导致S指数膨胀* 如果k太大就会导致chunk跨块传递S导致数值爆炸**L2Norm 把 K 固定模长 1**→ K 大小永远不变→ S 只随遗忘 g、递归更新**不会随向量长度膨胀**→ Chunk 分块递归数值完全稳定**2代码介绍**q l2norm_fwd(q)k l2norm_fwd(k)##### 3.2.1.3.2 chunk_gated_delta_rule_fwd**1代码介绍**def chunk_gated_delta_rule_fwd(q: torch.Tensor,k: torch.Tensor,v: torch.Tensor,g: torch.Tensor,beta: torch.Tensor,scale: float,initial_state: torch.Tensor,output_final_state: bool,cu_seqlens: torch.LongTensor | None None,prebuilt_metaNone,):......g chunk_local_cumsum(g,chunk_sizechunk_size,cu_seqlenscu_seqlens,block_indicesblock_indices_cumsum,)# obtain WY representation. u is actually the new v.A chunk_scaled_dot_kkt_fwd(kk,betabeta,g_cumsumg,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,output_dtypetorch.float32,)A solve_tril(AA,cu_seqlenscu_seqlens,chunk_indices_large_blockchunk_indices_large_block,chunk_indices_btchunk_indices_chunk64,output_dtypek.dtype,)w, u recompute_w_u_fwd(kk,vv,betabeta,AA,g_cumsumg,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,)h, v_new, final_state chunk_gated_delta_rule_fwd_h(kk,ww,uu,gg,initial_stateinitial_state,output_final_stateoutput_final_state,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,chunk_offsetschunk_offsets_chunk64,)......o chunk_fwd_o(qq,kk,vv_new,hh,gg,scalescale,cu_seqlenscu_seqlens,chunk_offsetschunk_offsets_chunk64,)......**2入参介绍**这里入参就是上面经过因果卷积ab转换成g和beta还有l2 norm之后的输出q shape[batch_size,seq_len, key_dim//tp]k shape[batch_size,seq_len, key_dim//tp]v shape[batch_size,seq_len, value_dim//tp]beta shape: [batch_size,seq_len, num_v_heads//tp]g shape: [batch_size,seq_len, num_v_heads//tp]**3API调用逻辑**###### 1准备工作forward_context get_forward_context()num_decodes 0chunk_size 64 # 固定分块大小64 token 一块根据seq_len切分###### 2 chunk_local_cumsumg chunk_local_cumsum(g, chunk_size64, ...)对每个 chunk 内的 g 做累计和用于快速计算历史状态衰减这里的 g 必须做 cumsum是为了在下面公式中后面的fwd_h让chunk内的递归公式变成**可并行计算**更详细地原本的GDN状态更新是其中ht必须依赖ht-1,必须串行现在对于g做累计和代入原公式可以得到**展开后St 不再依赖前一个 S_t-1了**只依赖* 初始状态 S0* 累积和 G_t、G_i输入输出shape不变g shape: [batch_size,seq_len, num_v_heads//tp]###### 3 chunk_scaled_dot_kkt_fwdA chunk_scaled_dot_kkt_fwd(k, beta, g_cumsumg, ...)* **chunk_scaled_dot_kkt_fwd_kernel 做一件事**在每个 chunk 内计算带门控、带缩放、带因果掩码的 k·k.T 矩阵公式A[i,j] β[i] * exp(g[i]-g[j]) * k[i]·k[j].T且 i j 才有效A shape:[batch_size,seq_len, num_v_heads//tp, chunk_size]###### 4 solve_tril三角求解 (I A)⁻¹A solve_tril(A, ...)**目的**让后面 state 递推变成 **O (1) 计算**###### 5recompute_w_u_fwd重新计算 w, u新 K, Vw, u recompute_w_u_fwd(k, v, beta, A, g_cumsumg, ...)* **输入**k, v, beta, A* **输出*** w新 key* u新 value* **u A · (v · β)** → 新 value* **w A ** ·** (k · β · exp(g))** → 新 key* **目的**让记忆 state 更新更快、更稳定v 原始 shape [batch_size,seq_len, num_v_heads//tp, head_v_dim] beta 进 kernel 前被转置成 [num_v_heads//tp,batch_size,seq_len] [16, 1, 8192]对于每个batch_size和每个head_num的chunk的v和k进行recompute┌─────────────┐ ┌─────────────┐│ A [64,64] │ │ V [64,128] │└──────┬──────┘ └──────┬──────┘│ ││ beta [64] → [64,1]│ ││ ▼│ V * beta [64,128]│ │└─────────┬─────────┘│▼U A (V*beta)[64,128]chunk A和chunk V和chunk beta做运算得到chunk u最终按照batch size和head num放到最终u不同的位置上| 数据 | shape | 含义 || ----------------- | ------------------------------ | --------------------------------------------- || 单个 chunk 输出 | **[64, 128]** | 一个 chunk 一个 head 的结果 || 最终完整 U | **[1, 8192, 48, 128]** | 全部 batch、全部 token、全部 head、全部维度 |对于k和beta和g的num_k_head和num_v_head不一致的情况这里用了如下方法ptr_k k (bos * Hg i_h // (H // Hg)) * K offs_t_2d * (Hg * K) offs_k * 1举例num_k_heads 48num_v_heads 16这里head的映射关系如下每 3 个 V head 共享 1 个 K headi_h (V head) | 对应的 K head0 → 0 //3 01 → 1 //3 02 → 2 //3 03 → 3 //3 14 → 4 //3 15 → 5 //3 1...45,46,47 → 15类似的w输出的shape是[batch_size,seq_len, num_v_heads//tp, head_v_dim]###### 6chunk_gated_delta_rule_fwd_h这里是在计算隐状态S是一个串行运行各个chunk的函数每个后面的St都和前面的St-1相关chunk内是并行操作最重要的步骤h, v_new, final_state chunk_gated_delta_rule_fwd_h(kk,ww,uu,gg,initial_stateinitial_state,output_final_stateoutput_final_state,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,chunk_offsetschunk_offsets_chunk64,)**输入*** kkey* w优化后的 key* u优化后的 value* g门控累计和控制更新强度* initial_state初始记忆 S₀上文已讲后续h全部用s代替来增加可读性本质不变可以是none这时就会跳过内部如下命令 如果有 shape[real_batch_size, num_v_heads//tp, head_k_dim, head_v_dim],后续每次都有前一次chunk做完的隐状态作为这次的initial stateif USE_INITIAL_STATE:h0_ptr h0 i_nh * K * Vptr_h0_bv1 h0_ptr offs_k * V offs_v1 * 1b_h1_bv1 tl.load(ptr_h0_bv1, maskmask_kv1, other0.0).to(tl.float32)ptr_h0_bv2 h0_ptr offs_k * V offs_v2 * 1b_h1_bv2 tl.load(ptr_h0_bv2, maskmask_kv2, other0.0).to(tl.float32)**输出*** **S** (每个 chunk 的状态)shape: [batch_size,chunk_num, num_v_heads//tp, head_k_dim, head_v_dim]**chunk_num 序列被切成多少个 chunk*** **St (final_state)** (最终状态)shape: [real_batch_size, num_v_heads//tp, head_k_dim, head_v_dim]真正在运行的代码是这一段b_v_new1 b_v1 - tl.dot(b_w, b_h1_bv1)b_v_new1 b_v_new1 * b_gb_h1_bv1 b_h1_bv1 * b_g_lastb_h1_bv1 tl.dot(b_k, b_v_new1)固对于第t个chunk写成数学公式为* St-1:第 t−1 块处理完毕后整条序列到目前为止所有历史信息压缩状态* St:第 t 块处理完之后新的全局累积状态* Kt:当前分块的键* Vt:当前分块的值这个v是前一步骤转换后的u* Wt:可以当成全局固定化权重,是前一步k转换后的w* 增量Vt-W*Ht-1这就是 “delta rule” 名字来源* gt第t个chunk的门控衰减系数指数衰减控制历史信息遗忘* e^gt对历史状态Ht-1做指数衰减 e^gt*Ht-1:遗忘旧记忆* e^(gL−gt)对当前增量做衰减* gL当前 chunk 最后一个 token 的门控也就是chunk内原始g的累加和更准确地公式可以如下###### 7 chunk_fwd_o算最终输出 o。这个函数对于所有chunk可以并行运行根据如下公式可以看到每个chunk直接选取H对应的位置即可公式如下输出由两部分组成**inter-chunk**当前 query 与历史累积状态的交互和 **intra-chunk**当前 chunk 内部 query 和 key 的注意力交互。**inter-chunk**当前 chunk 的 query 与前面所有 chunk 的累积隐状态的交互。对应代码b_o tl.dot(b_q, b_h)b_o b_o * tl.exp(b_g)[:, None]**intra-chunk**当前 chunk 内部注意力对应代码b_A tl.dot(b_q, b_k)b_A b_A * safe_exp(b_g[:, None] - b_g[None, :])b_A tl.where(m_A, b_A, 0)tl.dot(b_A, b_v)mask只保留下三角上三角全部置 0禁止看到未来 token# 门控衰减矩阵 Gamma[i,j] exp(g[i] - g[j])g 是 log-space 累积和# 举例g_cumsum [-0.1, -0.3, -0.5, -0.7]## Gamma # j0 j1 j2 j3# i0 [exp(0) 0 0 0 ] - 对角线无衰减# i1 [exp(-0.2) exp(0) 0 0 ] - token 0 对 token 1 衰减 exp(-0.2)# i2 [exp(-0.4) exp(-0.2) exp(0) 0 ]# i3 [exp(-0.6) exp(-0.4) exp(-0.2) exp(0)] - 越远的过去衰减越多为了更清楚地解释上面3-5步骤对于6-7步骤的逻辑这里引入WY表示这里通过deltanet的公式进行解释gated deltanet是一样的deltanet比较容易理解该公式可以用简单数学归纳法证明。首先定义n表示第n个token。当n1时公式固然成立。假设对于n-1也成立那么我们证明对于n也成立证明过程如下该证明不仅证明了公式的正确性也提供了w_n的计算公式通过Sn的递推公式我们可以证明通过归纳法那么对于在一个chunk i内部第r个位置我们通过递归到S[i]会有如下公式w和u在这里都是通过WY表示计算的但是从每一个块的第一个位置开始不是从序列起始位置开始从而第r个位置有如下w和u对于输出计算结合矩阵乘法形式可以得到最终公式#### 3.2.1.4 recurrent gated delta rule开启投机推理会运行该算子替代原先的fused_sigmoid_gating_delta_ruleRecurrent 模式是直接按时间步逐个计算递推过程适用于推理解码场景逐 token 生成。和prefill阶段的区别是把chunk_fwd_h和chunk_fwd_o都合并在一步里公式如下是对于逐个token递归这个和原生的gated delta rule是等价的门控衰减对应保证了 delta 项中擦除的是衰减后的旧值。代码实现如下b_h * exp(b_g) # h alpha * hb_v b_beta * (b_v - tl.sum(b_h * b_k[:, None], 0)) # v_new beta * (v - alpha*S k)b_h b_k[:, None] * b_v # h k * v_new^T另外在 kernel 中也做了GVA 的支持Q/K 每时间步前进 num_k_heads*head_k_dim//tpV/O 每时间步前进 num_v_heads*head_v_dim//tp。g 的步长也是 num_v_heads//tp或者是num_v_heads*head_k_dim//tp取决于是否是KDAKDA 模式**一个头K 个 g每个 key 维度一个**beta 的步长取决于是否为 headwise 模式。p_q H * Kp_k H * Kp_o HV * Vp_v HV * Vif not IS_KDA:p_g HVelse:p_gk HV * Kp_beta HV * (V if IS_BETA_HEADWISE else 1)#### 3.2.1.5 fused sigmoid gating delta rule和fused recurrent gated delta rule的区别是输入sigmoid输入是a,b recurrent输入是g和beta公式和fused recurrent gated delta rule一样只是在运行这段代码前又做了a,b到g,beta的转换可参考fused_gdn_gating_patch## 3.3 Qwen3.5-27B Profiling算子分析| **profiling算子名称** | **对应代码调用位置** | **代码路径** || -------------------------------------------------------- | -------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- | ---------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- || _causal_conv1d_update_kernel_npu_tiled_5 | _causal_conv1d_update_kernel_npu_tiled[grid]](x,weight,bias,conv_state,conv_state_indices,num_accepted_tokens,query_start_loc,block_idx_last_scheduled_token,initial_state_idx,out,batch,dim,seqlen,...) | [vllm-ascend/vllm\_ascend/ops/triton/mamba/causal\_conv1d.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/mamba/causal_conv1d.py#L527) || l2norm_fwd_kernel2_loop | q l2norm_fwd(q)k l2norm_fwd(k) | [vllm-ascend/vllm\_ascend/ops/triton/fla/l2norm.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/l2norm.py#L34) || chunk_local_cumsum_scalar_kernel | chunk_local_cumsum(g,chunk_sizechunk_size,cu_seqlenscu_seqlens,block_indicesblock_indices_cumsum,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/cumsum.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/cumsum.py#L116) || chunk_scaled_dot_kkt_fwd_kernel | chunk_scaled_dot_kkt_fwd(kk,betabeta,g_cumsumg,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,output_dtypetorch.float32,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/chunk\_scaled\_dot\_kkt.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/chunk_scaled_dot_kkt.py#L83) || solve_tril_16x16_kernel | A solve_tril(AA,cu_seqlenscu_seqlens,chunk_indices_large_blockchunk_indices_large_block,chunk_indices_btchunk_indices_chunk64,output_dtypek.dtype,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/solve\_tril.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/solve_tril.py#L330) || recompute_w_u_fwd_kernel | w, u recompute_w_u_fwd(kk,vv,betabeta,AA,g_cumsumg,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/wy\_fast.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/wy_fast.py#L98) || chunk_gated_delta_rule_fwd_kernel_h_blockdim64 | h, v_new, final_state chunk_gated_delta_rule_fwd_h(kk,ww,uu,gg,initial_stateinitial_state,output_final_stateoutput_final_state,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,chunk_offsetschunk_offsets_chunk64,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/chunk\_delta\_h.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/chunk_delta_h.py#L179) || chunk_fwd_kernel_o | o chunk_fwd_o(qq,kk,vv_new,hh,gg,scalescale,cu_seqlenscu_seqlens,chunk_offsetschunk_offsets_chunk64,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/chunk\_o.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/chunk_o.py#L112) || fused_recurrent_gated_delta_rule_fwd_kernel_11 | core_attn_out_non_spec, last_recurrent_state fused_recurrent_gated_delta_rule(qquery_non_spec,kkey_non_spec,vvalue_non_spec,gg_non_spec,betabeta_non_spec,initial_statessm_state,inplace_final_stateTrue,cu_seqlensnon_spec_query_start_loc[: attn_metadata.num_decodes 1],ssm_state_indicesnon_spec_state_indices_tensor,use_qk_l2norm_in_kernelTrue,) | [vllm/vllm/model\_executor/layers/fla/ops/fused\_recurrent.py at main · vllm-project/vllm (github.com)](https://github.com/vllm-project/vllm/blob/main/vllm/model_executor/layers/fla/ops/fused_recurrent.py#L481) || fused_sigmoid_gating_delta_rule_update_kernel_0 | core_attn_out_non_spec fused_sigmoid_gating_delta_rule_update(A_logself.A_log.contiguous(),dt_biasself.dt_bias.contiguous(),qquery_non_spec.contiguous(),kkey_non_spec.contiguous(),vvalue_non_spec.contiguous(),aa.contiguous(),bb.contiguous(),initial_state_sourcessm_state,initial_state_indicesnon_spec_state_indices_tensor,cu_seqlensnon_spec_query_start_loc,use_qk_l2norm_in_kernelTrue,softplus_beta1.0,softplus_threshold20.0,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/sigmoid\_gating.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/sigmoid_gating.py#L180) |# 四、总结本文总结了 Qwen3.5-27B 模型结构并针对 GDN 网络结构做了完整拆解与原理分析。GDN 模块预填充Prefill阶段最核心算子为chunk_gated_delta_rule_fwd_kernel_h和chunk_fwd_kernel_o自回归解码Decode阶段开启 MTP 投机解码时核心算子为fused_recurrent_gated_delta_rule关闭 MTP 常规单 Token 解码时核心算子为fused_sigmoid_gating_delta_rule。在 Prefill 分块 Chunk 计算流程中矩阵 A 构建、门控 g/beta 计算、Q/K 归一化、初始状态加载等全部前置步骤均是为后续**fwd_h 隐状态计算**与**fwd_o 输出计算**两个核心步骤提供输入支撑。其中chunk_gated_delta_rule_fwd_kernel_h负责完成 GDN 内部时序隐状态的分块并行求解chunk_fwd_kernel_o基于求解完成的隐状态进一步计算注意力最终输出。同时本文梳理了投机解码模式下的 GDN 执行逻辑投机草稿 Token 分支与已验证正常 Token 分支分开计算、分开执行因果卷积与时序递归仅在输出端按照位置索引合并结果既保证投机解码加速能力又保证模型因果正确性。整体 GDN 架构采用**Causal Conv1D 局部时序建模 线性注意力全局长依赖**的互补设计通过 Chunk 并行预填充、Recurrent 串行解码实现线性复杂度 O (N) 超长上下文推理能力相比传统 Transformer 注意力具备显著效率优势。
Qwen3.5-27B Dense模型网络结构以及关键算子分析
**作者**昇腾实战派# 一、背景介绍2026年千问发布了Qwen3.5系列模型模型结构与Qwen3-Next类似出现了Gated Attention模块该系列模型可通过vllm-ascend框架在昇腾平台上部署。本文基于transform和vllm-ascend框架中的代码分析,提取了Qwen3.5-27B的Profiling将对Qwen3.5-27B模型代码展开讲解详细介绍每个API的内容和作用。vllm-ascend相关代码链接[vllm-ascend/vllm\_ascend/ops/gdn.py at e14b89cf3021639ae2f4093a3052ad96c13e8b52 · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/e14b89cf3021639ae2f4093a3052ad96c13e8b52/vllm_ascend/ops/gdn.py#L253)# 二、Qwen3.5-27B模型结构## 2.1 Qwen3.5-27B模型结构介绍Qwen3.5-27B模型结构如图1所示主要包含几个模块**13层Linear Attention主要包含Gated DeltaNet和FNN模块****21层Full Attention主要包含Gated Attention和FNN模型其中Gated Attention与Qwen3系列模型Attention基本一致仍然沿用GQA的结构****3MTP**图1 Qwen3.5-27B模型结构示意图## 2.2 GDN简介### 2.2.1 GDN背景介绍通过注意力核线性化linear attention改造传统自注意力将复杂度从O(n2)降至O(n)采用业界顶尖的**GDNGated Delta Rule线性注意力架构继承 Delta Net 原生动态 Delta 离散化规则并吸收 Mamba 选择性 SSM 的全局遗忘与门控设计在实现跨序列全局平滑衰减的同时以单键值隐状态累积替换**机制完成序列建模无需全局 token 两两交互。Linear Attention主要优势有如下3点* Linear Attention 将标准 Attention 的 softmax 替换为近似核函数并改变 Q/K/V 的计算顺序先算再乘 Q将复杂度从O(L^2d)降为O(Ld^2)* 在自回归场景下Linear Attention 可以转化为 RNN 递推形式有状态 h、时序递归每一步仅需维护一个**固定尺寸的隐状态矩阵 h**无需保存历史全部 KV。* 递推形式每步串行且无法利用 tensor core因此采用** chunkwise 分块并行算法**兼顾效率。随着序列长度增加Flash Attention 的O(L^2d)计算量导致 kernel 逐渐进入 compute-bound而 Linear Attention 的O(Ld^2)复杂度优势开始显现。### 2.2.2 GDN算法简介Linear attention: 每个时间步 t 通过线性投影产生三个列向量: qt, vt, kt则当前的状态St可以表示为其中通过交换律可以推导到原来的公式原始 Linear Attention 在模型效果上远弱于标准 Transformer根本原因是**固定大小的隐状态容量有限导致”记忆碰撞”且缺乏遗忘机制**。为解决这个问题有以下两种改进方向* **Mamba2/GLA**引入标量门控衰减  可以快速擦除全局记忆但不能选择性更新单个键值对。* **DeltaNet**使用 delta 规则精确替换特定键值对但缺乏全局记忆的快速清除机制。**Gated Delta Networks (GDN)** GDN结构如图2所示将两者结合在线性递推的基础上同时引入门控衰减和 delta 更新强度 兼具全局遗忘和定向更新能力但需要用chunkwise算法实现。图2 Gated DeltaNet网络结构示意图### 2.2.3 GDN关键步骤总览**主要流程参考如下**1. 输入投影in_proj_qkv / in_proj_z / in_proj_ba→ 得到 QKV、z、b、a2. 因果卷积causal_conv1d→ 局部特征提取3. 生成 g(alpha)、beta→ 门控参数4. Gated Delta Rule 核心计算循环状态机→ 长序列建模5. 门控归一化 输出投影→ norm(z * out) → out_proj# 三、Qwen3.5-27B模型结构代码走读## 3.1 Qwen3.5-27B模型代码结构总览## 3.2 GatedDeltaNetAttention分析### 3.2.1 GatedDeltaNetAttention初始化self.hidden_size 模型维度 (如 2048 / 4096)self.num_k_heads Q/K 头数self.num_v_heads V 头数 (通常是 K 的整数倍)self.head_k_dim Q/K 每头维度self.head_v_dim V 每头维度self.key_dim num_k_heads * head_k_dimself.value_dim num_v_heads * head_v_dimself.conv_kernel_size 因果卷积核大小### 3.2.2 QKV、z、b、a输入映射if hasattr(self, in_proj_qkv):# LoRA path (Qwen3.5 only): separate in_proj_qkv and in_proj_zmixed_qkv, _ self.in_proj_qkv(hidden_states)ba, _ self.in_proj_ba(hidden_states)z, _ self.in_proj_z(hidden_states)z z.reshape(z.size(0), -1, self.head_v_dim)b, a ba.chunk(2, dim-1)b b.contiguous()a a.contiguous()else:mixed_qkvz, _ self.in_proj_qkvz(hidden_states)ba, _ self.in_proj_ba(hidden_states)if self.gqa_interleaved_layout:# Qwen3-Next: unpack the interleaved GQA layoutquery, key, value, z, b, a self.fix_query_key_value_ordering(mixed_qkvz, ba)query, key, value map(lambda x: rearrange(x, l p d - l (p d)), (query, key, value))mixed_qkv torch.cat((query, key, value), dim-1)else:# Qwen3.5: weights are already in [q, k, v, z] and [b, a] orderqkv_size (self.key_dim * 2 self.value_dim) // self.tp_sizez_size self.value_dim // self.tp_sizemixed_qkv, z mixed_qkvz.split([qkv_size, z_size], dim-1)z z.reshape(z.size(0), -1, self.head_v_dim)b, a ba.chunk(2, dim-1)b b.contiguous()a a.contiguous()* **in_proj_qkv**: 把输入映射成 QKV shape: hidden_size → key_dim key_dim value_dim输入shape[batch_size,seq_len, hidden_size//tp]输出shape[batch_size,seq_len, (key_dimkey_dimvalue_dim)//tp]* **in_proj_z**: 门控信号 shape: hidden_size → value_dim 输入shape[batch_size,seq_len, hidden_size//tp]输出shape[batch_size,seq_len, value_dim//tp]--- before reshape输出shape[batch_size,seq_len, num_v_heads//tp, head_v_dim]--- after reshape* **in_proj_ba**: 循环门控参数输入shape[batch_size,seq_len, hidden_size//tp]输出shape[batch_size,seq_len, 2*num_v_heads//tp]后续每个输入会根据上面输出进行拆解q shape[batch_size,seq_len, key_dim//tp]k shape[batch_size,seq_len, key_dim//tp]v shape[batch_size,seq_len, value_dim//tp]b shape: [batch_size,seq_len, num_v_heads//tp]a shape: [batch_size,seq_len, num_v_heads//tp]### 3.2.3 gdn_attention_core计算整个GatedDeltaNetAttention网络最核心的部分是调用了 gdn_attention_core core_attn_out torch.zeros((num_tokens, self.num_v_heads // self.tp_size, self.head_v_dim),dtypehidden_states.dtype,devicehidden_states.device,)torch.ops.vllm.gdn_attention_core(mixed_qkv,b,a,core_attn_out,_encode_layer_name(self.prefix),)**关键步骤**1. 申请一块全是0的输入内存 vLLM 算子是**in-place 写入**模式必须提前给空间。2. 跑完整的GatedDeltaNet计算结果写入上面空内存**【gdn_attention_core计算核心逻辑】**#### 3.2.3.1 causal_conv1d_update_npu**1API功能介绍****2API入参介绍**decode阶段对应的causal_conv1d备注prefill阶段是调用到了Ascend C固在此不展开描述* KERNEL_WIDTH w 卷积核长度qwen3.5固定width4卷积核大小较小属于small kernel* d通道维度 dim* t时间步 token 位置* xt[d]第 t 个 token、第 d 通道输入* w[d, k]第 d 通道、卷积核第 k 个权重代码 weight: (dim, width)* b[d]偏置 bias* yt[d]卷积原始输出* ot[d]最终输出代码里的 o* 因果规则**只能用 t, t-1, t-2,... 历史不能用未来 t1****3API核心公式**拆开逐段对应代码1. 卷积求和内核循环 j2. SiLU 激活源码 SILU_ACTIVATION最终* **xₜ**当前 token你有* **xₜ₋₁, xₜ₋₂ ...**过去的 token**必须存在 cache 里**decode推理的时候是**一个 token 一个 token 生成**的不可能每次都把整个历史序列重新传一遍。所以**cache_t 保存最近的w个历史 x供下一次卷积直接用**输出是mixed_qkv后续再拆分成q,k,v作为gated delta rule的输入公式如下query_spec, key_spec, value_spec self.rearrange_mixed_qkv(mixed_qkv_spec)query_non_spec, key_non_spec, value_non_spec self.rearrange_mixed_qkv(mixed_qkv_non_spec)#### 3.2.3.2 fused_gdn_gating_patch1API 功能介绍什么时候调用在prefill阶段和decode开启mtp阶段会用到2API 核心公式这里输出是g和这里的就是后面的关于g的操作* Softplus 保证输出**恒正**softplus(x)≥0β是温度系数缩放时间步threshold是溢出保护当βx过大时直接线性截断避免  爆炸。* 负号 - → 保证 exp(g) 永远 ** 1**对于后续GDN作用是让历史状态不断遗忘;* A_log → 每个头固定衰减率因为 beta_output 是**输入门控**必须满足* 范围 **(0,1)*** 对于后续GDN作用**控制当前 token 对状态 h 的贡献强度**也就是**控制**新信息注入强弱****#### **3.2.2.3 chunk_gated_delta_rule****1为什么在prefill阶段要用chunk**把长序列切成大小固定的块 BT64**长序列尤其 Prefill 阶段在 NPU/GPU 上**直接算会极慢、显存爆炸、无法并行且NPU/GPU 不擅长逐 token 循环对于Prefill 长序列无法一次性载入算子**串行依赖**对于每个时间步骤St依赖St-1序列方向完全串行无法利用序列维度的并行性。1. **块内**64 个 token 一次性并行算完内部递归2. **块间**只传递块首尾的状态 S* 块开始时载入上一块的最终状态 S_prev* 块内跑完所有 token 后输出本块最终状态 S_curr* 传给下一个块当初始状态**计算复杂度的下降**对于chunk计算可以写成如下公式这里N代表新项o代表旧项的门控衰减系数均可看成常量进行移项可以得到如下公式得到的因果下三角核如下也就是对应的公式S计算可以通过计算得到。由于这里的A是对应一个chunk大小的矩阵通过等比求和公式可以得到如下公式假设chunk大小是K对于下三角矩阵A^K0,固对于k做截断只需要计算IA...A^k-1即可复杂度降到了O(K^2)总的计算复杂度从如果序列长度是L)O(L^2)降到了O(LK)**为了保持全文的连贯性与可读性后续隐状态Ht均用St来表示****2输入shape介绍**q shape[batch_size,seq_len, num_k_heads//tp, head_k_dim]k shape[batch_size,seq_len, num_k_heads//tp, head_k_dim]v shape[batch_size,seq_len, num_v_heads//tp, head_v_dim]beta shape: [batch_size,seq_len, num_v_heads//tp]g shape: [batch_size,seq_len, num_v_heads//tp]**3调用逻辑**chunk_gated_delta_rule() [入口函数参数校验 格式转换]↓ChunkGatedDeltaRuleFunction.apply()↓forward() [自动求导的前向传播]↓ (可选)l2norm_fwd(q) L2归一化l2norm_fwd(k)↓chunk_gated_delta_rule_fwd() [真正的核心前向计算]↓┌─────────────────────────────────────────┐│ 核心计算流水线分块并行计算 ││ 1. chunk_local_cumsum(g) 门控累积和 ││ 2. chunk_scaled_dot_kkt_fwd() 计算A矩阵 ││ 3. solve_tril() 下三角求解 ││ 4. recompute_w_u_fwd() 计算w, u ││ 5. chunk_gated_delta_rule_fwd_h() 计算h ││ 6. chunk_fwd_o() 计算最终输出o │└─────────────────────────────────────────┘↓返回 o, final_state##### 3.2.2.3.1 l2norm_fwd**1功能介绍**这段代码中主要是通过triton实现L2 归一化实现的公式可以写成如下y x / sqrt( sum(x²) eps )目标**让变成单位向量**作用对于fwd_h里的公式后续会提到有大量的指数衰减项exp)* 如果k很大每次更新的S就会很大导致S指数膨胀* 如果k太大就会导致chunk跨块传递S导致数值爆炸**L2Norm 把 K 固定模长 1**→ K 大小永远不变→ S 只随遗忘 g、递归更新**不会随向量长度膨胀**→ Chunk 分块递归数值完全稳定**2代码介绍**q l2norm_fwd(q)k l2norm_fwd(k)##### 3.2.1.3.2 chunk_gated_delta_rule_fwd**1代码介绍**def chunk_gated_delta_rule_fwd(q: torch.Tensor,k: torch.Tensor,v: torch.Tensor,g: torch.Tensor,beta: torch.Tensor,scale: float,initial_state: torch.Tensor,output_final_state: bool,cu_seqlens: torch.LongTensor | None None,prebuilt_metaNone,):......g chunk_local_cumsum(g,chunk_sizechunk_size,cu_seqlenscu_seqlens,block_indicesblock_indices_cumsum,)# obtain WY representation. u is actually the new v.A chunk_scaled_dot_kkt_fwd(kk,betabeta,g_cumsumg,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,output_dtypetorch.float32,)A solve_tril(AA,cu_seqlenscu_seqlens,chunk_indices_large_blockchunk_indices_large_block,chunk_indices_btchunk_indices_chunk64,output_dtypek.dtype,)w, u recompute_w_u_fwd(kk,vv,betabeta,AA,g_cumsumg,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,)h, v_new, final_state chunk_gated_delta_rule_fwd_h(kk,ww,uu,gg,initial_stateinitial_state,output_final_stateoutput_final_state,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,chunk_offsetschunk_offsets_chunk64,)......o chunk_fwd_o(qq,kk,vv_new,hh,gg,scalescale,cu_seqlenscu_seqlens,chunk_offsetschunk_offsets_chunk64,)......**2入参介绍**这里入参就是上面经过因果卷积ab转换成g和beta还有l2 norm之后的输出q shape[batch_size,seq_len, key_dim//tp]k shape[batch_size,seq_len, key_dim//tp]v shape[batch_size,seq_len, value_dim//tp]beta shape: [batch_size,seq_len, num_v_heads//tp]g shape: [batch_size,seq_len, num_v_heads//tp]**3API调用逻辑**###### 1准备工作forward_context get_forward_context()num_decodes 0chunk_size 64 # 固定分块大小64 token 一块根据seq_len切分###### 2 chunk_local_cumsumg chunk_local_cumsum(g, chunk_size64, ...)对每个 chunk 内的 g 做累计和用于快速计算历史状态衰减这里的 g 必须做 cumsum是为了在下面公式中后面的fwd_h让chunk内的递归公式变成**可并行计算**更详细地原本的GDN状态更新是其中ht必须依赖ht-1,必须串行现在对于g做累计和代入原公式可以得到**展开后St 不再依赖前一个 S_t-1了**只依赖* 初始状态 S0* 累积和 G_t、G_i输入输出shape不变g shape: [batch_size,seq_len, num_v_heads//tp]###### 3 chunk_scaled_dot_kkt_fwdA chunk_scaled_dot_kkt_fwd(k, beta, g_cumsumg, ...)* **chunk_scaled_dot_kkt_fwd_kernel 做一件事**在每个 chunk 内计算带门控、带缩放、带因果掩码的 k·k.T 矩阵公式A[i,j] β[i] * exp(g[i]-g[j]) * k[i]·k[j].T且 i j 才有效A shape:[batch_size,seq_len, num_v_heads//tp, chunk_size]###### 4 solve_tril三角求解 (I A)⁻¹A solve_tril(A, ...)**目的**让后面 state 递推变成 **O (1) 计算**###### 5recompute_w_u_fwd重新计算 w, u新 K, Vw, u recompute_w_u_fwd(k, v, beta, A, g_cumsumg, ...)* **输入**k, v, beta, A* **输出*** w新 key* u新 value* **u A · (v · β)** → 新 value* **w A ** ·** (k · β · exp(g))** → 新 key* **目的**让记忆 state 更新更快、更稳定v 原始 shape [batch_size,seq_len, num_v_heads//tp, head_v_dim] beta 进 kernel 前被转置成 [num_v_heads//tp,batch_size,seq_len] [16, 1, 8192]对于每个batch_size和每个head_num的chunk的v和k进行recompute┌─────────────┐ ┌─────────────┐│ A [64,64] │ │ V [64,128] │└──────┬──────┘ └──────┬──────┘│ ││ beta [64] → [64,1]│ ││ ▼│ V * beta [64,128]│ │└─────────┬─────────┘│▼U A (V*beta)[64,128]chunk A和chunk V和chunk beta做运算得到chunk u最终按照batch size和head num放到最终u不同的位置上| 数据 | shape | 含义 || ----------------- | ------------------------------ | --------------------------------------------- || 单个 chunk 输出 | **[64, 128]** | 一个 chunk 一个 head 的结果 || 最终完整 U | **[1, 8192, 48, 128]** | 全部 batch、全部 token、全部 head、全部维度 |对于k和beta和g的num_k_head和num_v_head不一致的情况这里用了如下方法ptr_k k (bos * Hg i_h // (H // Hg)) * K offs_t_2d * (Hg * K) offs_k * 1举例num_k_heads 48num_v_heads 16这里head的映射关系如下每 3 个 V head 共享 1 个 K headi_h (V head) | 对应的 K head0 → 0 //3 01 → 1 //3 02 → 2 //3 03 → 3 //3 14 → 4 //3 15 → 5 //3 1...45,46,47 → 15类似的w输出的shape是[batch_size,seq_len, num_v_heads//tp, head_v_dim]###### 6chunk_gated_delta_rule_fwd_h这里是在计算隐状态S是一个串行运行各个chunk的函数每个后面的St都和前面的St-1相关chunk内是并行操作最重要的步骤h, v_new, final_state chunk_gated_delta_rule_fwd_h(kk,ww,uu,gg,initial_stateinitial_state,output_final_stateoutput_final_state,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,chunk_offsetschunk_offsets_chunk64,)**输入*** kkey* w优化后的 key* u优化后的 value* g门控累计和控制更新强度* initial_state初始记忆 S₀上文已讲后续h全部用s代替来增加可读性本质不变可以是none这时就会跳过内部如下命令 如果有 shape[real_batch_size, num_v_heads//tp, head_k_dim, head_v_dim],后续每次都有前一次chunk做完的隐状态作为这次的initial stateif USE_INITIAL_STATE:h0_ptr h0 i_nh * K * Vptr_h0_bv1 h0_ptr offs_k * V offs_v1 * 1b_h1_bv1 tl.load(ptr_h0_bv1, maskmask_kv1, other0.0).to(tl.float32)ptr_h0_bv2 h0_ptr offs_k * V offs_v2 * 1b_h1_bv2 tl.load(ptr_h0_bv2, maskmask_kv2, other0.0).to(tl.float32)**输出*** **S** (每个 chunk 的状态)shape: [batch_size,chunk_num, num_v_heads//tp, head_k_dim, head_v_dim]**chunk_num 序列被切成多少个 chunk*** **St (final_state)** (最终状态)shape: [real_batch_size, num_v_heads//tp, head_k_dim, head_v_dim]真正在运行的代码是这一段b_v_new1 b_v1 - tl.dot(b_w, b_h1_bv1)b_v_new1 b_v_new1 * b_gb_h1_bv1 b_h1_bv1 * b_g_lastb_h1_bv1 tl.dot(b_k, b_v_new1)固对于第t个chunk写成数学公式为* St-1:第 t−1 块处理完毕后整条序列到目前为止所有历史信息压缩状态* St:第 t 块处理完之后新的全局累积状态* Kt:当前分块的键* Vt:当前分块的值这个v是前一步骤转换后的u* Wt:可以当成全局固定化权重,是前一步k转换后的w* 增量Vt-W*Ht-1这就是 “delta rule” 名字来源* gt第t个chunk的门控衰减系数指数衰减控制历史信息遗忘* e^gt对历史状态Ht-1做指数衰减 e^gt*Ht-1:遗忘旧记忆* e^(gL−gt)对当前增量做衰减* gL当前 chunk 最后一个 token 的门控也就是chunk内原始g的累加和更准确地公式可以如下###### 7 chunk_fwd_o算最终输出 o。这个函数对于所有chunk可以并行运行根据如下公式可以看到每个chunk直接选取H对应的位置即可公式如下输出由两部分组成**inter-chunk**当前 query 与历史累积状态的交互和 **intra-chunk**当前 chunk 内部 query 和 key 的注意力交互。**inter-chunk**当前 chunk 的 query 与前面所有 chunk 的累积隐状态的交互。对应代码b_o tl.dot(b_q, b_h)b_o b_o * tl.exp(b_g)[:, None]**intra-chunk**当前 chunk 内部注意力对应代码b_A tl.dot(b_q, b_k)b_A b_A * safe_exp(b_g[:, None] - b_g[None, :])b_A tl.where(m_A, b_A, 0)tl.dot(b_A, b_v)mask只保留下三角上三角全部置 0禁止看到未来 token# 门控衰减矩阵 Gamma[i,j] exp(g[i] - g[j])g 是 log-space 累积和# 举例g_cumsum [-0.1, -0.3, -0.5, -0.7]## Gamma # j0 j1 j2 j3# i0 [exp(0) 0 0 0 ] - 对角线无衰减# i1 [exp(-0.2) exp(0) 0 0 ] - token 0 对 token 1 衰减 exp(-0.2)# i2 [exp(-0.4) exp(-0.2) exp(0) 0 ]# i3 [exp(-0.6) exp(-0.4) exp(-0.2) exp(0)] - 越远的过去衰减越多为了更清楚地解释上面3-5步骤对于6-7步骤的逻辑这里引入WY表示这里通过deltanet的公式进行解释gated deltanet是一样的deltanet比较容易理解该公式可以用简单数学归纳法证明。首先定义n表示第n个token。当n1时公式固然成立。假设对于n-1也成立那么我们证明对于n也成立证明过程如下该证明不仅证明了公式的正确性也提供了w_n的计算公式通过Sn的递推公式我们可以证明通过归纳法那么对于在一个chunk i内部第r个位置我们通过递归到S[i]会有如下公式w和u在这里都是通过WY表示计算的但是从每一个块的第一个位置开始不是从序列起始位置开始从而第r个位置有如下w和u对于输出计算结合矩阵乘法形式可以得到最终公式#### 3.2.1.4 recurrent gated delta rule开启投机推理会运行该算子替代原先的fused_sigmoid_gating_delta_ruleRecurrent 模式是直接按时间步逐个计算递推过程适用于推理解码场景逐 token 生成。和prefill阶段的区别是把chunk_fwd_h和chunk_fwd_o都合并在一步里公式如下是对于逐个token递归这个和原生的gated delta rule是等价的门控衰减对应保证了 delta 项中擦除的是衰减后的旧值。代码实现如下b_h * exp(b_g) # h alpha * hb_v b_beta * (b_v - tl.sum(b_h * b_k[:, None], 0)) # v_new beta * (v - alpha*S k)b_h b_k[:, None] * b_v # h k * v_new^T另外在 kernel 中也做了GVA 的支持Q/K 每时间步前进 num_k_heads*head_k_dim//tpV/O 每时间步前进 num_v_heads*head_v_dim//tp。g 的步长也是 num_v_heads//tp或者是num_v_heads*head_k_dim//tp取决于是否是KDAKDA 模式**一个头K 个 g每个 key 维度一个**beta 的步长取决于是否为 headwise 模式。p_q H * Kp_k H * Kp_o HV * Vp_v HV * Vif not IS_KDA:p_g HVelse:p_gk HV * Kp_beta HV * (V if IS_BETA_HEADWISE else 1)#### 3.2.1.5 fused sigmoid gating delta rule和fused recurrent gated delta rule的区别是输入sigmoid输入是a,b recurrent输入是g和beta公式和fused recurrent gated delta rule一样只是在运行这段代码前又做了a,b到g,beta的转换可参考fused_gdn_gating_patch## 3.3 Qwen3.5-27B Profiling算子分析| **profiling算子名称** | **对应代码调用位置** | **代码路径** || -------------------------------------------------------- | -------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- | ---------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- || _causal_conv1d_update_kernel_npu_tiled_5 | _causal_conv1d_update_kernel_npu_tiled[grid]](x,weight,bias,conv_state,conv_state_indices,num_accepted_tokens,query_start_loc,block_idx_last_scheduled_token,initial_state_idx,out,batch,dim,seqlen,...) | [vllm-ascend/vllm\_ascend/ops/triton/mamba/causal\_conv1d.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/mamba/causal_conv1d.py#L527) || l2norm_fwd_kernel2_loop | q l2norm_fwd(q)k l2norm_fwd(k) | [vllm-ascend/vllm\_ascend/ops/triton/fla/l2norm.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/l2norm.py#L34) || chunk_local_cumsum_scalar_kernel | chunk_local_cumsum(g,chunk_sizechunk_size,cu_seqlenscu_seqlens,block_indicesblock_indices_cumsum,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/cumsum.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/cumsum.py#L116) || chunk_scaled_dot_kkt_fwd_kernel | chunk_scaled_dot_kkt_fwd(kk,betabeta,g_cumsumg,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,output_dtypetorch.float32,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/chunk\_scaled\_dot\_kkt.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/chunk_scaled_dot_kkt.py#L83) || solve_tril_16x16_kernel | A solve_tril(AA,cu_seqlenscu_seqlens,chunk_indices_large_blockchunk_indices_large_block,chunk_indices_btchunk_indices_chunk64,output_dtypek.dtype,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/solve\_tril.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/solve_tril.py#L330) || recompute_w_u_fwd_kernel | w, u recompute_w_u_fwd(kk,vv,betabeta,AA,g_cumsumg,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/wy\_fast.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/wy_fast.py#L98) || chunk_gated_delta_rule_fwd_kernel_h_blockdim64 | h, v_new, final_state chunk_gated_delta_rule_fwd_h(kk,ww,uu,gg,initial_stateinitial_state,output_final_stateoutput_final_state,cu_seqlenscu_seqlens,chunk_indiceschunk_indices_chunk64,chunk_offsetschunk_offsets_chunk64,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/chunk\_delta\_h.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/chunk_delta_h.py#L179) || chunk_fwd_kernel_o | o chunk_fwd_o(qq,kk,vv_new,hh,gg,scalescale,cu_seqlenscu_seqlens,chunk_offsetschunk_offsets_chunk64,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/chunk\_o.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/chunk_o.py#L112) || fused_recurrent_gated_delta_rule_fwd_kernel_11 | core_attn_out_non_spec, last_recurrent_state fused_recurrent_gated_delta_rule(qquery_non_spec,kkey_non_spec,vvalue_non_spec,gg_non_spec,betabeta_non_spec,initial_statessm_state,inplace_final_stateTrue,cu_seqlensnon_spec_query_start_loc[: attn_metadata.num_decodes 1],ssm_state_indicesnon_spec_state_indices_tensor,use_qk_l2norm_in_kernelTrue,) | [vllm/vllm/model\_executor/layers/fla/ops/fused\_recurrent.py at main · vllm-project/vllm (github.com)](https://github.com/vllm-project/vllm/blob/main/vllm/model_executor/layers/fla/ops/fused_recurrent.py#L481) || fused_sigmoid_gating_delta_rule_update_kernel_0 | core_attn_out_non_spec fused_sigmoid_gating_delta_rule_update(A_logself.A_log.contiguous(),dt_biasself.dt_bias.contiguous(),qquery_non_spec.contiguous(),kkey_non_spec.contiguous(),vvalue_non_spec.contiguous(),aa.contiguous(),bb.contiguous(),initial_state_sourcessm_state,initial_state_indicesnon_spec_state_indices_tensor,cu_seqlensnon_spec_query_start_loc,use_qk_l2norm_in_kernelTrue,softplus_beta1.0,softplus_threshold20.0,) | [vllm-ascend/vllm\_ascend/ops/triton/fla/sigmoid\_gating.py at main · vllm-project/vllm-ascend (github.com)](https://github.com/vllm-project/vllm-ascend/blob/main/vllm_ascend/ops/triton/fla/sigmoid_gating.py#L180) |# 四、总结本文总结了 Qwen3.5-27B 模型结构并针对 GDN 网络结构做了完整拆解与原理分析。GDN 模块预填充Prefill阶段最核心算子为chunk_gated_delta_rule_fwd_kernel_h和chunk_fwd_kernel_o自回归解码Decode阶段开启 MTP 投机解码时核心算子为fused_recurrent_gated_delta_rule关闭 MTP 常规单 Token 解码时核心算子为fused_sigmoid_gating_delta_rule。在 Prefill 分块 Chunk 计算流程中矩阵 A 构建、门控 g/beta 计算、Q/K 归一化、初始状态加载等全部前置步骤均是为后续**fwd_h 隐状态计算**与**fwd_o 输出计算**两个核心步骤提供输入支撑。其中chunk_gated_delta_rule_fwd_kernel_h负责完成 GDN 内部时序隐状态的分块并行求解chunk_fwd_kernel_o基于求解完成的隐状态进一步计算注意力最终输出。同时本文梳理了投机解码模式下的 GDN 执行逻辑投机草稿 Token 分支与已验证正常 Token 分支分开计算、分开执行因果卷积与时序递归仅在输出端按照位置索引合并结果既保证投机解码加速能力又保证模型因果正确性。整体 GDN 架构采用**Causal Conv1D 局部时序建模 线性注意力全局长依赖**的互补设计通过 Chunk 并行预填充、Recurrent 串行解码实现线性复杂度 O (N) 超长上下文推理能力相比传统 Transformer 注意力具备显著效率优势。
