ub->gm传递l2_cache_ctl参数#920
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}
There was a problem hiding this comment.
Code Review
This pull request introduces support for GM store-side L2 cache control in the pto.mte_ub_gm and pto.copy_ubuf_to_gm operations by replacing a previously reserved field with the l2_cache_ctl parameter. The changes include updates to the dialect definition, parser, printer, verifier (ensuring constant values fit in the [0, 15] range), LLVM emitters, and documentation. Existing tests and examples have been updated to include the new parameter, and a new test file has been added to verify its roundtrip and expansion. There are no review comments, so I have no feedback to provide.
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.
and0d0
force-pushed
the
489-ub-gm-l2-cache
branch
from
7d57f84 to
f8686b4
Compare
zhangstevenunity
left a comment
zhangstevenunity
left a comment
There was a problem hiding this comment.
Reviewed the full cross-layer change: the ODS operand add (l2_cache_ctl as the 3rd operand of pto.mte_ub_gm), the reserved -> l2_cache_ctl rename on low-level pto.copy_ubuf_to_gm, both LLVM emitters (bit 60), ExpandDmaStorePattern forwarding, the [0,15] verifier, PTODSL _ops.py, docs, and the golden .pto updates.
Most of it is sound and backward compatible: default l2_cache_ctl=0 keeps the emitted config bits identical to the old reserved=0; bit 60 matches the load-direction copy_gm_to_ubuf packing; the docs "4 bits" field is consistent with the [0,15] check and shl 60; C++ has only two MteUbGmOp::build callers (both updated) and one CopyUbufToGmOp creator (forwards getL2CacheCtl()); PTODSL threads it as keyword-only, preserving existing positional callers, and the new regexes assert the operand position for both default (%c0) and explicit (%c7) cases.
One blocking issue: a second cross-layer emitter of pto.mte_ub_gm was missed.
[Blocker] TileLang DSL renderer still emits the old 3-operand form -> generated IR no longer parses
tilelang-dsl/python/tilelang_dsl/lowering.py _render_mte_ub_gm (around line 2528) reads args[2] as len_burst and renders:
pto.mte_ub_gm %src, %dst, %len_burst nburst(...) : !pto.ptr<..,ub>, !pto.ptr<..,gm>, i64, i64, i64, i64
(3 operands, 6 types). With l2_cache_ctl now a required 3rd operand, the updated parser puts %len_burst into the l2_cache_ctl slot and then fails at parseOperand(lenBurst) on the nburst keyword; the type list is also one short (6 vs the required 7). So any TileLang DSL v1 program using pto.mte_ub_gm(...) now produces IR that fails to parse/verify in ptoas.
The load counterpart in the same file, _render_mte_gm_ub, already emits ..., %l2_cache_ctl, %len_burst ... plus the matching type -- that is the pattern the store renderer should mirror. Since _analyze_mte_ub_gm only accepts 3 positional args + nburst/loops, the fix is to materialize a default 0 : i64 constant and emit it as the 3rd operand and 3rd type.
No test covers this: the only mte_ub_gm references under tilelang-dsl/tests/ are feature-tier/registration asserts, and tinsert_acc2vec.pto (correctly updated here) is a hand-authored ST input, not produced by this renderer -- so the break ships green.
Minor (non-blocking):
- The [0,15] range check is only on the store
MteUbGmOp::verify. The loadMteGmUbOp::verifyand low-levelCopyUbufToGmOp::verifystill do not range-checkl2_cache_ctl; the store path is safe because the wrapper forwards a verified value, but a hand-writtencopy_ubuf_to_gmwithl2_cache_ctl > 15would overflow into/past bit 63. A shared check would be more robust. - Surface asymmetry:
mte_gm_ub/mte_loadtakel2_cache_ctlpositionally (3rd arg), whilemte_ub_gm/mte_storetake it keyword-only at the tail. Reasonable for back-compat, but worth a doc note. docs/isa-legacy/02-dma-copy.mdstill labels the low-level 6th operand%reserved"Reserved field (set to 0)"; it is nowl2_cache_ctl.
| let arguments = (ins | ||
| PTO_BufferLikeType:$source, | ||
| PTO_BufferLikeType:$destination, | ||
| I64:$l2_cache_ctl, |
There was a problem hiding this comment.
Cross-layer emitter missed: TileLang DSL still renders the pre-l2_cache_ctl form.
Adding l2_cache_ctl as a required 3rd operand here also requires updating every emitter of the textual pto.mte_ub_gm. tilelang-dsl/python/tilelang_dsl/lowering.py:2528 _render_mte_ub_gm still emits:
pto.mte_ub_gm %src, %dst, %len_burst nburst(...) : ptr, ptr, i64, i64, i64, i64
(3 operands / 6 types). With this operand added, that IR fails to parse -- the parser puts %len_burst into the l2_cache_ctl slot and then hits the nburst keyword at parseOperand(lenBurst), and the type list is one short. Mirror _render_mte_gm_ub in the same file, which already emits ..., %l2_cache_ctl, %len_burst ... and the extra type; since the DSL surface does not expose l2_cache_ctl, materialize a default 0 : i64. No test exercises this render->parse path today, so it fails silently.
| let arguments = (ins | ||
| PTO_BufferLikeType:$source, | ||
| PTO_BufferLikeType:$destination, | ||
| I64:$l2_cache_ctl, |
There was a problem hiding this comment.
PR中提到：保持默认l2_cache_ctl=0，但是这里是强制要求参数？
There was a problem hiding this comment.
只在py引用的dsl接口层面非强制，后续没有值的话会补齐0，暂没考虑兼容旧.pto，可以改
| - **syntax:** | ||
| ```mlir | ||
| pto.mte_ub_gm %ub_src, %gm_dst, %len_burst | ||
| pto.mte_ub_gm %ub_src, %gm_dst, %l2_cache_ctl, %len_burst |
There was a problem hiding this comment.
通常，接口参数会按照参数热度排序。现在指令中的l2_cache_ctl参数，和上层接口/底层指令的参数位置都不太匹配：
- 上层dsl中是在末尾；
- 底层硬件指令
copy_ubuf_to_gm_align_v2中，是在nBurst, lenBurst之后； - 建议考虑下参数位置
and0d0
force-pushed
the
489-ub-gm-l2-cache
branch
from
74a4d0a to
20fafd5
Compare
4 participants
增加兼容性，保持默认l2_cache_ctl=0，显式传值时则使用显式值