[AMDGPU] Add DPP row_shl fast path for shuffle_down reduction offsets

[ZJLi2013]

Summary

amdgpu_shuffle_down_i32 currently always routes through the generic cross-half ds_bpermute + permlane64 helper, even for the small power-of-two offsets used by reduction trees. The existing FIXME calls this out:

Currently emulates shuffle_down via the cross-half ds_bpermute + permlane64 helper (~50-60 cycle latency). Should be upgraded to use DPP ROW_SHR instructions (~4-12 cycles) for reduction-pattern offsets (1, 2, 4, 8, 16); the cross-half case (offset >= 32) still needs the helper.

This PR implements that fast path for offsets {1, 2, 4, 8}: a single v_mov_b32_dpp per shuffle instead of the multi-instruction cross-half gather. Measured ~2.1x lower latency on RDNA4 and ~1.36x on CDNA3, with reduction consumers unchanged.

(One correction to the FIXME: shuffle_down reads lane i + offset, i.e. data moves left, so the correct instruction is row_shl, not row_shr.)

Change

For offsets {1, 2, 4, 8} emit DPP directly; everything else falls through to the existing helper:

i32 amdgpu_shuffle_down_i32(i32 offset, i32 value) {
#if defined(ARCH_amdgpu) && (defined(__x86_64__) || defined(__i386__) || defined(__amdgcn__))
  i32 r = value;
  if (offset == 1) { __asm__ volatile("v_mov_b32_dpp %0, %1 row_shl:1 row_mask:0xf bank_mask:0xf" : "+v"(r) : "v"(value)); return r; }
  if (offset == 2) { __asm__ volatile("v_mov_b32_dpp %0, %1 row_shl:2 row_mask:0xf bank_mask:0xf" : "+v"(r) : "v"(value)); return r; }
  if (offset == 4) { __asm__ volatile("v_mov_b32_dpp %0, %1 row_shl:4 row_mask:0xf bank_mask:0xf" : "+v"(r) : "v"(value)); return r; }
  if (offset == 8) { __asm__ volatile("v_mov_b32_dpp %0, %1 row_shl:8 row_mask:0xf bank_mask:0xf" : "+v"(r) : "v"(value)); return r; }
#endif
  return amdgpu_cross_half_shuffle_i32(amdgpu_lane_id() + offset, value);
}

Design notes:

• Offsets covered. row_shl:N shifts within each 16-lane row. Offsets {1,2,4,8} stay inside a row; offset 16 crosses a row and offset >= 32 crosses the wave half, so both keep the generic helper. offset is a compile-time constant after inlining, so the dead branches are eliminated.
• Boundary lanes. Lanes whose source falls outside the 16-lane row keep their own value (r is seeded with value). This matches how the reduction consumers use shuffle_down: reduce_*_tiled / block.reduce are decreasing-offset trees (32 -> 1), so by the time the offset is small the live lanes have contracted into a single 16-lane row and the boundary lanes are don't-care.
• shuffle_up intentionally unchanged. Its consumers are increasing-offset Hillis-Steele scans where every lane stays active and reads across a row boundary at offset 1 (lane 16 reads lane 15); row_shr returns the boundary lane's own value there, which would break the scan. shuffle_up keeps the generic helper.
• Why inline asm and not __builtin_amdgcn_update_dpp. runtime.cpp is compiled to bitcode by the host clang and only retargeted to amdgcn at JIT time, so the intrinsic is not available at compile time. The instruction text is carried verbatim into the IR — the same technique already used by the "+v" fence in amdgpu_cross_half_shuffle_i32.

Testing

Rebuilt the runtime bitcode from source and ran standalone quadrants kernels (no Genesis), single wave64, log2_size=6.

Correctness (reduce_add_tiled(6) lane 0 == sum(1..64) == 2080; interior shuffle_down checked per lane):

GPU	reduce consumers	shuffle_down interior
R9700 (gfx1201 / RDNA4)	PASS	PASS
MI300 (gfx942 / CDNA3)	PASS	PASS

Latency — single wave64, long data-dependent chain, ns per shuffle_down + fixed ALU (before = generic helper, after = DPP; offsets 16/32 stay on the helper in both and serve as an in-binary control):

R9700 (gfx1201 / RDNA4):

offset	path	before	after	speedup
1/2/4/8	DPP row_shl	~42.7 ns	~20 ns	~2.1x
16	helper	42.8	42.8	—
32	helper	33.3	33.3	—

MI300 (gfx942 / CDNA3):

offset	path	before	after	speedup
1/2/4/8	DPP row_shl	~68.2 ns	~50.2 ns	~1.36x
16	helper	68.25	68.19	—
32	helper	61.46	61.42	—

The CDNA3 speedup is smaller because there the cross-half helper is already a single wave-global ds_bpermute (permlane64 is identity on CDNA), so the baseline is cheaper; DPP still wins. The unchanged offset-16/32 numbers confirm only {1,2,4,8} take the new path and that the measurement is stable.

[chatgpt-codex-connector]

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 9998427dc4

ℹ️ About Codex in GitHub

Codex has been enabled to automatically review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

When you sign up for Codex through ChatGPT, Codex can also answer questions or update the PR, like "@codex address that feedback".

[chatgpt-codex-connector]

Preserve shuffle_down across DPP row boundaries

When subgroup.shuffle_down(value, 1) reaches this AMDGPU fast path, lanes at 16-lane DPP row boundaries (for example lane 15) keep their own value instead of reading lane 16, even though the public API/codegen route all subgroupShuffleDown calls here and document lane i + offset semantics for every in-range lane. This may be acceptable for the current reduction tree's dead lanes, but it corrupts any direct shuffle_down use or future algorithm that needs boundary lanes; keep the generic helper for boundary-sensitive calls or make this optimization reduction-specific.

Useful? React with 👍 / 👎.

[hughperkins]

@ZJLi2013 This codex comment seems serious and structural to me.

I don't think we can corrupt the meaning of 'shuffle down'. It is a well known operation, and should behave unchanged, across all platforms.

One option that occurs to me is to create a new operation 'shift_left', and implement this across all GPU platforms (Metal, Vulkan, CUDA, AMD).

Thoughts?