From 554199a656ac8d3ab2e18436f3ba8b06e3ea9fa0 Mon Sep 17 00:00:00 2001 From: Piotr Wilkin Date: Thu, 24 Sep 2026 16:33:32 +0200 Subject: [PATCH 1/2] perf(qwen-flash): fence each K block in the W8A8 WMMA GEMM The scheduler hoisted the next K block's LDS operand reads above the current block's WMMAs, which pushed every W8A8BlockedWmmaGEMMKernel instantiation to 256 VGPRs plus 128-140 bytes of scratch on gfx1151. One sched_barrier per K block removes the spills (176-184 VGPRs, no scratch) and roughly triples occupancy. Strix Halo, Qwen3.8-Flash-Next UD-Q4_K_XL, pp2048: - W8A8 GEMM: 45-51% less time per call (HC-down 1073 -> 525 us) - prefill: 1459.99 -> 1535.02 t/s (+5.1%) A captured replay of the HC-down dispatch is bitwise identical, and qwen38_flash_next_projection_ops_test passes. Co-Authored-By: Claude Opus 5.5 --- src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp b/src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp index bb3640cc67..934c61269e 100644 --- a/src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp +++ b/src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp @@ -3530,6 +3530,9 @@ __launch_bounds__(256) __global__ void W8A8BlockedWmmaGEMMKernel( } } } + // Hoisting the next K block's operands above this one's WMMAs spills + // registers on gfx1151 and halves occupancy. + __builtin_amdgcn_sched_barrier(0); } __syncthreads(); } From 893fb98117487018c33fdc91682f3c8d44318294 Mon Sep 17 00:00:00 2001 From: Piotr Wilkin Date: Thu, 24 Sep 2026 19:52:11 +0200 Subject: [PATCH 2/2] perf(qwen-flash): fence W8A8 K blocks only on LLVM 23+ ROCm 7.2.3 (LLVM 22) schedules W8A8BlockedWmmaGEMMKernel without spills (157-180 VGPRs), and there the fence costs about 3% per call. The spills come from LLVM 23 (TheRock ROCm 7.14 and 10.0 nightlies), so keep the fence only for that compiler. Checked with the ROCm 10.0 compiler: the real __clang_major__ (23) gives 159/159/178/173 VGPRs and no scratch; forcing it to 22 brings back 256 VGPRs and 128-264 bytes of scratch. qwen38_flash_next_projection_ops_test passes. Co-Authored-By: Claude Opus 5.5 --- src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp b/src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp index 934c61269e..4c94a7cd05 100644 --- a/src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp +++ b/src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp @@ -3530,9 +3530,12 @@ __launch_bounds__(256) __global__ void W8A8BlockedWmmaGEMMKernel( } } } - // Hoisting the next K block's operands above this one's WMMAs spills - // registers on gfx1151 and halves occupancy. +#if __clang_major__ >= 23 + // LLVM 23 hoists the next K block's operands above this one's WMMAs, + // which spills registers on gfx1151. ROCm 7.2.3 (LLVM 22) does not, and + // the fence costs it about 3%. __builtin_amdgcn_sched_barrier(0); +#endif } __syncthreads(); }