From c664e3fe5fc6282fee876a12aede9f7e98c71f02 Mon Sep 17 00:00:00 2001 From: HecreReed <821896444@qq.com> Date: Tue, 12 May 2026 09:48:48 +0800 Subject: [PATCH 1/6] feat: preserve debug name hints in ptoas emitc --- docs/designs/ptoas-debug-name-hints-design.md | 268 +++++++++++ test/lit/pto/debug_name_hints_cfg_emitc.pto | 33 ++ test/lit/pto/debug_name_hints_emitc.pto | 23 + tools/ptoas/ptoas.cpp | 425 +++++++++++++++++- 4 files changed, 747 insertions(+), 2 deletions(-) create mode 100644 docs/designs/ptoas-debug-name-hints-design.md create mode 100644 test/lit/pto/debug_name_hints_cfg_emitc.pto create mode 100644 test/lit/pto/debug_name_hints_emitc.pto 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..c7acd990a --- /dev/null +++ b/docs/designs/ptoas-debug-name-hints-design.md @@ -0,0 +1,268 @@ +# 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` 元数据 +- 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 只服务于调试命名，不参与语义。 + +### 6.3 前端接口约定 + +前端可以提供名字，也可以不提供。 + +- 提供时：PTOAS 尽量保留 +- 不提供时：PTOAS 继续使用现有 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 名 + +## 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` + +例如： + +- `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 时附带名字提示 +- PTOAS rewrite / lowering helper 在替换时传播名字 +- `PTOToEmitC` 中新建 `emitc::VariableOp`、`emitc::CastOp` 等值时继承或派生命名 +- 最终 `translateToCpp` 前或其包装层中应用名字分配规则 + +当前链路里已经存在大量重新物化变量的地方，因此不能只在 parser 或最终 printer 单点处理，必须贯穿传播和最终分配两个阶段。 + +## 11. 风险与限制 + +主要风险如下： + +- 某些 pass 新建值但没有传播名字，会导致局部退化回 `vN` +- 多结果 op 的 metadata 约定若不统一，前后端容易理解不一致 +- 名字传播若写成语义属性，容易污染 IR；因此必须坚持“调试元数据”定位 + +限制如下： + +- 该设计只能提供“尽量保留”，不能提供“逐号完全一致” +- 对 aggressive CSE 后的公共值，只能保留最终幸存值的名字 + +## 12. 测试建议 + +建议至少覆盖以下测试： + +- 单结果 op：前端名字能出现在最终 `kernel.cpp` +- 多结果 op：多个结果名能分别正确落到最终变量 +- 名字含非法字符：能被规范化 +- 重名冲突：能稳定追加后缀 +- 控制流 / `emitc.variable` / hoist：不会因为名字传播导致非法 C++ +- 未提供 hint：仍保持现有 `vN` 回退行为 + +## 13. 结论 + +本设计采用“名字作为调试元数据”的路线： + +- 不把名字当语义 +- 不要求逐号保真 +- 不把 IR 默认打印搞得很吵 +- 重点解决 `.pto`、前端 Python 和最终 `kernel.cpp` 之间“看得懂、对得上”的问题 + +这条路线对现有编译链路侵入较小，也最符合 issue `#337` 的真实诉求。 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..5dbc88801 --- /dev/null +++ b/test/lit/pto/debug_name_hints_cfg_emitc.pto @@ -0,0 +1,33 @@ +// 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 v1, int32_t v2, int32_t v3) { +// CHECK-NOT: PTOAS_NAME_HINTS +// CHECK: int32_t branch_sum; +// CHECK: int32_t branch_sum_1; +// CHECK: branch_sum = (int32_t) +// CHECK: branch_sum_1 = (int32_t) +// 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..e0d999644 --- /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 v1, int32_t v2) { +// CHECK-NOT: PTOAS_NAME_HINTS +// CHECK: int32_t sum_value = ( +// CHECK: int32_t sum_value_1 = ( +// CHECK: int32_t class_v = ( diff --git a/tools/ptoas/ptoas.cpp b/tools/ptoas/ptoas.cpp index f0e2b9cd9..58804bf43 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 #include +#include +#include #include using namespace mlir; @@ -283,6 +287,185 @@ static bool parseAutoSyncTailHint(llvm::StringRef hintStr, std::string &normaliz return false; } +static bool isCppIdentifierStart(char c) { + return std::isalpha(static_cast(c)) || c == '_'; +} + +static bool isCppIdentifierChar(char c) { + return std::isalnum(static_cast(c)) || c == '_'; +} + +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 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 &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 &hints) { + if (auto nameLoc = dyn_cast(loc)) { + std::string sanitized = sanitizeCppIdentifier(nameLoc.getName().getValue()); + if (!sanitized.empty()) + hints.push_back(std::move(sanitized)); + return; + } + + if (auto fusedLoc = dyn_cast(loc)) { + if (Attribute metadata = fusedLoc.getMetadata()) { + if (auto strAttr = dyn_cast(metadata)) { + std::string sanitized = sanitizeCppIdentifier(strAttr.getValue()); + if (!sanitized.empty()) + hints.push_back(std::move(sanitized)); + return; + } + if (auto arrayAttr = dyn_cast(metadata)) { + for (Attribute attr : arrayAttr) { + auto strAttr = dyn_cast(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(loc)) { + appendLocationNameHints(callSiteLoc.getCallee(), hints); + if (hints.empty()) + appendLocationNameHints(callSiteLoc.getCaller(), hints); + } +} + +static SmallVector +getResultNameHints(Operation *op) { + SmallVector 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 std::string buildNameHintMarker(llvm::ArrayRef hints) { + std::string marker = "/* PTOAS_NAME_HINTS:"; + 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 opsToAnnotate; + module.walk([&](Operation *op) { + if (op->getNumResults() == 0 || isa(op)) + return WalkResult::advance(); + if (op->getParentOfType()) + 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 hints = getResultNameHints(op); + if (hints.empty()) + continue; + builder.setInsertionPoint(op); + builder.create(op->getLoc(), + builder.getStringAttr(buildNameHintMarker(hints))); + } +} + // -------------------------------------------------------------------------- // Post-process C++ output: rewrite marker calls into Tile member calls. // @@ -524,11 +707,15 @@ static void materializeControlFlowOperands(Operation *rootOp) { if (!initAttr) continue; + Location valueLoc = op->getLoc(); + if (Operation *def = value.getDefiningOp()) + valueLoc = def->getLoc(); + Value tmp = - builder.create(op->getLoc(), value.getType(), + builder.create(valueLoc, value.getType(), initAttr) .getResult(); - builder.create(op->getLoc(), tmp, value); + builder.create(valueLoc, tmp, value); operand.set(tmp); } } @@ -688,6 +875,238 @@ static void rewriteHoistedGlobalTensorDecls(std::string &cpp) { cpp.swap(out); } +static std::optional> +parseNameHintMarker(llvm::StringRef markerBody) { + llvm::SmallVector 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> +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 matches; + if (kTieRegex.match(snippet, &matches) && matches.size() >= 2) { + llvm::SmallVector 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{matches[2].str()}; + + return std::nullopt; +} + +static void stripNameHintMarkers(std::string &cpp) { + constexpr llvm::StringLiteral kMarkerPrefix = "/* PTOAS_NAME_HINTS:"; + size_t searchPos = 0; + while (true) { + size_t markerPos = cpp.find(kMarkerPrefix.str(), searchPos); + if (markerPos == std::string::npos) + break; + + size_t markerEnd = cpp.find("*/", markerPos + kMarkerPrefix.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 rewriteNameHintMarkers(std::string &cpp) { + constexpr llvm::StringLiteral kMarkerPrefix = "/* PTOAS_NAME_HINTS:"; + llvm::StringMap replacements; + std::set usedNames; + + size_t searchPos = 0; + while (true) { + size_t markerPos = cpp.find(kMarkerPrefix.str(), searchPos); + if (markerPos == std::string::npos) + break; + + size_t bodyBegin = markerPos + kMarkerPrefix.size(); + size_t markerEnd = cpp.find("*/", bodyBegin); + if (markerEnd == std::string::npos) + break; + + auto hints = + parseNameHintMarker(llvm::StringRef(cpp).slice(bodyBegin, markerEnd)); + searchPos = markerEnd + 2; + if (!hints) + continue; + + size_t windowEnd = std::min(searchPos + static_cast(2048), + cpp.size()); + llvm::StringRef searchWindow = + llvm::StringRef(cpp).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) { + llvm::StringRef oldName = (*generatedNames)[i]; + if (replacements.count(oldName)) + continue; + std::string newName = makeUniqueCppIdentifier((*hints)[i], usedNames); + if (newName.empty() || newName == oldName) + continue; + replacements[oldName] = std::move(newName); + } + } + + stripNameHintMarkers(cpp); + 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); +} + namespace { struct ConstantDeclCandidate { size_t declLine = 0; @@ -1224,6 +1643,7 @@ int main(int argc, char **argv) { 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 +1665,7 @@ int main(int argc, char **argv) { rewriteAddPtrTraceMarkers(cppOutput, emitAddPtrTrace); rewriteScalarConstantDecls(cppOutput); rewriteHoistedGlobalTensorDecls(cppOutput); + rewriteNameHintMarkers(cppOutput); *outputOS << cppOutput; outputOS->flush(); From e936ce3833f6ba9c6b97f3e9450e6a92c13ed1df Mon Sep 17 00:00:00 2001 From: HecreReed <821896444@qq.com> Date: Tue, 12 May 2026 17:30:40 +0800 Subject: [PATCH 2/6] feat: preserve textual pto name hints in emitc output --- docs/designs/ptoas-debug-name-hints-design.md | 23 +- test/lit/pto/debug_name_hints_cfg_emitc.pto | 5 +- test/lit/pto/debug_name_hints_emitc.pto | 2 +- .../debug_name_hints_textual_cfg_emitc.pto | 36 + .../pto/debug_name_hints_textual_emitc.pto | 28 + tools/ptoas/ptoas.cpp | 832 ++++++++++++++++-- 6 files changed, 847 insertions(+), 79 deletions(-) create mode 100644 test/lit/pto/debug_name_hints_textual_cfg_emitc.pto create mode 100644 test/lit/pto/debug_name_hints_textual_emitc.pto diff --git a/docs/designs/ptoas-debug-name-hints-design.md b/docs/designs/ptoas-debug-name-hints-design.md index c7acd990a..8837ebfef 100644 --- a/docs/designs/ptoas-debug-name-hints-design.md +++ b/docs/designs/ptoas-debug-name-hints-design.md @@ -51,8 +51,9 @@ PTOAS 将“原始变量名”视为调试元数据，而不是语义属性。 具体做法： - 前端把名字提示写入 op 的 `Location` 元数据 +- 对直接输入的 textual `.pto`，PTOAS 额外从源码文本里提取 SSA 名、函数参数名和 block argument 名 - PTOAS 在 rewrite / lowering 时尽量传播该元数据 -- EmitC/C++ 输出阶段从该元数据恢复出可用的局部变量名 +- EmitC/C++ 输出阶段从该元数据和源码文本提示恢复出可用的局部变量名 这样做的核心收益是： @@ -110,7 +111,9 @@ loc(fused[...]) 前端可以提供名字，也可以不提供。 - 提供时：PTOAS 尽量保留 -- 不提供时：PTOAS 继续使用现有 fallback 命名 +- 不提供时： + - 若输入是 textual `.pto`，PTOAS 会先尝试回收源码里的 SSA / 参数 / block arg 名 + - 若仍然没有名字提示，再使用现有 fallback 命名 因此该能力是增量增强，不破坏现有前端。 @@ -162,6 +165,11 @@ loc(fused[...]) - 若来源明确，则使用来源名加后缀，如 `_phi`、`_cond`、`_tmp` - 若来源不明确，则使用稳定 fallback 名 +本次实现里，控制流合流出来的 block argument 还额外做了一层 C++ 末端修正： + +- 先尽量把 block argument 名字传播到 EmitC value loc +- 若 EmitC/C++ hoist 后仍退化成匿名 `vN`，则在函数级 C++ 后处理中把顶部 hoist 出来的匿名合流临时按 block arg 名回填 + ## 8. C++ 命名规则 ### 8.1 基本规则 @@ -226,9 +234,15 @@ loc(fused[...]) 该设计需要覆盖以下环节： - 前端生成 PTO IR 时附带名字提示 +- textual `.pto` 输入时，从源码文本恢复 SSA / 参数 / block arg 名 - PTOAS rewrite / lowering helper 在替换时传播名字 - `PTOToEmitC` 中新建 `emitc::VariableOp`、`emitc::CastOp` 等值时继承或派生命名 -- 最终 `translateToCpp` 前或其包装层中应用名字分配规则 +- 最终 `translateToCpp` 前后在 PTOAS 包装层中应用名字分配规则 + +特别说明两点： + +- 函数参数名没有直接依赖 EmitC 对 block argument loc 的最终打印行为，而是按函数名在最终 C++ 签名上做稳定重写。 +- 这一步是有意的，因为当前 CppEmitter 对参数名和局部 `vN` 命名是统一内部编号逻辑，单靠上游 emitter 不能满足 issue `#337` 的可读性要求。 当前链路里已经存在大量重新物化变量的地方，因此不能只在 parser 或最终 printer 单点处理，必须贯穿传播和最终分配两个阶段。 @@ -239,11 +253,13 @@ loc(fused[...]) - 某些 pass 新建值但没有传播名字，会导致局部退化回 `vN` - 多结果 op 的 metadata 约定若不统一，前后端容易理解不一致 - 名字传播若写成语义属性，容易污染 IR；因此必须坚持“调试元数据”定位 +- C++ 末端重写需要按函数作用域工作；若按整个文件全局替换，多个函数复用 `v1/v2/...` 会串名 限制如下： - 该设计只能提供“尽量保留”，不能提供“逐号完全一致” - 对 aggressive CSE 后的公共值，只能保留最终幸存值的名字 +- 对 textual `.pto` 的 SSA 名恢复依赖源码文本与解析顺序一致；它是调试增强，不是语义机制 ## 12. 测试建议 @@ -254,6 +270,7 @@ loc(fused[...]) - 名字含非法字符：能被规范化 - 重名冲突：能稳定追加后缀 - 控制流 / `emitc.variable` / hoist：不会因为名字传播导致非法 C++ +- textual `.pto`：函数参数名、局部 SSA 名、CFG block arg 名都能出现在最终 `kernel.cpp` - 未提供 hint：仍保持现有 `vN` 回退行为 ## 13. 结论 diff --git a/test/lit/pto/debug_name_hints_cfg_emitc.pto b/test/lit/pto/debug_name_hints_cfg_emitc.pto index 5dbc88801..9fec18d57 100644 --- a/test/lit/pto/debug_name_hints_cfg_emitc.pto +++ b/test/lit/pto/debug_name_hints_cfg_emitc.pto @@ -24,10 +24,13 @@ module { } } -// CHECK-LABEL: static AICORE int32_t name_hints_cfg_helper(bool v1, int32_t v2, int32_t v3) { +// 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 index e0d999644..75313b023 100644 --- a/test/lit/pto/debug_name_hints_emitc.pto +++ b/test/lit/pto/debug_name_hints_emitc.pto @@ -16,7 +16,7 @@ module { } } -// CHECK-LABEL: static AICORE int32_t name_hints_basic_helper(int32_t v1, int32_t v2) { +// 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 = ( 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/tools/ptoas/ptoas.cpp b/tools/ptoas/ptoas.cpp index 58804bf43..588f1a90d 100644 --- a/tools/ptoas/ptoas.cpp +++ b/tools/ptoas/ptoas.cpp @@ -287,6 +287,25 @@ static bool parseAutoSyncTailHint(llvm::StringRef hintStr, std::string &normaliz return false; } +struct ParsedTextualNameHints { + llvm::SmallVector, 8> functionArgHints; + llvm::SmallVector, 8> blockArgHints; + llvm::SmallVector, 32> opResultHints; +}; + +using FunctionArgHintMap = + llvm::StringMap>; +using FunctionBlockArgHintMap = + llvm::StringMap, 4>>; + +static bool isGeneratedValueName(llvm::StringRef name); +static SmallVector getValueNameHints(Value value); + +static bool isMlirValueNameChar(char c) { + return std::isalnum(static_cast(c)) || c == '_' || c == '.' || + c == '$' || c == '-'; +} + static bool isCppIdentifierStart(char c) { return std::isalpha(static_cast(c)) || c == '_'; } @@ -295,6 +314,164 @@ static bool isCppIdentifierChar(char c) { return std::isalnum(static_cast(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 +parseMlirNamedArguments(llvm::StringRef text) { + llvm::SmallVector 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(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 &names) { + size_t pos = 0; + while (pos < line.size() && + std::isspace(static_cast(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(line[pos]))) + ++pos; + if (pos < line.size() && line[pos] == ',') { + ++pos; + while (pos < line.size() && + std::isspace(static_cast(line[pos]))) + ++pos; + continue; + } + break; + } + + while (pos < line.size() && + std::isspace(static_cast(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 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); @@ -402,6 +579,176 @@ static void appendLocationNameHints(Location loc, } } +static bool hasLocationNameHints(Location loc) { + SmallVector hints; + appendLocationNameHints(loc, hints); + return !hints.empty(); +} + +static Location attachLocationNameHints(Location baseLoc, + llvm::ArrayRef hints, + MLIRContext *context) { + SmallVector 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(attrs.front()), baseLoc); + return FusedLoc::get(ArrayRef{baseLoc}, ArrayAttr::get(context, attrs), + context); +} + +static void applyValueNameHints(Value value, llvm::ArrayRef hints) { + if (!value || hints.empty() || hasLocationNameHints(value.getLoc())) + return; + value.setLoc(attachLocationNameHints(value.getLoc(), hints, value.getContext())); +} + +static void applyOperationResultNameHints(Operation *op, + llvm::ArrayRef hints) { + if (!op || op->getNumResults() == 0 || hints.empty() || + hasLocationNameHints(op->getLoc())) + return; + + SmallVector limitedHints; + limitedHints.reserve(std::min(op->getNumResults(), hints.size())); + for (size_t i = 0, e = std::min(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 &blocks) { + for (Region ®ion : 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()) { + 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{argHints[index]}); + } + } + ++funcHintIndex; + } + + SmallVector 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{blockHints[argIndex]}); + } + } + + size_t opHintIndex = 0; + module.walk([&](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()) { + SmallVector argHints; + bool hasAllHints = func.getNumArguments() != 0; + for (BlockArgument arg : func.getArguments()) { + SmallVector 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()) { + SmallVector nonEntryBlocks; + collectNonEntryBlocksInSourceOrder(func.getOperation(), nonEntryBlocks); + if (nonEntryBlocks.empty()) + continue; + + SmallVector, 4> blockHints; + blockHints.reserve(nonEntryBlocks.size()); + for (Block *block : nonEntryBlocks) { + SmallVector argHints; + bool hasAllHints = block->getNumArguments() != 0; + for (BlockArgument arg : block->getArguments()) { + SmallVector 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()) { + auto it = blockArgHints.find(func.getSymNameAttr()); + if (it == blockArgHints.end() || it->second.empty()) + continue; + + SmallVector 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{argHints[argIndex]}); + } + } + } +} + static SmallVector getResultNameHints(Operation *op) { SmallVector hints; @@ -431,8 +778,19 @@ getResultNameHints(Operation *op) { return hints; } -static std::string buildNameHintMarker(llvm::ArrayRef hints) { - std::string marker = "/* PTOAS_NAME_HINTS:"; +static SmallVector getValueNameHints(Value value) { + SmallVector 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 hints) { + std::string marker = ("/* " + prefix + ":").str(); for (size_t i = 0; i < hints.size(); ++i) { if (i != 0) marker.push_back(','); @@ -461,8 +819,9 @@ static void annotateEmitCNameHints(ModuleOp module) { if (hints.empty()) continue; builder.setInsertionPoint(op); - builder.create(op->getLoc(), - builder.getStringAttr(buildNameHintMarker(hints))); + builder.create( + op->getLoc(), + builder.getStringAttr(buildHintMarker("PTOAS_NAME_HINTS", hints))); } } @@ -707,9 +1066,7 @@ static void materializeControlFlowOperands(Operation *rootOp) { if (!initAttr) continue; - Location valueLoc = op->getLoc(); - if (Operation *def = value.getDefiningOp()) - valueLoc = def->getLoc(); + Location valueLoc = value.getLoc(); Value tmp = builder.create(valueLoc, value.getType(), @@ -900,12 +1257,84 @@ parseNameHintMarker(llvm::StringRef markerBody) { return hints; } +static std::optional> +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 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(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> 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"); + R"re((^|[^[:alnum:]_])(v[0-9]+)[[:space:]]*(=|;))re"); llvm::SmallVector matches; if (kTieRegex.match(snippet, &matches) && matches.size() >= 2) { @@ -934,73 +1363,10 @@ findNextHintedGeneratedNames(llvm::StringRef snippet) { return std::nullopt; } -static void stripNameHintMarkers(std::string &cpp) { - constexpr llvm::StringLiteral kMarkerPrefix = "/* PTOAS_NAME_HINTS:"; - size_t searchPos = 0; - while (true) { - size_t markerPos = cpp.find(kMarkerPrefix.str(), searchPos); - if (markerPos == std::string::npos) - break; - - size_t markerEnd = cpp.find("*/", markerPos + kMarkerPrefix.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 rewriteNameHintMarkers(std::string &cpp) { - constexpr llvm::StringLiteral kMarkerPrefix = "/* PTOAS_NAME_HINTS:"; - llvm::StringMap replacements; - std::set usedNames; - - size_t searchPos = 0; - while (true) { - size_t markerPos = cpp.find(kMarkerPrefix.str(), searchPos); - if (markerPos == std::string::npos) - break; - - size_t bodyBegin = markerPos + kMarkerPrefix.size(); - size_t markerEnd = cpp.find("*/", bodyBegin); - if (markerEnd == std::string::npos) - break; - - auto hints = - parseNameHintMarker(llvm::StringRef(cpp).slice(bodyBegin, markerEnd)); - searchPos = markerEnd + 2; - if (!hints) - continue; - - size_t windowEnd = std::min(searchPos + static_cast(2048), - cpp.size()); - llvm::StringRef searchWindow = - llvm::StringRef(cpp).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) { - llvm::StringRef oldName = (*generatedNames)[i]; - if (replacements.count(oldName)) - continue; - std::string newName = makeUniqueCppIdentifier((*hints)[i], usedNames); - if (newName.empty() || newName == oldName) - continue; - replacements[oldName] = std::move(newName); - } - } - - stripNameHintMarkers(cpp); +static void rewriteIdentifiersWithMap( + std::string &cpp, const llvm::StringMap &replacements) { if (replacements.empty()) return; - std::string rewritten; rewritten.reserve(cpp.size()); enum class LexState { @@ -1107,6 +1473,311 @@ static void rewriteNameHintMarkers(std::string &cpp) { 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 +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(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 +parseGeneratedDeclarationName(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(")) + 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 (!isGeneratedValueName(name)) + return std::nullopt; + return name.str(); +} + +static llvm::SmallVector +findTopLevelGeneratedDeclarations(llvm::StringRef segment) { + llvm::SmallVector 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 functionParamHints, + llvm::ArrayRef> blockArgHints) { + static constexpr llvm::StringLiteral kResultMarkerPrefix = + "/* PTOAS_NAME_HINTS:"; + + llvm::StringMap replacements; + std::set usedNames; + + 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) { + llvm::StringRef oldName = (*generatedParams)[i]; + std::string newName = + makeUniqueCppIdentifier(functionParamHints[i], usedNames); + if (newName.empty() || newName == oldName) + continue; + replacements[oldName] = std::move(newName); + } + } + } + + if (!blockArgHints.empty()) { + llvm::SmallVector generatedDecls = + findTopLevelGeneratedDeclarations(segment); + llvm::SmallVector 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) { + llvm::StringRef oldName = generatedDecls[startIndex + i]; + if (replacements.count(oldName)) + continue; + std::string newName = + makeUniqueCppIdentifier(flattenedBlockHints[i], usedNames); + if (newName.empty() || newName == oldName) + continue; + replacements[oldName] = std::move(newName); + } + } + } + + auto processMarker = [&](llvm::StringRef markerPrefix, + auto generatedNameFinder) { + size_t searchPos = 0; + while (true) { + size_t markerPos = segment.find(markerPrefix.str(), searchPos); + if (markerPos == std::string::npos) + break; + + size_t bodyBegin = markerPos + markerPrefix.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(2048), + segment.size()); + llvm::StringRef searchWindow = + llvm::StringRef(segment).slice(searchPos, windowEnd); + auto generatedNames = generatedNameFinder(searchWindow); + if (!generatedNames) + continue; + + size_t pairCount = std::min(hints->size(), generatedNames->size()); + for (size_t i = 0; i < pairCount; ++i) { + llvm::StringRef oldName = (*generatedNames)[i]; + if (replacements.count(oldName)) + continue; + std::string newName = makeUniqueCppIdentifier((*hints)[i], usedNames); + if (newName.empty() || newName == oldName) + continue; + replacements[oldName] = std::move(newName); + } + } + }; + + processMarker(kResultMarkerPrefix, findNextHintedGeneratedNames); + + stripAllHintMarkers(segment); + rewriteIdentifiersWithMap(segment, replacements); +} + +static void rewriteNameHintMarkers(std::string &cpp, + const FunctionArgHintMap &functionArgHints, + const FunctionBlockArgHintMap &functionBlockArgHints) { + llvm::SmallVector 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 paramHints; + llvm::SmallVector, 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; @@ -1397,6 +2068,7 @@ int main(int argc, char **argv) { OwningOpRef module; llvm::StringRef buf = (*fileOrErr)->getBuffer(); + std::optional 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(); @@ -1446,6 +2118,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( @@ -1458,6 +2131,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. @@ -1637,6 +2320,7 @@ int main(int argc, char **argv) { return 1; } + applyFunctionBlockArgNameHintsToEmitC(*module, functionBlockArgHints); dropEmptyEmitCExpressions(module.get()); materializeControlFlowOperands(module.get()); if (failed(reorderEmitCFunctions(module.get()))) { @@ -1665,7 +2349,7 @@ int main(int argc, char **argv) { rewriteAddPtrTraceMarkers(cppOutput, emitAddPtrTrace); rewriteScalarConstantDecls(cppOutput); rewriteHoistedGlobalTensorDecls(cppOutput); - rewriteNameHintMarkers(cppOutput); + rewriteNameHintMarkers(cppOutput, functionArgHints, functionBlockArgHints); *outputOS << cppOutput; outputOS->flush(); From a48fca7b6686c1241e99d95d6963731d342b967a Mon Sep 17 00:00:00 2001 From: HecreReed <821896444@qq.com> Date: Tue, 12 May 2026 19:30:57 +0800 Subject: [PATCH 3/6] fix: avoid identifier collisions in name hint rewrite --- docs/designs/ptoas-debug-name-hints-design.md | 1 + ..._name_hints_identifier_collision_emitc.pto | 21 ++ tools/ptoas/ptoas.cpp | 274 ++++++++++++------ 3 files changed, 211 insertions(+), 85 deletions(-) create mode 100644 test/lit/pto/debug_name_hints_identifier_collision_emitc.pto diff --git a/docs/designs/ptoas-debug-name-hints-design.md b/docs/designs/ptoas-debug-name-hints-design.md index 8837ebfef..734530970 100644 --- a/docs/designs/ptoas-debug-name-hints-design.md +++ b/docs/designs/ptoas-debug-name-hints-design.md @@ -203,6 +203,7 @@ loc(fused[...]) - 首个值保留原名 - 后续值追加 `_1`、`_2`、`_3` +- 分配前先预留函数里已经存在的参数名、局部声明名和 `using` 引入的标识符，避免重定义 例如： 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/tools/ptoas/ptoas.cpp b/tools/ptoas/ptoas.cpp index 588f1a90d..6c2245ead 100644 --- a/tools/ptoas/ptoas.cpp +++ b/tools/ptoas/ptoas.cpp @@ -1534,8 +1534,178 @@ static bool isTopLevelFunctionStartLine(llvm::StringRef trimmed) { return true; } +static std::optional +parseAnyDeclaredIdentifierName(llvm::StringRef line); + +static llvm::SmallVector +findTopLevelGeneratedDeclarations(llvm::StringRef segment); + static std::optional parseGeneratedDeclarationName(llvm::StringRef line) { + auto declaredName = parseAnyDeclaredIdentifierName(line); + if (!declaredName || !isGeneratedValueName(*declaredName)) + return std::nullopt; + return declaredName; +} + +static std::set collectDeclaredIdentifiersInFunctionSegment( + llvm::StringRef segment) { + std::set 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(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 +collectPendingIdentifierRenames( + llvm::StringRef segment, llvm::ArrayRef functionParamHints, + llvm::ArrayRef> blockArgHints) { + static constexpr llvm::StringLiteral kResultMarkerPrefix = + "/* PTOAS_NAME_HINTS:"; + llvm::SmallVector 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 generatedDecls = + findTopLevelGeneratedDeclarations(segment); + llvm::SmallVector 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(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 +parseAnyDeclaredIdentifierName(llvm::StringRef line) { llvm::StringRef trimmed = line.trim(); if (trimmed.empty() || trimmed.starts_with("#") || trimmed.starts_with("//") || !trimmed.ends_with(";")) @@ -1545,7 +1715,8 @@ parseGeneratedDeclarationName(llvm::StringRef line) { 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("while ") || body.starts_with("while(") || + body.starts_with("using namespace ")) return std::nullopt; llvm::StringRef lhs = body; @@ -1555,7 +1726,8 @@ parseGeneratedDeclarationName(llvm::StringRef line) { if (lastWs == llvm::StringRef::npos) return std::nullopt; llvm::StringRef name = lhs.drop_front(lastWs + 1).trim(); - if (!isGeneratedValueName(name)) + if (name.empty() || !isCppIdentifierStart(name.front()) || + !llvm::all_of(name, isCppIdentifierChar)) return std::nullopt; return name.str(); } @@ -1601,92 +1773,24 @@ static int countBraceDelta(llvm::StringRef line) { static void rewriteNameHintsInFunctionSegment( std::string &segment, llvm::ArrayRef functionParamHints, llvm::ArrayRef> blockArgHints) { - static constexpr llvm::StringLiteral kResultMarkerPrefix = - "/* PTOAS_NAME_HINTS:"; - llvm::StringMap replacements; - std::set usedNames; - - 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) { - llvm::StringRef oldName = (*generatedParams)[i]; - std::string newName = - makeUniqueCppIdentifier(functionParamHints[i], usedNames); - if (newName.empty() || newName == oldName) - continue; - replacements[oldName] = std::move(newName); - } - } - } - - if (!blockArgHints.empty()) { - llvm::SmallVector generatedDecls = - findTopLevelGeneratedDeclarations(segment); - llvm::SmallVector 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) { - llvm::StringRef oldName = generatedDecls[startIndex + i]; - if (replacements.count(oldName)) - continue; - std::string newName = - makeUniqueCppIdentifier(flattenedBlockHints[i], usedNames); - if (newName.empty() || newName == oldName) - continue; - replacements[oldName] = std::move(newName); - } - } + llvm::SmallVector pendingRenames = + collectPendingIdentifierRenames(segment, functionParamHints, blockArgHints); + std::set 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); } - auto processMarker = [&](llvm::StringRef markerPrefix, - auto generatedNameFinder) { - size_t searchPos = 0; - while (true) { - size_t markerPos = segment.find(markerPrefix.str(), searchPos); - if (markerPos == std::string::npos) - break; - - size_t bodyBegin = markerPos + markerPrefix.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(2048), - segment.size()); - llvm::StringRef searchWindow = - llvm::StringRef(segment).slice(searchPos, windowEnd); - auto generatedNames = generatedNameFinder(searchWindow); - if (!generatedNames) - continue; - - size_t pairCount = std::min(hints->size(), generatedNames->size()); - for (size_t i = 0; i < pairCount; ++i) { - llvm::StringRef oldName = (*generatedNames)[i]; - if (replacements.count(oldName)) - continue; - std::string newName = makeUniqueCppIdentifier((*hints)[i], usedNames); - if (newName.empty() || newName == oldName) - continue; - replacements[oldName] = std::move(newName); - } - } - }; - - processMarker(kResultMarkerPrefix, findNextHintedGeneratedNames); - stripAllHintMarkers(segment); rewriteIdentifiersWithMap(segment, replacements); } From eaa6f457049188a771011842c712f94054c4d300 Mon Sep 17 00:00:00 2001 From: HecreReed <821896444@qq.com> Date: Tue, 12 May 2026 20:19:02 +0800 Subject: [PATCH 4/6] test: relax emitc checks for preserved debug names --- test/lit/pto/async_put_get_emitc.pto | 28 +++++++++---------- test/lit/pto/eventid_array_dyn_sync.pto | 14 +++++----- test/lit/pto/eventid_array_get_set_get.pto | 10 +++---- test/lit/pto/eventid_array_no_cse.pto | 8 +++--- .../pto/issue481_addptr_gm_slot_buffer.pto | 16 +++++------ .../issue481_addptr_gm_slot_buffer_gss.pto | 16 +++++------ .../issue556_tpop_live_values_no_alias.pto | 12 ++++---- test/lit/pto/set_validshape_if.pto | 6 ++-- .../lit/pto/set_validshape_local_lowering.pto | 6 ++-- test/lit/pto/tassign_level3_loop_rebind.pto | 2 +- .../pto/tassign_level3_loop_rebind_gss.pto | 2 +- test/lit/pto/tdivs_dual_order_emitc.pto | 2 +- test/lit/pto/tdivs_dual_order_emitc_gss.pto | 2 +- test/lit/pto/tgather_three_forms_emitc.pto | 2 +- .../tgetval_level3_index_cast_regression.pto | 2 +- test/lit/pto/tile_compact_mode_emitc.pto | 6 ++-- test/lit/pto/tprint_alloc_tile_no_rebind.pto | 4 +-- .../pto/tpush_tpop_dynamic_validshape_a5.pto | 16 +++++------ ...ush_tpop_dynamic_validshape_default_a5.pto | 4 +-- test/lit/pto/tpush_tpop_emitc.pto | 12 ++++---- .../pto/tpush_tpop_frontend_lowering_a3.pto | 12 ++++---- .../tpush_tpop_frontend_lowering_a3_gss.pto | 12 ++++---- .../pto/tpush_tpop_frontend_lowering_a5.pto | 14 +++++----- .../tpush_tpop_globaltensor_frontend_a3.pto | 8 +++--- .../tpush_tpop_globaltensor_internal_a3.pto | 2 +- test/lit/pto/trem_emitc.pto | 2 +- test/lit/pto/trsqrt_emitc.pto | 4 +-- test/lit/pto/tsort32_emitc.pto | 4 +-- test/lit/pto/tstore_forms_emitc.pto | 6 ++-- 29 files changed, 117 insertions(+), 117 deletions(-) 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 [[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 [[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([[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; -// 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( -// A3: pto::comm::AsyncEvent [[GET_EVT:v[0-9]+]] = pto::comm::TGET_ASYNC( -// 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( +// A3: pto::comm::AsyncEvent [[GET_EVT:[_A-Za-z][_A-Za-z0-9]*]] = pto::comm::TGET_ASYNC( +// 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/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/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 [[POP0:v[0-9]+]]; -// CHECK: TPOP, Tile, TileSplitAxis::TILE_UP_DOWN>({{v[0-9]+}}, [[POP0]]); -// CHECK: Tile [[POP1:v[0-9]+]]; -// CHECK: TPOP, Tile, TileSplitAxis::TILE_UP_DOWN>({{v[0-9]+}}, [[POP1]]); -// CHECK: TMOV({{v[0-9]+}}, [[POP0]]); -// CHECK: TMOV({{v[0-9]+}}, [[POP1]]); +// CHECK: Tile [[POP0:[_A-Za-z][_A-Za-z0-9]*]]; +// CHECK: TPOP, Tile, TileSplitAxis::TILE_UP_DOWN>({{[_A-Za-z][_A-Za-z0-9]*}}, [[POP0]]); +// CHECK: Tile [[POP1:[_A-Za-z][_A-Za-z0-9]*]]; +// CHECK: TPOP, Tile, 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/set_validshape_if.pto b/test/lit/pto/set_validshape_if.pto index c70476efb..42106aff4 100644 --- a/test/lit/pto/set_validshape_if.pto +++ b/test/lit/pto/set_validshape_if.pto @@ -31,11 +31,11 @@ module { } } -// CHECK: Tile [[TILE:v[0-9]+]]; +// CHECK: Tile [[TILE:[_A-Za-z][_A-Za-z0-9]*]]; // 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 e247d8548..a5decb870 100644 --- a/test/lit/pto/set_validshape_local_lowering.pto +++ b/test/lit/pto/set_validshape_local_lowering.pto @@ -15,6 +15,6 @@ module { } } -// CHECK: Tile [[TILE:v[0-9]+]]; -// CHECK: TASSIGN([[TILE]], [[ADDR:v[0-9]+]]); -// CHECK: [[TILE]].SetValidShape([[ROW:v[0-9]+]], [[COL:v[0-9]+]]) +// CHECK: Tile [[TILE:[_A-Za-z][_A-Za-z0-9]*]]; +// 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 [[T:v[0-9]+]]; +// CHECK: Tile [[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 [[T:v[0-9]+]]; +// CHECK: Tile [[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 e9e7c8707..0a73142d4 100644 --- a/test/lit/pto/tgather_three_forms_emitc.pto +++ b/test/lit/pto/tgather_three_forms_emitc.pto @@ -21,7 +21,7 @@ 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, MaskPattern::P1111>( // A3-NOT: reinterpret_cast< // A3-NOT: TGATHER, Tile, Tile, Tile, CmpMode::EQ, 7>( 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> // A3-DAG: memref.alloc() : memref<1x16xf16, strided<[16, 1]>, #pto.address_space> // A3-DAG: memref.alloc() : memref<1x16xf16, strided<[17, 1]>, #pto.address_space> -// A3: Tile [[DEFAULT:v[0-9]+]]; -// A3: Tile [[COMPACT:v[0-9]+]]; -// A3: Tile [[ROWP1:v[0-9]+]]; +// A3: Tile [[DEFAULT:[_A-Za-z][_A-Za-z0-9]*]]; +// A3: Tile [[COMPACT:[_A-Za-z][_A-Za-z0-9]*]]; +// A3: Tile [[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 [[TILE:v[0-9]+]]; -// CHECK: TASSIGN([[TILE]], [[ADDR:v[0-9]+]]); +// CHECK: Tile [[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 diff --git a/test/lit/pto/tpush_tpop_dynamic_validshape_a5.pto b/test/lit/pto/tpush_tpop_dynamic_validshape_a5.pto index f6d3de8a5..f8ac389da 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 [[CUBE_TILE:v[0-9]+]]; -// A5: [[CUBE_TILE]].SetValidShape({{v[0-9]+}}, {{v[0-9]+}}); -// A5: TPOP, Tile, 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 [[CUBE_TILE:[_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, Tile, 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 [[VEC_TILE:v[0-9]+]]; -// A5: [[VEC_TILE]].SetValidShape({{v[0-9]+}}, {{v[0-9]+}}); -// A5: TPOP, Tile, 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 [[VEC_TILE:[_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, Tile, 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 {{v[0-9]+}} = Tile({{v[0-9]+}}, {{v[0-9]+}}); +// A5: Tile {{[_A-Za-z][_A-Za-z0-9]*}} = Tile({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}); // A5-NOT: SetValidShape // A5: TPOP, Tile, TileSplitAxis::TILE_LEFT_RIGHT>( // A5-LABEL: AICORE void vector_kernel( -// A5: Tile {{v[0-9]+}} = Tile({{v[0-9]+}}, {{v[0-9]+}}); +// A5: Tile {{[_A-Za-z][_A-Za-z0-9]*}} = Tile({{[_A-Za-z][_A-Za-z0-9]*}}, {{[_A-Za-z][_A-Za-z0-9]*}}); // A5-NOT: SetValidShape // A5: TPOP, Tile, 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, Tile, 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 {{v[0-9]+}}; +// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_C2V, 1024, 8, 8, false>( +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TPOP, Tile, TileSplitAxis::TILE_UP_DOWN>( // A3: TFREE, 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 -// A3: Tile {{v[0-9]+}}; +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TPOP, Tile, TileSplitAxis::TILE_NO_SPLIT>( -// A3: Tile {{v[0-9]+}}; +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TMOV( // A3: TFREE, 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 {{v[0-9]+}}; +// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>( +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TPUSH, Tile, TileSplitAxis::TILE_NO_SPLIT>( // A3: TPOP, Tile, TileSplitAxis::TILE_NO_SPLIT>( -// A3: Tile {{v[0-9]+}}; +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TNEG( // A3: TFREE, 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 -// A3: Tile {{v[0-9]+}}; +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TPOP, Tile, TileSplitAxis::TILE_NO_SPLIT>( -// A3: Tile {{v[0-9]+}}; +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TMOV( // A3: TFREE, 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 {{v[0-9]+}}; +// A3: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 4, true>( +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TPUSH, Tile, TileSplitAxis::TILE_NO_SPLIT>( // A3: TPOP, Tile, TileSplitAxis::TILE_NO_SPLIT>( -// A3: Tile {{v[0-9]+}}; +// A3: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A3: TNEG( // A3: TFREE, 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 -// A5: Tile {{v[0-9]+}}; +// A5: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A5: TPOP, Tile, TileSplitAxis::TILE_NO_SPLIT>( -// A5: Tile {{v[0-9]+}}; +// A5: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A5: TMOV( // A5: TFREE, 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 {{v[0-9]+}}; -// A5: Tile {{v[0-9]+}}; +// A5: auto {{[_A-Za-z][_A-Za-z0-9]*}} = TPipe<0, Direction::DIR_BOTH, 1024, 4, 2, true>( +// A5: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; +// A5: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A5: TMOV( // A5: TPUSH, Tile, TileSplitAxis::TILE_NO_SPLIT>( // A5: TPOP, Tile, TileSplitAxis::TILE_NO_SPLIT>( -// A5: Tile {{v[0-9]+}}; +// A5: Tile {{[_A-Za-z][_A-Za-z0-9]*}}; // A5: TNEG( // A5: TFREE, 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, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND> [[CUBE_ENTRY:v[0-9]+]](nullptr); +// CHECK: GlobalTensor, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND> [[CUBE_ENTRY:[_A-Za-z][_A-Za-z0-9]*]](nullptr); // CHECK: TALLOC, GlobalTensor, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>({{.*}}, [[CUBE_ENTRY]]); // CHECK: TSTORE // CHECK: TPUSH, GlobalTensor, 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, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND> [[VEC_ENTRY:v[0-9]+]](nullptr); +// CHECK: GlobalTensor, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND> [[VEC_ENTRY:[_A-Za-z][_A-Za-z0-9]*]](nullptr); // CHECK: TPOP, GlobalTensor, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>( // CHECK: TLOAD // CHECK: TFREE, GlobalTensor, 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, GlobalTensor, pto::Stride<256, 256, 256, 16, 1>, pto::Layout::ND>, TileSplitAxis::TILE_NO_SPLIT>( // CHECK: TPUSH, GlobalTensor, 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([[DST0]], [[VEC]]); // A3: TSTORE<{{.*}}AtomicType::AtomicAdd>([[DST0]], [[VEC]]); // A3: TSTORE([[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([[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([[DST1]], [[ACC]], [[PRE]]); From ca1265fd6c3de784912f37a5fec3b510af01944c Mon Sep 17 00:00:00 2001 From: HecreReed <821896444@qq.com> Date: Tue, 12 May 2026 20:19:03 +0800 Subject: [PATCH 5/6] test: use tile-native input in issue226 sync regression --- test/lit/pto/issue226_remove_redundant_pipe_pair.pto | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/test/lit/pto/issue226_remove_redundant_pipe_pair.pto b/test/lit/pto/issue226_remove_redundant_pipe_pair.pto index 7b8f292f1..a1805abf0 100644 --- a/test/lit/pto/issue226_remove_redundant_pipe_pair.pto +++ b/test/lit/pto/issue226_remove_redundant_pipe_pair.pto @@ -15,15 +15,15 @@ module { func.func @remove_redundant_pipe_pair( - %arg0: memref<64x1xf16, strided<[1, 1]>, #pto.address_space>) { + %arg0: !pto.ptr) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c2 = arith.constant 2 : index %c64 = arith.constant 64 : index - %vbuf0 = pto.bind_tile %arg0, %c64, %c1 - {config = #pto.tile_buf_config} - : memref<64x1xf16, strided<[1, 1]>, #pto.address_space> - -> memref<64x1xf16, strided<[1, 1], offset: ?>, #pto.address_space> + %vview = pto.make_tensor_view %arg0, shape = [%c64, %c1], strides = [%c1, %c1] + : !pto.tensor_view<64x1xf16> + %vbuf0 = pto.partition_view %vview, offsets = [%c0, %c0], sizes = [%c64, %c1] + : !pto.tensor_view<64x1xf16> -> !pto.partition_tensor_view<64x1xf16> pto.section.cube { %mat_a = pto.alloc_tile : !pto.tile_buf @@ -33,7 +33,7 @@ module { %acc = pto.alloc_tile : !pto.tile_buf scf.for %i = %c0 to %c2 step %c1 { - pto.tload ins(%vbuf0 : memref<64x1xf16, strided<[1, 1], offset: ?>, #pto.address_space>) + pto.tload ins(%vbuf0 : !pto.partition_tensor_view<64x1xf16>) outs(%mat_a : !pto.tile_buf) pto.tmov ins(%mat_a : !pto.tile_buf) outs(%left : !pto.tile_buf) From 1a1738c8f47268967e7b75d8a012b9af7d4b3a5c Mon Sep 17 00:00:00 2001 From: HecreReed <821896444@qq.com> Date: Wed, 13 May 2026 16:38:03 +0800 Subject: [PATCH 6/6] fix: keep qwen3 custom golden compatible with hinted names --- .../Qwen3DecodeA3/qwen3_decode_golden_lib.py | 30 +++++++++++++++++++ .../Qwen3DecodeA5/qwen3_decode_golden_lib.py | 30 +++++++++++++++++++ 2 files changed, 60 insertions(+) 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)