=== 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: Gated DeltaNet Decode Step. SPEC: - Op: Single token SSM update. 1. Conv1d (kernel=4) over the fused QKV [10240] input using a ring buffer. 2. Split into Q[2048], K[2048], V[6144]. 3. Map 16 QK heads to 48 V heads using `repeat_interleave(3)`. 4. Compute gates: beta = sigmoid(ssm_beta_proj). alpha = compute_delta_net_alpha(ssm_a, ssm_alpha_proj). (Leave `compute_delta_net_alpha` as a mathematically sound stub if the exact softplus/sigmoid identity is unknown). 5. Update state S[48, 128, 128]: e = v - S^T k; S <- alpha*S + beta * outer(k, e). 6. Output o = S^T q. - Shapes: Exactly as defined in the dossier. HARD CONSTRAINTS: - State S must be stored in fp32 to prevent drift. - One CTA per V-head (48 CTAs). - Roofline Target: The state R+W is 6.3 MB. At 260 GB/s, the physical floor is ~24 µs. Target >= 70% of measured HBM BW for the state traffic. Do not claim < 20 µs. OUTPUT: Complete .cu. Include the exact algebraic shortcut to avoid a second full state pass for the output `o`.
