=== 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 GB/s measured sustained HBM BW. T4 HAS: __dp4a (INT8 dot), __hfma2, FP16 tensor cores (WMMA m16n16k16), CUDA Graphs. T4 LACKS: cp.async, BF16, FP8, TMA, Ampere mma.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 | vocab_size: 248320 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 (GQA + Gated) --- blk.X.attn_q.weight: [5120, 12288] (24 Q heads @ 256 + 24 Gate heads @ 256 FUSED) blk.X.attn_k.weight: [5120, 1024] (4 KV heads @ 256) blk.X.attn_v.weight: [5120, 1024] (4 KV heads @ 256) blk.X.attn_output.weight: [6144, 5120] (24 heads * 256 dim = 6144) RoPE: partial, first 64 dims only. freq_base=10000000. --- SSM / GATED DELTANET LAYERS --- 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] (Base decay scalar per V-head) blk.X.ssm_alpha.weight: [5120, 48] (Input-dependent alpha gate projection) blk.X.ssm_beta.weight: [5120, 48] (Beta gate projection) blk.X.ssm_out.weight: [6144, 5120] (Output projection) Head Mapping: 16 QK heads (dim 128) map to 48 V heads (dim 128) via repeat_interleave(3). Recurrence Math: S <- alpha*S + beta * outer(k, v - S^T k). WARNING: alpha is a combination of ssm_a and ssm_alpha(x). Do not invent the formula; use a stub `compute_delta_net_alpha(ssm_a, alpha_proj)` if the exact identity is unknown. --- MTP / NEXTN HEAD (Layer 64) --- blk.64.nextn.eh_proj.weight: [10240, 5120] (GGUF shape: in=10240, out=5120) blk.64.nextn.enorm.weight: [5120] | blk.64.nextn.hnorm.weight: [5120] Math: out = proj @ eh_proj (No transpose! GGUF is [in, out]). Then RMSNorm -> shared LM Head. --- 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. WARNING: IQ4_XS uses a nonlinear 16-entry codebook, NOT the Q4_K superblock formula. Do not apply Q4_K math to IQ4_XS. === PHYSICS & CONTRACT === - Decode reads ~15.3-17.6 GB/token. Layer-split ceiling ~14-18 tok/s base. - Each WEIGHT byte loaded from DRAM at most once per token. - GEMV Accumulation: Dequant to fp16/fp32 in registers, accumulate in fp32, store fp16. Do NOT accumulate in fp16. - DeltaNet State Roofline: S is 3.15 MB fp32 per layer. R+W is 6.3 MB. At 260 GB/s, the absolute physical floor is ~24 Β΅s/layer. Target >= 70% of HBM BW for state traffic. - MTP Verify: SSM layers CANNOT batch verify. They require a sequential scan of length gamma+1, or a specialized chunk-verify. GQA layers can batch verify. Using the dossier, write a PRODUCTION CUDA kernel for SM 7.5: MTP NextN Head. SPEC: - Op: Forward through `blk.64.nextn` for the accepted prefix hidden state. - Inputs: Hidden states H[1, 5120], Embeddings E[1, 5120]. - Math: 1. h_norm = RMSNorm(H) 2. e_norm = RMSNorm(E) 3. proj = concat(h_norm, e_norm) -> shape [1, 10240] 4. out = proj @ eh_proj (Note: GGUF shape is [10240, 5120], so NO transpose). 5. Final RMSNorm -> multiply by shared LM Head (vocab 248320) to get logits. - Shapes: `eh_proj` is [10240, 5120] (Q6_K quantized). HARD CONSTRAINTS: - This kernel ONLY computes the NextN logits. It does NOT verify the draft. - Use shared memory to hold the 10240-dim intermediate projection. OUTPUT: Complete .cu with self-test.
