昇腾与GPU的开发差异

多核任务并行策略

NPU在Triton多核并行中是物理核强绑定模式，与GPU逻辑维度并行+硬件自动物理映射的模式形成核心差异

• 核心对比

  维度	GPU（NVIDIA）	昇腾（Ascend）
  grid 本质	逻辑任务维度（和物理核解耦）	物理核组映射（绑定 AI Core 拓扑）
  核数 / 维度限制	grid 维度 / 大小无硬限制	grid 大小≤AI Core 总数，2D 需匹配拓扑

GPU:可绑定多个维度轴（三维grid=[n,m,l] 等同于乘积n×m×l个并行线程），每个线程仅对应一次kernel执行，且仅执行一次。
NPU:Vector核，Cube核属于多个物理核，不同代际硬件核数不同，每个核仅执行一次Block,且支持对该Block重复调度执行。

充分利用核数

昇腾NPU具备多个计算核心，合理分配并充分利用所有可用核心，是提升算子性能的关键因素之一。 在调用Triton内核函数时，通过设置launch参数控制使用的核数量。以GELU算子为例：

triton_gelu[n, 1, 1](...)  # 第一个参数表示使用的核数，n表示使用n个核

通过对核数的调优，可实现对所有计算资源的充分调度和利用，从而最大化并行度与吞吐量。注意，当前版本核数需小于等于65535。

单核数据搬运策略

数据切分Tiling

写Triton内核函数时，合理的数据切分策略对性能优化至关重要。通过调整不同的切分粒度参数，可以在不同维度上平衡计算负载与内存访问效率。

常见的切分参数包括：

ncore：使用的核数（跨核切分）
xblock：核间数据块大小（核间切分）
xblock_sub：核内切分粒度（核内细粒度划分）

开发者可根据实际场景手动选择最优的切分配置，使得每次计算尽可能充分利用片上内存（On-chip Memory），避免频繁访问全局内存（Global Memory）造成的 性能瓶颈。

以GELU算子为例，通过调整切分参数，可以有效适配片上缓存容量限制，从而提升执行效率。

注：Atlas 800T/I A2产品的片上内存容量为192KB，因此在设计切分策略时需考虑该限制，确保每轮计算的数据量不超过片上内存容量。

GELU算子示例

GELU算子开发示例，使用3种方式计算结果。

standard_unary 为标准Torch计算。

triton_easy_kernel 为简单Triton实现。

triton_better_kernel为更高效的Triton实现。

标准Torch写法

输入tensor x0，经过torch计算实现 GELU 算子，返回结果值。

def standard_unary(x0):
    res = x0 * 0.5 * (1.0 + torch.erf(x0 / torch.sqrt(torch.tensor(2.0))))
    return res

简单Triton写法

以下是一个使用 Triton 编写的简单内核示例，用于展示如何定义和调用一个基本的Triton内核函数。此示例实现了一个简单的数学运算（GELU 激活函数）。

# 定义triton_kernel核函数
@triton.jit
def triton_easy_kernel(in_ptr0, out_ptr0, NUMEL: tl.constexpr):
    idx_block = tl.arange(0, NUMEL)
    x = tl.load(in_ptr0 + idx_block)
    ret = x * 0.5 * (1.0 + tl.erf(x / tl.sqrt(2.0)))
    tl.store(out_ptr0 + idx_block, ret)

注意事项

1. 内存限制：上述写法中，所有输入数据一次性被加载到内存中进行计算。如果输入张量过大，可能会超出单个内核的片上内存容量，导致内存溢出错误。 因此，这种简单的写法更适合于小规模张量的计算或用于理解 Triton 内核的基本写法和调用方式。

2. 适用场景：尽管这种方法有助于快速理解和入门 Triton 编程，但对于大规模数据集或高性能要求的应用场景，建议采用更复杂的数据切分策略（如 Tiling）， 以充分利用硬件资源并避免内存溢出问题。通过这种方式，开发者可以快速上手 Triton 编程，同时了解如何定义、调用以及优化 Triton 内核函数。

更高效triton写法

在昇腾 NPU 上使用 Triton 编写高性能算子时，为了充分利用硬件资源、避免内存溢出并提升执行效率，通常需要采用数据切分（Tiling）策略。 下面是一个经过优化的 Triton 内核实现示例，适用于大规模张量计算。

# 定义triton_kernel核函数
@triton.jit
def triton_better_kernel(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    for xoffset_sub in range(0, XBLOCK, XBLOCK_SUB):
        x_index = xoffset + xoffset_sub + tl.arange(0, XBLOCK_SUB)[:]
        xmask = x_index < xnumel
        x = tl.load(in_ptr0 + x_index, xmask)
        ret = x * 0.5 * (1.0 + tl.erf(x / tl.sqrt(2.0)))
        tl.store(out_ptr0 + x_index, ret, xmask)

# 调用triton_kernel核函数
ncore = 32
xblock = 32768
xblock_sub = 8192
triton_better_kernel[ncore, 1, 1](x0, out1, x0.numel(), xblock, xblock_sub)

关键代码解释

# 计算当前核处理数据块的起始偏移地址，实现核间切分。每个核仅负责 XBLOCK 大小的数据范围。
xoffset = tl.program_id(0) * XBLOCK

# 在单个核内部进一步细分数据块，每次处理 XBLOCK_SUB 大小的数据，实现核内切分。
for xoffset_sub in range(0, XBLOCK, XBLOCK_SUB):

# 构造当前迭代的数据索引数组，用于访问输入和输出张量。
x_index = xoffset + xoffset_sub + tl.arange(0, XBLOCK_SUB)[:]

# 设置掩码以防止越界访问，确保只处理合法范围内的数据。
xmask = x_index < xnumel

# 分别用于从全局内存加载数据到片上内存，以及将计算结果写回全局内存。
tl.load() 和 tl.store()

编译优化能力

AscendNPU IR优化

针对昇腾软硬件特性，适配了AscendNPU IR优化的编译选项，如下表所示。 使用方法：在autotune的配置阶段，传入编译选项的值 以开启multibuffer选项举例，在autotune的配置阶段，即triton.Config中，传入'multibuffer': True，详见autotune示例：

    def get_autotune_config():
        return [
            triton.Config({'XS': 1 * 128, 'multibuffer': True}),]

选项	能力	是否开启
multibuffer	开启流水并行数据搬运	默认true； true , false。 autotune中可配置
unit_flag	cube搬出的一个优化项	默认None；true , false。 autotune中可配置
limit_auto_multi_buffer_only_for_local_buffer	CV算子一个优化项，cube搬出的一个优化项	默认None；true , false。 autotune中可配置
limit_auto_multi_buffer_of_local_buffer	cube算子开启double buffer具体的scope	默认None；["no-limit", "no-l0c"]， autotune中可配置
set_workspace_multibuffer	只有在limit_auto_multi_buffer_only_for_local_buffer=false场景下生效	默认None；如[2,4]，autotune中可配置
enable_hivm_auto_cv_balance	set_workspace_multibuffer只有在limit_auto_multi_buffer_only_for_local_buffer=false场景下生效	默认None；true , false。 autotune中可配置
tile_mix_vector_loop	CV算子的一个优化项，当前vector可以切几份	默认None；如 [2,4,8]，autotune中可配置
tile_mix_cube_loop	CV算子一个优化项，当前cube可以切几份	默认None；如 [2,4,8]，autotune中可配置
auto_blockify_size	TRITON_ALL_BLOCKS_PARALLEL优化项，用于指定扩展的左起第一个维度的大小。	默认1；如 [2,4,8]，autotune中可配置

• 注：优化编译选项在ascend/backend/compiler.py代码中。
• 注：CV算子表示该算子运算过程中既使用了AI Core又使用了Vector Core。