Skip to content

[temp] Narrow VPTO changes in LLVM21 upgrade#825

Draft
mouliangyu wants to merge 33 commits into
hw-native-sys:mainfrom
mouliangyu:codex/vpto-llvm21-api-only-temp
Draft

[temp] Narrow VPTO changes in LLVM21 upgrade#825
mouliangyu wants to merge 33 commits into
hw-native-sys:mainfrom
mouliangyu:codex/vpto-llvm21-api-only-temp

Conversation

@mouliangyu

Copy link
Copy Markdown
Contributor

Temporary PR for reviewing the LLVM21 upgrade shape relative to main.

Focus:

  • keep VPTO changes scoped to LLVM21 API/interface migration
  • remove extra low-precision vcvt/convert carrier logic
  • remove known-unsupported SIM skip plumbing
  • delete the incorrect low_precision_hivm_llvm_ir_abi lit check

Validation run locally:

  • git diff --check
  • ninja -C .work/ptoas-build-llvm21-gcc ptoas
  • llvm-lit -v .work/ptoas-build-llvm21-gcc/test/lit --filter "vcvt_low_precision|vreg_low_precision"
  • ninja -C .work/ptoas-build-llvm21-gcc check-pto
  • VPTO SIM: vcvt-low-precision-roundtrip, vcvt-low-precision-special, vreg-low-precision-ldst

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request upgrades the LLVM/MLIR dependency from LLVM 19 to LLVM 21, adapting the codebase to various API changes such as using memTy.getStridesAndOffset, applyPatternsGreedily, and updating EmitC variable generation to handle LValueType and LoadOp. It also updates build documentation, Dockerfiles, and Python dependencies. The review feedback highlights Python 3.8 compatibility issues in type annotations within test_docs_as_test.py, an O(N^2) complexity bottleneck in isPTOEntryFunction that can be optimized with an early-exit loop, and a potential crash in kernel.py due to an uncaught ValueError during file decoding.

Important

The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.

from dataclasses import dataclass
from pathlib import Path
from typing import Iterable
from typing import Iterable, Optional

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

Since the project supports Python 3.8+ (as documented in README.md), using standard collection types like dict and list as generics in type annotations (PEP 585) will cause a runtime TypeError on Python 3.8. Import Dict and List from typing to maintain compatibility.

Suggested change
from typing import Iterable, Optional
from typing import Dict, Iterable, List, Optional

compile_kwargs: dict[str, object] | None = None
fixture: str | None = None
symbol: Optional[str] = None
compile_kwargs: Optional[dict[str, object]] = None

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

Use Dict from typing instead of dict for Python 3.8 compatibility.

Suggested change
compile_kwargs: Optional[dict[str, object]] = None
compile_kwargs: Optional[Dict[str, object]] = None



def find_block_metadata(path: Path, lines: list[str], fence_line: int) -> DocBlockMetadata | None:
def find_block_metadata(path: Path, lines: list[str], fence_line: int) -> Optional[DocBlockMetadata]:

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

Use List from typing instead of list for Python 3.8 compatibility.

Suggested change
def find_block_metadata(path: Path, lines: list[str], fence_line: int) -> Optional[DocBlockMetadata]:
def find_block_metadata(path: Path, lines: List[str], fence_line: int) -> Optional[DocBlockMetadata]:

Comment on lines +262 to 263
extra_namespace: Optional[dict[str, object]] = None,
) -> dict[str, object]:

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

Use Dict from typing instead of dict for Python 3.8 compatibility.

Suggested change
extra_namespace: Optional[dict[str, object]] = None,
) -> dict[str, object]:
extra_namespace: Optional[Dict[str, object]] = None
) -> Dict[str, object]:

Comment thread lib/PTO/IR/PTO.cpp
Comment on lines 2535 to 2548
bool mlir::pto::isPTOEntryFunction(func::FuncOp func) {
if (!func || func.isDeclaration())
return false;
return hasExplicitPTOEntryAttr(func);
if (auto attr = func->getAttrOfType<BoolAttr>(kEffectivePTOEntryAttrName))
return attr.getValue();
if (hasExplicitPTOEntryAttr(func))
return true;

ModuleOp module = func->getParentOfType<ModuleOp>();
if (!module)
return false;
SmallVector<func::FuncOp> defs = getPTOFunctionDefinitions(module);
return defs.size() == 1 && defs.front() == func;
}

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

Calling getPTOFunctionDefinitions inside isPTOEntryFunction performs a full $O(N)$ scan of the module's functions. When isPTOEntryFunction is called in a loop over all functions (such as in validatePTOEntryFunctions), this results in $O(N^2)$ complexity. Since we only care if there is exactly one definition, we can optimize this with an early-exit loop that returns false as soon as a second definition is encountered, reducing the average complexity to $O(1)$ for modules with multiple functions.

bool mlir::pto::isPTOEntryFunction(func::FuncOp func) {
  if (!func || func.isDeclaration())
    return false;
  if (auto attr = func->getAttrOfType<BoolAttr>(kEffectivePTOEntryAttrName))
    return attr.getValue();
  if (hasExplicitPTOEntryAttr(func))
    return true;

  ModuleOp module = func->getParentOfType<ModuleOp>();
  if (!module)
    return false;

  func::FuncOp singleFunc = nullptr;
  for (auto f : module.getOps<func::FuncOp>()) {
    if (!f.isDeclaration()) {
      if (singleFunc)
        return false;
      singleFunc = f;
    }
  }
  return singleFunc == func;
}

except SyntaxError:
try:
module = ast.parse(Path(path).read_text(encoding="utf-8"))
except (OSError, IOError, SyntaxError):

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

If the file at path contains invalid UTF-8 bytes, Path(path).read_text(encoding="utf-8") will raise a UnicodeDecodeError (which inherits from ValueError). Since ValueError is not caught by except (OSError, IOError, SyntaxError), this will cause a compiler crash. Add ValueError to the caught exceptions to handle this gracefully.

Suggested change
except (OSError, IOError, SyntaxError):
except (OSError, IOError, SyntaxError, ValueError):

@reedhecre

reedhecre commented Jun 16, 2026

Copy link
Copy Markdown

Codex Review

该评论由 review 机器人自动更新。

  • PR: [temp] Narrow VPTO changes in LLVM21 upgrade #825 [temp] Narrow VPTO changes in LLVM21 upgrade
  • Author: mouliangyu
  • Base/Head: main / codex/vpto-llvm21-api-only-temp
  • Head SHA: c9a83ca45ff9
  • Trigger: PR 有新提交
  • Generated At: 2026-06-16T17:21:10Z
  • Previous Head SHA: 0ca5522669f0
  • Status: failed at codex-review (exit=1)

Summary

Review failed at stage codex-review: exit=1

Findings

未生成结构化 findings,因为 review 过程提前失败。

Log Tail

 test/lit/pto/syncfinder_zero_loop_if_probe_gss.pto |   4 +-
 test/lit/pto/tassign_level3_loop_rebind.pto        |   3 +-
 test/lit/pto/tassign_level3_loop_rebind_gss.pto    |   3 +-
 test/lit/pto/tci_i16_emitc.pto                     |   4 +-
 test/lit/pto/tci_ui32_emitc.pto                    |   4 +-
 test/lit/pto/tprint_alloc_tile_no_rebind.pto       |   5 +-
 .../pto/tpush_tpop_globaltensor_frontend_a3.pto    |   6 +-
 .../pto/tpush_tpop_globaltensor_frontend_a5.pto    |  10 +-
 test/lit/pto/treshape_static_valid_shape_emitc.pto |   6 +-
 .../tile_fusion/mark_last_use_slot_mask_level2.pto |  10 +-
 .../a5/src/st/smoke/testcase/tmrgsort/tmrgsort.pto |   2 +-
 .../npu/a5/src/st/testcase/tmrgsort/tmrgsort.pto   |   2 +-
 tilelang-dsl/python/tilelang_dsl/kernel.py         |  21 +-
 tools/ptoas/ptoas.cpp                              | 261 +++++++++++++++--
 tools/ptobc/src/mlir_encode.cpp                    |  35 ++-
 tools/ptobc/src/ptobc_decode_print.cpp             |   3 +-
 .../ptobc/testdata/recent_mx_ops_v0_roundtrip.pto  |  26 ++
 tools/ptobc/testdata/recent_ops_v0_roundtrip.pto   |  12 -
 tools/ptobc/tests/CMakeLists.txt                   |   1 +
 tools/ptobc/tests/recent_ops_v0_encode.sh          |  11 +-
 .../ptobc/tests/tscatter_maskpattern_v0_encode.sh  |  14 +-
 93 files changed, 1066 insertions(+), 534 deletions(-)
===== END STAGE clone rc=0 @ 2026-06-17 01:20:35 =====

===== STAGE codex-review @ 2026-06-17 01:20:35 =====
set -euo pipefail
cd '/tmp/ptoas-pr-review-monitor/runs/20260617_012027_pr825/repo'
'codex' exec -C '/tmp/ptoas-pr-review-monitor/runs/20260617_012027_pr825/repo' -s read-only -c 'model_provider="codereview"' -c 'model="gpt-5.4"' -c 'model_reasoning_effort="xhigh"' --output-schema '/tmp/ptoas-pr-review-monitor/runs/20260617_012027_pr825/review_schema.json' -o '/tmp/ptoas-pr-review-monitor/runs/20260617_012027_pr825/codex_last_message.json' --color never - < '/tmp/ptoas-pr-review-monitor/runs/20260617_012027_pr825/review_prompt.txt'
OpenAI Codex v0.115.0 (research preview)
--------
workdir: /tmp/ptoas-pr-review-monitor/runs/20260617_012027_pr825/repo
model: gpt-5.4
provider: codereview
approval: never
sandbox: read-only
reasoning effort: xhigh
reasoning summaries: none
session id: 019ed173-03a2-7cb1-bb1a-91663c03885d
--------
user
你现在在审查 GitHub PR。

仓库:hw-native-sys/PTOAS
PR:#825 [temp] Narrow VPTO changes in LLVM21 upgrade
作者:mouliangyu
base branch:origin/main
head branch:HEAD(当前已 checkout 到 PR head)

要求:
1. 只审查这个 PR 相对 origin/main 的改动,必要时可以看上下文文件。
2. 重点找真实的 correctness / regression / contract mismatch / CI / runtime / compatibility 问题。
3. 不要提纯风格建议,不要提低价值猜测。
4. 严格按优先级输出:
   - P1:高概率会导致错误结果、编译/运行失败、严重回归、发布阻断
   - P2:重要缺陷、行为回归、遗漏校验/测试、较大兼容性问题
   - P3:次要但明确可改的问题
5. 如果没有问题,summary 直接写:未检查到 PR #825 存在问题,并返回 findings=[]。
6. 如果有问题,summary 简洁概括,findings 里每条都要给出:
   - severity
   - title
   - body(说明为什么是问题,尽量具体)
   - file(尽量给相对路径)
   - line(能确定就填整数,否则 null)

建议先查看:
- git status --short
- git diff --stat origin/main...HEAD
- git diff --unified=80 origin/main...HEAD

最终输出必须严格匹配 JSON schema。

mcp startup: no servers
Reconnecting... 1/5 (unexpected status 503 Service Unavailable: Service temporarily unavailable, url: https://codex.0u0o.com/responses, request id: 945908b9-4edf-41a9-9770-adb0c49c1d92)
Reconnecting... 2/5 (unexpected status 503 Service Unavailable: Service temporarily unavailable, url: https://codex.0u0o.com/responses, request id: 19d74b78-0f7a-4fc3-83c6-cf00b43638ca)
Reconnecting... 3/5 (unexpected status 503 Service Unavailable: Service temporarily unavailable, url: https://codex.0u0o.com/responses, request id: 81e25a7d-ad92-4b54-b748-d563080054b2)
Reconnecting... 4/5 (unexpected status 503 Service Unavailable: Service temporarily unavailable, url: https://codex.0u0o.com/responses, request id: fcc83cb3-e32f-404f-86d3-957fde3429c6)
Reconnecting... 5/5 (unexpected status 503 Service Unavailable: Service temporarily unavailable, url: https://codex.0u0o.com/responses, request id: 621363ce-1658-4517-a696-2d554834fe79)
ERROR: unexpected status 503 Service Unavailable: Service temporarily unavailable, url: https://codex.0u0o.com/responses, request id: 4c25c438-cdac-4c6b-8f81-ff55f2f4bcab
Warning: no last agent message; wrote empty content to /tmp/ptoas-pr-review-monitor/runs/20260617_012027_pr825/codex_last_message.json
===== END STAGE codex-review rc=1 @ 2026-06-17 01:21:10 =====

@mouliangyu mouliangyu force-pushed the codex/vpto-llvm21-api-only-temp branch 9 times, most recently from fec267a to 0ca5522 Compare June 16, 2026 16:57
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants