Skip to content

Commit 8f5984c

Browse files
authored
[AMD] Enable automatic loop fusion (#5876)
Pick up #5726 for the AMD backend.
1 parent 8f1c71d commit 8f5984c

2 files changed

Lines changed: 5 additions & 2 deletions

File tree

python/test/unit/language/test_core.py

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -6732,8 +6732,6 @@ def test_tl_range_num_stages(device):
67326732

67336733

67346734
def test_tl_range_fuse():
6735-
if is_hip():
6736-
pytest.skip("loop fusion is not enabled on AMD")
67376735

67386736
@triton.jit
67396737
def kernel(ub):

third_party/amd/backend/compiler.py

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -232,6 +232,11 @@ def make_ttgir(mod, metadata, options):
232232
passes.ttgpuir.add_optimize_dot_operands(pm, True)
233233
amd.passes.ttgpuir.add_hoist_layout_conversions(pm)
234234

235+
passes.ttgpuir.add_fuse_nested_loops(pm)
236+
passes.common.add_canonicalizer(pm)
237+
passes.common.add_licm(pm)
238+
passes.common.add_canonicalizer(pm)
239+
235240
global_prefetch = int(os.getenv("TRITON_HIP_GLOBAL_PREFETCH", "0"))
236241
local_prefetch = int(os.getenv("TRITON_HIP_LOCAL_PREFETCH", "0"))
237242

0 commit comments

Comments
 (0)