القائمة
双内核调用模式的 Triton Ascend 实现示例。展示在 forward 中先后调用两个 kernel 的标准写法：中间结果缓冲区分配、两次 kernel 启动。适用于需要分阶段计算的融合算子。
التثبيت باستخدام Codex أو Claude انسخ هذا Prompt والصقه في Codex أو Claude أو مساعد آخر ليراجع صفحة Skill ويثبّتها لك.
استنادا إلى تصنيف SOC المهني
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 中的类型别名。
| name | triton-ascend-example-double-kernel |
| description | 双内核调用模式的 Triton Ascend 实现示例。展示在 forward 中先后调用两个 kernel 的标准写法：中间结果缓冲区分配、两次 kernel 启动。适用于需要分阶段计算的融合算子。 |
| category | example |
| version | 1.0.0 |
| metadata | {"backend":"ascend","dsl":"triton_ascend","hardware":"Atlas A2, Atlas A3","framework":"torch"} |
当一个算子需要分两步计算（如先做变换再做归约），可在 forward 中依次启动两个 kernel：
class ModelNew(torch.nn.Module):
def __init__(self):
super().__init__()
try:
import torch
import triton
device = torch.npu.current_device()
properties = triton.runtime.driver.active.utils.get_device_properties(device)
self.VEC_CORE_NUM = properties.get("num_vectorcore", 40)
except:
self.VEC_CORE_NUM = 40
def forward(self, x):
intermediate = torch.empty_like(x)
output = torch.empty(out_shape, dtype=x.dtype, device=x.device)
grid = (self.VEC_CORE_NUM,)
kernel_stage1[grid](x, intermediate, ..., CORE_NUM=self.VEC_CORE_NUM)
kernel_stage2[grid](intermediate, output, ..., CORE_NUM=self.VEC_CORE_NUM)
return output
要点：
torch.empty_like 或指定 shape 分配