（来源：华为计算）

近期，上海人工智能实验室DeepLink团队推出扩展 Triton的深度学习编译器DLCompiler，以及面向大模型训练与推理、异构硬件适配的高性能算子库DLBlas。面向昇腾，研究团队通过扩展DSL深度融合，在性能保持无损的同时，突破了跨代迁移难题。并联合昇腾团队协同优化，通过扩展DSL深度融合，基于AscendNPU IR实现Triton OP在昇腾NPU上的高性能优化和性能无损运行，并应用在上海人工智能实验室多模态大模型训练中，提升训练性能和效率。

亮点速览

跨架构DSL扩展：通过扩展DSL，让DSA芯片（昇腾NPU）也能享受极致的编程体验和性能，成为“跨架构AI Kernel DSL”。

智能自动优化：实现智能核间调度，充分释放多核算力；结合创新的访存合并优化，将离散访问自动重组为高速连续访问，大幅提升算子性能与带宽利用率。

大模型瓶颈算子极致优化：提供高效Attention、GroupGemm、FuseMoe等关键算子实现，性能可达理论峰值80%+。GroupGemm、Matmul在昇腾NPU的Cube计算效率可达82%。

提供统一便捷的MoE接口：集成DeepEP、DeepGemm等算子，并根据token变化动态调整冗余专家分布解决MoE负载不均衡问题，并提供统一接口兼容不同推理框架。

当前仓库均已开源，可直接下载使用：

DLCompiler Github: 

https://github.com/DeepLink-org/DLCompiler

DLBlas Github: 

https://github.com/DeepLink-org/DLBlas

DeepLink 官网：

https://deeplink.org.cn/home

AscendNPU IR Gitcode: 

https://gitcode.com/Ascend/AscendNPU-IR

通过DSL扩展能力边界多款芯片上获最优性能表现

片上缓存分配（dl.alloc）

根据DSA硬件架构和内存排布要求，申请片上缓存大小并标定申请片上缓存层级（L0A，L0B，...），相比原生Triton，可以让算子开发者更细粒度描述算子tiling和流水排布。

Python@triton.jitdef custom_func_kernel(x_ptr,  output_ptr,                         n_elements, BLOCK_SIZE: tl.constexpr):    pid = tl.program_id(axis=0)    block_start = pid * BLOCK_SIZE    offsets = block_start + tl.arange(0, BLOCK_SIZE)    mask = offsets < n_elements    x = tl.load(x_ptr + offsets, mask=mask)    # 申请片上SRAM UB的buffer    y = dl.alloc([BLOCK_SIZE], 1.68, dtype=tl.float32, layout=dl.ND, scope=dl.L1)    # 实现L1->UB数据搬移    tl.store(y, x)

多处理单元流水抽象（dl.parallel）

针对DSA架构单个核内的Cube核计算单元和Vector核计算单元算力不一致的问题，dl.parallel可以更细粒度控制计算单元的并行。相比原生Triton，算子开发者可以直接控制Cube核计算单元和Vector核计算单元的并行行为。通过高效利用计算资源，获得更高的性能收益。

PythonSUB_BLK_M: tl.constexpr = BLOCK_SIZE_M // 2# 两个vector核计算单元并行计算for s in dl.parallel(0, 2, bind_sub_block=True):    left = s * SUB_BLK_M    right = (s + 1) * SUB_BLK_M    # 取slice切分local tensor    vec_sub_blk = accumulator[left:right, :]    if ACTIVATION == "leaky_relu_custom":        vec_sub_blk = leaky_relu_custom(vec_sub_blk)    c_sub_blk = vec_sub_blk.to(tl.float16)    # Write back the block of the output matrix C.

灵活编译提示（dl.compile_hint）

给编译器特定提示以执行特定的编译行为。原生Triton DSL抽象度高，某些场景编译器无法知晓完整语义，扩展dl.compile_hint以提示编译器执行特定行为，可以获得更好的性能收益。

Pythonfor k in tl.range(0, tl.cdiv(K, BLOCK_K)):    a_ptrs = a_ptrs_base + k * BLOCK_K    b_ptrs = b_ptrs_base + k * BLOCK_K    a = tl.load(a_ptrs,                mask=msk_m[:, None] and (offs_k[None, :] < K - k * BLOCK_K),                other=0.0)    # 提供类似pragma指导编译器优化  dl.compile_hint(a, "dot_pad_only_k")    b = tl.load(b_ptrs,                mask=msk_n[:, None] and (offs_k[None, :] < K - k * BLOCK_K),                other=0.0) dl.compile_hint(b, "dot_pad_only_k")    accumulator = tl.dot(a, b.T, acc=accumulator)

细粒度缓存切分（dl.extract_slice）

支持对SharedMemory、UB、L1等缓存层级中的localTensor取slice。可以与dl.parallel配合，细粒度控制计算资源，也可以配合访存，将非连续访存替换成连续大块访存，然后在LocalMemory中使用slice切分后分别进行计算。

Python@triton.jitdef triton_kernel(x_ptr, y_ptr, output_ptr, POS: tl.constexpr, N: tl.constexpr, BLOCK_SIZE_N: tl.constexpr):    pid = tl.program_id(axis=0)    start = pid * N    offsets = tl.arange(0, BLOCK_SIZE_N)    mask = offsets < N    x = tl.load(x_ptr + start + offsets, mask=mask)    y = tl.load(y_ptr + start + offsets, mask=mask)    # 片上缓存切分out_left = x[:POS] + y[:POS]   out_right = x[POS:] - y[POS:]    out_left_offsets = tl.arange(0, POS)    tl.store(output_ptr + start + out_left_offsets, out_left)    out_right_offsets = POS + out_left_offsets    tl.store(output_ptr + start + out_right_offsets, out_right, mask=out_right_offsets < N)

全新编程模型助力昇腾驱动Cube和 Vecto加速比迈向新高

Produce/Consumer编程模型

昇腾Cube核和Vector核分离式硬件架构，Cube核和Vector核分别拥有相互独立的指令队列。针对FA及其变种算子等需要Cube核和Vector核协同计算的场景，Cube计算单元和Vector计算单元互为生产者消费者关系。DLCompiler最新提出的编程模型采用生产者消费者的思想，分别描述Cube核和Vector核的计算逻辑，并在各自计算逻辑中分别使用dl.set_cross_flag(dl.SyncFlag.C2V, 0)和dl.wait_cross_flag(dl.SyncFlag.V2C, 1)描述Cube核和Vector核的协作关系。

Python@triton.jit_attn_fwd_kernel(Q, K, V, O, workspace_1, workspace_2, workspace_3):    for block_idx in range(pid, NUM_BLOCKS, NUM_CORES): # 绑核优化     with dl.async_task(scope=dl.async_task.cube):            q = tl.load(Q_block_ptr)            for start_n in range(lo, hi, BLOCK_N):                  k = tl.load(K_block_ptr)                qk = tl.dot(q, tl.trans(k))                  tl.store(workspace_1_ptr, qk)             dl.set_cross_flag(dl.SyncFlag.C2V, 0)           dl.wait_cross_flag(dl.SyncFlag.V2C, 1)                p_cast = tl.load(workspace_2_ptr)                  v = tl.load(V_block_ptr)                  acc_l0c = tl.dot(p_cast, v)                  tl.store(workspace_3_ptr, acc_l0c)   dl.set_cross_flag(dl.SyncFlag.C2V, 2)      with dl.async_task(scope=dl.async_task.vector):            for start_n in range(lo, hi, BLOCK_N):                dl.wait_cross_flag(dl.SyncFlag.C2V, 0)                qk = tl.load(workspace_1_ptr)                # compute                tl.store(workspace_2_ptr, p)               dl.set_cross_flag(dl.SyncFlag.V2C, 1)                # compute               dl.wait_cross_flag(dl.SyncFlag.C2V, 2)                acc_o_ub = tl.load(workspace_3_ptr)                acc_ptr = acc_ptr + acc_o_ub            # compute and store

完整代码可参考：

https://github.com/DeepLink-org/DLCompiler/blob/main/test/ascend/passed_tests/test_cv_flash_attention.py

DLCompiler提出的新编程模型有以下特点：

前端明确将Cube计算逻辑和Vector计算逻辑分开描述，分别使用with dl.async_task(scope=cube)和with dl.async_task(scope=vector)

前端申请workspace空间，控制Cube核和Vector核之间数据传递的workspace，workspace的DoubleBuffer逻辑由DLCompiler的编译pass自动生成；

前端描述Cube核和Vector核之间的同步指令，只指明event_id，其他参数由DLCompiler的编译pass自动生成。

无需了解硬件细节，编译优化释放芯片算力

核内调度

原生Triton中，Program ID表示一个block内启动的线程数。在DSA架构下，DLCompiler将其和任务类型对齐，通过自动分离代码映射到Cube核、Vector核资源，在考虑Cube核与Vector核负载均衡且保障Cube核与Vector核流水并行前提下，用满核资源。

访存合并

原生Triton中可将Gemm/Attention连续访存模式做到极致，对于stride不等于1或者非连续内存访问，基本就直接生成scalar load或warp内分散访问。这对于DSA架构采用SIMD指令，性能将会损失很厉害。DLCompiler通过把innermost loop对齐到memory-contiguous dimension，同时通过自动插入tile copy将非连续块转换为scratchedpad，再由compute core从scratchedpad进行消费。

助力芯片算力高效释放，高性能最佳实践

优化L2缓存

针对Matmul、Grouped_Matmul等矩阵乘算子，传统水平分核实现方式是优先完成结果矩阵的一行基本块计算，之后再计算下一行，以此类推。

当参与计算的张量shape比较大时使用传统水平分核方式会有如下问题：

同一时间多个计算核心都需要访问同一块左矩阵内存，产生Bank冲突，导致硬件访问效率降低。

当完成一整行分块矩阵乘运算时，已经将所有右矩阵数据全部使用上，右矩阵较大时会超过L2缓存的容量上限，此后每行运算都会产生缓存未命中，导致L2缓存命中率较低，影响算子执行效率。

使用对角线分核计算可以很大程度优化上面两点，此处以使用8*8对角线分核为例（可以Autotune）。8*8对角线分核方式中，每8*8分格内任务块编号如下：

在昇腾Atlas A2训练/推理系列产品中，以24个计算核并行执行任务为例进行分析。水平分核时，同一时间所有的核都在使用同一块左矩阵，导致理论访问bank冲突高达24，降低了左矩阵搬运效率。而8*8对角线分核分核的任务块内，行方向同一时刻只有3个核（0，8，16）在读左矩阵，列方向同一时刻也只有3个核（0，22，15）在读右矩阵，明显减小了Bank冲突。

L2缓存是所有计算核心共享的，所以理论上应尽可能使用L2缓存中的数据计算按计算。水平分核计算一行就需要使用到整个右矩阵。而对角线分核使用整个右矩阵理论可以执行8行数据的计算，数据局部性更优，对于L2缓存利用率较高。当L2缓存不足以放下整个右矩阵时，水平分核存在更加频繁的L2缓存换入换出。对角线分核作为Swizzle分核的变种，在DSA架构芯片上可以获得更高的性能收益。

高效访存

增加访存连续性：DLCompiler支持通过load/store原语使用地址偏移和mask实现灵活访存，但是连续大段访存时IO利用率更高。例如当读写二维Tensor的一个Block时，对低维的读写是连续的，对高维的读写是间断的，所以Autotune时，适当增加低维BlockSize可以增加访存连续性，性能更好。

使用块指针访存：块指针从语义上可以明确数据的排布形式，在DSL里直接提供了访存的shape、stride、offset、order等信息，比使用load/store原语更有利于编译优化，生成更优的访存指令。

使用编译器提示：访存时在DSL中增加max_constancy、max_contiguous、multiple_of三种编译器提示原语可以辅助编译器进行地址连续性分析，更有助于生成高效访存指令。

大段读取、切片计算：连续大段访存时IO利用率更高。某些场景中（例如数据不连续或者计算逻辑需要多次读取小数据），可以一次性读取大段数据，然后使用DLCompiler扩展的slice原语，在片上对大段数据取切片后分别计算，可以达到更高性能。

组合数据、大段写出：为了更好利用设备带宽。某些场景中，也可以使用DLCompiler扩展的slice原语，在片上将数据组合起来再整体写到GlobalMemory，可以达到更高性能。

大模型瓶颈算子精准优化

在大模型训练与推理中，针对计算瓶颈的关键算子在昇腾硬件平台上已实现显著的性能加速。通过深入优化算子实现、充分利用硬件特性，并结合软件协同设计取得突出进展：

针对昇腾NPU，通过扩展领域专用语言（DSL）并与昇腾毕昇编译器团队开展深度融合与协同优化，实现了关键算子的高效部署与性能提升。相比社区通用实现，优化后的算子性能提升达1.2至1.9倍。此外，针对部分社区现有算子在昇腾特定张量加速器（DSA）架构上支持不足的情况，完成了功能适配与基本通路验证，确保了其可在昇腾平台上正常运行。相关优化算子已在DLBlas仓库中集成发布。

MoE模型终极适配与生态融合

DLBlas为MoE系列模型提供统一接口层，支持SGLang、LMDeploy、vLLM等推理框架的即插即用，集成DeepEP，DeepGemm等接口，并针对MoE中的大EP场景优化实现two batch overlap和融合算子fusedMoE，并根据token动态选择expert均衡计算负载实现以缓解木桶效应。此外，针对FlashAttention的精度与效率问题，通过Triton Autotune参数优化与L2 cache策略调整，在保持精度的同时实现最高1.108倍加速。

DLBlas在Qwen3、DeepSeekV3等模型中助力开源框架提升性能：

上海人工智能实验室与昇腾强强合作，基于DeepLink和昇腾NPU，深耕编译优化等核心技术，全力推进AI工具链研发与安全高效的方案落地。针对昇腾NPU特性优化算子性能、构建统一计算通信中间表达，打造兼容主流框架的开发环境，为开发者提供更加便利的开发工具、更极致的计算性能。通过构建多场景支撑平台和垂域标杆应用，为科研和产业界提供更多技术创新的支持和应用参考，加速AI全行业自主技术创新与应用。