메뉴
比较两侧本身多为 `tl.int32`（offset / attn_arg 等），`arith.cmpi` 产出的是 `i1`；多段结果在 **`i1` 张量上做 `&`/`|`** 时， lowering 会在每条逻辑附近插入 **`extui`/`trunci`** 与 `select` 对齐；**每段比较后立刻 `.to(tl.int32)`**，让整条 mask 在 **`i32` 0/1 上 `&`/`|`**，后端更易连续处理 **`vand.i32`/`vor.i32`**。
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-case-vector-mask-i32 |
| description | 比较两侧本身多为 `tl.int32`（offset / attn_arg 等），`arith.cmpi` 产出的是 `i1`；多段结果在 **`i1` 张量上做 `&`/`|`** 时， lowering 会在每条逻辑附近插入 **`extui`/`trunci`** 与 `select` 对齐；**每段比较后立刻 `.to(tl.int32)`**，让整条 mask 在 **`i32` 0/1 上 `&`/`|`**，后端更易连续处理 **`vand.i32`/`vor.i32`**。 |
| category | improvement |
| version | 1.0.0 |
| metadata | {"case_type":"improvement","backend":"ascend","dsl":"triton_ascend","hardware":"Atlas A5"} |
tl.int32 参与比较（如 q_off/k_off、q_attn_arg/k_attn_arg），得到的 bool 再用 & / | 拼成 mask，最后交给 tl.where。让 arith.cmpi 的 i1 结果在tensor<…xi1> 上反复 andi/ori，Ascend 后端会在每条逻辑与最后的 arith.select 之间额外插入arith.extui / arith.trunci等，把向量的宽度对齐前后使用的tl.int32，显式 .to(tl.int32)中间 不再在 i1 上对齐宽度，vector 上更易使用指令vand.i32 / vor.i32。
@triton.jit
def mask(...):
triu = (q_off[:, None] <= k_off[None, :]).to(tl.int32)
return (
(triu & ((q_arg[:, None] == k_arg[None, :]).to(tl.int32)
| (k_arg[None, :] == 0).to(tl.int32)))
| (q_off[:, None] == k_off[None, :]).to(tl.int32))
tl.int32 的数据类型来源于比较操作符两边的数据类型。