diff --git a/docs/designs/ptoas-debug-name-hints-design.md b/docs/designs/ptoas-debug-name-hints-design.md
new file mode 100644
index 000000000..734530970
--- /dev/null
+++ b/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` 的真实诉求。
diff --git a/test/lit/pto/async_put_get_emitc.pto b/test/lit/pto/async_put_get_emitc.pto
index b191dbd0b..cf2c685cf 100644
--- a/test/lit/pto/async_put_get_emitc.pto
+++ b/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]]);
diff --git a/test/lit/pto/backedge_nested_same_pipe_prune_regression.pto b/test/lit/pto/backedge_nested_same_pipe_prune_regression.pto
index 426f60ecb..8b0444271 100644
--- a/test/lit/pto/backedge_nested_same_pipe_prune_regression.pto
+++ b/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]]);
diff --git a/test/lit/pto/debug_name_hints_cfg_emitc.pto b/test/lit/pto/debug_name_hints_cfg_emitc.pto
new file mode 100644
index 000000000..9fec18d57
--- /dev/null
+++ b/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)
diff --git a/test/lit/pto/debug_name_hints_emitc.pto b/test/lit/pto/debug_name_hints_emitc.pto
new file mode 100644
index 000000000..75313b023
--- /dev/null
+++ b/test/lit/pto/debug_name_hints_emitc.pto
@@ -0,0 +1,23 @@
+// RUN: ptoas %s | FileCheck %s
+
+module {
+  func.func @name_hints_basic(%lhs: i32, %rhs: i32) {
+    %0 = func.call @name_hints_basic_helper(%lhs, %rhs) : (i32, i32) -> i32
+    return
+  }
+
+  func.func private @name_hints_basic_helper(%lhs: i32, %rhs: i32) -> i32 {
+    %sum0 = arith.addi %lhs, %rhs : i32 loc("sum.value")
+    %sum1 = arith.addi %sum0, %sum0 : i32 loc("sum value")
+    %sum2 = arith.addi %sum1, %sum0 : i32 loc("class")
+    %sum3 = arith.addi %sum2, %sum1 : i32
+    %sum4 = arith.addi %sum2, %sum3 : i32
+    return %sum4 : i32
+  }
+}
+
+// CHECK-LABEL: static AICORE int32_t name_hints_basic_helper(int32_t lhs, int32_t rhs) {
+// CHECK-NOT: PTOAS_NAME_HINTS
+// CHECK: int32_t sum_value = (
+// CHECK: int32_t sum_value_1 = (
+// CHECK: int32_t class_v = (
diff --git a/test/lit/pto/debug_name_hints_identifier_collision_emitc.pto b/test/lit/pto/debug_name_hints_identifier_collision_emitc.pto
new file mode 100644
index 000000000..cb937c961
--- /dev/null
+++ b/test/lit/pto/debug_name_hints_identifier_collision_emitc.pto
@@ -0,0 +1,21 @@
+// RUN: ptoas %s | FileCheck %s
+
+module {
+  func.func @identifier_collision(%lhs: i32, %rhs: i32) {
+    %0 = func.call @identifier_collision_helper(%lhs, %rhs) : (i32, i32) -> i32
+    return
+  }
+
+  func.func private @identifier_collision_helper(%lhs: i32, %rhs: i32) -> i32 {
+    %sum = arith.addi %lhs, %rhs : i32 loc("T")
+    %sum2 = arith.addi %sum, %lhs : i32
+    %sum3 = arith.addi %sum2, %sum : i32
+    return %sum3 : i32
+  }
+}
+
+// CHECK-LABEL: static AICORE int32_t identifier_collision_helper(int32_t lhs, int32_t rhs) {
+// CHECK: using T = float;
+// CHECK-NOT: int32_t T = (
+// CHECK: int32_t T_1 = (
+
diff --git a/test/lit/pto/debug_name_hints_textual_cfg_emitc.pto b/test/lit/pto/debug_name_hints_textual_cfg_emitc.pto
new file mode 100644
index 000000000..c93e6c7d0
--- /dev/null
+++ b/test/lit/pto/debug_name_hints_textual_cfg_emitc.pto
@@ -0,0 +1,36 @@
+// RUN: ptoas %s | FileCheck %s
+
+module {
+  func.func @textual_name_hints_cfg(%cond: i1, %lhs: i32, %rhs: i32) {
+    %0 = func.call @textual_name_hints_cfg_helper(%cond, %lhs, %rhs) : (i1, i32, i32) -> i32
+    return
+  }
+
+  func.func private @textual_name_hints_cfg_helper(%cond: i1, %lhs: i32, %rhs: i32) -> i32 {
+    cf.cond_br %cond, ^bb1, ^bb2
+
+  ^bb1:
+    %then_value = arith.addi %lhs, %rhs : i32
+    cf.br ^bb3(%then_value : i32)
+
+  ^bb2:
+    %else_value = arith.subi %lhs, %rhs : i32
+    cf.br ^bb3(%else_value : i32)
+
+  ^bb3(%merged_value: i32):
+    %result_value = arith.addi %merged_value, %rhs : i32
+    return %result_value : i32
+  }
+}
+
+// CHECK-NOT: PTOAS_NAME_HINTS
+// CHECK-NOT: PTOAS_PARAM_HINTS
+// CHECK-LABEL: static AICORE int32_t textual_name_hints_cfg_helper(bool cond, int32_t lhs, int32_t rhs) {
+// CHECK: int32_t then_value;
+// CHECK: int32_t else_value;
+// CHECK: int32_t merged_value;
+// CHECK: then_value = (
+// CHECK: merged_value = then_value;
+// CHECK: else_value = (
+// CHECK: merged_value = else_value;
+// CHECK: return (int32_t) ((uint32_t) merged_value + (uint32_t) rhs);
diff --git a/test/lit/pto/debug_name_hints_textual_emitc.pto b/test/lit/pto/debug_name_hints_textual_emitc.pto
new file mode 100644
index 000000000..0bcd430c5
--- /dev/null
+++ b/test/lit/pto/debug_name_hints_textual_emitc.pto
@@ -0,0 +1,28 @@
+// RUN: ptoas %s | FileCheck %s
+
+module {
+  func.func @textual_name_hints_caller(%lhs: i32, %rhs: i32) {
+    %call_result = func.call @textual_name_hints_add(%lhs, %rhs) : (i32, i32) -> i32
+    %call_result_2 = func.call @textual_name_hints_sub(%lhs, %rhs) : (i32, i32) -> i32
+    return
+  }
+
+  func.func private @textual_name_hints_add(%lhs: i32, %rhs: i32) -> i32 {
+    %sum = arith.addi %lhs, %rhs : i32
+    return %sum : i32
+  }
+
+  func.func private @textual_name_hints_sub(%lhs: i32, %rhs: i32) -> i32 {
+    %diff = arith.subi %lhs, %rhs : i32
+    return %diff : i32
+  }
+}
+
+// CHECK-NOT: PTOAS_NAME_HINTS
+// CHECK-NOT: PTOAS_PARAM_HINTS
+// CHECK-LABEL: static AICORE int32_t textual_name_hints_add(int32_t lhs, int32_t rhs) {
+// CHECK-LABEL: static AICORE int32_t textual_name_hints_sub(int32_t lhs, int32_t rhs) {
+// CHECK-LABEL: AICORE void textual_name_hints_caller(int32_t lhs, int32_t rhs) {
+// CHECK: int32_t call_result = textual_name_hints_add(lhs, rhs);
+// CHECK: int32_t call_result_2 = textual_name_hints_sub(lhs, rhs);
+
diff --git a/test/lit/pto/eventid_array_dyn_sync.pto b/test/lit/pto/eventid_array_dyn_sync.pto
index f8fed9630..b9fa38d0f 100644
--- a/test/lit/pto/eventid_array_dyn_sync.pto
+++ b/test/lit/pto/eventid_array_dyn_sync.pto
@@ -14,10 +14,10 @@ module {
 }
 
 // CHECK-LABEL: __global__ AICORE void eventid_array_dyn_sync() {
-// CHECK: const int32_t {{v[0-9]+}} = 0;
-// CHECK: PTOAS_EventIdArray<4> {{v[0-9]+}};
-// CHECK: {{v[0-9]+}}[{{v[0-9]+}}] = {{v[0-9]+}};
-// CHECK: event_t {{v[0-9]+}} = (event_t) {{v[0-9]+}}[{{v[0-9]+}}];
-// CHECK: set_flag(PIPE_MTE2, PIPE_MTE3, {{v[0-9]+}});
-// CHECK: event_t {{v[0-9]+}} = (event_t) {{v[0-9]+}}[{{v[0-9]+}}];
-// CHECK: wait_flag(PIPE_MTE2, PIPE_MTE3, {{v[0-9]+}});
+// CHECK: const int32_t {{[_A-Za-z][_A-Za-z0-9]*}} = 0;
+// CHECK: PTOAS_EventIdArray<4> {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: {{[_A-Za-z][_A-Za-z0-9]*}}[{{[_A-Za-z][_A-Za-z0-9]*}}] = {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: event_t {{[_A-Za-z][_A-Za-z0-9]*}} = (event_t) {{[_A-Za-z][_A-Za-z0-9]*}}[{{[_A-Za-z][_A-Za-z0-9]*}}];
+// CHECK: set_flag(PIPE_MTE2, PIPE_MTE3, {{[_A-Za-z][_A-Za-z0-9]*}});
+// CHECK: event_t {{[_A-Za-z][_A-Za-z0-9]*}} = (event_t) {{[_A-Za-z][_A-Za-z0-9]*}}[{{[_A-Za-z][_A-Za-z0-9]*}}];
+// CHECK: wait_flag(PIPE_MTE2, PIPE_MTE3, {{[_A-Za-z][_A-Za-z0-9]*}});
diff --git a/test/lit/pto/eventid_array_get_set_get.pto b/test/lit/pto/eventid_array_get_set_get.pto
index 5aa826993..cf6c8da9e 100644
--- a/test/lit/pto/eventid_array_get_set_get.pto
+++ b/test/lit/pto/eventid_array_get_set_get.pto
@@ -17,10 +17,10 @@ module {
 }
 
 // CHECK-LABEL: __global__ AICORE void eventid_array_get_set_get() {
-// CHECK: PTOAS_EventIdArray<4> [[ARR:v[0-9]+]];
-// CHECK: [[ARR]][{{v[0-9]+}}] = {{v[0-9]+}};
-// CHECK: [[ARR]][{{v[0-9]+}}] = {{v[0-9]+}};
-// CHECK: event_t [[FIRST:v[0-9]+]] = (event_t) [[ARR]][{{v[0-9]+}}];
+// CHECK: PTOAS_EventIdArray<4> [[ARR:[_A-Za-z][_A-Za-z0-9]*]];
+// CHECK: [[ARR]][{{[_A-Za-z][_A-Za-z0-9]*}}] = {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: [[ARR]][{{[_A-Za-z][_A-Za-z0-9]*}}] = {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: event_t [[FIRST:[_A-Za-z][_A-Za-z0-9]*]] = (event_t) [[ARR]][{{[_A-Za-z][_A-Za-z0-9]*}}];
 // CHECK: set_flag(PIPE_MTE2, PIPE_MTE3, [[FIRST]]);
-// CHECK: event_t [[SECOND:v[0-9]+]] = (event_t) [[ARR]][{{v[0-9]+}}];
+// CHECK: event_t [[SECOND:[_A-Za-z][_A-Za-z0-9]*]] = (event_t) [[ARR]][{{[_A-Za-z][_A-Za-z0-9]*}}];
 // CHECK: wait_flag(PIPE_MTE2, PIPE_MTE3, [[SECOND]]);
diff --git a/test/lit/pto/eventid_array_no_cse.pto b/test/lit/pto/eventid_array_no_cse.pto
index eecea2ab1..d18f4f2d6 100644
--- a/test/lit/pto/eventid_array_no_cse.pto
+++ b/test/lit/pto/eventid_array_no_cse.pto
@@ -18,7 +18,7 @@ module {
 }
 
 // CHECK-LABEL: __global__ AICORE void eventid_array_no_cse() {
-// CHECK: PTOAS_EventIdArray<4> [[ARR0:v[0-9]+]];
-// CHECK: PTOAS_EventIdArray<4> [[ARR1:v[0-9]+]];
-// CHECK: [[ARR0]][{{v[0-9]+}}] = {{v[0-9]+}};
-// CHECK: [[ARR1]][{{v[0-9]+}}] = {{v[0-9]+}};
+// CHECK: PTOAS_EventIdArray<4> [[ARR0:[_A-Za-z][_A-Za-z0-9]*]];
+// CHECK: PTOAS_EventIdArray<4> [[ARR1:[_A-Za-z][_A-Za-z0-9]*]];
+// CHECK: [[ARR0]][{{[_A-Za-z][_A-Za-z0-9]*}}] = {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: [[ARR1]][{{[_A-Za-z][_A-Za-z0-9]*}}] = {{[_A-Za-z][_A-Za-z0-9]*}};
diff --git a/test/lit/pto/issue428_cube_sync_regression.pto b/test/lit/pto/issue428_cube_sync_regression.pto
index 9f129c670..69a5d49ea 100644
--- a/test/lit/pto/issue428_cube_sync_regression.pto
+++ b/test/lit/pto/issue428_cube_sync_regression.pto
@@ -5,98 +5,98 @@
 // - Preserve the tail drain before the auto tail barrier helper.
 //
 // CHECK-LABEL: tri_inv_block2x2_fp16(
-// CHECK: for (size_t v71 = v20; v71 < v22; v71 += v19) {
+// CHECK: for (size_t {{[_A-Za-z][_A-Za-z0-9]*}} = {{[_A-Za-z][_A-Za-z0-9]*}}; {{[_A-Za-z][_A-Za-z0-9]*}} < {{[_A-Za-z][_A-Za-z0-9]*}}; {{[_A-Za-z][_A-Za-z0-9]*}} += {{[_A-Za-z][_A-Za-z0-9]*}}) {
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID2);
-// CHECK: TMOV(v65, v49);
-// CHECK: TMOV(v67, v61);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID2);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID2);
 // CHECK: wait_flag(PIPE_FIX, PIPE_M, EVENT_ID2);
-// CHECK: TMATMUL(v69, v65, v67);
+// CHECK: TMATMUL({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_MTE1, EVENT_ID3);
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID3);
-// CHECK: TMOV(v67, v51);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID3);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID3);
-// CHECK: TMATMUL_ACC(v69, v69, v65, v67);
+// CHECK: TMATMUL_ACC({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_MTE1, EVENT_ID4);
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID2);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID2);
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID4);
-// CHECK: if ((int32_t) ((uint32_t) ((int32_t) v71) + (uint32_t) v8) < v21) {
-// CHECK: TMOV(v49, v69);
+// CHECK: if ((int32_t) ((uint32_t) ((int32_t) {{[_A-Za-z][_A-Za-z0-9]*}}) + (uint32_t) {{[_A-Za-z][_A-Za-z0-9]*}}) < {{[_A-Za-z][_A-Za-z0-9]*}}) {
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_FIX, PIPE_M, EVENT_ID3);
-// CHECK: TMOV(v65, v51);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID4);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID4);
 // CHECK: wait_flag(PIPE_FIX, PIPE_M, EVENT_ID3);
-// CHECK: TMATMUL(v69, v65, v67);
+// CHECK: TMATMUL({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID3);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID3);
-// CHECK: TMOV(v51, v69);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_MTE1, EVENT_ID2);
 // CHECK: set_flag(PIPE_FIX, PIPE_M, EVENT_ID2);
 // CHECK: }
-// CHECK: for (size_t v72 = v20; v72 < v22; v72 += v19) {
+// CHECK: for (size_t {{[_A-Za-z][_A-Za-z0-9]*}} = {{[_A-Za-z][_A-Za-z0-9]*}}; {{[_A-Za-z][_A-Za-z0-9]*}} < {{[_A-Za-z][_A-Za-z0-9]*}}; {{[_A-Za-z][_A-Za-z0-9]*}} += {{[_A-Za-z][_A-Za-z0-9]*}}) {
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID7);
-// CHECK: TMOV(v65, v53);
-// CHECK: TMOV(v67, v61);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID6);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID6);
 // CHECK: wait_flag(PIPE_FIX, PIPE_M, EVENT_ID6);
-// CHECK: TMATMUL(v69, v65, v67);
+// CHECK: TMATMUL({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_MTE1, EVENT_ID0);
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID0);
-// CHECK: TMOV(v67, v55);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID7);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID7);
-// CHECK: TMATMUL_ACC(v69, v69, v65, v67);
+// CHECK: TMATMUL_ACC({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_MTE1, EVENT_ID0);
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID6);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID6);
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID0);
-// CHECK: if ((int32_t) ((uint32_t) ((int32_t) v72) + (uint32_t) v8) < v21) {
-// CHECK: TMOV(v53, v69);
+// CHECK: if ((int32_t) ((uint32_t) ((int32_t) {{[_A-Za-z][_A-Za-z0-9]*}}) + (uint32_t) {{[_A-Za-z][_A-Za-z0-9]*}}) < {{[_A-Za-z][_A-Za-z0-9]*}}) {
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_FIX, PIPE_M, EVENT_ID7);
-// CHECK: TMOV(v65, v55);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
 // CHECK: wait_flag(PIPE_FIX, PIPE_M, EVENT_ID7);
-// CHECK: TMATMUL(v69, v65, v67);
+// CHECK: TMATMUL({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID7);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID7);
-// CHECK: TMOV(v55, v69);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_MTE1, EVENT_ID7);
 // CHECK: set_flag(PIPE_FIX, PIPE_M, EVENT_ID6);
 // CHECK: }
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
-// CHECK: TMOV(v53, v69);
-// CHECK: TSTORE(v48, v69);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
+// CHECK: TSTORE({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_FIX, PIPE_M, EVENT_ID0);
-// CHECK: TLOAD(v57, v36);
+// CHECK: TLOAD({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID3);
 // CHECK: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0);
-// CHECK: TMOV(v65, v53);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID3);
-// CHECK: TMOV(v67, v57);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
 // CHECK: wait_flag(PIPE_FIX, PIPE_M, EVENT_ID0);
-// CHECK: TMATMUL(v69, v65, v67);
+// CHECK: TMATMUL({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
-// CHECK: TMOV(v63, v69);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID1);
 // CHECK: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID1);
-// CHECK: TMOV(v65, v63);
-// CHECK: TMOV(v67, v49);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
-// CHECK: TMATMUL(v69, v65, v67);
+// CHECK: TMATMUL({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
-// CHECK: TSTORE(v45, v69);
+// CHECK: TSTORE({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID2);
 // CHECK: wait_flag(PIPE_FIX, PIPE_M, EVENT_ID2);
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID7);
diff --git a/test/lit/pto/issue454_loop_if_else_loop_carried_sync_regression.pto b/test/lit/pto/issue454_loop_if_else_loop_carried_sync_regression.pto
index 35e66d925..25ed502b9 100644
--- a/test/lit/pto/issue454_loop_if_else_loop_carried_sync_regression.pto
+++ b/test/lit/pto/issue454_loop_if_else_loop_carried_sync_regression.pto
@@ -6,27 +6,27 @@
 //
 // CHECK-LABEL: __global__ AICORE void loop_if_else_loop_carried_sync()
 // CHECK: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID1);
-// CHECK: TMOV(v10, v6);
-// CHECK: TMOV(v12, v8);
+// CHECK: TMOV([[LEFT:[_A-Za-z][_A-Za-z0-9]*]], [[M0:[_A-Za-z][_A-Za-z0-9]*]]);
+// CHECK: TMOV([[RIGHT:[_A-Za-z][_A-Za-z0-9]*]], [[M1:[_A-Za-z][_A-Za-z0-9]*]]);
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
-// CHECK: TMATMUL(v14, v10, v12);
+// CHECK: TMATMUL([[ACC:[_A-Za-z][_A-Za-z0-9]*]], [[LEFT]], [[RIGHT]]);
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
-// CHECK: TMOV(v6, v14);
+// CHECK: TMOV([[M0]], [[ACC]]);
 // CHECK: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0);
 // CHECK: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID0);
-// 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: wait_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID1);
-// CHECK: TMOV(v10, v6);
+// CHECK: TMOV([[LEFT]], [[M0]]);
 // CHECK: set_flag(PIPE_MTE1, PIPE_M, EVENT_ID1);
 // CHECK: wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID1);
-// CHECK: TMATMUL(v14, v10, v12);
+// CHECK: TMATMUL([[ACC]], [[LEFT]], [[RIGHT]]);
 // CHECK: set_flag(PIPE_M, PIPE_MTE1, EVENT_ID0);
 // CHECK: set_flag(PIPE_M, PIPE_FIX, EVENT_ID1);
 // CHECK: wait_flag(PIPE_M, PIPE_FIX, EVENT_ID1);
 // CHECK: wait_flag(PIPE_M, PIPE_MTE1, EVENT_ID0);
-// CHECK: if ((int32_t) v16 == v1) {
+// CHECK: if ((int32_t) {{[_A-Za-z][_A-Za-z0-9]*}} == {{[_A-Za-z][_A-Za-z0-9]*}}) {
 // CHECK: } else {
 // CHECK: set_flag(PIPE_FIX, PIPE_MTE1, EVENT_ID1);
 // CHECK: ptoas_auto_sync_tail(PTOAutoSyncTailMode::kBarrierAll);
diff --git a/test/lit/pto/issue481_addptr_gm_slot_buffer.pto b/test/lit/pto/issue481_addptr_gm_slot_buffer.pto
index f609a7cde..5bbe4a628 100644
--- a/test/lit/pto/issue481_addptr_gm_slot_buffer.pto
+++ b/test/lit/pto/issue481_addptr_gm_slot_buffer.pto
@@ -34,13 +34,13 @@ module {
   }
 }
 
-// CHECK-LABEL: AICORE void cube_kernel(__gm__ float* {{v[0-9]+}}, __gm__ float* {{v[0-9]+}})
-// CHECK: int64_t {{v[0-9]+}} = get_block_idx();
-// CHECK: __gm__ float* {{v[0-9]+}} = {{v[0-9]+}} + {{v[0-9]+}};
-// CHECK: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
+// CHECK-LABEL: AICORE void cube_kernel(__gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}}, __gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}})
+// CHECK: int64_t {{[_A-Za-z][_A-Za-z0-9]*}} = get_block_idx();
+// CHECK: __gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}} = {{[_A-Za-z][_A-Za-z0-9]*}} + {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
 
-// CHECK-LABEL: AICORE void vector_kernel(__gm__ float* {{v[0-9]+}}, __gm__ float* {{v[0-9]+}})
-// CHECK: int64_t {{v[0-9]+}} = get_block_idx();
-// CHECK: __gm__ float* {{v[0-9]+}} = {{v[0-9]+}} + {{v[0-9]+}};
-// CHECK: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
+// CHECK-LABEL: AICORE void vector_kernel(__gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}}, __gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}})
+// CHECK: int64_t {{[_A-Za-z][_A-Za-z0-9]*}} = get_block_idx();
+// CHECK: __gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}} = {{[_A-Za-z][_A-Za-z0-9]*}} + {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
 // CHECK-NOT: pto.addptr
diff --git a/test/lit/pto/issue481_addptr_gm_slot_buffer_gss.pto b/test/lit/pto/issue481_addptr_gm_slot_buffer_gss.pto
index f01256606..ffd82412a 100644
--- a/test/lit/pto/issue481_addptr_gm_slot_buffer_gss.pto
+++ b/test/lit/pto/issue481_addptr_gm_slot_buffer_gss.pto
@@ -34,13 +34,13 @@ module {
   }
 }
 
-// CHECK-LABEL: AICORE void cube_kernel(__gm__ float* {{v[0-9]+}}, __gm__ float* {{v[0-9]+}})
-// CHECK: int64_t {{v[0-9]+}} = get_block_idx();
-// CHECK: __gm__ float* {{v[0-9]+}} = {{v[0-9]+}} + {{v[0-9]+}};
-// CHECK: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
+// CHECK-LABEL: AICORE void cube_kernel(__gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}}, __gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}})
+// CHECK: int64_t {{[_A-Za-z][_A-Za-z0-9]*}} = get_block_idx();
+// CHECK: __gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}} = {{[_A-Za-z][_A-Za-z0-9]*}} + {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
 
-// CHECK-LABEL: AICORE void vector_kernel(__gm__ float* {{v[0-9]+}}, __gm__ float* {{v[0-9]+}})
-// CHECK: int64_t {{v[0-9]+}} = get_block_idx();
-// CHECK: __gm__ float* {{v[0-9]+}} = {{v[0-9]+}} + {{v[0-9]+}};
-// CHECK: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
+// CHECK-LABEL: AICORE void vector_kernel(__gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}}, __gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}})
+// CHECK: int64_t {{[_A-Za-z][_A-Za-z0-9]*}} = get_block_idx();
+// CHECK: __gm__ float* {{[_A-Za-z][_A-Za-z0-9]*}} = {{[_A-Za-z][_A-Za-z0-9]*}} + {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
 // CHECK-NOT: pto.addptr
diff --git a/test/lit/pto/issue556_tpop_live_values_no_alias.pto b/test/lit/pto/issue556_tpop_live_values_no_alias.pto
index de1734a5c..5ac1e26b4 100644
--- a/test/lit/pto/issue556_tpop_live_values_no_alias.pto
+++ b/test/lit/pto/issue556_tpop_live_values_no_alias.pto
@@ -57,9 +57,9 @@ module {
 }
 
 // CHECK-LABEL: AICORE void issue556_tpop_live_values_no_alias(
-// CHECK: Tile<TileType::Mat, half, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[POP0:v[0-9]+]];
-// CHECK: TPOP<TPipe<0, Direction::DIR_V2C, 512, 8, 8, false>, Tile<TileType::Mat, half, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_UP_DOWN>({{v[0-9]+}}, [[POP0]]);
-// CHECK: Tile<TileType::Mat, half, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[POP1:v[0-9]+]];
-// CHECK: TPOP<TPipe<0, Direction::DIR_V2C, 512, 8, 8, false>, Tile<TileType::Mat, half, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_UP_DOWN>({{v[0-9]+}}, [[POP1]]);
-// CHECK: TMOV({{v[0-9]+}}, [[POP0]]);
-// CHECK: TMOV({{v[0-9]+}}, [[POP1]]);
+// CHECK: Tile<TileType::Mat, half, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[POP0:[_A-Za-z][_A-Za-z0-9]*]];
+// CHECK: TPOP<TPipe<0, Direction::DIR_V2C, 512, 8, 8, false>, Tile<TileType::Mat, half, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_UP_DOWN>({{[_A-Za-z][_A-Za-z0-9]*}}, [[POP0]]);
+// CHECK: Tile<TileType::Mat, half, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[POP1:[_A-Za-z][_A-Za-z0-9]*]];
+// CHECK: TPOP<TPipe<0, Direction::DIR_V2C, 512, 8, 8, false>, Tile<TileType::Mat, half, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_UP_DOWN>({{[_A-Za-z][_A-Za-z0-9]*}}, [[POP1]]);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, [[POP0]]);
+// CHECK: TMOV({{[_A-Za-z][_A-Za-z0-9]*}}, [[POP1]]);
diff --git a/test/lit/pto/issue660_trowexpandmul_set_validshape_preserves_alloc_valid.pto b/test/lit/pto/issue660_trowexpandmul_set_validshape_preserves_alloc_valid.pto
index 313e1c6a7..d719ca2e9 100644
--- a/test/lit/pto/issue660_trowexpandmul_set_validshape_preserves_alloc_valid.pto
+++ b/test/lit/pto/issue660_trowexpandmul_set_validshape_preserves_alloc_valid.pto
@@ -39,9 +39,9 @@ module {
   }
 }
 
-// CHECK-DAG: const int32_t [[ROW:v[0-9]+]] = 16;
-// CHECK-DAG: const int32_t [[FULL_COL:v[0-9]+]] = 256;
-// CHECK: Tile<TileType::Vec, float, 16, 256, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[DST:v[0-9]+]] = Tile<TileType::Vec, float, 16, 256, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>([[ROW]], [[FULL_COL]]);
+// CHECK-DAG: const int32_t [[ROW:[_A-Za-z][_A-Za-z0-9]*]] = 16;
+// CHECK-DAG: const int32_t [[FULL_COL:[_A-Za-z][_A-Za-z0-9]*]] = 256;
+// CHECK: Tile<TileType::Vec, float, 16, 256, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[DST:[_A-Za-z][_A-Za-z0-9]*]] = Tile<TileType::Vec, float, 16, 256, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>([[ROW]], [[FULL_COL]]);
 // CHECK: TROWEXPANDMUL([[DST]],
 // CHECK: [[DST]].SetValidShape([[ROW]], [[ROW]]);
-// CHECK: TFILLPAD({{v[0-9]+}}, [[DST]]);
+// CHECK: TFILLPAD({{[_A-Za-z][_A-Za-z0-9]*}}, [[DST]]);
diff --git a/test/lit/pto/local_array_1d_emitc.pto b/test/lit/pto/local_array_1d_emitc.pto
index a423f6f0c..448fb09cc 100644
--- a/test/lit/pto/local_array_1d_emitc.pto
+++ b/test/lit/pto/local_array_1d_emitc.pto
@@ -23,6 +23,6 @@ module {
 }
 
 // CHECK-LABEL: local_array_1d
-// CHECK: int32_t [[A:v[0-9]+]][16];
-// CHECK: [[A]][{{v[0-9]+}}] = {{v[0-9]+}};
-// CHECK: [[A]][{{v[0-9]+}}] = [[A]][{{v[0-9]+}}];
+// CHECK: int32_t [[A:[_A-Za-z][_A-Za-z0-9]*]][16];
+// CHECK: [[A]][{{[_A-Za-z][_A-Za-z0-9]*}}] = {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: [[A]][{{[_A-Za-z][_A-Za-z0-9]*}}] = [[A]][{{[_A-Za-z][_A-Za-z0-9]*}}];
diff --git a/test/lit/pto/local_array_2d_emitc.pto b/test/lit/pto/local_array_2d_emitc.pto
index 7d419176a..81dcbfd7a 100644
--- a/test/lit/pto/local_array_2d_emitc.pto
+++ b/test/lit/pto/local_array_2d_emitc.pto
@@ -23,6 +23,6 @@ module {
 }
 
 // CHECK-LABEL: local_array_2d
-// CHECK: float [[M:v[0-9]+]][8][8];
-// CHECK: [[M]][{{v[0-9]+}}][{{v[0-9]+}}] = {{v[0-9]+}};
-// CHECK: [[M]][{{v[0-9]+}}][{{v[0-9]+}}] = [[M]][{{v[0-9]+}}][{{v[0-9]+}}];
+// CHECK: float [[M:[_A-Za-z][_A-Za-z0-9]*]][8][8];
+// CHECK: [[M]][{{[_A-Za-z][_A-Za-z0-9]*}}][{{[_A-Za-z][_A-Za-z0-9]*}}] = {{[_A-Za-z][_A-Za-z0-9]*}};
+// CHECK: [[M]][{{[_A-Za-z][_A-Za-z0-9]*}}][{{[_A-Za-z][_A-Za-z0-9]*}}] = [[M]][{{[_A-Za-z][_A-Za-z0-9]*}}][{{[_A-Za-z][_A-Za-z0-9]*}}];
diff --git a/test/lit/pto/local_array_get_rvalue_emitc.pto b/test/lit/pto/local_array_get_rvalue_emitc.pto
index e0d020c66..3f2e4621a 100644
--- a/test/lit/pto/local_array_get_rvalue_emitc.pto
+++ b/test/lit/pto/local_array_get_rvalue_emitc.pto
@@ -24,8 +24,8 @@ module {
 }
 
 // CHECK-LABEL: local_array_get_rvalue
-// CHECK: int32_t [[A:v[0-9]+]][16];
+// CHECK: int32_t [[A:[_A-Za-z][_A-Za-z0-9]*]][16];
 // `a[i] = ...a[j]... + ...;` — confirms the get result flows through arith as
 // a scalar rvalue. The exact cast/paren shape comes from arith.addi's signless
 // overflow-aware lowering; we only pin the subscript-rvalue flow here.
-// CHECK: [[A]][{{v[0-9]+}}] = {{.*}}[[A]][{{v[0-9]+}}]{{.*}};
+// CHECK: [[A]][{{[_A-Za-z][_A-Za-z0-9]*}}] = {{.*}}[[A]][{{[_A-Za-z][_A-Za-z0-9]*}}]{{.*}};
diff --git a/test/lit/pto/set_validshape_if.pto b/test/lit/pto/set_validshape_if.pto
index 671b43d4b..ee4be32c1 100644
--- a/test/lit/pto/set_validshape_if.pto
+++ b/test/lit/pto/set_validshape_if.pto
@@ -31,12 +31,12 @@ module {
   }
 }
 
-// CHECK-DAG: const int32_t [[C32:v[0-9]+]] = 32;
-// CHECK: Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[TILE:v[0-9]+]] = Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>([[C32]], [[C32]]);
+// CHECK-DAG: const int32_t [[C32:[_A-Za-z][_A-Za-z0-9]*]] = 32;
+// CHECK: Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[TILE:[_A-Za-z][_A-Za-z0-9]*]] = Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>([[C32]], [[C32]]);
 // CHECK: TASSIGN([[TILE]],
 // CHECK: if (
-// CHECK: [[TILE]].SetValidShape([[ROW1:v[0-9]+]], [[COL1:v[0-9]+]])
+// CHECK: [[TILE]].SetValidShape([[ROW1:[_A-Za-z][_A-Za-z0-9]*]], [[COL1:[_A-Za-z][_A-Za-z0-9]*]])
 // CHECK: } else {
-// CHECK: [[TILE]].SetValidShape([[ROW2:v[0-9]+]], [[COL2:v[0-9]+]])
+// CHECK: [[TILE]].SetValidShape([[ROW2:[_A-Za-z][_A-Za-z0-9]*]], [[COL2:[_A-Za-z][_A-Za-z0-9]*]])
 // CHECK: }
 // CHECK: TADD([[TILE]], [[TILE]], [[TILE]]);
diff --git a/test/lit/pto/set_validshape_local_lowering.pto b/test/lit/pto/set_validshape_local_lowering.pto
index 3adbc248e..c72c92b58 100644
--- a/test/lit/pto/set_validshape_local_lowering.pto
+++ b/test/lit/pto/set_validshape_local_lowering.pto
@@ -15,7 +15,7 @@ module {
   }
 }
 
-// CHECK-DAG: const int32_t [[C32:v[0-9]+]] = 32;
-// CHECK: Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[TILE:v[0-9]+]] = Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>([[C32]], [[C32]]);
-// CHECK: TASSIGN([[TILE]], [[ADDR:v[0-9]+]]);
-// CHECK: [[TILE]].SetValidShape([[ROW:v[0-9]+]], [[COL:v[0-9]+]])
+// CHECK-DAG: const int32_t [[C32:[_A-Za-z][_A-Za-z0-9]*]] = 32;
+// CHECK: Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[TILE:[_A-Za-z][_A-Za-z0-9]*]] = Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>([[C32]], [[C32]]);
+// CHECK: TASSIGN([[TILE]], [[ADDR:[_A-Za-z][_A-Za-z0-9]*]]);
+// CHECK: [[TILE]].SetValidShape([[ROW:[_A-Za-z][_A-Za-z0-9]*]], [[COL:[_A-Za-z][_A-Za-z0-9]*]])
diff --git a/test/lit/pto/tassign_level3_loop_rebind.pto b/test/lit/pto/tassign_level3_loop_rebind.pto
index d947bc60a..3d8428155 100644
--- a/test/lit/pto/tassign_level3_loop_rebind.pto
+++ b/test/lit/pto/tassign_level3_loop_rebind.pto
@@ -35,7 +35,7 @@ module {
 }
 
 // CHECK-LABEL: __global__ AICORE void tassign_loop_rebind() {
-// CHECK: Tile<TileType::Vec, half, 32, 32, BLayout::RowMajor, 32, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[T:v[0-9]+]];
+// CHECK: Tile<TileType::Vec, half, 32, 32, BLayout::RowMajor, 32, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[T:[_A-Za-z][_A-Za-z0-9]*]];
 // CHECK: for (
 // CHECK: TASSIGN([[T]],
 // CHECK: TPRINT([[T]]);
diff --git a/test/lit/pto/tassign_level3_loop_rebind_gss.pto b/test/lit/pto/tassign_level3_loop_rebind_gss.pto
index 9c07f4c1d..562691003 100644
--- a/test/lit/pto/tassign_level3_loop_rebind_gss.pto
+++ b/test/lit/pto/tassign_level3_loop_rebind_gss.pto
@@ -35,7 +35,7 @@ module {
 }
 
 // CHECK-LABEL: __global__ AICORE void tassign_loop_rebind() {
-// CHECK: Tile<TileType::Vec, half, 32, 32, BLayout::RowMajor, 32, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[T:v[0-9]+]];
+// CHECK: Tile<TileType::Vec, half, 32, 32, BLayout::RowMajor, 32, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[T:[_A-Za-z][_A-Za-z0-9]*]];
 // CHECK: for (
 // CHECK: TASSIGN([[T]],
 // CHECK: TPRINT([[T]]);
diff --git a/test/lit/pto/tdivs_dual_order_emitc.pto b/test/lit/pto/tdivs_dual_order_emitc.pto
index e036a1be8..84c70d0d7 100644
--- a/test/lit/pto/tdivs_dual_order_emitc.pto
+++ b/test/lit/pto/tdivs_dual_order_emitc.pto
@@ -16,5 +16,5 @@ module {
   }
 }
 
-// A3: TDIVS([[VDST:v[0-9]+]], [[VSRC:v[0-9]+]], [[VSCALAR:v[0-9]+]]);
+// A3: TDIVS([[VDST:[_A-Za-z][_A-Za-z0-9]*]], [[VSRC:[_A-Za-z][_A-Za-z0-9]*]], [[VSCALAR:[_A-Za-z][_A-Za-z0-9]*]]);
 // A3: TDIVS([[VDST]], [[VSCALAR]], [[VSRC]]);
diff --git a/test/lit/pto/tdivs_dual_order_emitc_gss.pto b/test/lit/pto/tdivs_dual_order_emitc_gss.pto
index bba9b3f84..716e07530 100644
--- a/test/lit/pto/tdivs_dual_order_emitc_gss.pto
+++ b/test/lit/pto/tdivs_dual_order_emitc_gss.pto
@@ -16,5 +16,5 @@ module {
   }
 }
 
-// A3: TDIVS([[VDST:v[0-9]+]], [[VSRC:v[0-9]+]], [[VSCALAR:v[0-9]+]]);
+// A3: TDIVS([[VDST:[_A-Za-z][_A-Za-z0-9]*]], [[VSRC:[_A-Za-z][_A-Za-z0-9]*]], [[VSCALAR:[_A-Za-z][_A-Za-z0-9]*]]);
 // A3: TDIVS([[VDST]], [[VSCALAR]], [[VSRC]]);
diff --git a/test/lit/pto/tgather_three_forms_emitc.pto b/test/lit/pto/tgather_three_forms_emitc.pto
index 56a6c4128..d96ff714e 100644
--- a/test/lit/pto/tgather_three_forms_emitc.pto
+++ b/test/lit/pto/tgather_three_forms_emitc.pto
@@ -21,9 +21,9 @@ module {
   }
 }
 
-// A3: TGATHER({{v[0-9]+}}, {{v[0-9]+}}, {{v[0-9]+}}, {{v[0-9]+}});
+// A3: TGATHER({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // A3: TGATHER<Tile<TileType::Vec, half, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, Tile<TileType::Vec, half, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, MaskPattern::P1111>(
 // A3-NOT: reinterpret_cast<
 // A3-NOT: TGATHER<Tile<TileType::Vec, int32_t, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, Tile<TileType::Vec, half, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, Tile<TileType::Vec, uint8_t, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, Tile<TileType::Vec, int32_t, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, CmpMode::EQ, 7>(
 // A3: TGATHER<Tile<TileType::Vec, int32_t, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, Tile<TileType::Vec, half, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, Tile<TileType::Vec, uint8_t, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, Tile<TileType::Vec, int32_t, 1, 32, BLayout::RowMajor, 1, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, CmpMode::EQ>(
-// A3-SAME: {{v[0-9]+}}, {{v[0-9]+}}, {{v[0-9]+}}, {{v[0-9]+}}, {{v[0-9]+}}, {{v[0-9]+}});
+// A3-SAME: {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
diff --git a/test/lit/pto/tgetval_level3_index_cast_regression.pto b/test/lit/pto/tgetval_level3_index_cast_regression.pto
index 7ae49373a..c6145c29c 100644
--- a/test/lit/pto/tgetval_level3_index_cast_regression.pto
+++ b/test/lit/pto/tgetval_level3_index_cast_regression.pto
@@ -45,5 +45,5 @@ module {
 }
 
 // CHECK-LABEL: __global__ AICORE void tgetval_level3_index_cast_regression() {
-// CHECK: int32_t [[ROW:v[0-9]+]] = [[IDX:v[0-9]+]].GetValue(
+// CHECK: int32_t [[ROW:[_A-Za-z][_A-Za-z0-9]*]] = [[IDX:[_A-Za-z][_A-Za-z0-9]*]].GetValue(
 // CHECK: TINSERT(
diff --git a/test/lit/pto/tile_compact_mode_emitc.pto b/test/lit/pto/tile_compact_mode_emitc.pto
index b1ba469a6..f03d81b79 100644
--- a/test/lit/pto/tile_compact_mode_emitc.pto
+++ b/test/lit/pto/tile_compact_mode_emitc.pto
@@ -28,6 +28,6 @@ module {
 // A3-DAG: memref.alloc() : memref<1x16xf16, strided<[16, 1]>, #pto.address_space<vec>>
 // A3-DAG: memref.alloc() : memref<1x16xf16, strided<[16, 1]>, #pto.address_space<vec>>
 // A3-DAG: memref.alloc() : memref<1x16xf16, strided<[17, 1]>, #pto.address_space<vec>>
-// A3: Tile<TileType::Vec, half, 1, 16, BLayout::RowMajor, 1, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[DEFAULT:v[0-9]+]];
-// A3: Tile<TileType::Vec, half, 1, 16, BLayout::RowMajor, 1, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Normal> [[COMPACT:v[0-9]+]];
-// A3: Tile<TileType::Vec, half, 1, 16, BLayout::RowMajor, 1, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::RowPlusOne> [[ROWP1:v[0-9]+]];
+// A3: Tile<TileType::Vec, half, 1, 16, BLayout::RowMajor, 1, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[DEFAULT:[_A-Za-z][_A-Za-z0-9]*]];
+// A3: Tile<TileType::Vec, half, 1, 16, BLayout::RowMajor, 1, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Normal> [[COMPACT:[_A-Za-z][_A-Za-z0-9]*]];
+// A3: Tile<TileType::Vec, half, 1, 16, BLayout::RowMajor, 1, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::RowPlusOne> [[ROWP1:[_A-Za-z][_A-Za-z0-9]*]];
diff --git a/test/lit/pto/tprint_alloc_tile_no_rebind.pto b/test/lit/pto/tprint_alloc_tile_no_rebind.pto
index a3cb04b08..e5d01eabb 100644
--- a/test/lit/pto/tprint_alloc_tile_no_rebind.pto
+++ b/test/lit/pto/tprint_alloc_tile_no_rebind.pto
@@ -13,8 +13,8 @@ module {
 }
 
 // CHECK-LABEL: __global__ AICORE void print_kernel() {
-// CHECK: Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, 32, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[TILE:v[0-9]+]];
-// CHECK: TASSIGN([[TILE]], [[ADDR:v[0-9]+]]);
+// CHECK: Tile<TileType::Vec, float, 32, 32, BLayout::RowMajor, 32, 32, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[TILE:[_A-Za-z][_A-Za-z0-9]*]];
+// CHECK: TASSIGN([[TILE]], [[ADDR:[_A-Za-z][_A-Za-z0-9]*]]);
 // CHECK-NOT: TASSIGN(
 // CHECK-NOT: .data()
 // CHECK-NOT: reinterpret_cast<uint64_t>
diff --git a/test/lit/pto/tpush_tpop_dynamic_validshape_a5.pto b/test/lit/pto/tpush_tpop_dynamic_validshape_a5.pto
index a39b43af6..812ef476d 100644
--- a/test/lit/pto/tpush_tpop_dynamic_validshape_a5.pto
+++ b/test/lit/pto/tpush_tpop_dynamic_validshape_a5.pto
@@ -47,13 +47,13 @@ module {
 }
 
 // A5-LABEL: AICORE void cube_kernel(
-// A5: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>(
-// A5: Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> [[CUBE_TILE:v[0-9]+]] = Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>({{v[0-9]+}}, {{v[0-9]+}});
-// A5: [[CUBE_TILE]].SetValidShape({{v[0-9]+}}, {{v[0-9]+}});
-// A5: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>, Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_LEFT_RIGHT>({{v[0-9]+}}, [[CUBE_TILE]]);
+// A5: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>(
+// A5: Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> [[CUBE_TILE:[_A-Za-z][_A-Za-z0-9]*]] = Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
+// A5: [[CUBE_TILE]].SetValidShape({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
+// A5: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>, Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_LEFT_RIGHT>({{[_A-Za-z][_A-Za-z0-9]*}}, [[CUBE_TILE]]);
 
 // A5-LABEL: AICORE void vector_kernel(
-// A5: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>(
-// A5: Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[VEC_TILE:v[0-9]+]] = Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>({{v[0-9]+}}, {{v[0-9]+}});
-// A5: [[VEC_TILE]].SetValidShape({{v[0-9]+}}, {{v[0-9]+}});
-// A5: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>, Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_LEFT_RIGHT>({{v[0-9]+}}, [[VEC_TILE]]);
+// A5: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>(
+// A5: Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> [[VEC_TILE:[_A-Za-z][_A-Za-z0-9]*]] = Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
+// A5: [[VEC_TILE]].SetValidShape({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
+// A5: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>, Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_LEFT_RIGHT>({{[_A-Za-z][_A-Za-z0-9]*}}, [[VEC_TILE]]);
diff --git a/test/lit/pto/tpush_tpop_dynamic_validshape_default_a5.pto b/test/lit/pto/tpush_tpop_dynamic_validshape_default_a5.pto
index 0b96de604..f1c64f643 100644
--- a/test/lit/pto/tpush_tpop_dynamic_validshape_default_a5.pto
+++ b/test/lit/pto/tpush_tpop_dynamic_validshape_default_a5.pto
@@ -45,11 +45,11 @@ module {
 }
 
 // A5-LABEL: AICORE void cube_kernel(
-// A5: Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}} = Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>({{v[0-9]+}}, {{v[0-9]+}});
+// A5: Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}} = Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // A5-NOT: SetValidShape
 // A5: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>, Tile<TileType::Mat, float, 16, 64, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_LEFT_RIGHT>(
 
 // A5-LABEL: AICORE void vector_kernel(
-// A5: Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}} = Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>({{v[0-9]+}}, {{v[0-9]+}});
+// A5: Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}} = Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}});
 // A5-NOT: SetValidShape
 // A5: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, false>, Tile<TileType::Vec, float, 16, 64, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_LEFT_RIGHT>(
diff --git a/test/lit/pto/tpush_tpop_emitc.pto b/test/lit/pto/tpush_tpop_emitc.pto
index f976c8d67..523d8bedf 100644
--- a/test/lit/pto/tpush_tpop_emitc.pto
+++ b/test/lit/pto/tpush_tpop_emitc.pto
@@ -32,20 +32,20 @@ module {
 }
 
 // A3-LABEL: AICORE void cube_push_gm(
-// A3: const int32_t {{v[0-9]+}} = 0;
-// A3: const int64_t {{v[0-9]+}} = 0;
+// A3: const int32_t {{[_A-Za-z][_A-Za-z0-9]*}} = 0;
+// A3: const int64_t {{[_A-Za-z][_A-Za-z0-9]*}} = 0;
 // A3: #if defined(__DAV_CUBE__)
-// A3: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>(
+// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>(
 // A3: TPUSH<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>, Tile<TileType::Acc, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 1024, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
 // A3: #endif // __DAV_CUBE__
 
 // A3-LABEL: AICORE void vector_pop_gm(
-// A3: const int32_t {{v[0-9]+}} = 0;
+// A3: const int32_t {{[_A-Za-z][_A-Za-z0-9]*}} = 0;
 // A3: #if defined(__DAV_VEC__)
 // A3: set_mask_norm();
 // A3: set_vector_mask(-1, -1);
-// A3: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
-// A3: Tile<TileType::Vec, float, 8, 16, BLayout::RowMajor, 8, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>(
+// A3: Tile<TileType::Vec, float, 8, 16, BLayout::RowMajor, 8, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TPOP<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>, Tile<TileType::Vec, float, 8, 16, BLayout::RowMajor, 8, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_UP_DOWN>(
 // A3: TFREE<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>, TileSplitAxis::TILE_LEFT_RIGHT>(
 // A3: #endif // __DAV_VEC__
diff --git a/test/lit/pto/tpush_tpop_frontend_lowering_a3.pto b/test/lit/pto/tpush_tpop_frontend_lowering_a3.pto
index 0bee9339f..7aaa06dfd 100644
--- a/test/lit/pto/tpush_tpop_frontend_lowering_a3.pto
+++ b/test/lit/pto/tpush_tpop_frontend_lowering_a3.pto
@@ -61,20 +61,20 @@ module {
 }
 
 // A3-LABEL: AICORE void cube_kernel(__gm__ float*
-// A3: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>(
+// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>(
 // A3: TPUSH<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>
-// A3: Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
-// A3: Tile<TileType::Left, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: Tile<TileType::Left, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TMOV(
 // A3: TFREE<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, TileSplitAxis::TILE_NO_SPLIT>(
 
 // A3-LABEL: AICORE void vector_kernel(__gm__ float*
-// A3: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>(
-// A3: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>(
+// A3: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TPUSH<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
 // A3: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
-// A3: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TNEG(
 // A3: TFREE<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, TileSplitAxis::TILE_NO_SPLIT>(
 
diff --git a/test/lit/pto/tpush_tpop_frontend_lowering_a3_gss.pto b/test/lit/pto/tpush_tpop_frontend_lowering_a3_gss.pto
index 979975bb6..9d45c628e 100644
--- a/test/lit/pto/tpush_tpop_frontend_lowering_a3_gss.pto
+++ b/test/lit/pto/tpush_tpop_frontend_lowering_a3_gss.pto
@@ -61,20 +61,20 @@ module {
 }
 
 // A3-LABEL: AICORE void cube_kernel(__gm__ float*
-// A3: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>(
+// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>(
 // A3: TPUSH<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>
-// A3: Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
-// A3: Tile<TileType::Left, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: Tile<TileType::Left, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TMOV(
 // A3: TFREE<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, TileSplitAxis::TILE_NO_SPLIT>(
 
 // A3-LABEL: AICORE void vector_kernel(__gm__ float*
-// A3: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>(
-// A3: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>(
+// A3: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TPUSH<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
 // A3: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
-// A3: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A3: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A3: TNEG(
 // A3: TFREE<TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>, TileSplitAxis::TILE_NO_SPLIT>(
 
diff --git a/test/lit/pto/tpush_tpop_frontend_lowering_a5.pto b/test/lit/pto/tpush_tpop_frontend_lowering_a5.pto
index dc4f0625a..413abb902 100644
--- a/test/lit/pto/tpush_tpop_frontend_lowering_a5.pto
+++ b/test/lit/pto/tpush_tpop_frontend_lowering_a5.pto
@@ -57,23 +57,23 @@ module {
 }
 
 // A5-LABEL: AICORE void cube_kernel(
-// A5: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>(
+// A5: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>(
 // A5: TPUSH<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>
-// A5: Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A5: Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A5: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>, Tile<TileType::Mat, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
-// A5: Tile<TileType::Left, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A5: Tile<TileType::Left, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A5: TMOV(
 // A5: TFREE<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>, TileSplitAxis::TILE_NO_SPLIT>(
 
 // A5-LABEL: AICORE void vector_kernel(
 // A5: if (get_subblockid() == 0) {
-// A5: auto {{v[0-9]+}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>(
-// A5: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
-// A5: Tile<TileType::Vec, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A5: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>(
+// A5: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
+// A5: Tile<TileType::Vec, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A5: TMOV(
 // A5: TPUSH<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>, Tile<TileType::Vec, float, 16, 16, BLayout::ColMajor, 16, 16, SLayout::RowMajor, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
 // A5: TPOP<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>, Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null>, TileSplitAxis::TILE_NO_SPLIT>(
-// A5: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{v[0-9]+}};
+// A5: Tile<TileType::Vec, float, 16, 16, BLayout::RowMajor, 16, 16, SLayout::NoneBox, 512, PadValue::Null, CompactMode::Null> {{[_A-Za-z][_A-Za-z0-9]*}};
 // A5: TNEG(
 // A5: TFREE<TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>, TileSplitAxis::TILE_NO_SPLIT>(
 // A5: }
diff --git a/test/lit/pto/tpush_tpop_globaltensor_frontend_a3.pto b/test/lit/pto/tpush_tpop_globaltensor_frontend_a3.pto
index c192d75cd..4d729a5f9 100644
--- a/test/lit/pto/tpush_tpop_globaltensor_frontend_a3.pto
+++ b/test/lit/pto/tpush_tpop_globaltensor_frontend_a3.pto
@@ -52,16 +52,16 @@ module {
 }
 
 // CHECK-LABEL: AICORE void cube_kernel
-// CHECK-SAME: (__gm__ float* [[CUBE_GM:v[0-9]+]],
+// CHECK-SAME: (__gm__ float* [[CUBE_GM:[_A-Za-z][_A-Za-z0-9]*]],
 // CHECK: TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>([[CUBE_GM]], {{.*}}, {{.*}});
-// CHECK: GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND> [[CUBE_ENTRY:v[0-9]+]](nullptr);
+// CHECK: GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND> [[CUBE_ENTRY:[_A-Za-z][_A-Za-z0-9]*]](nullptr);
 // CHECK: TALLOC<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>, GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>({{.*}}, [[CUBE_ENTRY]]);
 // CHECK: TSTORE
 // CHECK: TPUSH<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>, GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>({{.*}}, [[CUBE_ENTRY]]);
 // CHECK-LABEL: AICORE void vector_kernel
-// CHECK-SAME: (__gm__ float* [[VEC_GM:v[0-9]+]],
+// CHECK-SAME: (__gm__ float* [[VEC_GM:[_A-Za-z][_A-Za-z0-9]*]],
 // CHECK: TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>([[VEC_GM]], {{.*}}, {{.*}});
-// CHECK: GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND> [[VEC_ENTRY:v[0-9]+]](nullptr);
+// CHECK: GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND> [[VEC_ENTRY:[_A-Za-z][_A-Za-z0-9]*]](nullptr);
 // CHECK: TPOP<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>, GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>(
 // CHECK: TLOAD
 // CHECK: TFREE<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>, GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>(
diff --git a/test/lit/pto/tpush_tpop_globaltensor_internal_a3.pto b/test/lit/pto/tpush_tpop_globaltensor_internal_a3.pto
index 0b8072002..c56c19ed8 100644
--- a/test/lit/pto/tpush_tpop_globaltensor_internal_a3.pto
+++ b/test/lit/pto/tpush_tpop_globaltensor_internal_a3.pto
@@ -20,7 +20,7 @@ module {
 }
 
 // CHECK-LABEL: AICORE void cube_globaltensor_internal(
-// CHECK: __gm__ float* [[GM_DATA:v[0-9]+]] = PTOAS__GLOBAL_TENSOR_DATA(
+// CHECK: __gm__ float* [[GM_DATA:[_A-Za-z][_A-Za-z0-9]*]] = PTOAS__GLOBAL_TENSOR_DATA(
 // CHECK: TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>([[GM_DATA]], {{.*}}, {{.*}});
 // CHECK: TALLOC<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>, GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>(
 // CHECK: TPUSH<TPipe<0, Direction::DIR_C2V, 1024, 8, 8, true>, GlobalTensor<float, pto::Shape<1, 1, 1, 16, 16>, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>(
diff --git a/test/lit/pto/trem_emitc.pto b/test/lit/pto/trem_emitc.pto
index 073aaa7a6..9f5959bda 100644
--- a/test/lit/pto/trem_emitc.pto
+++ b/test/lit/pto/trem_emitc.pto
@@ -13,4 +13,4 @@ module {
 
 // A3: pto.trem ins([[SRC0:%[0-9]+]], [[SRC1:%[0-9]+]], [[TMP:%[0-9]+]] : memref<1x16xf32
 // A3: outs([[DST:%[0-9]+]] : memref<1x16xf32
-// A3: TREM([[VDST:v[0-9]+]], [[VSRC0:v[0-9]+]], [[VSRC1:v[0-9]+]], [[VTMP:v[0-9]+]]);
+// A3: TREM([[VDST:[_A-Za-z][_A-Za-z0-9]*]], [[VSRC0:[_A-Za-z][_A-Za-z0-9]*]], [[VSRC1:[_A-Za-z][_A-Za-z0-9]*]], [[VTMP:[_A-Za-z][_A-Za-z0-9]*]]);
diff --git a/test/lit/pto/trsqrt_emitc.pto b/test/lit/pto/trsqrt_emitc.pto
index 2287bfe29..b9b1b55bf 100644
--- a/test/lit/pto/trsqrt_emitc.pto
+++ b/test/lit/pto/trsqrt_emitc.pto
@@ -17,5 +17,5 @@ module {
 // A3: outs([[DST0:%[0-9]+]] : memref<1x16xf16
 // A3: pto.trsqrt ins([[SRC1:%[0-9]+]], [[TMP:%[0-9]+]] : memref<1x16xf16
 // A3: outs([[DST1:%[0-9]+]] : memref<1x16xf16
-// A3: TRSQRT([[VDST0:v[0-9]+]], [[VSRC0:v[0-9]+]]);
-// A3: TRSQRT([[VDST1:v[0-9]+]], [[VSRC1:v[0-9]+]], [[VTMP:v[0-9]+]]);
+// A3: TRSQRT([[VDST0:[_A-Za-z][_A-Za-z0-9]*]], [[VSRC0:[_A-Za-z][_A-Za-z0-9]*]]);
+// A3: TRSQRT([[VDST1:[_A-Za-z][_A-Za-z0-9]*]], [[VSRC1:[_A-Za-z][_A-Za-z0-9]*]], [[VTMP:[_A-Za-z][_A-Za-z0-9]*]]);
diff --git a/test/lit/pto/tsort32_emitc.pto b/test/lit/pto/tsort32_emitc.pto
index a49768169..9dceba7fd 100644
--- a/test/lit/pto/tsort32_emitc.pto
+++ b/test/lit/pto/tsort32_emitc.pto
@@ -18,5 +18,5 @@ module {
 // A3: outs([[DST0:%.*]] : memref<1x64xf16
 // A3: pto.tsort32 ins([[SRC]], [[IDX]], [[TMP:%.*]] : memref<1x32xf16
 // A3: outs([[DST1:%.*]] : memref<1x64xf16
-// A3: TSORT32([[VDST0:v[0-9]+]], [[VSRC:v[0-9]+]], [[VIDX:v[0-9]+]]);
-// A3: TSORT32([[VDST1:v[0-9]+]], [[VSRC]], [[VIDX]], [[VTMP:v[0-9]+]]);
+// A3: TSORT32([[VDST0:[_A-Za-z][_A-Za-z0-9]*]], [[VSRC:[_A-Za-z][_A-Za-z0-9]*]], [[VIDX:[_A-Za-z][_A-Za-z0-9]*]]);
+// A3: TSORT32([[VDST1:[_A-Za-z][_A-Za-z0-9]*]], [[VSRC]], [[VIDX]], [[VTMP:[_A-Za-z][_A-Za-z0-9]*]]);
diff --git a/test/lit/pto/tstore_forms_emitc.pto b/test/lit/pto/tstore_forms_emitc.pto
index e425ab57b..5c7dbeaa1 100644
--- a/test/lit/pto/tstore_forms_emitc.pto
+++ b/test/lit/pto/tstore_forms_emitc.pto
@@ -51,11 +51,11 @@ module {
   }
 }
 
-// A3: TSTORE([[DST0:v[0-9]+]], [[VEC:v[0-9]+]]);
+// A3: TSTORE([[DST0:[_A-Za-z][_A-Za-z0-9]*]], [[VEC:[_A-Za-z][_A-Za-z0-9]*]]);
 // A3: TSTORE<STPhase::Final>([[DST0]], [[VEC]]);
 // A3: TSTORE<{{.*}}AtomicType::AtomicAdd>([[DST0]], [[VEC]]);
 // A3: TSTORE<STPhase::Final, {{.*}}AtomicType::AtomicAdd>([[DST0]], [[VEC]]);
-// A3: TSTORE<{{.*}}AtomicType::AtomicAdd, ReluPreMode::NormalRelu>([[DST1:v[0-9]+]], [[ACC:v[0-9]+]]);
+// A3: TSTORE<{{.*}}AtomicType::AtomicAdd, ReluPreMode::NormalRelu>([[DST1:[_A-Za-z][_A-Za-z0-9]*]], [[ACC:[_A-Za-z][_A-Za-z0-9]*]]);
 // A3: TSTORE<STPhase::Final, {{.*}}AtomicType::AtomicAdd, ReluPreMode::NormalRelu>([[DST1]], [[ACC]]);
-// A3: TSTORE<{{.*}}AtomicType::AtomicAdd, ReluPreMode::NormalRelu>([[DST1]], [[ACC]], [[PRE:v[0-9]+]]);
+// A3: TSTORE<{{.*}}AtomicType::AtomicAdd, ReluPreMode::NormalRelu>([[DST1]], [[ACC]], [[PRE:[_A-Za-z][_A-Za-z0-9]*]]);
 // A3: TSTORE<STPhase::Final, {{.*}}AtomicType::AtomicAdd, ReluPreMode::NormalRelu>([[DST1]], [[ACC]], [[PRE]]);
diff --git a/test/samples/Qwen3DecodeA3/qwen3_decode_golden_lib.py b/test/samples/Qwen3DecodeA3/qwen3_decode_golden_lib.py
index 913317f85..08fdbc91d 100644
--- a/test/samples/Qwen3DecodeA3/qwen3_decode_golden_lib.py
+++ b/test/samples/Qwen3DecodeA3/qwen3_decode_golden_lib.py
@@ -8,6 +8,7 @@
 # See LICENSE in the root of the software repository for the full text of the License.
 
 import numpy as np
+from dataclasses import replace
 
 from validation_runtime import (
     bf16_to_float32,
@@ -477,10 +478,39 @@ def build_case_16(meta, generator, ints):
 }
 
 
+def _build_legacy_name_map(meta):
+    ordered = list(meta.read_order)
+    return {f'v{idx}': name for idx, name in enumerate(ordered, start=1)}
+
+
+def _with_legacy_meta_aliases(meta, legacy_to_actual):
+    elem_counts = dict(meta.elem_counts)
+    np_types = dict(meta.np_types)
+    for legacy, actual in legacy_to_actual.items():
+        if legacy in elem_counts:
+            continue
+        if actual not in elem_counts or actual not in np_types:
+            continue
+        elem_counts[legacy] = elem_counts[actual]
+        np_types[legacy] = np_types[actual]
+    return replace(meta, elem_counts=elem_counts, np_types=np_types)
+
+
+def _rewrite_legacy_buffer_names(entries, legacy_to_actual):
+    rewritten = {}
+    for name, value in entries.items():
+        rewritten[legacy_to_actual.get(name, name)] = value
+    return rewritten
+
+
 def run_case(case_name: str):
     meta = load_case_meta()
+    legacy_to_actual = _build_legacy_name_map(meta)
+    meta = _with_legacy_meta_aliases(meta, legacy_to_actual)
     generator = rng()
     ints = load_int32_assignments()
     buffers, golden = BUILDERS[case_name](meta, generator, ints)
+    buffers = _rewrite_legacy_buffer_names(buffers, legacy_to_actual)
+    golden = _rewrite_legacy_buffer_names(golden, legacy_to_actual)
     write_buffers(meta, buffers)
     write_golden(meta, golden)
diff --git a/test/samples/Qwen3DecodeA5/qwen3_decode_golden_lib.py b/test/samples/Qwen3DecodeA5/qwen3_decode_golden_lib.py
index 913317f85..08fdbc91d 100644
--- a/test/samples/Qwen3DecodeA5/qwen3_decode_golden_lib.py
+++ b/test/samples/Qwen3DecodeA5/qwen3_decode_golden_lib.py
@@ -8,6 +8,7 @@
 # See LICENSE in the root of the software repository for the full text of the License.
 
 import numpy as np
+from dataclasses import replace
 
 from validation_runtime import (
     bf16_to_float32,
@@ -477,10 +478,39 @@ def build_case_16(meta, generator, ints):
 }
 
 
+def _build_legacy_name_map(meta):
+    ordered = list(meta.read_order)
+    return {f'v{idx}': name for idx, name in enumerate(ordered, start=1)}
+
+
+def _with_legacy_meta_aliases(meta, legacy_to_actual):
+    elem_counts = dict(meta.elem_counts)
+    np_types = dict(meta.np_types)
+    for legacy, actual in legacy_to_actual.items():
+        if legacy in elem_counts:
+            continue
+        if actual not in elem_counts or actual not in np_types:
+            continue
+        elem_counts[legacy] = elem_counts[actual]
+        np_types[legacy] = np_types[actual]
+    return replace(meta, elem_counts=elem_counts, np_types=np_types)
+
+
+def _rewrite_legacy_buffer_names(entries, legacy_to_actual):
+    rewritten = {}
+    for name, value in entries.items():
+        rewritten[legacy_to_actual.get(name, name)] = value
+    return rewritten
+
+
 def run_case(case_name: str):
     meta = load_case_meta()
+    legacy_to_actual = _build_legacy_name_map(meta)
+    meta = _with_legacy_meta_aliases(meta, legacy_to_actual)
     generator = rng()
     ints = load_int32_assignments()
     buffers, golden = BUILDERS[case_name](meta, generator, ints)
+    buffers = _rewrite_legacy_buffer_names(buffers, legacy_to_actual)
+    golden = _rewrite_legacy_buffer_names(golden, legacy_to_actual)
     write_buffers(meta, buffers)
     write_golden(meta, golden)
diff --git a/tools/ptoas/ptoas.cpp b/tools/ptoas/ptoas.cpp
index f0e2b9cd9..6c2245ead 100644
--- a/tools/ptoas/ptoas.cpp
+++ b/tools/ptoas/ptoas.cpp
@@ -42,7 +42,11 @@
 #include "llvm/ADT/StringRef.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/ADT/StringMap.h"
+#include "llvm/ADT/SmallString.h"
+#include <algorithm>
 #include <memory>
+#include <optional>
+#include <set>
 #include <string>
 
 using namespace mlir;
@@ -283,6 +287,544 @@ static bool parseAutoSyncTailHint(llvm::StringRef hintStr, std::string &normaliz
   return false;
 }
 
+struct ParsedTextualNameHints {
+  llvm::SmallVector<llvm::SmallVector<std::string, 4>, 8> functionArgHints;
+  llvm::SmallVector<llvm::SmallVector<std::string, 4>, 8> blockArgHints;
+  llvm::SmallVector<llvm::SmallVector<std::string, 4>, 32> opResultHints;
+};
+
+using FunctionArgHintMap =
+    llvm::StringMap<llvm::SmallVector<std::string, 4>>;
+using FunctionBlockArgHintMap =
+    llvm::StringMap<llvm::SmallVector<llvm::SmallVector<std::string, 4>, 4>>;
+
+static bool isGeneratedValueName(llvm::StringRef name);
+static SmallVector<std::string, 4> getValueNameHints(Value value);
+
+static bool isMlirValueNameChar(char c) {
+  return std::isalnum(static_cast<unsigned char>(c)) || c == '_' || c == '.' ||
+         c == '$' || c == '-';
+}
+
+static bool isCppIdentifierStart(char c) {
+  return std::isalpha(static_cast<unsigned char>(c)) || c == '_';
+}
+
+static bool isCppIdentifierChar(char c) {
+  return std::isalnum(static_cast<unsigned char>(c)) || c == '_';
+}
+
+static bool parseMlirValueName(llvm::StringRef text, size_t &pos,
+                               std::string &name) {
+  if (pos >= text.size() || text[pos] != '%')
+    return false;
+  size_t begin = ++pos;
+  while (pos < text.size() && isMlirValueNameChar(text[pos]))
+    ++pos;
+  if (pos == begin)
+    return false;
+  name = text.slice(begin, pos).str();
+  return true;
+}
+
+static llvm::SmallVector<std::string, 4>
+parseMlirNamedArguments(llvm::StringRef text) {
+  llvm::SmallVector<std::string, 4> names;
+  for (size_t pos = 0; pos < text.size();) {
+    if (text[pos] != '%') {
+      ++pos;
+      continue;
+    }
+    std::string name;
+    size_t namePos = pos;
+    if (!parseMlirValueName(text, namePos, name)) {
+      ++pos;
+      continue;
+    }
+    while (namePos < text.size() &&
+           std::isspace(static_cast<unsigned char>(text[namePos])))
+      ++namePos;
+    if (namePos < text.size() && text[namePos] == ':')
+      names.push_back(std::move(name));
+    pos = namePos;
+  }
+  return names;
+}
+
+static bool parseLeadingOpResultNames(
+    llvm::StringRef line, llvm::SmallVectorImpl<std::string> &names) {
+  size_t pos = 0;
+  while (pos < line.size() &&
+         std::isspace(static_cast<unsigned char>(line[pos])))
+    ++pos;
+  if (pos >= line.size() || line[pos] != '%')
+    return false;
+
+  while (true) {
+    std::string name;
+    if (!parseMlirValueName(line, pos, name))
+      return false;
+    names.push_back(std::move(name));
+
+    while (pos < line.size() &&
+           std::isspace(static_cast<unsigned char>(line[pos])))
+      ++pos;
+    if (pos < line.size() && line[pos] == ',') {
+      ++pos;
+      while (pos < line.size() &&
+             std::isspace(static_cast<unsigned char>(line[pos])))
+        ++pos;
+      continue;
+    }
+    break;
+  }
+
+  while (pos < line.size() &&
+         std::isspace(static_cast<unsigned char>(line[pos])))
+    ++pos;
+  return pos < line.size() && line[pos] == '=';
+}
+
+static std::string stripMlirLineComments(llvm::StringRef text) {
+  std::string stripped;
+  stripped.reserve(text.size());
+  llvm::StringRef remaining = text;
+  while (!remaining.empty()) {
+    auto split = remaining.split('\n');
+    llvm::StringRef line = split.first;
+    llvm::StringRef rest = split.second;
+    llvm::StringRef body = line;
+    if (size_t commentPos = line.find("//"); commentPos != llvm::StringRef::npos)
+      body = line.take_front(commentPos);
+    stripped.append(body.begin(), body.end());
+    if (!rest.empty())
+      stripped.push_back('\n');
+    remaining = rest;
+  }
+  return stripped;
+}
+
+static ParsedTextualNameHints extractTextualNameHints(llvm::StringRef text) {
+  ParsedTextualNameHints hints;
+  std::string stripped = stripMlirLineComments(text);
+  llvm::StringRef source(stripped);
+
+  size_t searchPos = 0;
+  while (true) {
+    size_t funcPos = source.find("func.func", searchPos);
+    if (funcPos == llvm::StringRef::npos)
+      break;
+
+    size_t atPos = source.find('@', funcPos);
+    if (atPos == llvm::StringRef::npos)
+      break;
+    size_t lParenPos = source.find('(', atPos);
+    if (lParenPos == llvm::StringRef::npos)
+      break;
+
+    int depth = 0;
+    size_t rParenPos = llvm::StringRef::npos;
+    for (size_t i = lParenPos; i < source.size(); ++i) {
+      if (source[i] == '(') {
+        ++depth;
+      } else if (source[i] == ')') {
+        --depth;
+        if (depth == 0) {
+          rParenPos = i;
+          break;
+        }
+      }
+    }
+    if (rParenPos == llvm::StringRef::npos)
+      break;
+
+    hints.functionArgHints.push_back(
+        parseMlirNamedArguments(source.slice(lParenPos + 1, rParenPos)));
+    searchPos = rParenPos + 1;
+  }
+
+  llvm::StringRef remaining(source);
+  while (!remaining.empty()) {
+    auto split = remaining.split('\n');
+    llvm::StringRef line = split.first;
+    llvm::StringRef rest = split.second;
+    llvm::StringRef trimmed = line.ltrim();
+
+    if (trimmed.starts_with("^")) {
+      size_t lParenPos = trimmed.find('(');
+      size_t rParenPos = trimmed.rfind(')');
+      size_t colonPos = trimmed.rfind(':');
+      if (lParenPos != llvm::StringRef::npos &&
+          rParenPos != llvm::StringRef::npos && colonPos != llvm::StringRef::npos &&
+          lParenPos < rParenPos && rParenPos < colonPos) {
+        hints.blockArgHints.push_back(
+            parseMlirNamedArguments(trimmed.slice(lParenPos + 1, rParenPos)));
+      }
+    }
+
+    llvm::SmallVector<std::string, 4> resultNames;
+    if (parseLeadingOpResultNames(trimmed, resultNames))
+      hints.opResultHints.push_back(std::move(resultNames));
+
+    remaining = rest;
+  }
+
+  return hints;
+}
+
+static std::string sanitizeCppIdentifier(llvm::StringRef name) {
+  std::string sanitized;
+  sanitized.reserve(name.size() + 4);
+
+  auto appendUnderscore = [&]() {
+    if (sanitized.empty() || sanitized.back() != '_')
+      sanitized.push_back('_');
+  };
+
+  for (char c : name) {
+    if (isCppIdentifierChar(c))
+      sanitized.push_back(c);
+    else
+      appendUnderscore();
+  }
+
+  while (!sanitized.empty() && sanitized.back() == '_')
+    sanitized.pop_back();
+
+  if (sanitized.empty())
+    return {};
+  if (!isCppIdentifierStart(sanitized.front()))
+    sanitized.insert(sanitized.begin(), '_');
+  return sanitized;
+}
+
+static bool isReservedCppIdentifier(llvm::StringRef name) {
+  static const std::set<std::string> kReserved = {
+      "alignas",   "alignof",   "asm",       "auto",       "bool",
+      "break",     "case",      "catch",     "char",       "char8_t",
+      "char16_t",  "char32_t",  "class",     "const",      "consteval",
+      "constexpr", "constinit", "const_cast","continue",   "co_await",
+      "co_return", "co_yield",  "decltype",  "default",    "delete",
+      "do",        "double",    "dynamic_cast", "else",    "enum",
+      "explicit",  "export",    "extern",    "false",      "float",
+      "for",       "friend",    "goto",      "if",         "inline",
+      "int",       "long",      "mutable",   "namespace",  "new",
+      "noexcept",  "nullptr",   "operator",  "private",    "protected",
+      "public",    "register",  "reinterpret_cast", "requires",
+      "return",    "short",     "signed",    "sizeof",     "static",
+      "static_assert", "static_cast", "struct", "switch",  "template",
+      "this",      "thread_local", "throw",   "true",      "try",
+      "typedef",   "typeid",    "typename",  "union",      "unsigned",
+      "using",     "virtual",   "void",      "volatile",   "wchar_t",
+      "while"};
+  return kReserved.count(name.str()) != 0;
+}
+
+static std::string makeUniqueCppIdentifier(llvm::StringRef baseName,
+                                           std::set<std::string> &usedNames) {
+  std::string candidate = sanitizeCppIdentifier(baseName);
+  if (candidate.empty())
+    return {};
+  if (isReservedCppIdentifier(candidate))
+    candidate.append("_v");
+
+  std::string unique = candidate;
+  unsigned suffix = 1;
+  while (!usedNames.insert(unique).second) {
+    unique = candidate + "_" + std::to_string(suffix++);
+  }
+  return unique;
+}
+
+static void appendLocationNameHints(Location loc,
+                                    SmallVectorImpl<std::string> &hints) {
+  if (auto nameLoc = dyn_cast<NameLoc>(loc)) {
+    std::string sanitized = sanitizeCppIdentifier(nameLoc.getName().getValue());
+    if (!sanitized.empty())
+      hints.push_back(std::move(sanitized));
+    return;
+  }
+
+  if (auto fusedLoc = dyn_cast<FusedLoc>(loc)) {
+    if (Attribute metadata = fusedLoc.getMetadata()) {
+      if (auto strAttr = dyn_cast<StringAttr>(metadata)) {
+        std::string sanitized = sanitizeCppIdentifier(strAttr.getValue());
+        if (!sanitized.empty())
+          hints.push_back(std::move(sanitized));
+        return;
+      }
+      if (auto arrayAttr = dyn_cast<ArrayAttr>(metadata)) {
+        for (Attribute attr : arrayAttr) {
+          auto strAttr = dyn_cast<StringAttr>(attr);
+          if (!strAttr)
+            continue;
+          std::string sanitized = sanitizeCppIdentifier(strAttr.getValue());
+          if (!sanitized.empty())
+            hints.push_back(std::move(sanitized));
+        }
+        if (!hints.empty())
+          return;
+      }
+    }
+
+    for (Location childLoc : fusedLoc.getLocations())
+      appendLocationNameHints(childLoc, hints);
+    return;
+  }
+
+  if (auto callSiteLoc = dyn_cast<CallSiteLoc>(loc)) {
+    appendLocationNameHints(callSiteLoc.getCallee(), hints);
+    if (hints.empty())
+      appendLocationNameHints(callSiteLoc.getCaller(), hints);
+  }
+}
+
+static bool hasLocationNameHints(Location loc) {
+  SmallVector<std::string, 4> hints;
+  appendLocationNameHints(loc, hints);
+  return !hints.empty();
+}
+
+static Location attachLocationNameHints(Location baseLoc,
+                                        llvm::ArrayRef<std::string> hints,
+                                        MLIRContext *context) {
+  SmallVector<Attribute, 4> attrs;
+  attrs.reserve(hints.size());
+  for (llvm::StringRef hint : hints) {
+    if (!hint.empty())
+      attrs.push_back(StringAttr::get(context, hint));
+  }
+  if (attrs.empty())
+    return baseLoc;
+  if (attrs.size() == 1)
+    return NameLoc::get(cast<StringAttr>(attrs.front()), baseLoc);
+  return FusedLoc::get(ArrayRef<Location>{baseLoc}, ArrayAttr::get(context, attrs),
+                       context);
+}
+
+static void applyValueNameHints(Value value, llvm::ArrayRef<std::string> hints) {
+  if (!value || hints.empty() || hasLocationNameHints(value.getLoc()))
+    return;
+  value.setLoc(attachLocationNameHints(value.getLoc(), hints, value.getContext()));
+}
+
+static void applyOperationResultNameHints(Operation *op,
+                                          llvm::ArrayRef<std::string> hints) {
+  if (!op || op->getNumResults() == 0 || hints.empty() ||
+      hasLocationNameHints(op->getLoc()))
+    return;
+
+  SmallVector<std::string, 4> limitedHints;
+  limitedHints.reserve(std::min<size_t>(op->getNumResults(), hints.size()));
+  for (size_t i = 0, e = std::min<size_t>(op->getNumResults(), hints.size());
+       i < e; ++i)
+    limitedHints.push_back(hints[i]);
+  if (limitedHints.empty())
+    return;
+
+  op->setLoc(attachLocationNameHints(op->getLoc(), limitedHints, op->getContext()));
+}
+
+static void collectNonEntryBlocksInSourceOrder(
+    Operation *op, SmallVectorImpl<Block *> &blocks) {
+  for (Region &region : op->getRegions()) {
+    bool isEntryBlock = true;
+    for (Block &block : region) {
+      if (!isEntryBlock && block.getNumArguments() != 0)
+        blocks.push_back(&block);
+      isEntryBlock = false;
+      for (Operation &nestedOp : block)
+        collectNonEntryBlocksInSourceOrder(&nestedOp, blocks);
+    }
+  }
+}
+
+static void applyParsedTextualNameHints(ModuleOp module,
+                                        const ParsedTextualNameHints &hints) {
+  size_t funcHintIndex = 0;
+  for (func::FuncOp func : module.getOps<func::FuncOp>()) {
+    if (funcHintIndex < hints.functionArgHints.size()) {
+      auto argHints = hints.functionArgHints[funcHintIndex];
+      for (auto [index, arg] : llvm::enumerate(func.getArguments())) {
+        if (index < argHints.size())
+          applyValueNameHints(arg, llvm::ArrayRef<std::string>{argHints[index]});
+      }
+    }
+    ++funcHintIndex;
+  }
+
+  SmallVector<Block *, 16> nonEntryBlocks;
+  collectNonEntryBlocksInSourceOrder(module.getOperation(), nonEntryBlocks);
+  for (auto [index, block] : llvm::enumerate(nonEntryBlocks)) {
+    if (index >= hints.blockArgHints.size())
+      break;
+    auto blockHints = hints.blockArgHints[index];
+    for (auto [argIndex, arg] : llvm::enumerate(block->getArguments())) {
+      if (argIndex < blockHints.size())
+        applyValueNameHints(arg,
+                            llvm::ArrayRef<std::string>{blockHints[argIndex]});
+    }
+  }
+
+  size_t opHintIndex = 0;
+  module.walk<WalkOrder::PreOrder>([&](Operation *op) {
+    if (op->getNumResults() == 0)
+      return WalkResult::advance();
+    if (opHintIndex < hints.opResultHints.size())
+      applyOperationResultNameHints(op, hints.opResultHints[opHintIndex]);
+    ++opHintIndex;
+    return WalkResult::advance();
+  });
+}
+
+static FunctionArgHintMap collectFunctionArgNameHints(ModuleOp module) {
+  FunctionArgHintMap hintsByFunction;
+  for (func::FuncOp func : module.getOps<func::FuncOp>()) {
+    SmallVector<std::string, 4> argHints;
+    bool hasAllHints = func.getNumArguments() != 0;
+    for (BlockArgument arg : func.getArguments()) {
+      SmallVector<std::string, 4> hints = getValueNameHints(arg);
+      if (hints.empty()) {
+        hasAllHints = false;
+        break;
+      }
+      argHints.push_back(std::move(hints.front()));
+    }
+    if (hasAllHints)
+      hintsByFunction[func.getSymNameAttr()] = std::move(argHints);
+  }
+  return hintsByFunction;
+}
+
+static FunctionBlockArgHintMap collectFunctionBlockArgNameHints(ModuleOp module) {
+  FunctionBlockArgHintMap hintsByFunction;
+  for (func::FuncOp func : module.getOps<func::FuncOp>()) {
+    SmallVector<Block *, 8> nonEntryBlocks;
+    collectNonEntryBlocksInSourceOrder(func.getOperation(), nonEntryBlocks);
+    if (nonEntryBlocks.empty())
+      continue;
+
+    SmallVector<SmallVector<std::string, 4>, 4> blockHints;
+    blockHints.reserve(nonEntryBlocks.size());
+    for (Block *block : nonEntryBlocks) {
+      SmallVector<std::string, 4> argHints;
+      bool hasAllHints = block->getNumArguments() != 0;
+      for (BlockArgument arg : block->getArguments()) {
+        SmallVector<std::string, 4> hints = getValueNameHints(arg);
+        if (hints.empty()) {
+          hasAllHints = false;
+          break;
+        }
+        argHints.push_back(std::move(hints.front()));
+      }
+      if (hasAllHints)
+        blockHints.push_back(std::move(argHints));
+    }
+
+    if (!blockHints.empty())
+      hintsByFunction[func.getSymNameAttr()] = std::move(blockHints);
+  }
+  return hintsByFunction;
+}
+
+static void applyFunctionBlockArgNameHintsToEmitC(
+    ModuleOp module, const FunctionBlockArgHintMap &blockArgHints) {
+  for (emitc::FuncOp func : module.getOps<emitc::FuncOp>()) {
+    auto it = blockArgHints.find(func.getSymNameAttr());
+    if (it == blockArgHints.end() || it->second.empty())
+      continue;
+
+    SmallVector<Block *, 8> nonEntryBlocks;
+    collectNonEntryBlocksInSourceOrder(func.getOperation(), nonEntryBlocks);
+    for (size_t blockIndex = 0,
+                e = std::min(nonEntryBlocks.size(), it->second.size());
+         blockIndex < e; ++blockIndex) {
+      Block *block = nonEntryBlocks[blockIndex];
+      auto argHints = it->second[blockIndex];
+      for (auto [argIndex, arg] : llvm::enumerate(block->getArguments())) {
+        if (argIndex < argHints.size())
+          applyValueNameHints(arg, llvm::ArrayRef<std::string>{argHints[argIndex]});
+      }
+    }
+  }
+}
+
+static SmallVector<std::string, 4>
+getResultNameHints(Operation *op) {
+  SmallVector<std::string, 4> hints;
+  if (!op || op->getNumResults() == 0)
+    return hints;
+
+  appendLocationNameHints(op->getLoc(), hints);
+  if (hints.empty())
+    return hints;
+
+  hints.erase(std::remove_if(hints.begin(), hints.end(),
+                             [](const std::string &name) {
+                               return name.empty();
+                             }),
+              hints.end());
+  if (hints.empty())
+    return hints;
+
+  if (op->getNumResults() == 1) {
+    if (hints.size() > 1)
+      hints.resize(1);
+    return hints;
+  }
+
+  if (hints.size() > op->getNumResults())
+    hints.resize(op->getNumResults());
+  return hints;
+}
+
+static SmallVector<std::string, 4> getValueNameHints(Value value) {
+  SmallVector<std::string, 4> hints;
+  if (!value)
+    return hints;
+  appendLocationNameHints(value.getLoc(), hints);
+  if (hints.size() > 1)
+    hints.resize(1);
+  return hints;
+}
+
+static std::string buildHintMarker(llvm::StringRef prefix,
+                                   llvm::ArrayRef<std::string> hints) {
+  std::string marker = ("/* " + prefix + ":").str();
+  for (size_t i = 0; i < hints.size(); ++i) {
+    if (i != 0)
+      marker.push_back(',');
+    marker.append(hints[i]);
+  }
+  marker.append(" */\n");
+  return marker;
+}
+
+static void annotateEmitCNameHints(ModuleOp module) {
+  llvm::SmallVector<Operation *, 32> opsToAnnotate;
+  module.walk<WalkOrder::PreOrder>([&](Operation *op) {
+    if (op->getNumResults() == 0 || isa<emitc::VerbatimOp>(op))
+      return WalkResult::advance();
+    if (op->getParentOfType<emitc::ExpressionOp>())
+      return WalkResult::advance();
+    if (getResultNameHints(op).empty())
+      return WalkResult::advance();
+    opsToAnnotate.push_back(op);
+    return WalkResult::advance();
+  });
+
+  OpBuilder builder(module.getContext());
+  for (Operation *op : opsToAnnotate) {
+    SmallVector<std::string, 4> hints = getResultNameHints(op);
+    if (hints.empty())
+      continue;
+    builder.setInsertionPoint(op);
+    builder.create<emitc::VerbatimOp>(
+        op->getLoc(),
+        builder.getStringAttr(buildHintMarker("PTOAS_NAME_HINTS", hints)));
+  }
+}
+
 // --------------------------------------------------------------------------
 // Post-process C++ output: rewrite marker calls into Tile member calls.
 //
@@ -524,11 +1066,13 @@ static void materializeControlFlowOperands(Operation *rootOp) {
       if (!initAttr)
         continue;
 
+      Location valueLoc = value.getLoc();
+
       Value tmp =
-          builder.create<emitc::VariableOp>(op->getLoc(), value.getType(),
+          builder.create<emitc::VariableOp>(valueLoc, value.getType(),
                                             initAttr)
               .getResult();
-      builder.create<emitc::AssignOp>(op->getLoc(), tmp, value);
+      builder.create<emitc::AssignOp>(valueLoc, tmp, value);
       operand.set(tmp);
     }
   }
@@ -688,6 +1232,656 @@ static void rewriteHoistedGlobalTensorDecls(std::string &cpp) {
   cpp.swap(out);
 }
 
+static std::optional<llvm::SmallVector<std::string, 4>>
+parseNameHintMarker(llvm::StringRef markerBody) {
+  llvm::SmallVector<std::string, 4> hints;
+  markerBody = markerBody.trim();
+  if (markerBody.empty())
+    return std::nullopt;
+
+  size_t start = 0;
+  while (start <= markerBody.size()) {
+    size_t comma = markerBody.find(',', start);
+    llvm::StringRef token = markerBody.slice(
+        start, comma == llvm::StringRef::npos ? markerBody.size() : comma);
+    token = token.trim();
+    if (!token.empty())
+      hints.push_back(token.str());
+    if (comma == llvm::StringRef::npos)
+      break;
+    start = comma + 1;
+  }
+
+  if (hints.empty())
+    return std::nullopt;
+  return hints;
+}
+
+static std::optional<llvm::SmallVector<std::string, 4>>
+findNextHintedGeneratedParams(llvm::StringRef snippet) {
+  size_t lParenPos = snippet.find('(');
+  if (lParenPos == llvm::StringRef::npos)
+    return std::nullopt;
+
+  int parenDepth = 0;
+  size_t rParenPos = llvm::StringRef::npos;
+  for (size_t i = lParenPos; i < snippet.size(); ++i) {
+    char c = snippet[i];
+    if (c == '(') {
+      ++parenDepth;
+    } else if (c == ')') {
+      --parenDepth;
+      if (parenDepth == 0) {
+        rParenPos = i;
+        break;
+      }
+    }
+  }
+  if (rParenPos == llvm::StringRef::npos)
+    return std::nullopt;
+
+  llvm::StringRef params = snippet.slice(lParenPos + 1, rParenPos);
+  llvm::SmallVector<std::string, 4> names;
+  size_t partBegin = 0;
+  int angleDepth = 0;
+  int bracketDepth = 0;
+  parenDepth = 0;
+  for (size_t i = 0; i <= params.size(); ++i) {
+    char c = i < params.size() ? params[i] : ',';
+    if (c == '<') {
+      ++angleDepth;
+    } else if (c == '>' && angleDepth > 0) {
+      --angleDepth;
+    } else if (c == '[') {
+      ++bracketDepth;
+    } else if (c == ']' && bracketDepth > 0) {
+      --bracketDepth;
+    } else if (c == '(') {
+      ++parenDepth;
+    } else if (c == ')' && parenDepth > 0) {
+      --parenDepth;
+    }
+
+    bool atSeparator =
+        (i == params.size()) ||
+        (c == ',' && angleDepth == 0 && bracketDepth == 0 && parenDepth == 0);
+    if (!atSeparator)
+      continue;
+
+    llvm::StringRef param = params.slice(partBegin, i).trim();
+    partBegin = i + 1;
+    if (param.empty())
+      continue;
+
+    size_t end = param.size();
+    while (end > 0 && std::isspace(static_cast<unsigned char>(param[end - 1])))
+      --end;
+    size_t begin = end;
+    while (begin > 0 && isCppIdentifierChar(param[begin - 1]))
+      --begin;
+    llvm::StringRef token = param.slice(begin, end);
+    if (isGeneratedValueName(token))
+      names.push_back(token.str());
+  }
+
+  if (names.empty())
+    return std::nullopt;
+  return names;
+}
+
+static std::optional<llvm::SmallVector<std::string, 4>>
+findNextHintedGeneratedNames(llvm::StringRef snippet) {
+  static const llvm::Regex kTieRegex(
+      R"re(std::tie\(([[:space:]]*v[0-9]+([[:space:]]*,[[:space:]]*v[0-9]+)*)\)[[:space:]]*=)re");
+  static const llvm::Regex kSingleRegex(
+      R"re((^|[^[:alnum:]_])(v[0-9]+)[[:space:]]*(=|;))re");
+
+  llvm::SmallVector<llvm::StringRef, 4> matches;
+  if (kTieRegex.match(snippet, &matches) && matches.size() >= 2) {
+    llvm::SmallVector<std::string, 4> names;
+    llvm::StringRef tieNames = matches[1];
+    size_t start = 0;
+    while (start < tieNames.size()) {
+      size_t comma = tieNames.find(',', start);
+      llvm::StringRef token = tieNames.slice(
+          start, comma == llvm::StringRef::npos ? tieNames.size() : comma);
+      token = token.trim();
+      if (!token.empty())
+        names.push_back(token.str());
+      if (comma == llvm::StringRef::npos)
+        break;
+      start = comma + 1;
+    }
+    if (!names.empty())
+      return names;
+  }
+
+  matches.clear();
+  if (kSingleRegex.match(snippet, &matches) && matches.size() >= 3)
+    return llvm::SmallVector<std::string, 4>{matches[2].str()};
+
+  return std::nullopt;
+}
+
+static void rewriteIdentifiersWithMap(
+    std::string &cpp, const llvm::StringMap<std::string> &replacements) {
+  if (replacements.empty())
+    return;
+  std::string rewritten;
+  rewritten.reserve(cpp.size());
+  enum class LexState {
+    Normal,
+    LineComment,
+    BlockComment,
+    StringLiteral,
+    CharLiteral,
+  };
+  LexState state = LexState::Normal;
+
+  auto appendCurrent = [&](char c) { rewritten.push_back(c); };
+
+  for (size_t i = 0; i < cpp.size();) {
+    char c = cpp[i];
+    char next = i + 1 < cpp.size() ? cpp[i + 1] : '\0';
+
+    switch (state) {
+    case LexState::Normal:
+      if (c == '/' && next == '/') {
+        state = LexState::LineComment;
+        appendCurrent(c);
+        appendCurrent(next);
+        i += 2;
+        continue;
+      }
+      if (c == '/' && next == '*') {
+        state = LexState::BlockComment;
+        appendCurrent(c);
+        appendCurrent(next);
+        i += 2;
+        continue;
+      }
+      if (c == '"') {
+        state = LexState::StringLiteral;
+        appendCurrent(c);
+        ++i;
+        continue;
+      }
+      if (c == '\'') {
+        state = LexState::CharLiteral;
+        appendCurrent(c);
+        ++i;
+        continue;
+      }
+      if (isCppIdentifierStart(c)) {
+        size_t end = i + 1;
+        while (end < cpp.size() && isCppIdentifierChar(cpp[end]))
+          ++end;
+        llvm::StringRef token(cpp.data() + i, end - i);
+        auto it = replacements.find(token);
+        if (it != replacements.end())
+          rewritten.append(it->second);
+        else
+          rewritten.append(token.begin(), token.end());
+        i = end;
+        continue;
+      }
+      appendCurrent(c);
+      ++i;
+      continue;
+
+    case LexState::LineComment:
+      appendCurrent(c);
+      ++i;
+      if (c == '\n')
+        state = LexState::Normal;
+      continue;
+
+    case LexState::BlockComment:
+      appendCurrent(c);
+      ++i;
+      if (c == '*' && next == '/') {
+        appendCurrent(next);
+        ++i;
+        state = LexState::Normal;
+      }
+      continue;
+
+    case LexState::StringLiteral:
+      appendCurrent(c);
+      ++i;
+      if (c == '\\' && i < cpp.size()) {
+        appendCurrent(cpp[i]);
+        ++i;
+      } else if (c == '"') {
+        state = LexState::Normal;
+      }
+      continue;
+
+    case LexState::CharLiteral:
+      appendCurrent(c);
+      ++i;
+      if (c == '\\' && i < cpp.size()) {
+        appendCurrent(cpp[i]);
+        ++i;
+      } else if (c == '\'') {
+        state = LexState::Normal;
+      }
+      continue;
+    }
+  }
+
+  cpp.swap(rewritten);
+}
+
+static void stripHintMarkersWithPrefix(std::string &cpp,
+                                       llvm::StringRef markerPrefix) {
+  size_t searchPos = 0;
+  while (true) {
+    size_t markerPos = cpp.find(markerPrefix.str(), searchPos);
+    if (markerPos == std::string::npos)
+      break;
+
+    size_t markerEnd = cpp.find("*/", markerPos + markerPrefix.size());
+    if (markerEnd == std::string::npos)
+      break;
+    markerEnd += 2;
+    while (markerEnd < cpp.size() &&
+           (cpp[markerEnd] == '\r' || cpp[markerEnd] == '\n'))
+      ++markerEnd;
+
+    cpp.erase(markerPos, markerEnd - markerPos);
+    searchPos = markerPos;
+  }
+}
+
+static void stripAllHintMarkers(std::string &cpp) {
+  stripHintMarkersWithPrefix(cpp, "/* PTOAS_NAME_HINTS:");
+  stripHintMarkersWithPrefix(cpp, "/* PTOAS_PARAM_HINTS:");
+}
+
+static bool isHintMarkerLine(llvm::StringRef trimmed) {
+  return trimmed.starts_with("/* PTOAS_NAME_HINTS:") ||
+         trimmed.starts_with("/* PTOAS_PARAM_HINTS:");
+}
+
+static std::optional<std::string>
+extractFunctionNameFromSegment(llvm::StringRef segment) {
+  size_t lParenPos = segment.find('(');
+  if (lParenPos == llvm::StringRef::npos)
+    return std::nullopt;
+  size_t end = lParenPos;
+  while (end > 0 && std::isspace(static_cast<unsigned char>(segment[end - 1])))
+    --end;
+  size_t begin = end;
+  while (begin > 0 && isCppIdentifierChar(segment[begin - 1]))
+    --begin;
+  if (begin == end)
+    return std::nullopt;
+  return segment.slice(begin, end).str();
+}
+
+static bool isTopLevelFunctionStartLine(llvm::StringRef trimmed) {
+  if (trimmed.empty() || trimmed.starts_with("#") || !trimmed.ends_with("{"))
+    return false;
+  if (!trimmed.contains('(') || !trimmed.contains(')'))
+    return false;
+  if (trimmed.starts_with("if ") || trimmed.starts_with("if(") ||
+      trimmed.starts_with("for ") || trimmed.starts_with("for(") ||
+      trimmed.starts_with("while ") || trimmed.starts_with("while(") ||
+      trimmed.starts_with("switch ") || trimmed.starts_with("switch(") ||
+      trimmed.starts_with("catch ") || trimmed.starts_with("catch("))
+    return false;
+  return true;
+}
+
+static std::optional<std::string>
+parseAnyDeclaredIdentifierName(llvm::StringRef line);
+
+static llvm::SmallVector<std::string, 4>
+findTopLevelGeneratedDeclarations(llvm::StringRef segment);
+
+static std::optional<std::string>
+parseGeneratedDeclarationName(llvm::StringRef line) {
+  auto declaredName = parseAnyDeclaredIdentifierName(line);
+  if (!declaredName || !isGeneratedValueName(*declaredName))
+    return std::nullopt;
+  return declaredName;
+}
+
+static std::set<std::string> collectDeclaredIdentifiersInFunctionSegment(
+    llvm::StringRef segment) {
+  std::set<std::string> declaredNames;
+
+  if (size_t lParenPos = segment.find('('); lParenPos != llvm::StringRef::npos) {
+    int parenDepth = 0;
+    size_t rParenPos = llvm::StringRef::npos;
+    for (size_t i = lParenPos; i < segment.size(); ++i) {
+      char c = segment[i];
+      if (c == '(') {
+        ++parenDepth;
+      } else if (c == ')') {
+        --parenDepth;
+        if (parenDepth == 0) {
+          rParenPos = i;
+          break;
+        }
+      }
+    }
+    if (rParenPos != llvm::StringRef::npos) {
+      llvm::StringRef params = segment.slice(lParenPos + 1, rParenPos);
+      size_t partBegin = 0;
+      int angleDepth = 0;
+      int bracketDepth = 0;
+      parenDepth = 0;
+      for (size_t i = 0; i <= params.size(); ++i) {
+        char c = i < params.size() ? params[i] : ',';
+        if (c == '<') {
+          ++angleDepth;
+        } else if (c == '>' && angleDepth > 0) {
+          --angleDepth;
+        } else if (c == '[') {
+          ++bracketDepth;
+        } else if (c == ']' && bracketDepth > 0) {
+          --bracketDepth;
+        } else if (c == '(') {
+          ++parenDepth;
+        } else if (c == ')' && parenDepth > 0) {
+          --parenDepth;
+        }
+
+        bool atSeparator =
+            (i == params.size()) ||
+            (c == ',' && angleDepth == 0 && bracketDepth == 0 &&
+             parenDepth == 0);
+        if (!atSeparator)
+          continue;
+
+        llvm::StringRef param = params.slice(partBegin, i).trim();
+        partBegin = i + 1;
+        if (param.empty())
+          continue;
+
+        size_t end = param.size();
+        while (end > 0 &&
+               std::isspace(static_cast<unsigned char>(param[end - 1])))
+          --end;
+        size_t begin = end;
+        while (begin > 0 && isCppIdentifierChar(param[begin - 1]))
+          --begin;
+        llvm::StringRef name = param.slice(begin, end);
+        if (!name.empty() && isCppIdentifierStart(name.front()) &&
+            llvm::all_of(name, isCppIdentifierChar))
+          declaredNames.insert(name.str());
+      }
+    }
+  }
+
+  llvm::StringRef remaining = segment;
+  while (!remaining.empty()) {
+    auto split = remaining.split('\n');
+    llvm::StringRef line = split.first;
+    llvm::StringRef rest = split.second;
+    if (auto declaredName = parseAnyDeclaredIdentifierName(line))
+      declaredNames.insert(*declaredName);
+    remaining = rest;
+  }
+
+  return declaredNames;
+}
+
+struct PendingIdentifierRename {
+  std::string oldName;
+  std::string baseHint;
+};
+
+static llvm::SmallVector<PendingIdentifierRename, 8>
+collectPendingIdentifierRenames(
+    llvm::StringRef segment, llvm::ArrayRef<std::string> functionParamHints,
+    llvm::ArrayRef<llvm::SmallVector<std::string, 4>> blockArgHints) {
+  static constexpr llvm::StringLiteral kResultMarkerPrefix =
+      "/* PTOAS_NAME_HINTS:";
+  llvm::SmallVector<PendingIdentifierRename, 8> pendingRenames;
+
+  if (!functionParamHints.empty()) {
+    if (auto generatedParams = findNextHintedGeneratedParams(segment)) {
+      size_t pairCount =
+          std::min(functionParamHints.size(), generatedParams->size());
+      for (size_t i = 0; i < pairCount; ++i) {
+        pendingRenames.push_back(
+            PendingIdentifierRename{(*generatedParams)[i], functionParamHints[i]});
+      }
+    }
+  }
+
+  if (!blockArgHints.empty()) {
+    llvm::SmallVector<std::string, 4> generatedDecls =
+        findTopLevelGeneratedDeclarations(segment);
+    llvm::SmallVector<std::string, 4> flattenedBlockHints;
+    for (auto blockHints : blockArgHints)
+      flattenedBlockHints.append(blockHints.begin(), blockHints.end());
+    if (!flattenedBlockHints.empty() &&
+        generatedDecls.size() >= flattenedBlockHints.size()) {
+      size_t startIndex = generatedDecls.size() - flattenedBlockHints.size();
+      for (size_t i = 0; i < flattenedBlockHints.size(); ++i) {
+        pendingRenames.push_back(PendingIdentifierRename{
+            generatedDecls[startIndex + i], flattenedBlockHints[i]});
+      }
+    }
+  }
+
+  size_t searchPos = 0;
+  while (true) {
+    size_t markerPos = segment.find(kResultMarkerPrefix.str(), searchPos);
+    if (markerPos == std::string::npos)
+      break;
+
+    size_t bodyBegin = markerPos + kResultMarkerPrefix.size();
+    size_t markerEnd = segment.find("*/", bodyBegin);
+    if (markerEnd == std::string::npos)
+      break;
+
+    auto hints = parseNameHintMarker(
+        llvm::StringRef(segment).slice(bodyBegin, markerEnd));
+    searchPos = markerEnd + 2;
+    if (!hints)
+      continue;
+
+    size_t windowEnd =
+        std::min(searchPos + static_cast<size_t>(2048), segment.size());
+    llvm::StringRef searchWindow =
+        llvm::StringRef(segment).slice(searchPos, windowEnd);
+    auto generatedNames = findNextHintedGeneratedNames(searchWindow);
+    if (!generatedNames)
+      continue;
+
+    size_t pairCount = std::min(hints->size(), generatedNames->size());
+    for (size_t i = 0; i < pairCount; ++i) {
+      pendingRenames.push_back(
+          PendingIdentifierRename{(*generatedNames)[i], (*hints)[i]});
+    }
+  }
+
+  return pendingRenames;
+}
+
+static std::optional<std::string>
+parseAnyDeclaredIdentifierName(llvm::StringRef line) {
+  llvm::StringRef trimmed = line.trim();
+  if (trimmed.empty() || trimmed.starts_with("#") || trimmed.starts_with("//") ||
+      !trimmed.ends_with(";"))
+    return std::nullopt;
+  llvm::StringRef body = trimmed.drop_back().rtrim();
+  if (body.starts_with("return") || body.starts_with("goto ") ||
+      body.starts_with("if ") || body.starts_with("if(") ||
+      body.starts_with("switch ") || body.starts_with("switch(") ||
+      body.starts_with("for ") || body.starts_with("for(") ||
+      body.starts_with("while ") || body.starts_with("while(") ||
+      body.starts_with("using namespace "))
+    return std::nullopt;
+
+  llvm::StringRef lhs = body;
+  if (size_t eqPos = body.find('='); eqPos != llvm::StringRef::npos)
+    lhs = body.take_front(eqPos).rtrim();
+  size_t lastWs = lhs.find_last_of(" \t");
+  if (lastWs == llvm::StringRef::npos)
+    return std::nullopt;
+  llvm::StringRef name = lhs.drop_front(lastWs + 1).trim();
+  if (name.empty() || !isCppIdentifierStart(name.front()) ||
+      !llvm::all_of(name, isCppIdentifierChar))
+    return std::nullopt;
+  return name.str();
+}
+
+static llvm::SmallVector<std::string, 4>
+findTopLevelGeneratedDeclarations(llvm::StringRef segment) {
+  llvm::SmallVector<std::string, 4> names;
+  size_t lBracePos = segment.find('{');
+  if (lBracePos == llvm::StringRef::npos)
+    return names;
+
+  llvm::StringRef body = segment.drop_front(lBracePos + 1);
+  llvm::StringRef remaining = body;
+  while (!remaining.empty()) {
+    auto split = remaining.split('\n');
+    llvm::StringRef line = split.first;
+    llvm::StringRef rest = split.second;
+    llvm::StringRef trimmed = line.trim();
+    if (trimmed.empty()) {
+      remaining = rest;
+      continue;
+    }
+    if (trimmed.starts_with("using "))
+      break;
+    if (auto generatedName = parseGeneratedDeclarationName(trimmed))
+      names.push_back(*generatedName);
+    remaining = rest;
+  }
+  return names;
+}
+
+static int countBraceDelta(llvm::StringRef line) {
+  int delta = 0;
+  for (char c : line) {
+    if (c == '{')
+      ++delta;
+    else if (c == '}')
+      --delta;
+  }
+  return delta;
+}
+
+static void rewriteNameHintsInFunctionSegment(
+    std::string &segment, llvm::ArrayRef<std::string> functionParamHints,
+    llvm::ArrayRef<llvm::SmallVector<std::string, 4>> blockArgHints) {
+  llvm::StringMap<std::string> replacements;
+  llvm::SmallVector<PendingIdentifierRename, 8> pendingRenames =
+      collectPendingIdentifierRenames(segment, functionParamHints, blockArgHints);
+  std::set<std::string> usedNames =
+      collectDeclaredIdentifiersInFunctionSegment(segment);
+  for (const PendingIdentifierRename &rename : pendingRenames)
+    usedNames.erase(rename.oldName);
+
+  for (const PendingIdentifierRename &rename : pendingRenames) {
+    llvm::StringRef oldName = rename.oldName;
+    if (replacements.count(oldName))
+      continue;
+    std::string newName = makeUniqueCppIdentifier(rename.baseHint, usedNames);
+    if (newName.empty() || newName == oldName)
+      continue;
+    replacements[oldName] = std::move(newName);
+  }
+
+  stripAllHintMarkers(segment);
+  rewriteIdentifiersWithMap(segment, replacements);
+}
+
+static void rewriteNameHintMarkers(std::string &cpp,
+                                   const FunctionArgHintMap &functionArgHints,
+                                   const FunctionBlockArgHintMap &functionBlockArgHints) {
+  llvm::SmallVector<std::string, 0> lines;
+  for (llvm::StringRef ref(cpp); !ref.empty();) {
+    auto split = ref.split('\n');
+    lines.push_back(split.first.str());
+    ref = split.second;
+  }
+
+  std::string rewritten;
+  rewritten.reserve(cpp.size());
+  size_t cursor = 0;
+  int topLevelBraceDepth = 0;
+
+  auto appendLines = [&](size_t begin, size_t end, bool stripMarkers) {
+    if (begin >= end)
+      return;
+    std::string chunk;
+    for (size_t i = begin; i < end; ++i) {
+      chunk.append(lines[i]);
+      if (i + 1 != end || end != lines.size())
+        chunk.push_back('\n');
+    }
+    if (stripMarkers)
+      stripAllHintMarkers(chunk);
+    rewritten.append(chunk);
+  };
+
+  size_t i = 0;
+  while (i < lines.size()) {
+    llvm::StringRef trimmed = llvm::StringRef(lines[i]).trim();
+    if (topLevelBraceDepth == 0 && isTopLevelFunctionStartLine(trimmed)) {
+      size_t segmentBegin = i;
+      while (segmentBegin > cursor &&
+             isHintMarkerLine(llvm::StringRef(lines[segmentBegin - 1]).trim()))
+        --segmentBegin;
+
+      appendLines(cursor, segmentBegin, true);
+
+      size_t segmentEnd = i;
+      int segmentBraceDepth = 0;
+      bool sawOpeningBrace = false;
+      for (; segmentEnd < lines.size(); ++segmentEnd) {
+        int lineDelta = countBraceDelta(lines[segmentEnd]);
+        if (lines[segmentEnd].find('{') != std::string::npos)
+          sawOpeningBrace = true;
+        segmentBraceDepth += lineDelta;
+        if (sawOpeningBrace && segmentBraceDepth == 0) {
+          ++segmentEnd;
+          break;
+        }
+      }
+
+      std::string segment;
+      for (size_t lineIndex = segmentBegin; lineIndex < segmentEnd; ++lineIndex) {
+        segment.append(lines[lineIndex]);
+        if (lineIndex + 1 != segmentEnd || segmentEnd != lines.size())
+          segment.push_back('\n');
+      }
+      llvm::SmallVector<std::string, 4> paramHints;
+      llvm::SmallVector<llvm::SmallVector<std::string, 4>, 4> blockHints;
+      if (auto functionName = extractFunctionNameFromSegment(segment)) {
+        auto it = functionArgHints.find(*functionName);
+        if (it != functionArgHints.end())
+          paramHints = it->second;
+        auto blockIt = functionBlockArgHints.find(*functionName);
+        if (blockIt != functionBlockArgHints.end())
+          blockHints = blockIt->second;
+      }
+      rewriteNameHintsInFunctionSegment(segment, paramHints, blockHints);
+      rewritten.append(segment);
+
+      cursor = segmentEnd;
+      i = segmentEnd;
+      topLevelBraceDepth = 0;
+      continue;
+    }
+
+    topLevelBraceDepth += countBraceDelta(lines[i]);
+    ++i;
+  }
+
+  appendLines(cursor, lines.size(), true);
+  cpp.swap(rewritten);
+}
+
 namespace {
 struct ConstantDeclCandidate {
   size_t declLine = 0;
@@ -978,6 +2172,7 @@ int main(int argc, char **argv) {
 
   OwningOpRef<ModuleOp> module;
   llvm::StringRef buf = (*fileOrErr)->getBuffer();
+  std::optional<ParsedTextualNameHints> textualNameHints;
   const bool isPTOBC = (buf.size() >= 6 && std::memcmp(buf.data(), "PTOBC\0", 6) == 0);
   auto normalizeArch = [](llvm::StringRef archValue) {
     std::string normalized = archValue.str();
@@ -1027,6 +2222,7 @@ int main(int argc, char **argv) {
     }
   } else {
     // Parse textual MLIR (.pto).
+    textualNameHints = extractTextualNameHints(buf);
     llvm::SourceMgr sourceMgr;
     sourceMgr.AddNewSourceBuffer(std::move(*fileOrErr), llvm::SMLoc());
     pto::ScopedPTOParserTargetArch scopedParserArch(
@@ -1039,6 +2235,16 @@ int main(int argc, char **argv) {
     }
   }
 
+  if (module && textualNameHints)
+    applyParsedTextualNameHints(*module, *textualNameHints);
+
+  FunctionArgHintMap functionArgHints;
+  FunctionBlockArgHintMap functionBlockArgHints;
+  if (module)
+    functionArgHints = collectFunctionArgNameHints(*module);
+  if (module)
+    functionBlockArgHints = collectFunctionBlockArgNameHints(*module);
+
   // If the CLI explicitly requested an arch, it overrides the input module.
   // Otherwise, preserve the textual module's arch when present and only fall
   // back to the effective default.
@@ -1218,12 +2424,14 @@ int main(int argc, char **argv) {
     return 1;
   }
 
+  applyFunctionBlockArgNameHintsToEmitC(*module, functionBlockArgHints);
   dropEmptyEmitCExpressions(module.get());
   materializeControlFlowOperands(module.get());
   if (failed(reorderEmitCFunctions(module.get()))) {
     llvm::errs() << "Error: Failed to order emitted functions for C++ emission.\n";
     return 1;
   }
+  annotateEmitCNameHints(*module);
 
   // Emit C++ to string, then post-process, then write to output file.
   std::string cppOutput;
@@ -1245,6 +2453,7 @@ int main(int argc, char **argv) {
   rewriteAddPtrTraceMarkers(cppOutput, emitAddPtrTrace);
   rewriteScalarConstantDecls(cppOutput);
   rewriteHoistedGlobalTensorDecls(cppOutput);
+  rewriteNameHintMarkers(cppOutput, functionArgHints, functionBlockArgHints);
   
   *outputOS << cppOutput;
   outputOS->flush();