fix: remove device printf from GDN kernel — corrupts output on Blackwell SM 12.0

The debug printf inside gated_delta_net_step (added in cf54a0b) causes
garbage output on Blackwell GPUs. Device printf on SM 12.0 appears to
corrupt the kernel's register state or introduce synchronization
artifacts, producing blank lines, repeated characters, or echoed
prompt templates instead of coherent text.

Also reverts the kh_mode parameter to the simpler division-only
formula (head_id / k_heads_per_v) since the MLX format is the only
one used on the CUDA backend.

Verified: 'The capital of France is' → 'Paris' on both RTX 4090
(SM 8.9) and RTX PRO 6000 Blackwell (SM 12.0).

Author: ssubbotin

cuda_infer/infer.cu

@@ -1940,14 +1940,13 @@ static void layer_forward(Model *model, int layer_idx, int pos, int K) {
         //   MLX (g_quant_format != 1): division — V heads 0..3 share K head 0
         //   GGUF (g_quant_format == 1): modulo  — V heads 0,16,32,48 share K head 0 (matches llama.cpp)
         uint32_t khpv = LINEAR_NUM_V_HEADS / LINEAR_NUM_K_HEADS;
-        uint32_t kh_mode = (g_quant_format == 1) ? 1u : 0u;
         gated_delta_net_step<<<LINEAR_NUM_V_HEADS, 128>>>(
             model->delta_state[layer_idx],
             model->buf_conv_output,                           // q [2048] (scaled for GGUF)
             model->buf_conv_output + LINEAR_TOTAL_KEY,        // k [2048]
             model->buf_conv_output + 2 * LINEAR_TOTAL_KEY,    // v [8192]
             model->buf_g_decay, model->buf_beta_gate,
-            model->buf_delta_output, khpv, kh_mode);
+            model->buf_delta_output, khpv);
 
         if (layer_idx == 0 && g_dump_layer0) {
             float d5[5];

cuda_infer/kernels.cuh

@@ -731,16 +731,14 @@ __global__ void gated_delta_net_step(
     const float* __restrict__ g_decay, // [64]
     const float* __restrict__ beta_gate, // [64]
     float* __restrict__ output,        // [8192]
-    uint32_t k_heads_per_v,            // = 4
-    uint32_t kh_mode                   // 0 = division (MLX), 1 = modulo (llama.cpp/GGUF)
+    uint32_t k_heads_per_v             // = 4
 ) {
     uint32_t head_id = blockIdx.x;
     uint32_t vi = threadIdx.x;
     // Key head mapping is format-dependent:
     //   MLX:        kh = head_id / k_heads_per_v  (chunked: V heads 0..3 share K head 0)
     //   llama.cpp:  kh = head_id % num_k_heads    (interleaved: V heads 0,16,32,48 share K head 0)
-    uint32_t n_kh = gridDim.x / k_heads_per_v;
-    uint32_t kh = (kh_mode == 0) ? (head_id / k_heads_per_v) : (head_id % n_kh);
+    uint32_t kh = head_id / k_heads_per_v;
     float g = g_decay[head_id];
     float beta = beta_gate[head_id];
 
@@ -767,15 +765,6 @@ __global__ void gated_delta_net_step(
         out_val += state[state_base + ki] * q[k_base + ki];
     output[v_base + vi] = out_val;
 
-    // Debug: print head 1, vi=0 intermediate values
-    if (head_id == 1 && vi == 0) {
-        printf("[gdn] h1v0: g=%.6f beta=%.6f kv_mem=%.8f v=%.8f delta=%.8f out=%.8f\n",
-               g, beta, kv_mem, v[v_base], delta, out_val);
-        printf("[gdn] h1v0: q[0:3]=%.6f %.6f %.6f  k[0:3]=%.6f %.6f %.6f\n",
-               q[k_base], q[k_base+1], q[k_base+2], k[k_base], k[k_base+1], k[k_base+2]);
-        printf("[gdn] h1v0: state[0:3]=%.8f %.8f %.8f\n",
-               state[state_base], state[state_base+1], state[state_base+2]);
-    }
 }
 
 // ============================================================================