=== HARDWARE GROUND TRUTH: 2x NVIDIA Tesla T4 (Kaggle) === A...

Prompt

=== HARDWARE GROUND TRUTH: 2x NVIDIA Tesla T4 (Kaggle) === Architecture: SM 7.5 (Turing). 40 SMs/GPU. 16 GB GDDR6, ~260-280 GB/s measured sustained bandwidth. T4 HAS: __dp4a (INT8 dot), __hfma2, FP16 tensor cores (mma.sync m8n8k4), CUDA Graphs. T4 LACKS: cp.async, BF16, FP8, TMA, mma.sync m16n8k16, >64KB dynamic smem/block. Interconnect: PCIe 3.0 only. No NVLink. P2P usually disabled. Pipeline split is mandatory. === MODEL GROUND TRUTH: Qwen3.8-27B (GGUF Arch: `qwen35`) === File: unsloth/Qwen3.8-27B-GGUF (UD-Q4_K_XL / UD-Q4_K_M mix) general.architecture: qwen35 qwen35.block_count: 65 (Layers 0-63 are transformer/SSM, Layer 64 is MTP/NextN) qwen35.embedding_length: 5120 qwen35.feed_forward_length: 17408 qwen35.full_attention_interval: 4 (Layers 3,7,11...63 are Full GQA. Others are SSM/DeltaNet) --- FULL ATTENTION LAYERS (e.g., blk.3, blk.7) --- blk.X.attn_q.weight: [5120, 12288] (Note: 12288 implies 48 Q heads @ dim 256) blk.X.attn_k.weight: [5120, 1024] (4 KV heads @ dim 256) blk.X.attn_v.weight: [5120, 1024] (4 KV heads @ dim 256) blk.X.attn_output.weight: [6144, 5120] RoPE: partial, first 64 dims only. freq_base=10000000. --- SSM / GATED DELTANET LAYERS (e.g., blk.0, blk.1, blk.2) --- blk.X.attn_qkv.weight: [5120, 10240] (Fused Q=2048, K=2048, V=6144) blk.X.ssm_conv1d.weight: [4, 10240] (Conv1d kernel=4 over QKV) blk.X.ssm_a: [48] (Decay factor per V-head) blk.X.ssm_alpha.weight: [5120, 48] (Alpha gate projection) blk.X.ssm_beta.weight: [5120, 48] (Beta gate projection) blk.X.ssm_out.weight: [6144, 5120] (Output projection) Recurrence: S <- alpha*S + beta * k (v - S^T k)^T. State is [48, 128, 128] per layer. --- MTP / NEXTN HEAD (Layer 64) --- blk.64.nextn.eh_proj.weight: [10240, 5120] (Embedding+Hidden projection for draft token) blk.64.nextn.enorm.weight: [5120] blk.64.nextn.hnorm.weight: [5120] blk.64.nextn.shared_head_norm.weight: [5120] --- QUANTIZATION MIX (Unsloth Dynamic) --- Tensors use mixed quants: IQ3_S, IQ4_XS, IQ4_NL, Q3_K, Q4_K, Q5_K, Q6_K, Q8_0, F32. Q4_K superblock (256 weights = 144 bytes): { half d; half dmin; uint8 scales[12]; uint8 qs[128]; } Dequant: w = d * (sc * q4) - dmin * m. IQ4_XS uses 4-bit with specific superblock packing. === PHYSICS & CONTRACT === - Decode reads ~15.3-17.6 GB/token. Layer-split ceiling ~14-18 tok/s base. - MTP verify (Layer 64) amortizes weight reads. Target 30 tok/s requires MTP acceptance alpha >= 2.0. - Each WEIGHT byte loaded from DRAM at most once per token. - fp32 accumulate for all reductions and SSM state. fp16x2 ok for GEMV. Using the dossier, write a PRODUCTION CUDA kernel for SM 7.5: Mixed-Quant GEMV. SPEC: - Op: y[N] = dequant(W[N,K]) @ x[K], fp16 x/y. - Shapes: Must handle [5120, 10240] (SSM QKV), [5120, 12288] (Attn Q), [5120, 17408] (FFN Gate/Up). - Quantization: Must dynamically dequantize Q4_K, Q5_K, and IQ4_XS based on a tensor-type flag passed to the kernel. Use `__dp4a` for INT8/INT4 dot products where applicable, and `__hfma2` for fp16 accumulation. - Reference: Standard GGML dequantization logic for Q4_K and IQ4_XS. HARD CONSTRAINTS: - SM 7.5 only. No cp.async, no bf16. - One warp per 256-weight superblock. x broadcast via shared memory. - Fully coalesced 128-bit loads. OUTPUT: Complete .cu with kernel, host launcher, and a self-test that generates random Q4_K and IQ4_XS weights, runs the kernel, and diffs against a PyTorch fp32 reference.