=== 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 the ENGINE DESIGN DOC for a from-scratch dual-T4 inference engine for the `qwen35` architecture. Do NOT write kernels yet. Deliver: 1. Kernel dependency graph for one decode token across all 65 layers, explicitly handling the split between SSM layers (fused QKV) and Full Attention layers (separate Q/K/V). 2. Byte-balanced pipeline split: Calculate the exact byte weight of SSM vs Full Attention layers using the tensor shapes in the dossier. Assign layers to GPU0 and GPU1 to minimize pipeline bubble. 3. MTP Verify Plan: Detail how `blk.64.nextn` is used. Explain how to batch the verification of gamma drafted tokens through the main 64 layers, and then pass them through the `eh_proj` [10240, 5120] and NextN head to compute acceptance logits in a single batched forward pass. 4. State Management: Define the memory layout for the 48 SSM recurrent states (fp32) vs the 16 Full Attention KV caches (q8_0). 5. Quantization Strategy: How to handle the mixed UD quants (IQ4_XS, Q4_K, Q5_K, Q6_K) in a single GEMV kernel without branching overhead. Cite dossier tensor names for every claim.
