[Bugfix] Bugfix of pass order for hopper (#117)

524991fe · Yu Cheng · GitHub · 62843b88 · 524991fe
Commit 524991fe authored Feb 25, 2025 by Yu Cheng Committed by GitHub Feb 25, 2025
Show whitespace changes
Inline Side-by-side

Showing with 4 additions and 4 deletions

tilelang/engine/lower.py tilelang/engine/lower.py +4 -4

No files found.
--- a/tilelang/engine/lower.py
+++ b/tilelang/engine/lower.py
@@ -203,14 +203,14 @@ def lower(
    mod = tl.transform.ThreadPartialSync("shared.dyn")(mod)
    mod = tir.transform.InferFragment()(mod)
    mod = tir.transform.LowerThreadAllreduce()(mod)
+    mod = tl.transform.LowerHopperIntrin()(mod)
+    mod = tl.transform.ThreadSync("shared")(mod)
+    mod = tl.transform.ThreadSync("shared.dyn")(mod)
+    mod = tir.transform.InjectPTXAsyncCopy()(mod)

    mod = tl.transform.AnnotateDeviceRegions()(mod)
    mod = tir.transform.SplitHostDevice()(mod)
    mod = tir.transform.MergeSharedMemoryAllocations()(mod)
-    mod = tl.transform.ThreadSync("shared")(mod)
-    mod = tl.transform.ThreadSync("shared.dyn")(mod)
-    mod = tl.transform.LowerHopperIntrin()(mod)
-    mod = tir.transform.InjectPTXAsyncCopy()(mod)

    mod = tl.transform.MakePackedAPI()(mod)
    mod = tir.transform.LowerDeviceKernelLaunch()(mod)