=== 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 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 handle the split between SSM layers (fused QKV, Conv1D ring buffer state, recurrent S) and Full Attention layers (separate Q/K/V, Gated Q, RoPE). 2. Byte-balanced pipeline split: Do NOT calculate bytes from shapes. Instruct the engine to use the `gguf` python package to read `tensor.nbytes` and `tensor.type` to build an exact byte-weight table. Assign layers to GPU0 and GPU1 to balance the measured bytes, accounting for GPU0 holding the embeddings and GPU1 holding the LM Head/NextN. 3. MTP Verify Plan: Detail how `blk.64.nextn` is used. Explain the verify loop: - GQA layers: Batched prefill of length gamma+1. - SSM layers: Sequential scan of length gamma+1 through the recurrence. - State Management: How to checkpoint the 48 SSM states (3.15 MB each) and Conv1D ring buffers before verify, and rewind on rejection. 4. State Management: Define the memory layout for the 48 SSM recurrent states (fp32), the Conv1D ring buffers (3 x 10240 fp16), and the 16 Full Attention KV caches (q8_0). 5. Quantization Strategy: Explain the host-side template dispatch for the mixed UD quants. Detail how Q4_K, Q5_K, and IQ4_XS require entirely separate dequantization logic (especially the IQ4_XS nonlinear codebook) and must be dispatched at the host level to avoid warp divergence. Cite dossier tensor names and GGUF shape conventions for every claim.
