feat: preserve debug name hints in ptoas emitc

docs/designs/ptoas-debug-name-hints-design.md

@@ -0,0 +1,286 @@
+# PTOAS 变量名保留与调试名提示设计
+
+## 1. 文档范围
+
+本文定义 issue `#337` 对应的两项能力：
+
+- 前端为 PTO IR 附带“原始变量名”提示信息
+- PTOAS 在 `.pto -> .cpp` 编译链路中尽量沿用该名字，提升可读性与问题定位效率
+
+本文只讨论调试可读性，不改变任何 IR 语义、优化行为或生成代码功能。
+
+## 2. 背景问题
+
+当前 PTOAS 在 `level3` 编译链路中会经过多轮 rewrite、CSE 和 EmitC 降低，最终 `kernel.cpp`
+中的局部变量名通常被重新编号为 `v0`、`v1`、`v2`。
+
+这会带来两个直接问题：
+
+- 输入 `.pto` 与输出 `.cpp` 不容易一一对照
+- 前端 Python 里的业务变量名无法保留下来，定位问题时大量 `vXXX` 可读性很差
+
+issue `#337` 希望解决的是“可对照、可定位”，而不是改变编译结果本身。
+
+## 3. 设计目标
+
+本设计的目标如下：
+
+- 支持前端向 PTO IR 传递可选的变量名提示信息
+- 生成 `kernel.cpp` 时优先使用这些提示信息，减少无意义的 `vXXX`
+- 默认打印 IR 时不明显增加噪声，避免把 IR 变丑
+- 在发生 CSE、控制流合并、lowering 新造临时值时，仍能给出稳定、可理解的名字
+
+## 4. 非目标
+
+本设计明确不保证以下行为：
+
+- 不保证输入 `.pto` 中的 `%v37` 在输出 `.cpp` 中仍然严格叫 `v37`
+- 不保证所有中间优化后仍保留一一对应关系
+- 不为名字建立任何语义约束；名字仅用于调试与阅读
+- 不要求所有前端都必须提供名字提示
+
+原因很直接：当前链路存在 CSE、值合并、值拆分、新值物化与 `emitc.variable` 提升，最终 C++
+中的很多值并不存在与输入 SSA 的严格一一映射关系。
+
+## 5. 总体方案
+
+### 5.1 核心思路
+
+PTOAS 将“原始变量名”视为调试元数据，而不是语义属性。
+
+具体做法：
+
+- 前端把名字提示写入 op 的 `Location` 元数据
+- 对直接输入的 textual `.pto`，PTOAS 额外从源码文本里提取 SSA 名、函数参数名和 block argument 名
+- PTOAS 在 rewrite / lowering 时尽量传播该元数据
+- EmitC/C++ 输出阶段从该元数据和源码文本提示恢复出可用的局部变量名
+
+这样做的核心收益是：
+
+- IR 语义不受影响
+- 默认 textual IR 不需要把一堆 `name_hint` 属性直接打印出来
+- 只有在显式看调试信息时，这些名字才会显式出现
+
+### 5.2 为什么不会让 IR 变丑
+
+本设计不建议把名字直接做成普通 op 属性，比如：
+
+```mlir
+%0 = pto.foo ... { pto.result_name_hints = ["query_tile"] }
+```
+
+这种方案虽然直观，但会让所有 IR 都充满调试字段。
+
+本设计改为把名字放进 `Location`：
+
+- 默认 IR 打印时，location 不会成为主要视觉噪声
+- 只有在显式打开 debug info 打印时，名字才会展示出来
+- 语义属性区保持干净
+
+因此，IR 里是“多带了调试备注”，不是“把主体语法变复杂”。
+
+## 6. IR 名字承载方式
+
+### 6.1 单结果 op
+
+单结果 op 的名字提示记录在 `NameLoc` 或 `FusedLoc` 元数据中。
+
+示意：
+
+```mlir
+%0 = pto.tload ... loc("query_tile")
+```
+
+其含义是：该 op 的主结果推荐名字为 `query_tile`。
+
+### 6.2 多结果 op
+
+多结果 op 不能只靠单个 `NameLoc` 表达所有结果名，因此使用 `FusedLoc` 的 metadata
+携带结果名数组。
+
+逻辑示意：
+
+```text
+loc(fused<metadata=["lhs", "rhs", "acc"]>[...])
+```
+
+这里 metadata 只服务于调试命名，不参与语义。
+
+### 6.3 前端接口约定
+
+前端可以提供名字，也可以不提供。
+
+- 提供时：PTOAS 尽量保留
+- 不提供时：
+  - 若输入是 textual `.pto`，PTOAS 会先尝试回收源码里的 SSA / 参数 / block arg 名
+  - 若仍然没有名字提示，再使用现有 fallback 命名
+
+因此该能力是增量增强，不破坏现有前端。
+
+## 7. 名字传播规则
+
+### 7.1 直接透传
+
+如果一个 rewrite 基本是一对一替换：
+
+- 新值继承原值名字
+
+示例：
+
+- `pto.xxx -> emitc.cast`
+- `pto.xxx -> emitc.variable`
+- `pto.xxx -> 某个一对一的 helper op`
+
+### 7.2 派生命名
+
+如果 lowering 会从一个源值派生多个新值，则在源名字基础上追加稳定后缀。
+
+建议后缀包括：
+
+- `_cast`
+- `_addr`
+- `_tile`
+- `_shape`
+- `_stride`
+- `_tmp`
+
+例如：
+
+- `query_tile -> query_tile_addr`
+- `query_tile -> query_tile_cast`
+
+### 7.3 合并场景
+
+如果多个值被合并成一个值，例如 CSE 或公共表达式复用：
+
+- 优先保留支配值的已有名字
+- 若名字冲突或为空，回退到稳定生成名
+
+### 7.4 控制流与 hoist 场景
+
+`scf` / `cf` / `emitc.variable` 路径会引入额外临时变量，这些变量在源 `.pto` 中通常没有严格对应项。
+
+对这类值：
+
+- 若来源明确，则使用来源名加后缀，如 `_phi`、`_cond`、`_tmp`
+- 若来源不明确，则使用稳定 fallback 名
+
+本次实现里，控制流合流出来的 block argument 还额外做了一层 C++ 末端修正：
+
+- 先尽量把 block argument 名字传播到 EmitC value loc
+- 若 EmitC/C++ hoist 后仍退化成匿名 `vN`，则在函数级 C++ 后处理中把顶部 hoist 出来的匿名合流临时按 block arg 名回填
+
+## 8. C++ 命名规则
+
+### 8.1 基本规则
+
+生成 `kernel.cpp` 时，PTOAS 对每个可命名局部值执行如下规则：
+
+1. 先取调试名字提示
+2. 做 C++ 标识符规范化
+3. 处理关键字冲突
+4. 处理同名冲突
+5. 若仍不可用，则回退到 `vN`
+
+### 8.2 规范化规则
+
+建议规则如下：
+
+- 非 `[A-Za-z0-9_]` 字符替换为 `_`
+- 首字符若不是字母或下划线，则补前缀 `_`
+- C++ 关键字追加后缀 `_v`
+- 连续下划线可压缩
+
+例如：
+
+- `q/k` -> `q_k`
+- `32tmp` -> `_32tmp`
+- `class` -> `class_v`
+
+### 8.3 冲突处理
+
+若同一作用域下出现重名：
+
+- 首个值保留原名
+- 后续值追加 `_1`、`_2`、`_3`
+- 分配前先预留函数里已经存在的参数名、局部声明名和 `using` 引入的标识符，避免重定义
+
+例如：
+
+- `query_tile`
+- `query_tile_1`
+- `query_tile_2`
+
+## 9. 可见性与开关
+
+### 9.1 默认行为
+
+默认情况下：
+
+- IR 语义不变
+- 若前端提供了名字提示，PTOAS 生成 C++ 时优先使用
+- 普通 IR 打印不要求显式展示这些提示
+
+### 9.2 调试打印
+
+若开发者需要查看 IR 中实际承载的调试名字，可通过调试打印模式展示 location。
+
+也就是说：
+
+- 平时看 IR：保持干净
+- 排查名字传播问题时：打开 debug info 看 metadata
+
+## 10. 与现有链路的关系
+
+该设计需要覆盖以下环节：
+
+- 前端生成 PTO IR 时附带名字提示
+- textual `.pto` 输入时，从源码文本恢复 SSA / 参数 / block arg 名
+- PTOAS rewrite / lowering helper 在替换时传播名字
+- `PTOToEmitC` 中新建 `emitc::VariableOp`、`emitc::CastOp` 等值时继承或派生命名
+- 最终 `translateToCpp` 前后在 PTOAS 包装层中应用名字分配规则
+
+特别说明两点：
+
+- 函数参数名没有直接依赖 EmitC 对 block argument loc 的最终打印行为，而是按函数名在最终 C++ 签名上做稳定重写。
+- 这一步是有意的，因为当前 CppEmitter 对参数名和局部 `vN` 命名是统一内部编号逻辑，单靠上游 emitter 不能满足 issue `#337` 的可读性要求。
+
+当前链路里已经存在大量重新物化变量的地方，因此不能只在 parser 或最终 printer 单点处理，必须贯穿传播和最终分配两个阶段。
+
+## 11. 风险与限制
+
+主要风险如下：
+
+- 某些 pass 新建值但没有传播名字，会导致局部退化回 `vN`
+- 多结果 op 的 metadata 约定若不统一，前后端容易理解不一致
+- 名字传播若写成语义属性，容易污染 IR；因此必须坚持“调试元数据”定位
+- C++ 末端重写需要按函数作用域工作；若按整个文件全局替换，多个函数复用 `v1/v2/...` 会串名
+
+限制如下：
+
+- 该设计只能提供“尽量保留”，不能提供“逐号完全一致”
+- 对 aggressive CSE 后的公共值，只能保留最终幸存值的名字
+- 对 textual `.pto` 的 SSA 名恢复依赖源码文本与解析顺序一致；它是调试增强，不是语义机制
+
+## 12. 测试建议
+
+建议至少覆盖以下测试：
+
+- 单结果 op：前端名字能出现在最终 `kernel.cpp`
+- 多结果 op：多个结果名能分别正确落到最终变量
+- 名字含非法字符：能被规范化
+- 重名冲突：能稳定追加后缀
+- 控制流 / `emitc.variable` / hoist：不会因为名字传播导致非法 C++
+- textual `.pto`：函数参数名、局部 SSA 名、CFG block arg 名都能出现在最终 `kernel.cpp`
+- 未提供 hint：仍保持现有 `vN` 回退行为
+
+## 13. 结论
+
+本设计采用“名字作为调试元数据”的路线：
+
+- 不把名字当语义
+- 不要求逐号保真
+- 不把 IR 默认打印搞得很吵
+- 重点解决 `.pto`、前端 Python 和最终 `kernel.cpp` 之间“看得懂、对得上”的问题
+
+这条路线对现有编译链路侵入较小，也最符合 issue `#337` 的真实诉求。

test/lit/pto/async_put_get_emitc.pto

@@ -15,22 +15,22 @@ module {
 }
 
 // A3-LABEL: AICORE void async_put_get(
-// A3: Tile<TileType::Vec, int8_t, 1, 256, BLayout::RowMajor, 1, 256, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[SCRATCH:v[0-9]+]];
-// A3: TASSIGN([[SCRATCH]], [[SCRATCH_ADDR:v[0-9]+]]);
-// A3: pto::comm::AsyncSession [[SESSION:v[0-9]+]];
-// A3: pto::comm::sdma::SdmaBaseConfig [[CFG:v[0-9]+]] = {32768ULL, 0ULL, 1u};
+// A3: Tile<TileType::Vec, int8_t, 1, 256, BLayout::RowMajor, 1, 256, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[SCRATCH:[_A-Za-z][_A-Za-z0-9]*]];
+// A3: TASSIGN([[SCRATCH]], [[SCRATCH_ADDR:[_A-Za-z][_A-Za-z0-9]*]]);
+// A3: pto::comm::AsyncSession [[SESSION:[_A-Za-z][_A-Za-z0-9]*]];
+// A3: pto::comm::sdma::SdmaBaseConfig [[CFG:[_A-Za-z][_A-Za-z0-9]*]] = {32768ULL, 0ULL, 1u};
 // A3: pto::comm::BuildAsyncSession<pto::comm::DmaEngine::SDMA>([[SCRATCH]], {{.*}}, [[SESSION]], {{.*}}, [[CFG]], {{.*}});
 // A3: using [[SHAPETY:.*]] = pto::Shape<1, 1, 1, 1, 128>;
 // A3: using [[STRIDETY:.*]] = pto::Stride<128, 128, 128, 128, 1>;
 // A3: constexpr pto::Layout [[LAYOUT:.*]] = pto::Layout::ND;
-// A3: [[SHAPETY]] [[SHAPE0:v[0-9]+]] = [[SHAPETY]]();
-// A3: [[STRIDETY]] [[STRIDE0:v[0-9]+]] = [[STRIDETY]]();
+// A3: [[SHAPETY]] [[SHAPE0:[_A-Za-z][_A-Za-z0-9]*]] = [[SHAPETY]]();
+// A3: [[STRIDETY]] [[STRIDE0:[_A-Za-z][_A-Za-z0-9]*]] = [[STRIDETY]]();
 // A3: using [[GLTNSRTY:.*]] = GlobalTensor<float, [[SHAPETY]], [[STRIDETY]], [[LAYOUT]]>;
-// A3: [[GLTNSRTY]] [[GT0:v[0-9]+]] = [[GLTNSRTY]]({{.*}}, [[SHAPE0]], [[STRIDE0]]);
-// A3: [[SHAPETY]] [[SHAPE1:v[0-9]+]] = [[SHAPETY]]();
-// A3: [[STRIDETY]] [[STRIDE1:v[0-9]+]] = [[STRIDETY]]();
-// A3: [[GLTNSRTY]] [[GT1:v[0-9]+]] = [[GLTNSRTY]]({{.*}}, [[SHAPE1]], [[STRIDE1]]);
-// A3: pto::comm::AsyncEvent [[PUT_EVT:v[0-9]+]] = pto::comm::TPUT_ASYNC<pto::comm::DmaEngine::SDMA>(
-// A3: pto::comm::AsyncEvent [[GET_EVT:v[0-9]+]] = pto::comm::TGET_ASYNC<pto::comm::DmaEngine::SDMA>(
-// A3: bool [[PUT_DONE:v[0-9]+]] = [[PUT_EVT]].Wait([[SESSION]]);
-// A3: bool [[GET_DONE:v[0-9]+]] = [[GET_EVT]].Test([[SESSION]]);
+// A3: [[GLTNSRTY]] [[GT0:[_A-Za-z][_A-Za-z0-9]*]] = [[GLTNSRTY]]({{.*}}, [[SHAPE0]], [[STRIDE0]]);
+// A3: [[SHAPETY]] [[SHAPE1:[_A-Za-z][_A-Za-z0-9]*]] = [[SHAPETY]]();
+// A3: [[STRIDETY]] [[STRIDE1:[_A-Za-z][_A-Za-z0-9]*]] = [[STRIDETY]]();
+// A3: [[GLTNSRTY]] [[GT1:[_A-Za-z][_A-Za-z0-9]*]] = [[GLTNSRTY]]({{.*}}, [[SHAPE1]], [[STRIDE1]]);
+// A3: pto::comm::AsyncEvent [[PUT_EVT:[_A-Za-z][_A-Za-z0-9]*]] = pto::comm::TPUT_ASYNC<pto::comm::DmaEngine::SDMA>(
+// A3: pto::comm::AsyncEvent [[GET_EVT:[_A-Za-z][_A-Za-z0-9]*]] = pto::comm::TGET_ASYNC<pto::comm::DmaEngine::SDMA>(
+// A3: bool [[PUT_DONE:[_A-Za-z][_A-Za-z0-9]*]] = [[PUT_EVT]].Wait([[SESSION]]);
+// A3: bool [[GET_DONE:[_A-Za-z][_A-Za-z0-9]*]] = [[GET_EVT]].Test([[SESSION]]);

test/lit/pto/backedge_nested_same_pipe_prune_regression.pto

@@ -5,13 +5,13 @@
 // the wider same-pipe pair can be removed before event-id allocation.
 //
 // CHECK-LABEL: __global__ AICORE void backedge_nested_same_pipe_prune()
-// CHECK: for (size_t v16 = (size_t) v1; v16 < ((size_t) v3); v16 += (size_t) v2) {
+// CHECK: for (size_t {{[_A-Za-z][_A-Za-z0-9]*}} = (size_t) {{[_A-Za-z][_A-Za-z0-9]*}}; {{[_A-Za-z][_A-Za-z0-9]*}} < ((size_t) {{[_A-Za-z][_A-Za-z0-9]*}}); {{[_A-Za-z][_A-Za-z0-9]*}} += (size_t) {{[_A-Za-z][_A-Za-z0-9]*}}) {
 // CHECK-NEXT: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID[[INNER:[0-9]+]]);
-// CHECK-NEXT: TMOV(v10, v6);
-// CHECK-NEXT: TMOV(v12, v8);
-// CHECK: TMATMUL(v14, v10, v12);
-// CHECK: TMOV(v8, v14);
-// CHECK-NEXT: TMOV(v6, v14);
+// CHECK-NEXT: TMOV([[LEFT:[_A-Za-z][_A-Za-z0-9]*]], [[M0:[_A-Za-z][_A-Za-z0-9]*]]);
+// CHECK-NEXT: TMOV([[RIGHT:[_A-Za-z][_A-Za-z0-9]*]], [[M1:[_A-Za-z][_A-Za-z0-9]*]]);
+// CHECK: TMATMUL([[ACC:[_A-Za-z][_A-Za-z0-9]*]], [[LEFT]], [[RIGHT]]);
+// CHECK: TMOV([[M1]], [[ACC]]);
+// CHECK-NEXT: TMOV([[M0]], [[ACC]]);
 // CHECK-NEXT: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID[[INNER]]);
 // CHECK-NEXT: }
 // CHECK-NEXT: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID[[INNER]]);

test/lit/pto/debug_name_hints_cfg_emitc.pto

@@ -0,0 +1,36 @@
+// RUN: ptoas %s | FileCheck %s
+
+module {
+  func.func @name_hints_cfg(%cond: i1, %lhs: i32, %rhs: i32) {
+    %0 = func.call @name_hints_cfg_helper(%cond, %lhs, %rhs) : (i1, i32, i32) -> i32
+    return
+  }
+
+  func.func private @name_hints_cfg_helper(%cond: i1, %lhs: i32, %rhs: i32) -> i32 {
+    cf.cond_br %cond, ^bb1, ^bb2
+
+  ^bb1:
+    %then = arith.addi %lhs, %rhs : i32 loc("branch.sum")
+    cf.br ^bb3(%then : i32)
+
+  ^bb2:
+    %else = arith.subi %lhs, %rhs : i32 loc("branch-sum")
+    cf.br ^bb3(%else : i32)
+
+  ^bb3(%merged: i32):
+    %out = arith.addi %merged, %rhs : i32 loc("result")
+    %ret = arith.addi %out, %out : i32
+    return %ret : i32
+  }
+}
+
+// CHECK-LABEL: static AICORE int32_t name_hints_cfg_helper(bool cond, int32_t lhs, int32_t rhs) {
+// CHECK-NOT: PTOAS_NAME_HINTS
+// CHECK: int32_t branch_sum;
+// CHECK: int32_t branch_sum_1;
+// CHECK: int32_t merged;
+// CHECK: branch_sum = (int32_t)
+// CHECK: merged = branch_sum;
+// CHECK: branch_sum_1 = (int32_t)
+// CHECK: merged = branch_sum_1;
+// CHECK: result = (int32_t)