Skip to content

[Bug] PTODSL failed to compile when a variable carried by a while loop is conditionally reassigned before the loop #1558

Closed

Description

@liggest

Component

PTODSL

Description

PTODSL 中,触及全部下述条件的 kernel 在转换为 mlir 时报错

  1. 一个变量在 while 之前的 if 里被赋值(if n > 0: acc = ...)
  2. 该变量在 while 块内先读后写(acc = acc + ...)
  3. 该变量在 while 之后不再被使用

报错大致为

MLIRError: Verification failed:
error: unknown: operand #0 does not dominate this use

Reproduction (minimal)

  • 此用例报错
from ptodsl import pto


@pto.jit(name="while_carry_merge", kernel_kind="vector", target="a5", mode="explicit")
def kernel(out: pto.ptr(pto.i32, "gm"), n: pto.i32):
    buf = pto.castptr(pto.const(0, dtype=pto.i64), pto.ptr(pto.ui8, "ub"))
    acc = pto.const(0.0, dtype=pto.f32)
    if n > 0:
        acc = pto.load(pto.castptr(buf, pto.ptr(pto.f32, "ub")), 100)
    i = pto.const(0, dtype=pto.i32)
    while i < n:
        acc = acc + pto.load(pto.castptr(buf, pto.ptr(pto.f32, "ub")), pto.cast(i, pto.i64))
        i = pto.cast(i + 1, pto.i32)
    pto.store(i, out, pto.const(0, dtype=pto.i64))


if __name__ == "__main__":
    kernel.compile()
    print("compiled OK")
  • 只要在循环结束后仍使用变量,便不会报错
from ptodsl import pto


@pto.jit(name="while_carry_merge", kernel_kind="vector", target="a5", mode="explicit")
def kernel(out: pto.ptr(pto.i32, "gm"), n: pto.i32):
    buf = pto.castptr(pto.const(0, dtype=pto.i64), pto.ptr(pto.ui8, "ub"))
    acc = pto.const(0.0, dtype=pto.f32)
    if n > 0:
        acc = pto.load(pto.castptr(buf, pto.ptr(pto.f32, "ub")), 100)
    i = pto.const(0, dtype=pto.i32)
    while i < n:
        acc = acc + pto.load(pto.castptr(buf, pto.ptr(pto.f32, "ub")), pto.cast(i, pto.i64))
        i = pto.cast(i + 1, pto.i32)
    pto.store(i, out, pto.const(0, dtype=pto.i64))
+    pto.store(acc, pto.castptr(buf, pto.ptr(pto.f32, "ub")), 200)  # 依然使用 acc


if __name__ == "__main__":
    kernel.compile()
    print("compiled OK")

Expected behavior

PTODSL 编译成功

Actual behavior / error logs

Traceback (most recent call last):
  File "test_while_carry.py", line 18, in <module>
    while_carry_merge_repro.compile()
  File "ptodsl/ptodsl/_jit.py", line 343, in compile
    return self._compiler.compile(**constexpr_bindings)
  File "ptodsl/ptodsl/_kernel_compilation.py", line 142, in compile
    compiled.build()
  File "ptodsl/ptodsl/_tracing/artifacts.py", line 36, in build
    built = self._module_factory()
  File "ptodsl/ptodsl/_tracing/runtime.py", line 98, in build_module
    self.verify_module(module)
  File "ptodsl/ptodsl/_tracing/runtime.py", line 52, in verify_module
    module.operation.verify()
ptoas.mlir._mlir_libs._site_initialize.<locals>.MLIRError: Verification failed:
error: unknown: operand #0 does not dominate this use
 note: unknown: see current operation: 
  %6:2 = "scf.while"(%19, %5) ({
  ^bb0(%arg4: f32, %arg5: i32):
    %16 = "pto.cmpi"(%arg5, %arg1) <{predicate = #pto.scalar_cmp_predicate<lt>, signedness = #pto.signedness<signed>}> : (i32, i32) -> i1
    "scf.condition"(%16, %arg4, %arg5) : (i1, f32, i32) -> ()
  }, {
  ^bb0(%arg2: f32, %arg3: i32):
    %9 = "pto.castptr"(%1) : (!pto.ptr<ui8, ub>) -> !pto.ptr<f32, ub>
    %10 = "pto.exti"(%arg3) <{signedness = #pto.signedness<signed>}> : (i32) -> i64
    %11 = "pto.index_cast"(%10) <{signedness = #pto.signedness<signed>}> : (i64) -> index
    %12 = "pto.load"(%9, %11) : (!pto.ptr<f32, ub>, index) -> f32
    %13 = "pto.addf"(%arg2, %12) <{fastmath = #pto.fastmath<none>}> : (f32, f32) -> f32
    %14 = "pto.constant"() <{value = 1 : i32}> : () -> i32
    %15 = "pto.addi"(%arg3, %14) <{overflowFlags = #pto.overflow<none>}> : (i32, i32) -> i32
    "scf.yield"(%13, %15) : (f32, i32) -> ()
  }) : (f32, i32) -> (f32, i32)
 note: unknown: operand defined here (op in a child region)

实际场景生成的 IR 中,acc 的新值产生于 if 分支内部,if 本身没有返回任何结果;而紧跟其后的 while 把分支内部那个值拿来做循环初始值

"scf.if"(%4) ({
  %19 = "pto.load"(%17, %18) : (!pto.ptr<f32, ub>, index) -> f32
  "scf.yield"() : () -> ()
}, {}) : (i1) -> ()             // if 没有结果,acc 的新值留在分支内部
%6:2 = "scf.while"(%19, %5) ({  // while 用 %19 做初始值 → %19 不支配该使用
  ...

Git commit

48e67a2

Host platform

Linux (x86_64)

Target Ascend arch (if relevant)

a5

PTOAS build level (if relevant)

No response

Activity

  1. added
    bugSomething isn't working
    tilelangIssues from tilelang
    on Sep 18, 2026
  2. TaoTao-real commented on Sep 21, 2026

    @TaoTao-real
    Contributor

    已确认该问题属于 PTODSL AST 活跃变量分析缺陷,需要在 PTOAS 仓库中修复。

    while 分析错误地从循环入口活跃集合中移除了循环体内赋值的变量,导致循环后不再使用的 loop-carried 值没有成为前置 if 的返回值。随后 scf.while 直接引用分支内部定义,产生 dominance verification error。修复后保留这类循环入口活跃变量。

    GitCode PR:https://gitcode.com/cann/pto-as/merge_requests/237

    验证:原始模式在基线提交上可稳定重现支配关系错误;修复后编译通过,ptodsl/tests/test_scf_while_ast.py 通过。

  3. liggest commented on Sep 21, 2026

    @liggest
    ContributorAuthor

    已验证 PR 有效,感谢🎉

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    bugSomething isn't workingtilelangIssues from tilelang

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions