DeepInfer 源码拆解（三十九）：一个 TileLang 入口里的多段 Kernel——前缀和、PDL 与段间依赖的发射编排

传统 stream 的串行

图源：NVIDIA CUDA Programming Guide 4.5.1 “GPU activity timeline”。

在同一个 CUDA stream 里，primary_kernel 和 secondary_kernel 的执行区间完全不重叠。中间的空白同时包含两段时间：secondary_kernel 的 launch 延迟（驱动把 grid 配置、参数打包、下发到硬件队列的时间，通常几微秒量级）和 driver/硬件对 primary 完成态与显存写回可见性的等待。对于 decode 阶段一个 step 只有几十个 kernel、每个 kernel 本身可能只有一两微秒的场景，这段空白占比不小。

preamble 可以提前跑

图源：NVIDIA CUDA Programming Guide 4.5.1 “Preamble section of secondary_kernel”。

几乎每个 kernel 都有一段 preamble：读常驻参数到寄存器、算线程在全局中的位置、为共享内存清零、初始化归约器。这些工作和 primary 写入的结果无关。如果 secondary 的 block 能在 primary 结束之前就被调度起来、先把 preamble 跑完，那么当 primary 真正结束时，secondary 只剩下真正依赖数据的部分要跑。

图源：NVIDIA CUDA Programming Guide 4.5.1 “Concurrent execution of primary_kernel and secondary_kernel”。

这张图里的关键是重叠区。它既包含 secondary 的 preamble，也可能包含 secondary 在同步点之后访问不依赖 primary 结果的数据的那部分——只要这些访问在显式同步点之后没写，就可以提前做。PDL 的规则里明确允许这样做，但需要程序员自己保证不越界。

三条规则

按 CUDA Programming Guide 4.5.2 的描述，PDL 由三件事构成：

规则一：primary 显式触发。 primary kernel 在”它的结果已经对 secondary 可见”的时刻，所有 block 都应执行 cudaTriggerProgrammaticLaunchCompletion()。若 primary 不触发，则它所有 block 退出时会隐式触发。

规则二：secondary 用扩展 launch API 启动。 secondary 通过 cudaLaunchKernelEx 带上 cudaLaunchAttributeProgrammaticStreamSerialization = 1 启动。带上这个属性后，驱动可以不等待 primary 完成与显存写回就提前调度 secondary，条件是 primary 的所有 block 都已完成触发。

规则三：secondary 必须自己同步。 即便提前调度了，secondary 的 block 也可能在 primary 写出的数据可见之前就开始执行。所以 secondary 要在访问 primary 产物的地方调用 cudaGridDependencySynchronize()（或其它等价机制）。

文档里还有一句容易被忽略的话：

these methods provide the opportunity for the primary and secondary kernels to execute concurrently, however this behavior is opportunistic and not guaranteed to lead to concurrent kernel execution. Reliance on concurrent execution in this manner is unsafe and can lead to deadlock.

即重叠是机会性的，不能把它当保证来用；同步点之前提前跑的逻辑不能假设自己一定能拿到资源。

这套语义在 TileLang 里的入口

TileLang 把上面三件事分别映射成三个用户可见的对象：

• T.pdl_trigger() → codegen 成 cudaTriggerProgrammaticLaunchCompletion()（src/cuda/codegen/codegen_cuda.cc 中对 tl::pdl_trigger() 的处理）；
• T.pdl_sync() → codegen 成 cudaGridDependencySynchronize()；
• 带 pdl_sync 的 prim_func 被标记上 kHasGridSync，运行时把这个标记翻译成 launch 参数 kUseProgramaticDependentLaunch（src/cuda/codegen/rt_mod_cuda.cc），最终在 cudaLaunchKernelEx 里打开 cudaLaunchAttributeProgrammaticStreamSerialization。

值得注意的一个细节：只有含 pdl_sync 的 kernel 才会被以 PDL 方式发射。pdl_trigger 单独出现只说明”我准备好让别人提前起来”，但对方是否提前起，取决于对方自己的 launch 属性。这个不对称在编译器里还有一条硬约束——lower_pdl.cc 里写着：

NVCC has issues with __ldg when using PDL synchronization. Suppress the annotation when kHasGridSync is set. Cannot invoke __ldg function with pdl_sync

也就是说，开启了 pdl_sync 的 kernel 里不能用 __ldg（NVCC 的已知问题），同时 __restrict__ 参数限定符也会被一并抑制。工程上这意味着：给一个 kernel 加 pdl_sync，不只是加一行同步，还会改变它的编译产物——寄存器级只读缓存、指针别名优化都会受影响。这条约束在写性能敏感 kernel 时必须记着。

还有一条语义上的隐含规则：pdl_sync 同步的对象是同一 stream 里最近一次发射的、允许 programmatic 依赖的 kernel，而不是”任意前序 kernel”。在 TileLang 里由于多段 T.Kernel 会按顺序发射到同一 stream，实际效果就是”等上一段”。如果中间夹了一个普通 kernel（不带 PDL 属性），那么中间的普通 kernel 会严格等前一段完全结束，后一段的 pdl_sync 等的则是这个普通 kernel——整条链的 PDL 收益就被中间那一段吃掉了。这是多段组合时最常见的”看起来用了 PDL 但没快”的原因。

第一段：每个”前缀行”扫一遍 batch

1
T.assume(batch_size >= 0)
2
with T.Kernel(num_rows, threads=_THREADS) as row:
3
    tid = T.get_thread_binding()
4
    T.pdl_sync()
5
    running_total = T.alloc_var(T.int32, 0)
6
    for chunk in T.serial(num_chunks):
7
        vals = T.alloc_shared((_THREADS,), T.int32)
8
        for j in T.Parallel(_THREADS):
9
            elem = chunk * _THREADS + j
10
            if elem < batch_size:
11
                sq = cu_seqlens_q[elem + 1] - cu_seqlens_q[elem]
12
                sk = cu_seqlens_k[elem + 1] - cu_seqlens_k[elem]
13
                if row == num_ratios:
14
                    vals[j] = 2 * ((sq + cp2 - 1) // cp2)
15
                _ = [_write_ratio_len(vals, row, j, ri, ratio, sk) for ri, ratio in ratio_items]
16
            else:
17
                vals[j] = 0
18
        T.sync_threads()
19
        T.cumsum(vals, dim=0)
20
        T.sync_threads()
21
        ...

这段在做什么？

• grid 布局：num_rows 个 block，每 block 对应一张输出表的”行”。行的含义由 row 的下标决定——row < num_ratios 是各个压缩比对应的 cu_lens，row == num_ratios 是 CP mirror 索引表自己的前缀偏移。于是所有前缀和在同一段 kernel 里、不同 block 之间并行完成；行与行之间没有依赖。
• 每行内部：按 _THREADS 个元素一块切 batch。每块先用 T.Parallel 把这一块的”贡献”算进共享内存 vals，再 T.sync_threads() 保证可见，然后对共享内存做 T.cumsum（块内并行扫描），再把”不含自身的累计和”写进 cu_lens，最后把块总量叠加到 running_total。
• 关键点：块内用并行扫描，块间串行累加。 running_total 是一个变量级的状态，每个 block 都有一份，块间的串行走访是 num_chunks 次 T.cumsum + 一次标量加法。cu_seqlens 的 batch 一般不大（几十到几百条），这部分的代价远比”一个 block 串行扫一遍”低，又远比”整个前缀和做成跨 block 的两遍扫描（decoupled look-back）“简单。这是一个非常工程化的取舍：用”块内并行、块间串行”的折中，换掉一次 kernel 的复杂度。
• vals 被复用：注意 _write_ratio_len 会把同一份 vals[j] 写成不同压缩比的贡献值，然后一起做 T.cumsum——所以一行 cu_lens 存的是某个压缩比下、各序列起始偏移的独立前缀表，它们在 row 之间互不干扰。

cumsum 的细节也值得抠一下。TileLang 的 T.cumsum(vals, dim=0) 对共享内存里的 _THREADS 个元素做 inclusive scan，结果是 vals[j] 变成 vals[0] + ... + vals[j]。然后代码写：

1
if j == 0:
2
    cu_lens[row, elem] = running_total
3
else:
4
    cu_lens[row, elem] = running_total + vals[j - 1]

这两行得到的是 exclusive 前缀和（含 vals[j-1]、不含 vals[j]），而 running_total 加上 vals[_THREADS-1]，就是这一块的总量。最后 if tid == 0: cu_lens[row, batch_size] = running_total 补上末尾的完整总量，让 cu_lens 拥有 batch_size + 1 个元素，和 cu_seqlens 的布局对齐。

为什么第一段自己也有 T.pdl_sync()

第一段开头也写了 T.pdl_sync()。前缀和的第一段本身不依赖更早的 kernel 产物吗？从源码看，它需要读取 cu_seqlens_q、cu_seqlens_k，这两者在 DeepInfer 的元数据构造流水线里由上游 kernel（或者对 decode 的 prepare_v4_decode_metadata 路径）写入。把 pdl_sync 放在 kernel 最前面有两种效果：

1. 语义上：保证读到的是完整写回的数据——即使上游 kernel 已经在更早的时刻触发了 launch completion，这个读也可能发生在显存写回完成之前，pdl_sync 会等”上游所有被依赖的 primary 都已写完”。
2. 性能上：pdl_sync 之前的指令（本例中只是 T.get_thread_binding()）可以提前跑，相当于白送的一小段 preamble。在这个 kernel 里 preamble 本身很轻，收益不大；但对那些 preamble 重（比如要先清零一大块共享内存）的 kernel，收益就很明显。

顺带说一句：第一段自己的 pdl_sync 也是它被标记 kHasGridSync、从而以 PDL 方式启动的原因。如果想让一个 kernel 成为 PDL 链里的”被动方”，就必须在里面放一个 pdl_sync；只有 pdl_trigger 是不够的。这是一个非常容易踩的坑：有人以为只要上游加了 pdl_trigger，下游就会自动提前起来——不是的。

第一段末尾没有 pdl_trigger 会怎样

第一段代码中没有调用 T.pdl_trigger()。按 NVIDIA 文档的规则一，不触发就意味着隐式触发在”所有 block 退出时”发生。结果是第二段仍然可以用 PDL 启动，但它能提前覆盖的只有 launch 延迟本身——因为触发时间点实际被推到了和 kernel 完全结束几乎同时，下游 block 拿不到任何与上游后续工作重叠的机会。也就是说，pdl_trigger 是否显式写出，决定了重叠窗口的大小：

• 显式触发在”数据写好之后、剩余工作之前” → 重叠窗口大；
• 隐式触发 → 重叠窗口基本为零，只剩 launch 延迟。

那么对前缀和来说，能不能显式触发？理论上可以——把触发放在”写完 cu_lens 的最后一块之后”就行。但注意这段代码的结构：with T.Kernel 的整个 body 在每个 block 里各执行一遍，T.pdl_trigger() 写在哪里就意味着每个 block 都在那里调用一次。文档要求”所有 block 执行了触发才放行”，所以这里必须保证每个 block 都能到达触发点。如果把触发放进只在部分线程执行的分支里，就只有触及该分支的那些 block 会显式触发，其余 block 只能靠退出时的隐式触发兜底，显式的收益被稀释。这也是为什么在 fused_sample.py 里，触发被写在”最后一块的 finalize_reducer 之前”，而且写在所有线程都会执行的路径上：

1
if i0 == N_LSE_CHUNKS - 1:
2
    T.pdl_trigger()
3
T.finalize_reducer(blk_sum)

注释写得很直接：Release the next launch before the final block reduction.——把触发放在最后一次块内归约之前，可以让下游的 block 在归约期间就开始跑 preamble。

如果想让前缀和也受益于这种”提前触发”，可以把触发写在”每个 block 已经把自己那行写完全部 num_chunks 块”之后，也就是 cu_lens[row, batch_size] = running_total 之后、但在函数末尾。DeepInfer 这里没有这么做，原因值得猜一下：前缀和 kernel 本身极短，触发能省的时间可能还抵不上一次额外的触发指令；而且第二段的依赖是全部行的 cu_lens 都写完，不是单行写完——某个 block 提前触发对其它行没意义。每行并行、块内串行的设计让”是否触发”这件事从收益上接近零，所以不写是最佳工程选择。

第二段：并行散列

1
# The output-parallel scatter lets long individual sequences occupy
2
# many blocks instead of being serialized inside one sequence block.
3
with T.Kernel(nb_rev + nb_out, threads=_THREADS) as bx:
4
    tid = T.get_thread_binding()
5
    # Depends on PDL: T.pdl_sync() below waits for the prefix kernel to finish writing cu_lens.
6
    T.pdl_sync()
7
    if bx < nb_rev:
8
        t = bx * _THREADS + tid
9
        if t < total_q:
10
            _fill_reverse(cu_seqlens_q, cu_lens, mirror_cp_reverse_index, t)
11
    else:
12
        out_pos = (bx - nb_rev) * _THREADS + tid
13
        if out_pos < output_len:
14
            _fill_out(cu_seqlens_q, cu_seqlens_k, cu_lens, mirror_cp_index, ks_out, ke_out, out_pos, cp_rank)

第二段是典型的”以输出元素为并行单位”。grid 是两个任务段的拼接：前 nb_rev 个 block 负责填 mirror_cp_reverse_index（长度 total_q），后 nb_out 个 block 负责填 mirror_cp_index 与各压缩比的 ks/ke（长度 output_len）。每个 block 处理 _THREADS 个元素。

注释 The output-parallel scatter lets long individual sequences occupy many blocks instead of being serialized inside one sequence block 点出了一个很关键的并行策略选择：如果按序列切 block（一个 block 处理一个序列），长序列会拖垮这个 block，而短序列的 block 会空转；按全局输出位置切 block，长序列自然被切散到多个 block 上，负载更均匀。

具体到 _fill_reverse 的逻辑：给定全局 token 下标 t，用二分查找定位它属于哪个序列 seg（cu_seqlens_q[seg] <= t < cu_seqlens_q[seg+1]），然后算出这个 token 在 zigzag 切分下的源 rank 与局部下标，结合 cu_lens 里对应 rank 和序列的起始偏移，得到 mirror_cp_reverse_index[t]。_fill_out 的结构相似，只是它处理的是”输出位置”而不是”全局 token 位置”：每个输出位置要找到它对应的源位置，再把 ks/ke 的偏移（delta）写进对应的 ks_out/ke_out 行。

注意这两者都需要 cu_lens——这正是 PDL 依赖的来源：cu_lens 由第一段写、第二段读，而且第一段每行的 cu_lens 是串行累加出来的，不可能在第二段开始前就知道。PDL 把这段依赖表达成”我（第二段）被启动之前，你可以（第一段）先触发，但我在 pdl_sync 之后才读 cu_lens”。

把 prepare_cp_mirror 的两段依赖画成时间线，会更容易看清 pdl_sync 卡在哪里：

图中第一段没有显式 T.pdl_trigger()，因此触发点被推到每个 CTA 退出时——重叠窗口很小；第二段在 T.pdl_sync() 之前只做地址与常量计算，之后才读 cu_lens。这张图是按本文对源码的拆解画出的结构示意，不是论文原图。

段间依赖不是数据竞争，但需要显式同步

这里有一个容易误读的点。PDL 下的第二段可能在第一段的某些 block 还没写完全部行时就被调度起来了——触发在 block 级别，文档要求”所有 block 都触发”，但”触发”并不等于”写完”。所以第二段在 pdl_sync 之后读 cu_lens 是安全的，因为 pdl_sync 会等所有被依赖的 primary 的写回可见。但如果有人把 pdl_sync 去掉，或者把它放在读 cu_lens 之后，就会出现典型的”读到半截 cu_lens”的竞态——这种 bug 在单卡小 batch 的开发环境里可能一辈子都不触发，但在大 batch 上概率会显著上升。

再看第一段自己的 pdl_sync：它保证读 cu_seqlens_q/k 之前这些输入已写回。两个同步点夹住一段完全独立、对 cu_lens 没有依赖的计算，恰好就是 PDL 设计的目标形态。

一个旁证是 TileLang 自带的 PDL 测试（testing/python/language/test_tilelang_language_pdl.py）。其中 kernels_with_pdl_pipeline 和 kernels_with_pdl_overlap_window 都在同一个 @T.prim_func 里写了两个 with T.Kernel(...) 段，前一段以 T.pdl_trigger() 结尾，后一段以 T.pdl_sync() 开头，再配一堆”前置/后置工作”来制造重叠窗口。这说明”一个入口里多段 T.Kernel、用 PDL 串起来”是 TileLang 一等公民的写法，而不是 DeepInfer 自己绕出来的习惯用法。DeepInfer 的 prepare_cp_mirror_tl 只是把这套写法用到了真实推理负载上。

情形一：前置依赖段 + 依赖段（prepare_cp_mirror / yoco_cp_boundary_fused）

就是上面拆的那种。段之间是显式的数据依赖，靠 pdl_sync 建立顺序，靠 pdl_trigger 扩大重叠窗口。

情形二：存储段 + 计算段（compress）

compress.py 里的 compress_and_update_state_prefill_tl 有 4 个 T.Kernel。第一段的前缀是：

1
# Read the previous chunk's tail before store overwrites state_cache.
2
with T.Kernel(total_c, NUM_DIM_BLOCKS, threads=compute_threads) as (bx, bdim):
3
    ...

后面几段分别负责”把当前 chunk 的尾部写进环形 state_cache”和”计算本 chunk 的新压缩 KV”。注释那句话是核心：计算段要在存储段覆盖 state_cache 之前读取旧值。如果顺序搞反了，压缩窗口里”跨 chunk 边界的那部分”就会读到已经被新 chunk 覆盖掉的旧 token——压缩结果错一位。这种情形里段间依赖也是 pdl_sync，但它的重点不是”加速”，而是用顺序保证正确性：把”先读旧、后写新”的时序写进同一个 prim_func，比在 Python 里调两次 launch 再靠变量命名表达更强、更不容易被后续重构改坏。

compress_and_update_state_decode_tl 里则反过来——先存后算，然后用 if (klen % COMPRESS_RATIO) != 0: kv_compressed[idx, ...] = 0 处理”压缩窗口不满”的情况。两个函数的方向相反，正是 prefill 与 decode 对”窗口对齐”的不同需求。

情形三：两种线程映射（compress）

compress.py 的 prefill 版本里，计算段用 compute_threads、存储段用 store_threads，而且两段的 grid 形状也不一样。多段的另一种动机是同一批数据要经过两个不同的线程切分策略：第一段按 token × 维度块并行做 per-token 计算，第二段按维度 tile 并行做 store——两段的线程映射不一样，合到一起只会互相掣肘。把每段写成一个 T.Kernel，让 TileLang 为每段独立选线程数、独立做布局推断。

情形四：两阶段采样（fused_sample）

fused_sample.py 里的 fused_sample_tl 是一个入口两段 T.Kernel：第一段算每行的 LSE（log-sum-exp），第二段用 LSE 做 top-p/top-k 与采样。两段之间有一个显式 workspace：

1
# Per-row (-log(sum_exp), -max_logit), shared between the two launches.
2
lse_workspace: T.Tensor((total_q, 2), T.float),

这其实是把中间结果显式写回显存，是 PDL 链里唯一的跨段通道。第一段在两个 LSE 值算完后调用 T.pdl_trigger()，位置刻意选在最后一次 finalize_reducer 之前；第二段在 preamble 里先把温度、top-p/top-k 的常量、共享内存归约器都准备好，直到真正要读 lse_workspace 时才 T.pdl_sync()。两段还各自写了 T.annotate_min_blocks_per_sm(3) 和 (4)：这是一个 occupancy 提示，让编译器在寄存器/共享内存预算内尽量保证每个 SM 能驻留这么多 block，机会性并行下下游 block 才有位置和上游 block 同驻。

情形五：初始化段 + 原子更新段（finalize_verification）

sampler.py 里的 finalize_verification_tl 更极端：第一段只有一个 block、32 个线程，干的事就是清零三个统计量：

1
# A separate launch initializes the counters before any atomic updates,
2
# without a cooperative grid barrier that would prevent PDL launches.
3
with T.Kernel(1, threads=32):
4
    T.pdl_sync()
5
    T.clear(num_generated_tokens)
6
    T.clear(num_verified_tokens)
7
    if T.get_thread_binding() == 0:
8
        batch_confidence_error[0] = 0.0

注释把设计动机说得很白：如果不用一个独立的 launch 清零，而是让主 kernel 的每个 block 自己判断”我是不是第一个”，就得用 cooperative grid barrier 来保证所有 block 看到清零后的值——而 cooperative launch 要求所有 block 同时驻留，这与 PDL 的机会性并行在实现上难以共存。用一次极小的前置 launch 换掉一次 grid barrier，是这段代码最核心的取舍。

第二段才是真正的 per-row 归约与写回。注意第一段的 T.Kernel(1, threads=32) 只有 1 个 block：这既是”清零”这种全局任务的自然形态，也让它的触发（隐式，因为没有写 pdl_trigger）几乎无延迟。

元数据流：几个入口串成一条链

以 decode 路径为例，一个 step 的元数据要经过几步准备：

• compute_seqlens_k_zos_tl（compute_seqlens_k_zos.py）更新 seqlens_k：读 seqlens_q 与上一轮接受的 token 数，写出新的 KV 长度；
• prepare_v4_decode_metadata_tl（prepare_decode_metadata.py）按 seqlens_k、cu_seqlens_q 逐序列生成 positions、slot_mapping_store_swa、kv_swa_index_offset_len、各压缩比的 klen 与 positions_c/slot_mapping_store_ca；
• load_kv_index_tl（get_kv_index.py）和 kv_index_cat_tl（kv_index_cat.py）把这些 (offset, len) 对展开成内核要用的整行下标；
• convert_kv_index_tl（flash_mla_utils.py）在稀疏注意力之前把逻辑 token 下标换成物理槽位。

这几个入口之间每一个都是”后一个必须看到前一个写完的结果”。它们在每个入口的开头都写了 T.pdl_sync()：

1
with T.Kernel(T.ceildiv(n_seq, BLOCK_SIZE), threads=BLOCK_SIZE) as bx, T.Parallel(BLOCK_SIZE) as tx:
2
    sk = T.alloc_var("int32")
3
    seq_idx = bx * BLOCK_SIZE + tx
4
    T.pdl_sync()
5
    if seq_idx < n_seq:
6
        sk = seqlens_k[seq_idx] + seqlens_q[seq_idx]
7
        prev_idx = remap[seq_idx]
8
        if prev_idx != -1:
9
            sk = sk - nextn + prev_num_accepted_tokens[prev_idx]
10
        seqlens_k[seq_idx] = sk

和 prepare_cp_mirror 的多段写法对照：多段写法把两次 launch 压成一次调用，单段写法保留独立的入口但让 GPU 侧不必等前一个 kernel 完全退出。两者的取舍不同：

• 压成一次调用省的是 host 侧的 Python 调用和参数打包，代价是 grid 里出现的量不能在两个段之间各自自由地动态（见上一节）；
• 保留多个入口省的是”改一个算子不用动另一个”，代价是每个入口都要自己开头写 pdl_sync，依赖关系只体现在”谁先被调度器调用”这个隐含顺序里。

DeepInfer 在元数据链上选了后者。这符合它把元数据和计算分开设计的整体思路——这些入口的调用顺序由 C++ 侧调度器决定并固定下来，GPU 侧只需要”在每条边上尽量提前起”。

T.pdl_sync() 写在 Kernel 第一行意味着什么

上面 compute_seqlens_k_zos_tl 的写法在 DeepInfer 里极其普遍：with T.Kernel(...) 之后前三行是 T.alloc_var、算 seq_idx、然后 T.pdl_sync()。这看起来像”没有 preamble”，但其实不是：

• T.alloc_var 和索引计算在 lowering 之后大多会变成寄存器里的常量表达式，或者只需要读 kernel 参数里的标量；
• 真正有 preamble 价值的往往是共享内存的清零、归约器的初始化、常量表的装载，这些如果写在 pdl_sync 之前，就能在上游还在跑的时候完成。

fused_sample.py 的采样段把这一点做得很满：温度、top_p/top_k 激活判断、hist/topk_hist/ps_s/ks_s 这几块共享内存的填充，全部放在 T.pdl_sync() 之前：

1
T.fill(ps_s, ps)
2
T.fill(ks_s, ks)
3
T.fill(desired_s, 0)
4

5
# LSE's trigger follows its input sync, so the preamble can read inputs.
6
T.pdl_sync()
7
lse0 = T.alloc_var(T.float, lse_workspace[row, 0])

这几行共享内存写入是每行都要做的常数准备，和 LSE 无关，正好可以藏进上游 LSE 段的尾部。相比 compute_seqlens_k_zos_tl 这种”sync 放第一行”的写法，采样段的 PDL 收益要实打实得多。

什么时候不该开 PDL

把视角放到整个 decode step 上会看到一个容易忽略的问题：pdl_sync 会给 kernel 带上 kHasGridSync 属性，进而禁掉 __ldg/__restrict__（后面工程清单第 5 条）。所以对一个计算很重、但只读上游的一小部分元数据的 kernel（比如注意力本身），开 PDL 反而可能因为丢了只读缓存路径而变慢。

DeepInfer 的做法是分层：元数据构造、采样、下标展开这类轻 kernel 大量用 pdl_sync；真正重的注意力与 GEMM kernel 交给 FlashMLA、DeepGEMM 这些底层库自己决定是否开 PDL，MoE 的跨卡通信则交给 DeepEP，而不是在 TileLang 层强行套一层。也就是说，pdl_sync 的覆盖面在 DeepInfer 里是”薄计算、多依赖”的算子集合，而不是全站默认。这也解释了为什么 compress.py 里虽然有两段，但第一段（total_c, NUM_DIM_BLOCKS 的二维 grid）和后面几段的线程数、grid 形状都差别很大——它们不是在”为了节省 launch”而合并，而是因为必须共享一个由 TileLang 统一管理的编译单元。

1. 段间的依赖要画清楚：哪些是”必须等”、哪些是”等也无所谓”

每个 T.pdl_sync 都是对”必须等”的声明。如果两段之间没有真正的依赖（比如第二段只读输入、第一段也只读输入），那么加 pdl_sync 不会错，但会错失性能：本来可以把第二段的整个执行都藏在第一段背后。判断标准只有一个——第二段是否在同步点之后读取第一段写的数据；如果不是，同步点就应该尽量往后放（或者干脆不放）。

2. 触发点要选在”下游需要的数据已写好”之后、但”剩余工作越多越好”之前

fused_sample.py 里把 pdl_trigger 放在最后一次 finalize_reducer 之前，是这条的教科书用法：LSE 的值已经在 lse_workspace 里写好，剩下的只是块内归约，于是下游采样 kernel 可以在这段归约期间就开始准备常量与共享内存。prepare_cp_mirror 里不显式触发，是因为触发能省的时间极少。在写自己的 kernel 时，应该看每个 block 的 body 里”写下游产物”之后还剩多少工作——如果剩余量远小于触发带来的前置量，那就不值得。

3. 触发必须是”所有 block 都会到达的路径”

把触发写在 if tid == 0 或 if bx < nb_rev 之类的分支里、且只让部分 block 到达，会让”所有 block 都触发”这条条件无法满足，显式触发实际不生效，结果等价于只靠隐式触发——看起来加了 PDL，实际上没重叠。所以 kv_index_cat.py 里两个分支都写了：

1
if i == pipeline_r - 1:
2
    T.pdl_trigger()

这行在两个分支内各出现一次，且都不在 if tid 之下——每个 block 的最后一次流水迭代都会触发。

4. pdl_sync 的放置位置决定 preamble 的大小

gather.py 里把 T.pdl_sync() 放在循环开始前、而在分配 index/values fragment 之后：

1
index = T.alloc_fragment([buf_m], index_dtype)
2
values = T.alloc_fragment([buf_m, buf_n], dtype)
3
sm = bm * BM
4
sd = bd * BD
5
T.pdl_sync()

fragment 的分配本身是编译期的事，没有运行时开销；但更一般的原则是：把不依赖上游结果的地址计算、常量装载放在 pdl_sync 前面。TileLang 的 T.pdl_sync() 是一个普通表达式，编译器不会自动帮你把工作搬过去——所以必须在 DSL 里显式地排好顺序。

5. 一个 kernel 里有 T.pdl_sync，就不能用 __ldg 和 __restrict__

这是 lower_pdl.cc 的硬约束。如果你的 kernel 在开启 PDL 前后性能变化异常（尤其变差），第一件要查的就是”是否靠 __ldg/__restrict__ 拿到了额外的只读缓存或别名优化”。在 TileLang 里，这体现为 kHasGridSync 属性被传给 codegen，导致 PrintFunctionSignature 抑制 __restrict__。若想保留这些优化，可以把 pdl_sync 挪到一个”本来就以 PDL 方式启动的兄弟 kernel”里，或者把依赖降到 Python 侧、回到两次普通 launch。

6. 工厂参数决定”代码形状”，动态维度决定”每次的 grid”

prepare_cp_mirror_tl 的 compress_ratios 元组里有多少个 ratio，决定 num_rows、决定要展开几份 _write_ratio_len、决定 cu_lens 有几行——这些是代码形状，必须在工厂参数里固定。而 batch_size/total_q/output_len 只影响 grid 大小和循环边界，写成 T.dynamic 没问题。写多段 kernel 时要分清这两类量：出现在 T.macro 展开、T.alloc_* 形状、T.unroll 次数里的量必须静态；只出现在 grid、循环 upper bound、下标判断里的量可以动态。搞混的典型后果是 JIT 缓存被形状打爆，或者编译期报”无法在 trace 期求值”。

7. 多段之间的中间结果必须走显存

PDL 链里的两段属于不同的 kernel，各自的 T.alloc_shared 是段内私有的。唯一的跨段通道是显存。fused_sample.py 的 lse_workspace 就是这个通道的命名规范化版本：[total_q, 2] 的 fp32 显存区，第一段写、第二段读；prepare_cp_mirror 的 cu_lens 同理。这条约束意味着，多段组合不能用来”把中间结果藏在共享内存里”，跨段的中间结果必须显式规划，而且这个显存区在 CUDA Graph 捕获时必须作为 graph 的输入/输出张量被固定下来。

8. T.Pipelined 的流水与 PDL 的”机会性并行”会争 SM

kv_index_cat.py 里用了 T.Pipelined(pipeline_lr) 来流水化对 kv_index 的拷贝。每段的流水深度、线程数、SM 占用都是独立选的。需要注意的是，如果前一段在流水末尾触发、后一段在 preamble 里也开流水，两段会争同一批 SM 的驻留名额。PDL 文档那句”opportunistic”在这里就变得很具体：下游流水能开多深，取决于上游 block 退出得有多快，而不是你在源码里写的 num_stages。TileLang 自带的 test_tilelang_language_pdl.py 里用 work_iters=8192 给两段都制造”前置/后置工作”，正是为了在受控条件下观察这种资源竞争。