菜单
Grid/Block 配置策略，包括核数选择、并行度调优、二次切分和大 shape 算子处理方案。适用于需要确定 kernel 启动参数、优化多核并行效率、或处理超大规模数据的内核代码生成场景
用 Codex 或 Claude 帮你安装 复制这段 Prompt，粘贴到 Codex、Claude 或其他助手里，让它检查 Skill 页面并帮你完成安装。
基于 SOC 职业分类
| name | triton-ascend-grid-config |
| description | Grid/Block 配置策略，包括核数选择、并行度调优、二次切分和大 shape 算子处理方案。适用于需要确定 kernel 启动参数、优化多核并行效率、或处理超大规模数据的内核代码生成场景 |
| category | fundamental |
| version | 1.0.0 |
| metadata | {"backend":"ascend","dsl":"triton_ascend","hardware":"Atlas A2, Atlas A3"} |
(x,), (x, y), (x, y, z)适用于按行/按块独立处理的算子（Element-wise、Reduce、Normalization 等）。
@triton.jit
def kernel(
input_ptr, output_ptr, M, N,
stride_m, stride_n,
BLOCK_N: tl.constexpr,
CORE_NUM: tl.constexpr,
):
pid = tl.program_id(0)
# 交错处理：pid=0 处理第 0, CORE_NUM, 2*CORE_NUM, ... 行
for row_idx in range(pid, M, CORE_NUM):
row_ptr = input_ptr + row_idx * stride_m
out_ptr = output_ptr + row_idx * stride_m
for col_start in range(0, N, BLOCK_N):
offs = col_start + tl.arange(0, BLOCK_N)
mask = offs < N
data = tl.load(row_ptr + offs * stride_n, mask=mask)
result = compute(data)
tl.store(out_ptr + offs * stride_n, result, mask=mask)
必须在 __init__ 中获取，禁止在 forward 中调用（触发设备同步）。
import torch
import triton
class ModelNew(torch.nn.Module):
def __init__(self):
super().__init__()
try:
device = torch.npu.current_device()
properties = triton.runtime.driver.active.utils.get_device_properties(device)
self.VEC_CORE_NUM = properties.get("num_vectorcore", 40)
self.CUBE_CORE_NUM = properties.get("num_aicore", 20)
except:
self.VEC_CORE_NUM = 40
self.CUBE_CORE_NUM = 20
def forward(self, x):
M, N = x.shape
out = torch.empty_like(x)
grid = (self.VEC_CORE_NUM,)
kernel[grid](x, out, M, N, x.stride(0), x.stride(1),
BLOCK_N=256, CORE_NUM=self.VEC_CORE_NUM)
return out
VEC_CORE_NUMCUBE_CORE_NUM若 BLOCK_SIZE 超限或单次切分超硬件缓存，可嵌套循环做多层切分：
for m_start in range(pid_m * BLOCK_M, min((pid_m + 1) * BLOCK_M, M), SUB_BLOCK_M):
for n_start in range(pid_n * BLOCK_N, min((pid_n + 1) * BLOCK_N, N), SUB_BLOCK_N):
# 处理 SUB_BLOCK 大小的子块
AscendC direct-invoke 崩溃/挂起修复索引：Kernel timeout、hang、Segmentation Fault、aic error、buffer 死锁、plog/memcheck 调试。
AscendC direct-invoke 精度失败修复索引：输出全 0/随机值、DataCopy 对齐、EnQue/DeQue 同步、FP16/FP32 精度、Cast RoundMode、DumpTensor 分段定位。
AscendC direct-invoke 工程契约：WA 使用 kernel.py + ascendc_op/，ModelNew 调用 torch.ops.npu.*，adapter 负责复制工程、CMake 构建和 npu-arch patch。适用于 dsl=ascendc 的 autoresearch 任务。
把注册式 AscendC 算子迁移为 direct-invoke 工程的保真原则：kernel 算法和 tiling 公式不乱改，只替换注册框架胶水、入口 ABI、host launch 与 PyTorch extension。
CATLASS TileShape 与 on-chip 缓存容量约束：L1/L0A/L0B/L0C 预算公式、fp16/fp32 Pingpong 双缓冲、512B 对齐与排布对 Tile 选型的影响。调参前必读。
CATLASS Gemm 性能调优：DispatchPolicy、Tile 与分核负载均衡、Swizzle、何时用 padding/Split-K/Preload。面向 AR 修改 catlass_kernel.asc 中的类型别名。