AscendNPU-IR 自动分块（Auto Blockify）：逻辑块到物理块的高效调度优化全解析

• 编译器
• 人工智能
• Ascend
• CANN

本文以 AscendNPU-IR 开源仓库中 auto_blockify.md 文档为主体，结合AutoBlockifyParallelLoop与SIMTAutoBlockify两个核心 Pass 的源码实现与测试用例展开。读完本文，你将理解：为什么昇腾算子执行流水线需要自动分块；SIMD 与 SIMT 两种模式下自动分块的算法原理与接口用法；Super-blocking 性能扩展如何工作；以及如何通过TRITON_ALL_PARALLEL、--enable-auto-blockify-loop、--simt-auto-blockify、--super-block-factor=N等开关在 triton-ascend 与 bishengir-compile 中启用该特性。

一、背景：调度效率决定算子性能

Auto Blockify Pass 是昇腾（Ascend）兼容算子执行流水线中的一项核心技术优化，其目标是把逻辑块（logical block）高效映射到硬件物理块（physical block）。在当前架构下，调度效率直接决定算子性能：如果逻辑块与物理块之间能够实现一一对应，调度开销将完全消除，从而带来显著的性能提升。

在 AscendNPU-IR 的实际架构中，可用物理块数量通常远小于计算所需的逻辑块数量。典型场景下，物理块不足 50 个，而逻辑块可能达到 500 个以上。当两者存在约 10 倍的差距时，通过 Auto Blockify 消除调度开销后，加速效果可以超过原始速度的 2 倍。

Auto Blockify 的触发方式因使用场景而异：

• 通过 triton-ascend 运行 Triton kernel 时，在 Triton 前端设置环境变量TRITON_ALL_PARALLEL即可激活 Auto Blockify 逻辑；
• 面向 AscendNPU-IR 开发者，在bishengir-compile命令中加入--enable-auto-blockify-loop标志。

二、SIMD 模式：AutoBlockifyParallelLoop

SIMD 模式下的自动分块由 HIVM Dialect 侧的AutoBlockifyParallelLoopPass 实现，其定义位于 Passes.td：

def AutoBlockifyParallelLoop : Pass<"auto-blockify-parallel-loop", "func::FuncOp"> { let summary = "Enable auto loop on blocks when logical blocknum is larger than physical one"; let constructor = "mlir::hivm::createAutoBlockifyParallelLoopPass()"; let dependentDialects = ["hivm::HIVMDialect", "hacc::HACCDialect"]; }

2.1 算法原理

该 Pass 的核心思路是：在原有 IR 基础上额外引入一层外层循环，使每个物理块在单次调度中可以顺序处理多个逻辑块。抽象逻辑如下：

for outer from 0,...,ceildiv(logical_block_dim, physical_block_dim) for inner from 0,...,physical_block_dim <- Used as block.idx. use(min(outer * physical_block_dim + inner, logical_block_dim))

逻辑推演分三步理解：

1. 原始调度方式

   未做优化时，每个逻辑块对应一个独立的block.idx，等价于：

   block.idx = hivm.get_block_idx use(block.idx) -------Equivalent to-------------- for block.idx from 0,...,logical_block_num use(block.idx)

2. 使用TRITON_ALL_PARALLEL的场景

   用户在 triton 适配器中加入TRITON_ALL_PARALLEL标志后，kernel 被限制为只以最大物理块数量启动（假设逻辑块数量 > 物理块数量）。此时执行被限制为：

   for block.idx from 0,...,physical_block_num <- From get_block_idx use(block.idx)

   仅靠这一层循环逻辑无法覆盖所有计算索引，会造成索引缺失（missing indices）。这正是引入 Auto Blockify Pass 的原因——它自动补充一层外层循环（分块层），填补这一空缺。

   注意：如果没有通过 Triton Adapter 集成，必须自行保证 block 维度的设置与上述逻辑一致。

3. Auto Blockify 处理后的最终逻辑

   当 Pass 自动补全循环结构后，最终执行逻辑为：

   for outer from 0,...,ceildiv(logical_block_dim, physical_block_dim) for inner from 0,...,physical_block_dim <- used as block.idx use(min(outer * physical_block_dim + inner, logical_block_dim))

   其中min(outer * physical_block_dim + inner, logical_block_dim)的作用是：当最后一个外层迭代无法整除时，将逻辑块索引钳制到合法范围内，避免越界访问。

2.2 源码实现要点

从 AutoBlockifyParallelLoop.cpp 的loopOnLogicBlock函数可以看清整个改写流程：

• 识别逻辑块数：遍历入口块，寻找带kLogicalBlockNumAttr（IR 中的logical_block_num）标记的annotation.mark操作，取其操作数作为逻辑块总数；若找不到，Pass 报错"Logical Block number not found"并失败；
• 定位物理块索引：查找hivm::GetBlockIdxOp，其返回值（i64）经arith.trunci截断为 i32 作为scf.for循环的起始值（下界），即物理块从0到物理块维度；
• 推导物理块数：getPhysicalBlockNum通过queryFuncCoreType判断函数核心类型，再结合模块上的dlti.target_system_spec中hacc设备规格（DeviceSpec）读取物理块数——AIV（Vector）核心取VECTOR_CORE_COUNT，Cube 相关核心取CUBE_CORE_COUNT；若无法推断则报错"Physical block num cannot be inferred"；
• 构建外层循环：创建scf.for，步长为物理块数，上界为逻辑块数；把除标记操作、get_block_idx及其依赖之外的所有指令整体搬移到循环体内；
• 重写 block 索引使用：replaceBlockIdUsers把原先直接使用get_block_idx结果的地方，替换为outer_iv * physical_block_num + get_block_idx（即外层迭代变量乘以物理块数再加上内层物理块索引），得到完整的一维逻辑块 ID；
• 标记新循环：为生成的scf.for打上autoblockify.subloop属性（BlockifyLoopAttrName），便于后续 Pass 识别该循环已被分块处理。

值得一提的是，源码中还支持auto_blockify_size函数属性：当该属性大于 1 时，上界会先执行ceildiv(logical_block_num, blockify_size)，步长也相应调整，用于进一步细化分块粒度。

2.3 接口与使用条件

该特性在bishengir-compile中由--enable-auto-blockify-loop标志控制（选项定义见 Options.td，默认值为false）；同时也可以直接用bishengir-opt的--auto-blockify-parallel-loop标志单独调用。

要正确使用该特性，必须满足以下两个前置条件：

1. 逻辑块数必须可用：Pass 通过查找带有kLogicalBlockNumAttr属性（IR 中表现为annotation.mark %1 {logical_block_num}）的值来获取逻辑块数。必须确保该值存在，否则 Pass 调用会失败；
2. 必须存在hivm get_block_idx操作：该操作返回从 0 到 block 维度的块索引。使用 Auto Blockify 时，需要修改调用设备 kernel 时的 block 维度（以最大物理块维度启动，与上述算法一致），使get_block_idx返回 0 到physical_block_num范围内的值。

2.4 管线集成位置

在 HIVMPipelines.cpp 中，该 Pass 被明确放置在createInferFuncCoreTypePass（推断核心类型）之后，源码注释解释了原因：

AutoBlockifyParallelLoopPass needs to be after infer core type because num. of physical blocks we loop on is dependent on core type

也就是说，外层循环需要以"每个物理块处理多少逻辑块"为依据，而这个数量取决于核心类型（Vector 还是 Cube），因此必须先完成核心类型推断。同时该 Pass 仅在enableTritonKernelCompile && enableAutoBlockifyLoop同时为真时插入，说明它面向的是 Triton kernel 编译路径。

2.5 输入与输出示例

以下是一个完整的add_kernel示例。输入 IR 中，逻辑块数由arg6 * arg7 * arg8计算并由annotation.mark {logical_block_num}标记（约 1024 的倍数），hivm.hir.get_block_idx提供一维物理块索引：

module attributes {dlti.target_system_spec = #dlti.target_system_spec<"NPU" : #hacc.target_device_spec<#dlti.dl_entry<"AI_CORE_COUNT", 20 : i32>, #dlti.dl_entry<"CUBE_CORE_COUNT", 20 : i32>, #dlti.dl_entry<"VECTOR_CORE_COUNT", 40 : i32>, #dlti.dl_entry<"UB_SIZE", 1572864 : i32>, #dlti.dl_entry<"L1_SIZE", 4194304 : i32>, #dlti.dl_entry<"L0A_SIZE", 524288 : i32>, #dlti.dl_entry<"L0B_SIZE", 524288 : i32>, #dlti.dl_entry<"L0C_SIZE", 1048576 : i32>, #dlti.dl_entry<"UB_ALIGN_SIZE", 256 : i32>, #dlti.dl_entry<"L1_ALIGN_SIZE", 256 : i32>, #dlti.dl_entry<"L0C_ALIGN_SIZE", 4096 : i32>>>, hivm.module_core_type = #hivm.module_core_type<AIV>} { func.func @add_kernel(%arg0: i64 {hacc.arg_type = #hacc.arg_type<ffts_base_address>}, %arg1: memref<?xi8> {hacc.arg_type = #hacc.arg_type<workspace>}, %arg2: memref<?xf32> {tt.divisibility = 16 : i32}, %arg3: memref<?xf32> {tt.divisibility = 16 : i32}, %arg4: memref<?xf32> {tt.divisibility = 16 : i32}, %arg5: i32 {tt.divisibility = 16 : i32}, %arg6: i32, %arg7: i32, %arg8: i32) attributes {WorkspaceArgIdx = 0 : i64, func_dyn_memref_args = dense<[false, true, true, true, true, false, false, false, false]> : vector<9xi1>, hacc.entry, hacc.function_kind = #hacc.function_kind<DEVICE>, hivm.func_core_type = #hivm.func_core_type<AIV>} { %c1024 = arith.constant 1024 : index %c1024_i32 = arith.constant 1024 : i32 %c0 = arith.constant 0 : index hivm.hir.set_mask_norm %0 = arith.muli %arg6, %arg7 : i32 %1 = arith.muli %0, %arg8 : i32 annotation.mark %1 {logical_block_num} : i32 // This logical_block_num is the original large value. %2 = hivm.hir.get_block_idx -> i64 // for block.idx from 0,...,block_num %3 = arith.trunci %2 : i64 to i32 %4 = arith.muli %arg8, %arg7 : i32 %5 = arith.divsi %3, %4 : i32 %6 = arith.remsi %5, %arg6 : i32 %7 = arith.muli %6, %c1024_i32 : i32 %8 = arith.index_cast %7 : i32 to index %reinterpret_cast = memref.reinterpret_cast %arg2 to offset: [%8], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>> %alloc = memref.alloc() : memref<1024xf32> %9 = arith.addi %8, %c1024 : index %10 = arith.index_cast %arg5 : i32 to index %11 = arith.maxsi %8, %10 : index %12 = arith.minsi %9, %11 : index %13 = arith.subi %12, %8 : index %subview = memref.subview %reinterpret_cast[0] [%13] [1] : memref<1024xf32, strided<[1], offset: ?>> %subview_0 = memref.subview %alloc[0] [%13] [1] : memref<1024xf32> to memref<?xf32, strided<[1]>> hivm.hir.load ins(%subview : memref<?xf32, strided<[1], offset: ?>>) outs(%subview_0 : memref<?xf32, strided<[1]>>) left_padding_num = %c0 : index init_out_buffer = false %14 = bufferization.to_tensor %alloc restrict writable : memref<1024xf32> %reinterpret_cast_1 = memref.reinterpret_cast %arg3 to offset: [%8], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>> %alloc_2 = memref.alloc() : memref<1024xf32> %subview_3 = memref.subview %reinterpret_cast_1[0] [%13] [1] : memref<1024xf32, strided<[1], offset: ?>> %subview_4 = memref.subview %alloc_2[0] [%13] [1] : memref<1024xf32> to memref<?xf32, strided<[1]>> hivm.hir.load ins(%subview_3 : memref<?xf32, strided<[1], offset: ?>>) outs(%subview_4 : memref<?xf32, strided<[1]>>) left_padding_num = %c0 : index init_out_buffer = false %15 = bufferization.to_tensor %alloc_2 restrict writable : memref<1024xf32> %16 = tensor.empty() : tensor<1024xf32> %17 = hivm.hir.vadd ins(%14, %15 : tensor<1024xf32>, tensor<1024xf32>) outs(%16 : tensor<1024xf32>) -> tensor<1024xf32> %reinterpret_cast_5 = memref.reinterpret_cast %arg4 to offset: [%8], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>> %extracted_slice = tensor.extract_slice %17[0] [%13] [1] : tensor<1024xf32> to tensor<?xf32> %subview_6 = memref.subview %reinterpret_cast_5[0] [%13] [1] : memref<1024xf32, strided<[1], offset: ?>> hivm.hir.store ins(%extracted_slice : tensor<?xf32>) outs(%subview_6 : memref<?xf32, strided<[1], offset: ?>>) return } }

经过--auto-blockify-parallel-loop处理后的输出 IR 发生了如下关键变化：

• 外层新增scf.for %arg9 = %c0_i32 to %2 step %c1_i32，其中%2 = ceildivsi(%1, 40)，40来自VECTOR_CORE_COUNT，即ceildiv(logical_block_num, physical_block_dim)；
• 原函数体（hivm.hir.load、vadd、store等计算）整体搬入外层循环体内；
• get_block_idx仍提供内层物理块索引，但所有原本直接使用 block 索引做地址计算的地方，改为先计算outer * physical_block_num + inner再钳位min(..., logical_block_num)得到完整逻辑块 ID：

module attributes {dlti.target_system_spec = #dlti.target_system_spec<"NPU" : #hacc.target_device_spec<#dlti.dl_entry<"AI_CORE_COUNT", 20 : i32>, #dlti.dl_entry<"CUBE_CORE_COUNT", 20 : i32>, #dlti.dl_entry<"VECTOR_CORE_COUNT", 40 : i32>, #dlti.dl_entry<"UB_SIZE", 1572864 : i32>, #dlti.dl_entry<"L1_SIZE", 4194304 : i32>, #dlti.dl_entry<"L0A_SIZE", 524288 : i32>, #dlti.dl_entry<"L0B_SIZE", 524288 : i32>, #dlti.dl_entry<"L0C_SIZE", 1048576 : i32>, #dlti.dl_entry<"UB_ALIGN_SIZE", 256 : i32>, #dlti.dl_entry<"L1_ALIGN_SIZE", 256 : i32>, #dlti.dl_entry<"L0C_ALIGN_SIZE", 4096 : i32>>>, hivm.module_core_type = #hivm.module_core_type<AIV>} { func.func @add_kernel(%arg0: i64 {hacc.arg_type = #hacc.arg_type<ffts_base_address>}, %arg1: memref<?xi8> {hacc.arg_type = #hacc.arg_type<workspace>}, %arg2: memref<?xf32> {tt.divisibility = 16 : i32}, %arg3: memref<?xf32> {tt.divisibility = 16 : i32}, %arg4: memref<?xf32> {tt.divisibility = 16 : i32}, %arg5: i32 {tt.divisibility = 16 : i32}, %arg6: i32, %arg7: i32, %arg8: i32) attributes {WorkspaceArgIdx = 0 : i64, func_dyn_memref_args = dense<[false, true, true, true, true, false, false, false, false]> : vector<9xi1>, hacc.entry, hacc.function_kind = #hacc.function_kind<DEVICE>, hivm.func_core_type = #hivm.func_core_type<AIV>} { %0 = arith.muli %arg6, %arg7 : i32 %1 = arith.muli %0, %arg8 : i32 annotation.mark %1 {logical_block_num} : i32 // This logical_block_num is the original large value. %c0_i32 = arith.constant 0 : i32 %c40_i32 = arith.constant 40 : i32 // 40 is the number of physical blocks here. %2 = arith.ceildivsi %1, %c40_i32 : i32 // ceildiv(logical_block_num, physical_block_dim) %c1_i32 = arith.constant 1 : i32 scf.for %arg9 = %c0_i32 to %2 step %c1_i32 : i32 { // Outer loop %c1024 = arith.constant 1024 : index %c1024_i32 = arith.constant 1024 : i32 %c0 = arith.constant 0 : index hivm.hir.set_mask_norm %3 = hivm.hir.get_block_idx -> i64 // Inner loop (locked to physical blocks) %4 = arith.trunci %3 : i64 to i32 %5 = arith.muli %arg9, %c40_i32 : i32 // outer_i * physical_block_num %6 = arith.addi %5, %4 : i32 // outer_i * physical_block_num + inner %7 = arith.minsi %6, %1 : i32 // (min(outer*physical_block_dim + inner, logical_block_num)) %8 = arith.extsi %7 : i32 to i64 %9 = arith.trunci %8 : i64 to i32 %10 = arith.muli %arg8, %arg7 : i32 %11 = arith.divsi %9, %10 : i32 %12 = arith.remsi %11, %arg6 : i32 %13 = arith.muli %12, %c1024_i32 : i32 %14 = arith.index_cast %13 : i32 to index %reinterpret_cast = memref.reinterpret_cast %arg2 to offset: [%14], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>> %alloc = memref.alloc() : memref<1024xf32> %15 = arith.addi %14, %c1024 : index %16 = arith.index_cast %arg5 : i32 to index %17 = arith.maxsi %14, %16 : index %18 = arith.minsi %15, %17 : index %19 = arith.subi %18, %14 : index %subview = memref.subview %reinterpret_cast[0] [%19] [1] : memref<1024xf32, strided<[1], offset: ?>> %subview_0 = memref.subview %alloc[0] [%19] [1] : memref<1024xf32> to memref<?xf32, strided<[1]>> hivm.hir.load ins(%subview : memref<?xf32, strided<[1], offset: ?>>) outs(%subview_0 : memref<?xf32, strided<[1]>>) left_padding_num = %c0 : index %20 = bufferization.to_tensor %alloc restrict writable : memref<1024xf32> %reinterpret_cast_1 = memref.reinterpret_cast %arg3 to offset: [%14], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>> %alloc_2 = memref.alloc() : memref<1024xf32> %subview_3 = memref.subview %reinterpret_cast_1[0] [%19] [1] : memref<1024xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1], offset: ?>> %subview_4 = memref.subview %alloc_2[0] [%19] [1] : memref<1024xf32> to memref<?xf32, strided<[1]>> hivm.hir.load ins(%subview_3 : memref<?xf32, strided<[1], offset: ?>>) outs(%subview_4 : memref<?xf32, strided<[1]>>) left_padding_num = %c0 : index %21 = bufferization.to_tensor %alloc_2 restrict writable : memref<1024xf32> %22 = tensor.empty() : tensor<1024xf32> %23 = hivm.hir.vadd ins(%20, %21 : tensor<1024xf32>, tensor<1024xf32>) outs(%22 : tensor<1024xf32>) -> tensor<1024xf32> %reinterpret_cast_5 = memref.reinterpret_cast %arg4 to offset: [%14], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>> %extracted_slice = tensor.extract_slice %23[0] [%19] [1] : tensor<1024xf32> to tensor<?xf32> %subview_6 = memref.subview %reinterpret_cast_5[0] [%19] [1] : memref<1024xf32, strided<[1], offset: ?>> hivm.hir.store ins(%extracted_slice : tensor<?xf32>) outs(%subview_6 : memref<?xf32, strided<[1], offset: ?>>) } return } }

2.6 Triton Adapter 集成与前置准备

该 Pass 已广泛应用于 triton 适配器流水线中。正确用法是从前端（triton）通过TRITON_ALL_PARALLEL=1环境变量开启——这个环境变量同时完成了准备工作（锁定 block 数量），然后自动以正确的标志调用对应编译器命令。

为什么 Triton 路径能开箱即用？因为在 Triton 流水线中存在名为TritonGlobalKernelArgsToHIVMOpPass的 Pass（实现见 TritonGlobalKernelArgsToHIVMOp.cpp），它会自动保证logical_block_num标记值存在并创建所需的get_block_idx操作。其工作原理（源码 L56-L76 注释）是：Triton 的启动网格总是 3D 的，而hivm::get_block_idx是一维的，因此该 Pass 把 Triton kernel 函数末尾的 6 个 i32 参数（3 个 PROGRAM_NUM_ARGS + 3 个 PROGRAM_ID_ARGS）转换为一维逻辑块索引：用program_num_x * program_num_y * program_num_z计算总逻辑块数并打上logical_block_num标记，随后创建hivm::GetBlockIdxOp，并按 Triton 的 x-fastest 启动顺序（idx = pid_0 + pid_1 * num_x + pid_2 * num_x * num_y）将一维索引解码回 3D program_id。这样，后续的 AutoBlockifyParallelLoopPass 就能直接消费这些信息。

2.7 约束条件

• 可并行性（Parallelizability）：Auto Blockify 算法仅适用于完全可并行的代码。即每个逻辑块的计算与访存必须能够安全地并行执行，块与块之间不能存在依赖；
• 使用场景（Use scenario）：如果逻辑块数量非常少（接近或小于物理块数），该 Pass 无法带来收益，不应开启。

三、SIMT 模式：SIMTAutoBlockify

SIMT 模式下的自动分块，在功能和使用方式上与 SIMD 模式本质相同，本节重点介绍 SIMT 路径特有的部分。对应 Pass 为SIMTAutoBlockify，作用于mlir::triton::FuncOp（TTIR 层），定义见 Triton/Transforms/Passes.td：

def SIMTAutoBlockify : Pass<"simt-auto-blockify", "mlir::triton::FuncOp"> { let summary = "Blockify Triton SIMT kernels when logical grid is larger than the physical core count"; ... let options = [Option<"superBlockFactor", "superblock-factor", "unsigned", /*default=*/"1", "The super-blocking factor (must be a power of 2, " ">= 1). 1 means no super-blocking.">]; }

3.1 算法原理

与 SIMD 模式类似，该调度自动增加一层循环，并用"由新循环 IV 与物理核 ID 计算出的逻辑核 ID"替换 SIMT 版本的逻辑核 ID 指令（tt.get_program_id x/y/z）。计算逻辑如下：

Original kernel: pid_x = tt.get_program_id x pid_y = tt.get_program_id y pid_z = tt.get_program_id z <kernel body> After rewriting: logical = grid_x * grid_y * grid_z // Compute logical_block_num. chunk = ceildiv(logical, physical_block_dim) hw_idx = gpu.linear_block_id // Inner loop (locked to the physical block). start = hw_idx * chunk // The first logical block processed by this physical block. upper = min(start + chunk, logical) // The last logical block processed by this physical block. for iv in [start, upper) step 1: pid_x = iv % grid_x // Compute the 3D logical core ID. pid_y = (iv / grid_x) % grid_y pid_z = iv / (grid_x * grid_y) <kernel body, with the original tt.get_program_id instruction replaced by the pid_* resolved above>

源码实现（SIMTAutoBlockify.cpp）中有几个值得注意的实现细节：

• 硬件线性块 ID：gpu.linear_block_id最终会 lower 为ascend_dpx::BlockIdxOp，并进一步落到线性硬件 block-id 内建指令。由于tt.get_program_id x/y/z是 3D 索引、无法直接给出线性核 ID，Pass 必须显式重写所有tt.get_program_id操作；
• 网格元数据保留：tt.get_num_programs保持原样，供后续 lower 阶段使用；
• 物理核数来源：getPhysicalBlockNum从模块的 NPU TargetSpec 中读取VECTOR_CORE_COUNT作为物理块数（SIMT 面向 Vector 核心）；
• prologue 构建：将函数入口块拆分为 prologue（计算logical、chunk、lowerBound、upperBound并创建scf.for）和原函数体两个块，随后把原函数体搬入循环体；多块函数体则用scf.execute_region包裹；
• 3D 解码：buildLogicalProgramIds按 x-fastest 顺序解码——pid_x = iv % grid_x、pid_y = (iv / grid_x) % grid_y、pid_z = iv / (grid_x * grid_y)，并将所有tt.get_program_id替换为对应值；
• 跳过规则：非公开函数或带返回值的函数被跳过；kernel 内没有任何tt.program_id时静默跳过（对应测试 simt-auto-blockify-silent-skip.mlir）。

在 Triton 管线中（TritonPipelines.cpp），SIMTAutoBlockify位于AdaptTritonIRKernelPass（携带super-block-barrier选项）之后、OptimizeSIMTExpressionsPass之前，并以options.superBlockFactor作为superblock-factor参数传入。

3.2 接口描述

SIMT 模式下，该特性同样由bishengir-compile的--enable-auto-blockify-loop标志控制，也可直接用bishengir-opt的--simt-auto-blockify标志单独调用。

3.3 可选性能扩展：Super-blocking（超级分块）

在 GPU SIMT 编程中，常见做法是单个 kernel 函数配置较小的工作负载、并启动大量逻辑核。这种场景下Vector 核心的计算能力得不到充分利用，无法达到峰值性能。基于自动分块，Super-blocking 扩展让一个物理核并行处理多个相邻的逻辑核，以提升算力利用率。

接口描述：

该特性由bishengir-compile的--super-block-factor=N标志控制，其中N表示并行处理的逻辑核数量，默认值为 1（即关闭 super-blocking，只做常规自动分块）。也可以直接用bishengir-opt的 Pass 选项-simt-auto-blockify="superblock-factor=N"调用。选项定义见 Options.td。

逻辑描述：

常规自动分块引入的循环，其步长改为N，表示同时处理N个相邻逻辑块。若原函数每个逻辑块启动W个 warp，调度后的函数将启动N×W个 warp，并根据 warp ID 计算新的逻辑核 ID：

for iv in [start, upper) step N: warp_id = thread_id_x / 32 // Calculate the warp ID. local = warp_id % N // Calculate the ID of the logical core processed in parallel within the core. linear = iv + local // Calculate the one-dimensional logical core ID. if linear < upper: // Ensure that the logical core ID does not go out of bounds. <kernel body, compute the 3D logical core ID using linear and replace the original tt.get_program_id>

源码中的实现策略（SIMTAutoBlockify.cpp 及其注释）可以更精确地描述为：

warpId = tid.x / WARP (WARP = 32); blockId = gpu.linearBlockId; // super-block id newBlockId = blockId * SBF + (warpId % SBF); // recover original logical block id newWarpId = warpId / SBF; // recover original warp id tid.x = newWarpId * WARP + laneId;

即：把 warp ID 的低位（warpId % SBF）提取出来作为逻辑块 ID 的低位，恢复出真实的逻辑块 ID；warpId / SBF作为真实 warp 号。同时循环步长被调整为SBF，并插入if linear < upper边界检查（scf.if），防止超出的逻辑核 ID 越界。此外，Pass 还要求superblock-factor必须是2 的幂且 ≥ 1，否则直接报错signalPassFailure。

3.4 约束条件

• warp 数量：昇腾最多支持 64 个 warp，因此N×W的值必须小于等于 64；
• 共享内存：开启 super-blocking 后，N个并行逻辑核会把物理核的共享内存均分为N份，需要关注每个逻辑核的内存占用以避免溢出；
• 使用场景：如果逻辑块数量非常少，该 Pass 无法带来收益；如果每个逻辑核的工作负载本身较大，super-blocking 的性能增益也有限。

四、测试与验证

仓库为两个 Pass 提供了完整的 LIT 测试，可作为功能验证与理解参考：

• SIMD 路径：auto-blockify-parallel-loop.mlir 通过bishengir-opt %s -auto-blockify-parallel-loop运行，并用 FileCheck 校验：annotation.mark带logical_block_num标记、hivm.hir.get_block_idx生成下界、arith.constant物理块数作为步长、生成的scf.for循环以及autoblockify.subloop属性；
• SIMT 路径：simt-auto-blockify.mlir 覆盖了默认、superblock-factor=1与superblock-factor=2三种场景，校验gpu.linear_block_id的获取、chunk/start/upper的计算、scf.for循环以及 3D 逻辑核 ID 的解码（divui/remui组合）；此外还有 simt-auto-blockify-2.mlir、simt-auto-blockify-multi-block.mlir（多块函数体经scf.execute_region包裹）以及 negative/no-target/silent-skip 等边界用例。

五、总结与适用建议

Auto Blockify 通过在逻辑块与物理块之间引入"外层循环 + 内层物理块索引 + 索引钳位"的映射结构，在物理块远少于逻辑块（典型如 50 vs 500，相差约 10 倍）的场景下，用单波次调度覆盖全部计算索引，从而消除调度开销、实现超过 2 倍的加速。选择与使用该特性时，建议遵循以下要点：

1. SIMD（Vector 算子）路径：Triton 前端设TRITON_ALL_PARALLEL=1，或bishengir-compile加--enable-auto-blockify-loop；bishengir-opt可用--auto-blockify-parallel-loop单独调试；
2. SIMT 路径：同样由--enable-auto-blockify-loop控制，bishengir-opt用--simt-auto-blockify；需要进一步压榨 Vector 算力时，可加--super-block-factor=N（N 为 2 的幂，N×W ≤ 64，默认 1 关闭）；
3. 前置条件：IR 中必须存在logical_block_num标记与hivm get_block_idx操作（Triton 路径由 TritonGlobalKernelArgsToHIVMOpPass 自动补齐），且 kernel 启动的 block 维度需锁定为物理块数；
4. 适用前提：只适用于完全可并行、无块间依赖的 kernel；当逻辑块数很少或单逻辑核负载很大时，收益有限甚至为负，应结合具体算子特性决定是否开启。

• 编译器
• 人工智能
• Ascend
• CANN