From a02de7360f62f4e36d06330fb0bcce929bf06b28 Mon Sep 17 00:00:00 2001 From: Cary Palmer <24235924+professorpalmer@users.noreply.github.com> Date: Sat, 19 Sep 2026 09:20:25 -0500 Subject: [PATCH 1/2] cuda: use the 4-column GDN warp layout on all Ampere+ NVIDIA, not only GB10 (+6% prefill on Ada) The S_v == 128 GATED_DELTA_NET kernel has a 4-columns-per-warp variant that was gated to GGML_CUDA_CC_DGX_SPARK. The recurrence is serial over tokens, so prefill on a hybrid model lives or dies on this kernel's per-step efficiency, and the wider layout is a straight win on consumer Ampere/Ada as well. Host launch geometry and compiled kernel now derive the column count from one constexpr predicate, gdn_cols_per_warp(arch, S_v, KDA): the kernel evaluates it on __CUDA_ARCH__, the host on ggml_cuda_highest_compiled_arch(cc) (the arch the device will actually run, not the raw runtime capability), and only for GGML_CUDA_CC_IS_NVIDIA(cc). Under GGML_USE_HIP / GGML_USE_MUSA the predicate is 1 on both sides, so AMD capability values (which carry an offset above GGML_CUDA_CC_AMPERE) can no longer make the host launch 8 z-blocks of 4 columns against a 1-column kernel. Fixes the geometry mismatch raised in review. RTX 4070 12 GB, Bonsai 2 27B PTQ1_0: pp2048 1220 -> 1297 t/s, decode unchanged. Maintainer paired runs at the previous revision: pp512 +3.7% RTX 3090, +1.3% RTX 4090, +0.6% H100, +0.3% RTX 5090; tg128 within noise everywhere. test-backend-ops CUDA0 vs CPU: GATED_DELTA_NET 39/39 supported cases. --- ggml/src/ggml-cuda/gated_delta_net.cu | 26 ++++++++++++++++++++++---- 1 file changed, 22 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/gated_delta_net.cu b/ggml/src/ggml-cuda/gated_delta_net.cu index 5cf6968a6e1b..d089c36db873 100644 --- a/ggml/src/ggml-cuda/gated_delta_net.cu +++ b/ggml/src/ggml-cuda/gated_delta_net.cu @@ -1,6 +1,21 @@ #include "gated_delta_net.cuh" #include "ggml-cuda/common.cuh" +// Columns per warp for the S_v == 128, non-KDA kernel. The host launch geometry and the compiled +// kernel must agree, so both derive it from this one predicate: the kernel passes __CUDA_ARCH__, +// the host passes the highest compiled arch the device will actually run +// (ggml_cuda_highest_compiled_arch(cc)), not the raw runtime capability. NVIDIA Ampere and newer +// take 4 columns (tuned on GB10, a straight win on consumer Ampere/Ada too); HIP and MUSA keep 1, +// whatever their capability values (which carry vendor offsets above GGML_CUDA_CC_AMPERE) say. +static constexpr __host__ __device__ int gdn_cols_per_warp(int arch, int S_v, bool KDA) { +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + (void) arch; (void) S_v; (void) KDA; + return 1; +#else + return GGML_CUDA_CC_IS_NVIDIA(arch) && arch >= GGML_CUDA_CC_AMPERE && S_v == 128 && !KDA ? 4 : 1; +#endif +} + static __global__ void gdn_precompute_exp(const float * g, float * g_exp, int64_t n) { for (int64_t i = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; i < n; i += (int64_t) blockDim.x*gridDim.x) { @@ -44,10 +59,10 @@ gated_delta_net_cuda(const float * q, const uint32_t sequence = blockIdx.y; // Each warp owns one or more columns, using warp-level primitives to reduce across rows. const int lane = threadIdx.x; -#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK - constexpr int cols_per_warp = S_v == 128 && !KDA ? 4 : 1; +#if defined(__CUDA_ARCH__) + constexpr int cols_per_warp = gdn_cols_per_warp(__CUDA_ARCH__, S_v, KDA); #else - constexpr int cols_per_warp = 1; + constexpr int cols_per_warp = 1; // host pass only; never executed #endif const int col = (blockIdx.z * blockDim.y + threadIdx.y) * cols_per_warp; @@ -217,7 +232,10 @@ static void launch_gated_delta_net( const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size; const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; const int num_warps = 4; - const int cols_per_warp = cc == GGML_CUDA_CC_DGX_SPARK && S_v == 128 && !KDA ? 4 : 1; + // The recurrence is serial over tokens, so prefill lives or dies on this kernel's per-step + // efficiency (RTX 4070, Bonsai 2 27B: pp2048 1220 -> 1297 t/s, decode unchanged). Must match + // the kernel's compile-time choice: same predicate, fed the arch that was compiled for this cc. + const int cols_per_warp = GGML_CUDA_CC_IS_NVIDIA(cc) ? gdn_cols_per_warp(ggml_cuda_highest_compiled_arch(cc), S_v, KDA) : 1; dim3 grid_dims(H, n_seqs, (S_v + num_warps * cols_per_warp - 1) / (num_warps * cols_per_warp)); dim3 block_dims(warp_size <= S_v ? warp_size : S_v, num_warps, 1); From 937c4d857282f9c8fb91682443309559856d9b76 Mon Sep 17 00:00:00 2001 From: Cary Palmer <24235924+professorpalmer@users.noreply.github.com> Date: Sun, 20 Sep 2026 20:08:25 -0500 Subject: [PATCH 2/2] cuda: keep only the GDN cols-per-warp launch invariant in comments. --- ggml/src/ggml-cuda/gated_delta_net.cu | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml-cuda/gated_delta_net.cu b/ggml/src/ggml-cuda/gated_delta_net.cu index d089c36db873..f0d08152acb0 100644 --- a/ggml/src/ggml-cuda/gated_delta_net.cu +++ b/ggml/src/ggml-cuda/gated_delta_net.cu @@ -1,12 +1,7 @@ #include "gated_delta_net.cuh" #include "ggml-cuda/common.cuh" -// Columns per warp for the S_v == 128, non-KDA kernel. The host launch geometry and the compiled -// kernel must agree, so both derive it from this one predicate: the kernel passes __CUDA_ARCH__, -// the host passes the highest compiled arch the device will actually run -// (ggml_cuda_highest_compiled_arch(cc)), not the raw runtime capability. NVIDIA Ampere and newer -// take 4 columns (tuned on GB10, a straight win on consumer Ampere/Ada too); HIP and MUSA keep 1, -// whatever their capability values (which carry vendor offsets above GGML_CUDA_CC_AMPERE) say. +// Host and device must agree on columns/warp. Kernel uses __CUDA_ARCH__; host uses ggml_cuda_highest_compiled_arch(cc). NVIDIA Ampere+ and S_v==128 and !KDA -> 4; HIP/MUSA stay 1. static constexpr __host__ __device__ int gdn_cols_per_warp(int arch, int S_v, bool KDA) { #if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) (void) arch; (void) S_v; (void) KDA; @@ -232,9 +227,7 @@ static void launch_gated_delta_net( const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size; const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; const int num_warps = 4; - // The recurrence is serial over tokens, so prefill lives or dies on this kernel's per-step - // efficiency (RTX 4070, Bonsai 2 27B: pp2048 1220 -> 1297 t/s, decode unchanged). Must match - // the kernel's compile-time choice: same predicate, fed the arch that was compiled for this cc. + // Same predicate as the kernel, fed the arch compiled for this cc. const int cols_per_warp = GGML_CUDA_CC_IS_NVIDIA(cc) ? gdn_cols_per_warp(ggml_cuda_highest_compiled_arch(cc), S_v, KDA) : 1; dim3 grid_dims(H, n_seqs, (S_v + num_warps * cols_per_warp - 1) / (num_warps * cols_per_warp)); dim3 block_dims(warp_size <= S_v ? warp_size : S_v, num_warps, 1);