Fix all warnings + stack overflow + TQ_STATIC_ASSERT (closes #11)

- TQ_STATIC_ASSERT: C11 _Static_assert (was recursive no-op in C mode)
- Stack overflow: recon[256] → recon[512] for Gemma 4 head_dim=512
- Remove unused Metal shader constants (TG_SIZE, MAX_SHARED_DIM, TQ_INDICES_SIZE)
- Remove unused g_repack_count, n_tiles variables
- Replace MIN() macro with ternary to avoid GNU extension warning

Zero warnings, 34/34 tests, score 99.2%.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Author: unamedkr

include/turboquant/tq_types.h

@@ -8,7 +8,7 @@
 #ifdef __cplusplus
 #define TQ_STATIC_ASSERT(cond, msg) static_assert(cond, msg)
 #else
-#define TQ_STATIC_ASSERT(cond, msg) TQ_STATIC_ASSERT(cond, msg)
+#define TQ_STATIC_ASSERT(cond, msg) _Static_assert(cond, msg)
 #endif
 
 /* Cross-platform math constants (some platforms lack M_PI) */

src/backend/metal/tq_matmul.metal

@@ -118,11 +118,6 @@ constant uchar ksigns_iq2xs[128] = {
  * in threadgroup shared memory to avoid redundant global reads.
  * ============================================================ */
 
-/* Threadgroup size for matmul kernels */
-constant uint TG_SIZE = 256;
-
-/* Maximum input dimension cacheable in shared memory (32KB / 4 = 8192 floats) */
-constant uint MAX_SHARED_DIM = 8192;
 
 kernel void matmul_iq2_xxs(
     device const uchar*  weight    [[buffer(0)]],  /* [out_dim * row_bytes] */

src/backend/metal/tq_metal_dispatch.m

@@ -1725,7 +1725,7 @@ void tq_metal_repack_q4(const uint8_t* src_qs, const float* src_scales,
 #define TQ_REPACK_CACHE_SIZE 128
 static struct { const void* key; id<MTLBuffer> qs; id<MTLBuffer> sc; int out_dim; int in_dim; }
     g_repack_cache[TQ_REPACK_CACHE_SIZE];
-static int g_repack_count = 0;
+static int g_repack_count __attribute__((unused)) = 0;
 
 static void encode_q4_matmul(id<MTLComputeCommandEncoder> enc,
                               id<MTLBuffer> input_buf,
@@ -1736,8 +1736,6 @@ static void encode_q4_matmul(id<MTLComputeCommandEncoder> enc,
     if (!tq_pipe_matmul_tq_q4) return;
 
     int n_blocks = in_dim / 32;
-    const int TILE = 32;
-    int n_tiles = (out_dim + TILE - 1) / TILE;
 
     /* Fast Q4 kernel: llama.cpp-inspired uint16 mask trick + SIMD-group.
      * No repacking needed — reads original row-major Q4 layout.
@@ -2173,15 +2171,21 @@ int tq_metal_forward_layer(
             [enc setBuffer:g_gpu_k  offset:0 atIndex:1];
             [enc setBuffer:pos_buf  offset:0 atIndex:2];
             [enc setBuffer:kvd_buf  offset:0 atIndex:3];
-            [enc dispatchThreads:MTLSizeMake(kv_dim, 1, 1)
-               threadsPerThreadgroup:MTLSizeMake(MIN(kv_dim, 256), 1, 1)];
+            {
+                NSUInteger tg_w = (NSUInteger)(kv_dim < 256 ? kv_dim : 256);
+                [enc dispatchThreads:MTLSizeMake(kv_dim, 1, 1)
+                   threadsPerThreadgroup:MTLSizeMake(tg_w, 1, 1)];
+            }
             [enc memoryBarrierWithScope:MTLBarrierScopeBuffers];
 
             /* Write V to cache */
             [enc setBuffer:vc_buf   offset:0 atIndex:0];
             [enc setBuffer:g_gpu_v  offset:0 atIndex:1];
-            [enc dispatchThreads:MTLSizeMake(kv_dim, 1, 1)
-               threadsPerThreadgroup:MTLSizeMake(MIN(kv_dim, 256), 1, 1)];
+            {
+                NSUInteger tg_w = (NSUInteger)(kv_dim < 256 ? kv_dim : 256);
+                [enc dispatchThreads:MTLSizeMake(kv_dim, 1, 1)
+                   threadsPerThreadgroup:MTLSizeMake(tg_w, 1, 1)];
+            }
             [enc memoryBarrierWithScope:MTLBarrierScopeBuffers];
         }
 

src/backend/metal/tq_moe_kernel.metal

@@ -146,11 +146,6 @@ constant uchar ksigns_iq2xs[128] = {
  * Each thread processes a subset of blocks; caller must reduce.
  * ============================================================ */
 
-/* Threadgroup size */
-constant uint TG_SIZE = 256;
-
-/* Maximum input cacheable in shared memory (32KB / 4 = 8192 floats) */
-constant uint MAX_SHARED_DIM = 8192;
 
 /**
  * Partial IQ2_XXS dot for blocks assigned to this thread.

src/backend/metal/tq_polar.metal

@@ -15,7 +15,7 @@ using namespace metal;
 
 constant int TQ_BK           = 128;
 constant int TQ_PAIRS        = 64;   /* TQ_BK / 2 */
-constant int TQ_INDICES_SIZE = 32;   /* TQ_BK / 4  (pairs/2 bytes) */
+/* TQ_INDICES_SIZE = TQ_BK / 4 = 32 (pairs/2 bytes), used only in C host code */
 
 /* ============================================================
  * Block structures (matching C layout)

src/engine/tq_transformer.c

@@ -1313,7 +1313,7 @@ static void self_attn_forward(tq_model_t* model, tq_state_t* s, int l, int pos)
         const tq_type_traits_t* dbg_traits = &TQ_TRAITS[s->kv_quant_type];
         float mse = 0, cos_num = 0, cos_d1 = 0, cos_d2 = 0;
         uint8_t tmp_buf[1024];
-        float recon[256];
+        float recon[512]; /* max head_dim is 512 (Gemma 4 full layers) */
         for (int kh = 0; kh < 1; kh++) { /* first head only */
             const float* key_src = s->k + kh * head_dim;
             dbg_traits->quantize(key_src, tmp_buf, head_dim);
@@ -1890,7 +1890,7 @@ static void self_attn_forward(tq_model_t* model, tq_state_t* s, int l, int pos)
         } else if (s->value_quant_bits == 2) {
             /* Q2 value path: dequantize and accumulate.
              * Q2 has a more complex codebook, so we keep the dequant path. */
-            float v_tmp[512]; /* max head_dim is 256, safe with margin */
+            float v_tmp[512]; /* max head_dim is 512 (Gemma 4 full layers) */
             size_t layer_off_qs = (size_t)l * max_seq * s->value_stride_qs;
             size_t layer_off_sc = (size_t)l * max_seq * s->value_stride_scales;
             int n_blocks_per_head = (head_dim + 31) / 32;