From d09a120ec30b50de8328f456b8d0c05ec7adf5ba Mon Sep 17 00:00:00 2001 From: Robert Esclapez Garcia Date: Thu, 30 Jul 2026 08:06:09 -0700 Subject: [PATCH] CUDA: dispatch the chunked gated_delta_net in CU mode when it fills the device The kernel is bound by issue pressure on the LDS pipe: 0.00% bank conflicts, MemUnitBusy 38.6%, and only 3.03 VALU instructions per LDS instruction against 7-9 for the token-by-token kernel. CU mode confines a block to one CU, which halves the LDS contention domain, and it costs no residency here - a WGP holds floor(128 KB / 57.25 KB) = 2 of these blocks in WGP mode and 2*floor(64 KB / 57.25 KB) = 2 in CU mode. It only pays once every CU has a block to run. A block occupies one CU in CU mode but spreads over the whole WGP in WGP mode, so below one block per WGP the second CU of each WGP sits idle. Measured at 512 tokens, cost per op against block count, with nsm = 20 WGPs on gfx1151: blocks 4 8 16 20 24 32 40 48 delta +40% +40% +33% +19% -17% -9% -9% -9% The launch therefore picks between two entry points on H*n_seqs > nsm. The work-group processor mode is a function attribute rather than something a template parameter can select, so the two entry points share one inlined body and differ only in that attribute. For head counts that take the CU path the op is 6.4-10.6% cheaper across every sequence length from 64 to 4096 tokens, worth +0.67% end to end on Qwen3.5-35B-A3B at pp4096 - what a tenth off an op that is ~7% of prefill predicts. Head counts below the threshold keep the WGP path: Qwen3.5-0.8B, whose 16 heads take it, measures +0.00% at pp128 over ten paired iterations. Assisted-by: Claude Opus 4.7 --- ggml/src/ggml-cuda/gated_delta_net_chunked.cu | 63 +++++++++++++------ 1 file changed, 45 insertions(+), 18 deletions(-) diff --git a/ggml/src/ggml-cuda/gated_delta_net_chunked.cu b/ggml/src/ggml-cuda/gated_delta_net_chunked.cu index 767d030c3880..6e5a254fd3f9 100644 --- a/ggml/src/ggml-cuda/gated_delta_net_chunked.cu +++ b/ggml/src/ggml-cuda/gated_delta_net_chunked.cu @@ -109,23 +109,34 @@ bool ggml_cuda_gdn_chunked_supported(bool kda, bool keep_rs, int64_t S_v, int64_ #define GDN_MIN_WAVES_PER_SIMD 1 #endif -__global__ void __launch_bounds__(GDN_BLOCK_SIZE, GDN_MIN_WAVES_PER_SIMD) -gdn_chunked_f32(const float * __restrict__ q, - const float * __restrict__ k, - const float * __restrict__ v, - const float * __restrict__ g, - const float * __restrict__ beta, - const float * __restrict__ state_in, - float * __restrict__ dst, - float * __restrict__ state_out, - const int64_t n_tokens, - const int64_t n_heads, - const int64_t sq1, const int64_t sq2, const int64_t sq3, - const int64_t sv1, const int64_t sv2, const int64_t sv3, - const int64_t sb1, const int64_t sb2, const int64_t sb3, - const uint3 neqk1_magic, - const uint3 rq3_magic, - const float scale) { +// CU mode confines a block to one CU, which halves the LDS contention domain. It costs no residency: +// a WGP holds floor(128 KB / 57.25 KB) = 2 of these blocks in WGP mode and 2*floor(64 KB / 57.25 KB) +// = 2 in CU mode. It only pays once every CU has a block to run, hence the launch-time choice below. +#if defined(GGML_USE_HIP) && defined(RDNA3_5) +#define GDN_CU_MODE __attribute__((target("cumode"))) +#else +#define GDN_CU_MODE +#endif + +// The work-group processor mode is a function attribute, so it cannot be a template parameter: the +// two entry points below share this body and differ only in that attribute. +#define GDN_CHUNKED_PARAMS \ + const float * __restrict__ q, const float * __restrict__ k, \ + const float * __restrict__ v, const float * __restrict__ g, \ + const float * __restrict__ beta, const float * __restrict__ state_in, \ + float * __restrict__ dst, float * __restrict__ state_out, \ + const int64_t n_tokens, const int64_t n_heads, \ + const int64_t sq1, const int64_t sq2, const int64_t sq3, \ + const int64_t sv1, const int64_t sv2, const int64_t sv3, \ + const int64_t sb1, const int64_t sb2, const int64_t sb3, \ + const uint3 neqk1_magic, const uint3 rq3_magic, \ + const float scale + +#define GDN_CHUNKED_ARGS \ + q, k, v, g, beta, state_in, dst, state_out, n_tokens, n_heads, \ + sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale + +static __device__ __forceinline__ void gdn_chunked_body(GDN_CHUNKED_PARAMS) { constexpr int CHUNK = GDN_CHUNK; constexpr int HEAD_DIM = GDN_HEAD_DIM; constexpr int COL_LANES = GDN_COL_LANES; @@ -399,6 +410,16 @@ gdn_chunked_f32(const float * __restrict__ q, } } +__global__ void __launch_bounds__(GDN_BLOCK_SIZE, GDN_MIN_WAVES_PER_SIMD) +gdn_chunked_f32_wgp(GDN_CHUNKED_PARAMS) { + gdn_chunked_body(GDN_CHUNKED_ARGS); +} + +__global__ void GDN_CU_MODE __launch_bounds__(GDN_BLOCK_SIZE, GDN_MIN_WAVES_PER_SIMD) +gdn_chunked_f32_cu(GDN_CHUNKED_PARAMS) { + gdn_chunked_body(GDN_CHUNKED_ARGS); +} + void ggml_cuda_gdn_chunked(ggml_backend_cuda_context & ctx, const ggml_cuda_gdn_chunked_args & args) { GGML_ASSERT(args.S_v == GDN_HEAD_DIM); @@ -410,7 +431,13 @@ void ggml_cuda_gdn_chunked(ggml_backend_cuda_context & ctx, const ggml_cuda_gdn_ const ggml_cuda_kernel_launch_params launch_params(block_nums, block_dims, 0, ctx.stream()); - ggml_cuda_kernel_launch(gdn_chunked_f32, launch_params, + // nsm counts WGPs on RDNA, and a block occupies one CU in CU mode but spreads over the whole WGP + // in WGP mode. Below one block per WGP, CU mode leaves the second CU of each WGP idle and costs + // about a quarter of the op; above it every CU is busy either way and CU mode is worth ~10%. + const bool cu_mode = args.H * args.n_seqs > ggml_cuda_info().devices[ggml_cuda_get_device()].nsm; + const auto kernel = cu_mode ? gdn_chunked_f32_cu : gdn_chunked_f32_wgp; + + ggml_cuda_kernel_launch(kernel, launch_params, args.q, args.k, args.v, args.g, args.beta, args.state_in, args.dst, args.state_out, args.n_tokens, args.H, args.sq1, args.sq2, args.sq3, args.sv1, args.sv2, args.sv3, args.sb1, args.sb2, args.sb3,