[Triton] Integrate reorder-loop-loads into LoopUnroll pass

[isaname]

Background

After loop unrolling with tl.range(loop_unroll_factor=N), load and compute operations are interleaved in the loop body. This prevents optimal memory pipeline
utilization — ideally all loads should be issued early to overlap with subsequent computation.

Previously a standalone TritonReorderLoopLoads pass handled post-unroll load clustering. This PR merges that logic directly into the LoopUnroll pass, performing
reordering in-place immediately after unrolling, eliminating the extra pass and simplifying the pipeline.

Design

1. Reorder logic integrated into LoopUnroll

For loops with tt.reorder attribute:

• Full unroll: loads and their transitive deps are clustered to the front of the unrolled op range in the parent block.
• Partial unroll: same clustering is applied within the expanded loop body.

Key helpers: willFullyUnroll() (static trip count check), collectDepsInBlock() (transitive dep collection), reorderLoadsInRange() (load-first reordering
preserving relative order within each group).

2. User interface

Added reorder=False parameter to tl.range. When reorder=True with loop_unroll_factor set, the tt.reorder attribute is attached to the scf.for, triggering
load clustering during unroll.

for i in tl.range(0, N, loop_unroll_factor=4, reorder=True):
    a = tl.load(a_ptr + i * BLOCK)
    b = tl.load(b_ptr + i * BLOCK)
    c = a + b
    tl.store(c_ptr + i * BLOCK, c)

Performance

cp_gather_indexer_k_quant_cache

batch_size	seq_len	block_size	head_dim	quant_block_size	origin	pipeline	reorder
1	1	16	128	128	1.325	1.472	1.473
1	2	16	128	128	1.324	1.452	1.453
1	4	16	128	128	1.324	1.471	1.472
1	8	16	128	128	1.329	1.453	1.454
1	16	16	128	128	1.320	1.459	1.461
1	17	16	128	128	1.332	1.471	1.472
1	32	16	128	128	1.336	1.433	1.433
1	64	16	128	128	1.342	1.453	1.453
1	128	16	128	128	1.334	1.437	1.439
1	256	16	128	128	1.307	1.426	1.425
1	512	16	128	128	1.308	1.389	1.389
1	1024	16	128	128	1.309	1.389	1.387
1	2048	16	128	128	1.280	1.351	1.351
1	4096	16	128	128	1.298	1.358	1.360
1	8192	16	128	128	1.317	1.355	1.386
1	8192	64	128	128	1.318	1.382	1.369
1	16384	16	128	128	1.391	1.462	1.463
1	32768	16	128	128	1.715	1.720	1.719
1	32768	64	128	128	1.703	1.723	1.727
1	65536	16	128	128	2.015	2.001	1.995
1	65536	64	128	128	2.029	2.008	2.004
2	1	16	128	128	1.324	1.446	1.445
2	4096	16	128	128	1.249	1.241	1.246
4	1	16	128	128	1.293	1.428	1.429
4	4096	16	128	128	1.286	1.305	1.305
8	1	16	128	128	1.286	1.401	1.401
8	4096	16	128	128	1.189	1.212	1.213
16	1	16	128	128	1.394	1.502	1.503
16	2048	16	128	128	0.929	0.952	0.948
16	2048	64	128	128	0.925	0.948	0.947
32	1	16	128	128	1.323	1.386	1.386
32	1024	16	128	128	1.342	1.373	1.374
64	1	16	128	128	1.537	1.602	1.601
64	512	16	128	128	1.497	1.552	1.553
64	512	64	128	128	1.500	1.551	1.548
128	1	16	128	128	2.086	2.162	2.163
128	256	16	128	128	1.868	1.922	1.921
256	1	16	128	128	3.026	3.169	3.170
256	128	16	128	128	2.637	2.715	2.715
256	128	64	128	128	2.643	2.709	2.713

fused_inv_rope_fp8_quant

num_tokens	num_heads	n_groups	heads_per_group	tma_aligned_scales	origin	pipeline	reorder
256	128	8	16	False	1.063	1.163	1.200
256	128	8	16	True	1.003	1.117	1.141
512	128	8	16	False	1.057	1.142	1.189
512	128	8	16	True	1.003	1.126	1.131
1024	128	8	16	False	1.035	1.150	1.215
1024	128	8	16	True	0.996	1.140	1.158
2048	128	8	16	False	1.038	1.179	1.225
2048	128	8	16	True	0.999	1.163	1.178
4096	128	8	16	False	1.051	1.182	1.230
4096	128	8	16	True	1.001	1.159	1.180

[CLAassistant]

Thank you for your submission! We really appreciate it. Like many open source projects, we ask that you sign our Contributor License Agreement before we can accept your contribution.
_{You have signed the CLA already but the status is still pending? Let us recheck it.}