From 9998427dc400d94ba3c141f39df6687183a36d7e Mon Sep 17 00:00:00 2001 From: zjli2013 Date: Thu, 9 Jul 2026 18:42:40 +0800 Subject: [PATCH] [AMDGPU] Add DPP row_shl fast path for shuffle_down reduction offsets amdgpu_shuffle_down_i32 previously always routed through the generic ds_bpermute + permlane64 cross-half helper (~50-60 cycles). For the power-of-two reduction-step offsets {1,2,4,8} the source lane stays within the same 16-lane DPP row, so a single v_mov_b32_dpp row_shl:N computes the shuffle directly (~4-12 cycles). Lanes whose source falls outside the row keep their own value (r is seeded with value); reduction trees treat those boundary lanes as don't-care once the tree has contracted below 16 active lanes, so reduce_{add,min,max}_tiled and block.reduce stay correct. offset 16 crosses a row and offset >= 32 crosses the wave half, so both fall through to the generic helper unchanged. shuffle_up keeps the helper: its Hillis-Steele scan consumers keep all lanes active and read across row boundaries at offset 1, which row_shr cannot express. Emitted as inline asm (not __builtin_amdgcn_update_dpp) because runtime.cpp is compiled to bitcode by the host clang and only retargeted to amdgcn at JIT time; the instruction text is carried verbatim into the IR, same technique as the existing +v fence in the cross-half helper. Measured on R9700 (gfx1201 / RDNA4), single-wave dependent-chain latency microbenchmark: offsets {1,2,4,8} ~42.7ns -> ~20ns (~2.1x); offset 16/32 unchanged (still helper). reduce_add_tiled(log2_size=6) == 2080 (sum 1..64), reduction consumers PASS. --- .../runtime/llvm/runtime_module/runtime.cpp | 20 +++++++++++++++++++ 1 file changed, 20 insertions(+) diff --git a/quadrants/runtime/llvm/runtime_module/runtime.cpp b/quadrants/runtime/llvm/runtime_module/runtime.cpp index 36e10d871e..146e94973e 100644 --- a/quadrants/runtime/llvm/runtime_module/runtime.cpp +++ b/quadrants/runtime/llvm/runtime_module/runtime.cpp @@ -1163,6 +1163,22 @@ f64 amdgpu_shuffle_f64(i32 index, f64 value) { // 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. i32 amdgpu_shuffle_down_i32(i32 offset, i32 value) { +#if defined(ARCH_amdgpu) && (defined(__x86_64__) || defined(__i386__) || defined(__amdgcn__)) + // DPP fast path for the reduction-step offsets {1,2,4,8}. ``row_shl:N`` shifts data left within each + // 16-lane row (lane ``i`` reads lane ``i+N``); lanes whose source falls outside the row keep their own + // value (``r`` is seeded with ``value``), which reduction trees treat as don't-care once the tree has + // contracted below 16 active lanes. Emitted as inline asm rather than ``__builtin_amdgcn_update_dpp`` + // because runtime.cpp is compiled to bitcode by the host clang (the module is retargeted to amdgcn only + // at JIT time), so the intrinsic isn't available at compile time; the instruction text is carried + // verbatim into the IR, same technique as the ``+v`` fence above. ``offset`` is resolved at compile + // time via constant propagation after inlining, so the dead branches are eliminated. offset 16 crosses + // a row and >=32 crosses the wave half, so they fall through to the generic helper. + 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); } @@ -1197,6 +1213,10 @@ f64 amdgpu_shuffle_down_f64(i32 offset, f64 value) { // Mirrors `amdgpu_shuffle_down`: the cross-half helper is a generic gather, so `shuffle_up` is just `shuffle_down` // with the source lane index decremented instead of incremented. The same DPP fast-path FIXME applies here too. i32 amdgpu_shuffle_up_i32(i32 offset, i32 value) { + // No DPP fast path here (unlike shuffle_down): shuffle_up feeds Hillis-Steele scans, where every lane + // stays active and reads across 16-lane row boundaries starting at offset 1 (lane 16 reads lane 15). + // ``row_shr`` would return the boundary lane's own value there instead of the cross-row neighbour, so + // the scan miscompiles. shuffle_up therefore always uses the generic cross-half helper. return amdgpu_cross_half_shuffle_i32(amdgpu_lane_id() - offset, value); }