Replace tl.make_block_ptr with plain pointer arithmetic for Triton 3.8

[yuyzhang512]

Motivation

Triton 3.8 removed block pointers. tl.make_block_ptr still exists as a symbol but raises at trace time:

NotImplementedError: Block pointers have been removed in favor of the tensor descriptor API

so every kernel using it fails to compile. This is an API removal, not a GPU issue — it reproduces identically on gfx950 and gfx942. ATOM currently has 181 such sites across 15 files, all of which break the moment the Triton wheel moves to 3.8.

Technical Details

All 181 block accesses become plain pointer arithmetic with explicit bounds masks:

file	sites		file	sites
fla_ops/fused_merge_recompute.py	34		fla_ops/wy_fast.py	7
fla_ops/solve_tril.py	28		fla_ops/chunk_o.py	6
fla_ops/chunk_fused.py	24		fla_ops/chunk_o_vk.py	6
fla_ops/chunk_delta_h.py	23		minimax_m3/sparse_attn.py	6
fla_ops/chunk_delta_h_vk.py	23		fla_ops/fused_cumsum_kkt.py	5
fla_ops/cumsum.py	8		fla_ops/chunk_scaled_dot_kkt.py	4
sglang/.../minimax_m3_sparse.py	4		fla_ops/l2norm.py	2
minimax_m3/index_topk.py	1

# before
p = tl.make_block_ptr(base, (T, K), (H * K, 1), (i_t * BT, i_k * BK), (BT, BK), (1, 0))
b = tl.load(p, boundary_check=(0, 1))

# after
r = i_t * BT + tl.arange(0, BT)
c = i_k * BK + tl.arange(0, BK)
b = tl.load(base + r[:, None] * (H * K) + c[None, :],
            mask=(r < T)[:, None] & (c < K)[None, :], other=0.0)

Semantics are preserved exactly:

• Masks cover only the dims listed in the original boundary_check. Several sites here check just dim 0 (boundary_check=(0,), padding_option="zero"); a dim the author left unchecked was already known in-bounds, and masking it would change codegen for no reason.
• padding_option="zero" and the default both become other=0.0.
• <ptr>.dtype.element_ty is retargeted to the base pointer.
• There is no tl.advance anywhere in the repo, so no pointer-advancing loops needed restructuring.

cumsum.py was converted by hand rather than mechanically. Both of its kernels select between two alternative block pointers in an if/else and then issue a single load, so a naive "nearest preceding declaration" rewrite would silently bind the load to the else variant and drop the if one. The branch-dependent base pointer and stride are hoisted into the branches and the masked access is shared. No other file has this shape.

Test Plan

gfx950, triton 3.8.0+amd.rocm7.1.0.gitf6a045ff:

1. Confirm zero remaining tl.make_block_ptr sites and no syntax errors repo-wide.
2. Import all 14 affected modules.
3. Compile and run the kernels that have a checkable reference, comparing against torch.

Test Result

Zero make_block_ptr sites remain (was 181), no syntax errors, and all 14 affected modules import.

Kernels compiled and ran against torch references:

kernel	max |err|
chunk_local_cumsum_scalar (1-D, hand-converted branchy path)	2.4e-06
chunk_local_cumsum_vector (2-D, hand-converted branchy path)	0.0
l2norm_fwd	3.0e-08
solve_tril (vs torch.linalg.inv(I + A))	5.2e-04

Validation scope — please review with this in mind

The four kernels above are numerically verified. The remaining files — notably fused_merge_recompute.py (34 sites) and chunk_delta_h{,_vk}.py (23 each) — are converted by the same audited transform and import cleanly, but I could not find a test in this repo that exercises them numerically, so they are compile-validated only. Those are the places most worth a careful reviewer pass.

black could not be run here (it targets py3.15, the box has py3.12), so formatting has not been checked against the repo linter.

[github-actions]

🏷️ CI Guide

Runs automatically on every eligible PR before approval:

• ✅ Pre Checkin: Black, Ruff, catalog schema validation, non-GPU unit tests

Heavy model tests:

• ✅ Run after the PR is approved and Pre Checkin passes
• ✅ Run immediately when an approval review is submitted
• ✅ Can be requested before approval with labels

Label	Tests
ci:full	Run all heavy PR model tests: native ATOM, vLLM, and SGLang
ci:atom	Run native ATOM model accuracy tests
ci:vllm	Run ATOM vLLM OOT model accuracy tests
ci:sglang	Run ATOM SGLang model accuracy tests

Heavy jobs are skipped when the PR is not approved and no matching ci:* label is present.
Add labels via the sidebar or gh pr edit 2003 --add-label <label>