feat(ptodsl): Migrate TileLib from TilelangDSL to PTODSL: Daemon and metadata#894
Merged
Merged
Conversation
Codex Review该评论由 review 机器人自动更新。
SummaryReview failed at stage Findings未生成结构化 findings，因为 review 过程提前失败。 Log Tail |
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Code Review
This pull request introduces the PTODSL TileLib backend, adding the InsertTemplateAttributes pass to attach legal template candidates to tile operations and updating ExpandTileOp to support both TileLang and PTODSL backends. It also implements a Unix-socket daemon server and client in Python to handle template metadata queries and specialization rendering. The review feedback is highly constructive, pointing out critical issues such as a potential denial-of-service vulnerability in socket message framing, race conditions in the template registry under concurrent access, and incorrect cache eviction logic. Additionally, the feedback addresses cross-platform compilation issues on Windows caused by POSIX-specific headers and functions, and suggests using MLIR's diagnostic infrastructure instead of direct llvm::errs() calls for better error reporting.
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.
ManiSadati
changed the title
Zhendong404
suggested changes
Jul 3, 2026
ManiSadati
force-pushed
the
mani/ptodsl
branch
2 times, most recently
from
c14ecf2 to
fe8ff93
Compare
… and constraint support
…f the changed files
Contributor
Author
|
/run a5 |
Contributor
Author
|
/run a3 |
|
已接收
页面会自动刷新，可以直接看当前阶段、排队情况和最近结果。 |
|
已接收
页面会自动刷新，可以直接看当前阶段、排队情况和最近结果。 |
A5 板测失败
日志尾部 |
A3 板测完成（有跳过）
|
Contributor
|
/run a5 |
|
已接收
页面会自动刷新，可以直接看当前阶段、排队情况和最近结果。 |
A5 板测失败
失败用例
|
A5 板测失败详情：PR #894test_intercore_sync_a5
shr
sels
sel
scatter
rowexpandsub
rowexpandmul
rowexpanddiv
bitcast_inplace_cvt
rope_kv_cache
rmsnorm
post_rmsnorm
partmin
partition5d_dynamic_a5
mscatter
movfp_fixpipe_reuse
mgather
colexpandsub
colexpandmax
cmps
cmp
|
Contributor
|
/run a5 |
|
已接收
页面会自动刷新，可以直接看当前阶段、排队情况和最近结果。 |
A5 板测失败
失败用例
|
A5 板测失败详情：PR #894sels
sel
scatter
rowexpandsub
rowexpandmul
rowexpanddiv
rope_kv_cache
rmsnorm
post_rmsnorm
vector_example_dag_kernel_add
plan_memory_nested_loops
partmin
mscatter
movfp_fixpipe_reuse
mgather
cmps
cmp
|
zhangstevenunity
requested changes
Jul 14, 2026
Collaborator
There was a problem hiding this comment.
Review - feat(ptodsl): Migrate TileLib to PTODSL (Daemon + metadata)
Deep-dive review of the C++ integration (InsertTemplateAttributes, ExpandTileOp, ptoas.cpp, PTOViewToMemref, FoldTileBufIntrinsics, TilelangDaemon), the PTODSL serving layer (daemon / wire / registry / render / constraints), and the 115 A5 templates. This is a well-architected, substantial piece of work, and the earlier bot-flagged infra issues are genuinely resolved. My one blocking concern is about what this PR ships as the default, not the infrastructure itself.
Blocker: this PR silently makes PTODSL the default backend while the A5 board is red
The description says "Kept TileLang as the default backend" and lists "Making PTODSL the default TileLib backend" under Out of scope. The code does the opposite:
tools/ptoas/ptoas.cpp:419->llvm::cl::init(TileLibBackend::PTODSL)include/PTO/Transforms/Passes.td:545->tileLibBackenddefault"ptodsl"
(Commit c5e583a9 "make ptodsl default"; ptodsl/README.md was updated to match, so only the PR body is now stale.)
This matters because the now-default A5 path is failing the board. Latest run in this PR (69df138b84fb, 2026-07-14): OK 209 / FAIL 17 / SKIP 1, with genuine numeric miscompiles (not infra flakes):
sel,sels,cmp,cmpsrowexpandsub,rowexpandmul,rowexpanddivscatter,mscatter,mgatherpartmin,movfp_fixpipe_reusermsnorm,post_rmsnorm,rope_kv_cache- segfaults:
plan_memory_nested_loops,vector_example_dag_kernel_add(exit 139)
The new expand_tile_op_ptodsl_* lit tests are FileCheck-structural only (they verify candidate-metadata insertion and MLIR shape, not numeric output), so ninja check-pto is green while these kernels compute wrong answers. The board (remote-NPU) is the only numeric gate, and it is red. The failures are A5-specific because templates exist only under templates/a5/ (A3/EmitC is unaffected, hence A3's 221/0).
Recommendation: keep --tile-lib-backend=tilelang as the default in this PR (matching the stated scope), land the PTODSL infra + templates behind the opt-in flag, and flip the default in a focused follow-up once the A5 board is green. That also de-risks the deliberate "no fallback to TileLang" policy: PTODSL-default + no-fallback means any daemon/env misconfiguration becomes a hard compile failure for every A5 user.
The failing set is disparate (the shared row-expand body and surface vsub/vmul/vdiv are operand-order-correct; the sel/cmp f32 paths use intricate pintlv_b16/punpack predicate interleaving that the out=64.0 constant mismatch implicates), so these look like per-template bugs the team is already iterating on. I am not asking to fix them all here, only not to default onto them yet.
Non-blocking
Daemon lifecycle (confirmed):
- Multi-daemon process leak in mixed-backend fatobj compiles.
runPTOASJobsruns child jobs in a sequential in-process loop (driver.cpp:1195), and each VPTO child with tile ops callsDaemonManager::start, which uses a pid-only socket path (TilelangDaemon.cpp:28) and overwrites the single staticprocessInfoslot with noisRunning()guard (TilelangDaemon.cpp:87). With >=2 VPTO-tile-op children, only the last daemon pid is tracked;stop()kills only that one, leaking the earlierpython -m ...daemonprocess(es). Suggest guarding onisRunning()(reuse the daemon) and/or making the socket path unique per job. - No cleanup on signal/crash. Cleanup is registered only via
std::atexit(TilelangDaemon.cpp:142), which does not run on SIGINT (Ctrl-C), SIGTERM, or SIGSEGV, leaking the daemon process and/tmp/tilelib_daemon_{pid}.sock. Pre-existing, but PR #894 makes this the default path, so it is now the common case.
Daemon has no per-connection recv timeout. serving/daemon.py:428 + wire.py:recv_exactly: the server is now sequential, with no timeout on accepted sockets. A client that sends a 4-byte length prefix then fewer bytes and holds the connection wedges the single serve thread indefinitely; later real RPCs hang until each caller's 30s ExecuteAndWait fires, and the daemon never recovers. Same-user only (chmod 0o600), and a killed client EOFs cleanly, so only a live partial-send wedges it.
constraints.py:332 swallows all constraint-predicate exceptions as "not satisfied." A buggy predicate (typo / AttributeError) silently makes its candidate illegal instead of surfacing; if it is the only candidate, the op later fails with a confusing "no legal template candidates." Consider letting unexpected exceptions propagate (or log them).
Compile-time cost: 2N Python subprocess spawns. InsertTemplateAttributes spawns one get_metadata helper per tile op and ExpandTileOp one render helper per op (ExpandTileOp.cpp:1015), each re-importing the client stack; the metadata half is new work this PR adds. Worth batching or reusing a persistent client in a follow-up.
Minor daemon hardening: TOCTOU between bind() and chmod(0o600) (daemon.py:326, tighten via umask(0o077) before bind); a successful render >64 MB escapes send_message as an uncaught ValueError past dispatch's try; and context_attrs merge precedence differs between metadata_request (attrs override) and evaluate_candidate (setdefault, built context wins) - align them.
Verified good
- Wire DoS fixed: length prefix checked against the 64 MiB cap before allocation (
wire.py:45). - Registry race fixed: sequential
UnixStreamServer, all locks removed - claim holds. - Broken-symlink socket cleanup and cache-eviction-on-overwrite: both correct.
replaceOpWithClonedAttrs(PTOViewToMemref.cpp:230) correctly preservescandidates/precision attrs across lowering (same op type in/out).- Pass order correct:
InsertTemplateAttributesruns beforeFusionPlan; daemon started only for VPTO + has-tile-ops + PTODSL; both passes receive the same auto-started socket.
| "Use the legacy TileLang DSL TileLib"), | ||
| clEnumValN(TileLibBackend::PTODSL, "ptodsl", | ||
| "Use the PTODSL TileLib daemon")), | ||
| llvm::cl::init(TileLibBackend::PTODSL)); |
Collaborator
There was a problem hiding this comment.
Blocker: this flips the default backend to PTODSL. The PR description says TileLang is kept as the default and lists making PTODSL default as out-of-scope, but this line (and Passes.td:545) default to ptodsl. Combined with the current A5 board failures (OK 209 / FAIL 17, real numeric miscompiles in sel/sels/cmp/cmps/rowexpand{sub,mul,div}/partmin/mgather/mscatter/rmsnorm/rope_kv_cache + segfaults), this defaults every A5 user onto a numerically-broken path while ninja check-pto stays green (lit tests are structural, not numeric). Suggest keeping tilelang as the default here and flipping in a follow-up once the board is green.
| "Path to Unix domain socket for daemon RPC (if empty, uses subprocess)"> | ||
| "Path to Unix domain socket for daemon RPC">, | ||
| Option<"tileLibBackend", "tile-lib-backend", "std::string", | ||
| /*default=*/"\"ptodsl\"", |
Collaborator
There was a problem hiding this comment.
Pass-side default also flipped to ptodsl here (mirrors ptoas.cpp:419). Same concern as the CLI default: this makes PTODSL the default for --emit-vpto pipelines while the A5 board is red. Keep tilelang until the numeric failures are resolved.
| return false; | ||
| } | ||
|
|
||
| processInfo = std::make_pair(procInfo.Pid, socketPath); |
Collaborator
There was a problem hiding this comment.
processInfo is overwritten unconditionally with no isRunning() check, and generateSocketPath() (line 28) keys the socket only on ::getpid() (constant for the whole process). In a mixed-backend fatobj compile, runPTOASJobs (driver.cpp:1195) runs VPTO child jobs in a sequential in-process loop; each tile-op child calls start() again with the same socket path. The earlier daemon's pid is lost, so the exit-time stop() kills only the last one and the earlier python -m ...daemon process leaks. Suggest guarding on isRunning() (reuse the daemon) and/or a per-job unique socket path.
| try: | ||
| if not predicate(**kwargs): | ||
| return False | ||
| except Exception: |
Collaborator
There was a problem hiding this comment.
passes() swallows any exception from a constraint predicate as "not satisfied" (return False). A buggy predicate (typo, AttributeError, KeyError) therefore silently marks its candidate illegal rather than surfacing the error; if it is the only candidate, the op later fails with a confusing "InsertTemplateAttributes found no legal template candidates" far from the real cause. Consider letting unexpected exceptions propagate (or at least logging them) so template-authoring bugs are visible.
| class _Handler(socketserver.BaseRequestHandler): | ||
| def handle(self): | ||
| try: | ||
| request = recv_message(self.request) |
Collaborator
There was a problem hiding this comment.
The daemon is now a sequential UnixStreamServer with no timeout set on accepted connections, and recv_exactly blocks until it gets the full announced length. A client that sends the 4-byte prefix then fewer bytes and holds the socket open wedges the single serve thread forever; every subsequent real RPC then hangs until each C++ caller's 30s ExecuteAndWait times out, and the daemon never recovers (a regression from the old threaded model). Consider a per-handler socket timeout.
A5 板测成功
|
A3 板测完成（有跳过）
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
6 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
PTOAS currently relies on TileLang DSL to render TileOp templates for the VPTO backend.
This PR introduces a PTODSL-native TileLib path while preserving the existing TileLang backend. It establishes the infrastructure needed to migrate templates incrementally without changing the default backend or requiring performance-based version selection yet.
What changed
PTODSL TileLib
taddtsubtmultmaxtmintdivtcolmaxtaddandtmul.PTODSL daemon
PTOAS integration
--tile-lib-backend=tilelang|ptodsl.InsertTemplateAttributes, which runs before fusion and stores legal candidates on each TileOp as an attribute.idnameloop_depthpostupdatetailExpandTileOpto consume the attached candidates and render the first remaining candidate.Current selection behavior
This PR does not introduce fusion-aware or cost-model-based version selection.
When multiple legal candidates remain,
ExpandTileOpdeterministically renders candidate index zero. A future change can filter or reorder the candidate array before expansion.Out of scope
lib/TileOps.Validation
PTODSL TileLib daemon
Run the focused daemon tests:
This covers:
PTODSL TileLib Python suite
Run all TileLib Python tests:
python3 -m unittest discover \ -s ptodsl/tests \ -p 'test_tilelib_*.py'This covers template registration, legality constraints, candidate selection,
structured MLIR rendering, elementwise templates, and daemon behavior.
PTOAS integration
Run the focused PTODSL and legacy TileLang integration tests through the
generated lit configuration:
This covers:
InsertTemplateAttributesplacement before fusion;