fix(rocm): avoid pointer-select crash in DSV4 sparse attention

[nekomario28]

Review/delivery carrier

This is an upstream-ready review carrier, not a merge target for this fork's diverged main. The base branch review/upstream-main-f7c31e is an immutable snapshot of current FlashML-org/FreeToken main at f7c31e92dbf296de3a5bb1b9c5fcb58f988e3a07 so the diff remains exactly one upstream-reviewable file.

Summary

Avoid an AMD Triton compiler assertion in both DSV4 sparse-attention paths by preserving the existing single-load pointer-select path on CUDA and using separately masked window/compressed pool loads only on ROCm.

Changed file only:

python/freetoken/kernel/triton/dsv4/sparse_attn.py

Head commit: 51a2d2ad7d22b2ce05eee6fa80f11cf52fbee6c9

Root cause

On RX 7800 XT (gfx1101) with PyTorch 2.11.0+rocm7.2, HIP 7.2.26015, Triton 3.6.0, both decode and prefill reproduced the same AMD compiler assertion in TritonAMDGPUCanonicalizePointers:

ConvertArithSelectOp: expected can narrow to be the same for both fatPtrT and fatPtrF

The failing IR came from the pointer-valued tl.where(is_win, win_ptr, cmp_ptr). This matches the failure class tracked in triton-lang/triton#9859.

Fix shape

SEPARATE_POOL_LOADS: tl.constexpr is enabled only when torch.version.hip is not None.

• ROCm: two separately masked pool loads; their mutually exclusive results are summed.
• CUDA: retains the original pointer select and single gather tl.load, preserving the existing shared-memory design constraint for consumer Blackwell.

Physical RED → GREEN evidence

Baseline RED:

• Dormant Giant temporary carrier Support Qwen3.8 27B FlashML-org/FreeToken#106
• run 32932005040, job 98065895708
• RX 7800 XT / gfx1101
• decode m=1, split-k: RED_REPRODUCED
• prefill m=2: RED_REPRODUCED
• both localized to TritonAMDGPUCanonicalizePointers

Hypothesis discriminator:

• carrier FreeToken engine installation doesn't work in Fedora 43 using AppImage FlashML-org/FreeToken#107, run 32932483894, job 98067252224
• separate masked loads: decode + prefill both GREEN_NUMERICAL vs PyTorch reference

Exact clean commit validation:

• run 32933110257, job 98069026160
• exact head 51a2d2ad7d22b2ce05eee6fa80f11cf52fbee6c9
• parent f7c31e92dbf296de3a5bb1b9c5fcb58f988e3a07
• product blob 7dcefe54183b12864a64d4cde82f2bcc63fa87fb
• SHA-256 a6efd0490949a9ca8bac8c83d8bca9caa59596189676bd67809b9c442964fd37
• provenance gate: PASS_MAIN_DIRECT_ONE_FILE_VERIFIED_BLOB
• decode: GREEN_NUMERICAL, splits=3
• prefill: GREEN_NUMERICAL, splits=0
• final classification: PASS_MAIN_DIRECT_FIX_GFX1101

Claim boundary

This proves the exercised DSV4 decode/prefill paths on RX 7800 XT / gfx1101 / ROCm 7.2 / Triton 3.6.0. It does not claim performance parity, full DSV4 model serving, all ROCm/Triton versions, or physical NVIDIA regression.

A direct upstream PR was attempted but the connected GitHub integration returned 403 Resource not accessible by integration; this carrier preserves the exact upstream-ready diff and evidence until an authorized GitHub/CLI path can deliver it.