From 2795d2b7f275a19e96f3742efa83bfae26f681b7 Mon Sep 17 00:00:00 2001 From: muh-bot Date: Wed, 5 Aug 2026 03:20:38 +0000 Subject: [PATCH] =?UTF-8?q?[muh]=20tuning=5Freduce:=20=E4=BF=AE=E6=AD=A3?= =?UTF-8?q?=20SMEM=20=E6=A8=A1=E5=9E=8B=E9=94=99=E8=AF=AF=20(=E5=9F=BA?= =?UTF-8?q?=E4=BA=8E=20agent=5Freduce.cuh=20=E6=BA=90=E7=A0=81=E5=88=86?= =?UTF-8?q?=E6=9E=90)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 关键修正: 旧注释: 'tile = tpb * ipt * accum_size ≤ 48KB SMEM' 实际: reduce 不使用 BlockLoad 的 SMEM staging buffer! 数据直接从 global memory 加载到寄存器 (striped/vectorized) SMEM 仅用于 BlockReduce (warp shuffle scratch, << 1KB) 真正的约束是寄存器压力: items[ITEMS_PER_THREAD] 全在寄存器里 items=24 float32 → ~24 regs → 可以接受 (64K regs/SM) from agent_reduce.cuh: - ConsumeFullTile: 直接 striped load 或 VectorT 加载到 items[] 数组 - 没有 BlockLoad::TempStorage (不像 scan 有 union {load, store, scan}) - ATTEMPT_VECTORIZATION 条件: vec_size>1, items%vec==0, sizeof≤8 - 数值不变 (items=24 对 float32 是合理的, 只是理由从错误的 SMEM 改为正确的寄存器压力) from kernel_reduce.cuh: - atomic path (not_guaranteed): atomicAdd 聚合, 16 SMs 上 contention 极低 - LOAD_LDG 保留 (等 bench 数据, topk 显示 LOAD_DEFAULT 可能更优) --- muh/include/muh/tuning/tuning_reduce.cuh | 142 ++++++++++------------- 1 file changed, 64 insertions(+), 78 deletions(-) diff --git a/muh/include/muh/tuning/tuning_reduce.cuh b/muh/include/muh/tuning/tuning_reduce.cuh index 747ce03d..32196c19 100644 --- a/muh/include/muh/tuning/tuning_reduce.cuh +++ b/muh/include/muh/tuning/tuning_reduce.cuh @@ -7,25 +7,41 @@ // HARDWARE PROFILE (confirmed via ixsmi on Phanthy Cloud): // SM count: 16 (NOT 50 from spec sheet) // SMEM: 48KB (49152 bytes) per block -// L2 cache: 6MB (vs SM100's 50MB — 8.3× smaller) +// L2 cache: 6MB // HBM BW: 900 GB/s -// BW/SM: 900/16 = 56 GB/s (≈ B200 level, NOT A100's 18 GB/s) +// BW/SM: 900/16 = 56 GB/s // Warp size: 32 // -// SM=16 IMPACT ON TUNING: -// With only 16 SMs and max 2 CTAs/SM occupancy, there are at most 32 concurrent -// CTAs. Each CTA must process MORE data per tile to compensate for fewer CTAs. -// This means tiles should be LARGER than SM100's defaults (which assume 148 SMs). -// Target: fill SMEM to ≥ 70% where possible (current det paths use only 23%). +// CRITICAL ARCHITECTURE INSIGHT (from agent_reduce.cuh source): +// Reduce does NOT use BlockLoad — data is loaded directly to registers +// via striped access (ConsumeFullTile) or vectorized loads (VectorT). +// The SMEM is ONLY used by BlockReduce (warp shuffles + small scratch). +// BlockReduce SMEM ≈ threads * sizeof(AccumT) / 32 (one value per warp). // -// CCCL upstream structure (for reference): -// compute_capability >= 10.0 → sm100_tuning specializations (type-dispatch) -// compute_capability >= 6.0 → Policy600 {threads=256, items=16, vec=4} -// compute_capability >= 5.0 → Policy500 {threads=256, items=20, vec=4} -// Three determinism modes: run_to_run, gpu_to_gpu, not_guaranteed +// Previous muh tuning assumed "tile = tpb * ipt * accum_size ≤ 48KB SMEM" +// — THIS WAS WRONG. That formula applies to scan (which uses BlockLoad staging). +// For reduce, the real constraint is REGISTER PRESSURE: +// Each thread holds AccumT items[ITEMS_PER_THREAD] in registers. +// items=24 for float32 → 24 registers → acceptable (BI-V100 has 64K registers/SM) +// items=24 also means fewer CTAs needed → good for 16 SMs // -// NOTE: scale_mem_bound returns {items, threads} (items-first), matching -// CCCL's scaling_result struct. Destructure as auto [i, t] = ...; +// Vectorized load condition (agent_reduce.cuh line ~180): +// ATTEMPT_VECTORIZATION requires: vec_size > 1, items % vec == 0, +// is_pointer, is_trivially_relocatable, sizeof(InputT) <= 8 +// +// CCCL kernel_reduce.cuh two execution paths: +// 1. multi-tile (DeviceReduceKernel): each block reduces a tile partition +// - StableReductionOrder=false: atomicAdd aggregation (lowest latency) +// - StableReductionOrder=true: write to d_out[blockIdx.x], then single-tile pass +// 2. single-tile (DeviceReduceSingleTileKernel): one block for all data +// +// BI-V100: 16 SMs → max 32 CTAs → atomic contention is negligible +// → not_guaranteed determinism (atomic path) is optimal for decode +// +// PENDING REAL BENCHMARK: LOAD_LDG vs LOAD_DEFAULT +// topk bench showed LOAD_DEFAULT (ld=0) beat LOAD_LDG (ld=1). +// reduce may be similar — BI-V100's L1/L2 behavior differs from NVIDIA. +// Keep LOAD_LDG for now (CCCL default), switch if bench says otherwise. #pragma once @@ -56,70 +72,65 @@ enum class determinism_t { // ============================================================ // BI-V100 tuning values for plus<> operator // -// SM=16 strategy: maximize tile size within 48KB SMEM. -// With 32 concurrent CTAs (16 SMs × 2 occupancy), each CTA -// should process ≥ 49152/(accum_size) elements per tile. -// -// CCCL benchmark format reference: -// ipt_.tpb_.ipv_ -// where s_N = speedup vs TUNE_BASE at 2^N elements +// Tile sizing rationale (corrected): +// NOT SMEM-limited — reduce loads to registers, not SMEM staging. +// Constraint is register pressure + occupancy: +// items=24, float32 → 24 regs for data + overhead → ~40 regs/thread +// SM has 64K regs → 64K/40 = 1600 threads max → 3 CTAs @ 512 threads ✓ +// Also want large tiles (few CTAs) because 16 SMs can't launch many anyway. +// CCCL SM100 uses items=16 for float32 (with 148 SMs, more CTAs = more parallelism) +// BI-V100 uses items=24 (fewer CTAs, each processes more data) // ============================================================ // --- plus<> operator, two-phase (WARP_REDUCTIONS) --- struct bi100_plus_accum1_o4 { - // accum_size=1 (int8/uint8/bool), tile = 512*32*1 = 16384 (33% SMEM) - // Scaled from CCCL: nominal_4B_items=16 → items=16*4/1=64, clamped to 32 - // SM=16: want larger tile → threads=512, items=32 + // accum_size=1, reg pressure: 32 regs × 1B = trivial static constexpr int items = 32; static constexpr int threads = 512; - static constexpr int vec = 4; + static constexpr int vec = 4; // 4×1B = 4-byte vectorized load }; struct bi100_plus_accum2_o4 { - // accum_size=2 (int16/uint16/float16/bfloat16), tile = 512*24*2 = 24576 (50%) - // Qwen3.6 uses bfloat16 for KV cache — this is a hot path + // accum_size=2 (float16/bfloat16 — Qwen3.6 KV cache hot path) + // reg pressure: 24 regs × 2B accum overhead → moderate static constexpr int items = 24; static constexpr int threads = 512; - static constexpr int vec = 2; + static constexpr int vec = 2; // 2×2B = 4-byte vectorized load }; struct bi100_plus_float32_o4 { - // accum_size=4, tile = 512*24*4 = 49152 (100% SMEM — max utilization) - // SM100 ref: ipt_16.tpb_512.ipv_2 → tile=32768 (67% SMEM) - // SM=16 optimization: increase items from 16→24 to fill SMEM - // This gives each CTA 50% more data, compensating for fewer CTAs + // accum_size=4 (float32 — paged_attention score reduction) + // reg pressure: 24 × 4B = 96B per thread → ~24 regs just for data + // CCCL SM100: items=16, threads=512, vec=2 + // BI-V100: items=24 (fewer CTAs = each does more, compensating for 16 SMs) static constexpr int items = 24; static constexpr int threads = 512; - static constexpr int vec = 2; + static constexpr int vec = 2; // items%vec=0 ✓ → vectorized load enabled }; struct bi100_plus_float32_o8 { - // Same as o4 but with 8-byte offset — vec=1 for alignment static constexpr int items = 24; static constexpr int threads = 512; - static constexpr int vec = 1; + static constexpr int vec = 1; // 8-byte offset → unaligned risk, disable vectorize }; struct bi100_plus_float64_o4 { - // accum_size=8, SM100 uses threads=640 items=16 → tile=81920 > 49152 OVERFLOW! - // Max items at threads=384: 49152/(384*8) = 16 → tile = 49152 (100%) - // Alternatively threads=512 items=12 → tile = 49152 (100%) - // Choose 384×16: more items/thread = fewer loop iterations = better ILP + // accum_size=8, reg pressure: 16 × 2 regs (double) = 32 regs → fine + // CCCL SM100: items=16, threads=640 → 640 NOT multiple of 32 on all warps + // BI-V100: threads=384 (12 warps, clean), items=16 static constexpr int items = 16; static constexpr int threads = 384; - static constexpr int vec = 2; + static constexpr int vec = 2; // 2×8B = 16-byte vectorized load }; struct bi100_plus_float64_o8 { - // 8-byte offset + 8-byte accum: vec=1 static constexpr int items = 16; static constexpr int threads = 384; static constexpr int vec = 1; }; struct bi100_plus_int64_o4 { - // Same SMEM constraint as float64 (accum_size=8) static constexpr int items = 16; static constexpr int threads = 384; static constexpr int vec = 2; @@ -132,63 +143,51 @@ struct bi100_plus_int64_o8 { }; struct bi100_plus_accum16_o4 { - // accum_size=16 (int128/complex), tile = 192*16*16 = 49152 (100%) + // accum_size=16 (int128/complex) + // reg pressure: 16 × 4 regs = 64 regs → high, keep items=16 + // threads=192 (6 warps) to allow 2 CTAs per SM for occupancy static constexpr int items = 16; static constexpr int threads = 192; - static constexpr int vec = 1; + static constexpr int vec = 1; // sizeof(InputT)=16 > 8 → ATTEMPT_VECTORIZATION=false }; // --- Deterministic tunings: BLOCK_REDUCE_RAKING, vec=1 --- -// SM=16 fix: increase tile from ~23% to ≥50% SMEM utilization +// RAKING needs more SMEM than WARP_REDUCTIONS (scratch array per warp-step) +// but still far less than scan's BlockLoad staging struct bi100_det_float32 { - // OLD: threads=224 items=13 → tile=11648 (23% SMEM) — way too small for 16 SMs - // NEW: threads=384 items=32 → tile=49152 (100% SMEM) - // With only 32 concurrent CTAs, maxing SMEM per CTA is critical + // Deterministic: RAKING with large tile + // items=32, threads=384 → 32 regs for data → acceptable static constexpr int items = 32; static constexpr int threads = 384; }; struct bi100_det_float64 { - // OLD: threads=128 items=11 → tile=11264 (23% SMEM) - // NEW: threads=384 items=16 → tile=49152 (100% SMEM) static constexpr int items = 16; static constexpr int threads = 384; }; struct bi100_det_int32 { - // int32 deterministic: threads=384 items=32 → tile=49152 (100%) static constexpr int items = 32; static constexpr int threads = 384; }; struct bi100_det_int16 { - // int16/float16/bfloat16 deterministic - // threads=384 items=64 → tile=49152 (100%) + // int16/float16/bfloat16: items=64 → 64 regs × 2B/accum → moderate static constexpr int items = 64; static constexpr int threads = 384; }; -// --- Default fallback for unknown types/ops --- +// --- Default fallback --- struct bi100_default { - // SM60-equivalent but with SM=16 tile maximization - // threads=256 items=24 → at accum_size=4: tile=24576 (50% SMEM, safe margin) - // at accum_size=8: 256*24*8 = 49152 (100%) static constexpr int items = 24; static constexpr int threads = 256; static constexpr int vec = 4; }; // ============================================================ -// policy_selector — full dispatch matching CCCL structure -// -// Dispatch order: -// 1. determinism mode (gpu_to_gpu → RAKING, else → WARP_REDUCTIONS) -// 2. operator type (plus → specialized, min/max → same as plus for BI-V100) -// 3. accum_size (1B, 2B, 4B, 8B, 16B) -// 4. offset_size (4B vs 8B affects vec_size) -// 5. accum_type (float32/float64 get specific tunings) +// policy_selector // ============================================================ struct policy_selector { @@ -198,10 +197,8 @@ struct policy_selector { int accum_size; determinism_t determinism = determinism_t::run_to_run; - // --- Deterministic path: BLOCK_REDUCE_RAKING --- constexpr ReducePolicy get_deterministic(const hardware_capability& hw) const { if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) { - // Type-specific tunings for deterministic reduce if (accum_size <= 2) { auto [i, t] = scale_mem_bound(bi100_det_int16::threads, bi100_det_int16::items, accum_size); @@ -221,19 +218,15 @@ struct policy_selector { return {rp, rp}; } } - // Fallback for unknown hardware auto [i, t] = scale_mem_bound(256, 16, accum_size); ReducePassPolicy rp{t, i, 1, BLOCK_REDUCE_RAKING, LOAD_DEFAULT}; return {rp, rp}; } - // --- Two-phase path: BLOCK_REDUCE_WARP_REDUCTIONS --- constexpr ReducePolicy get_two_phase(const hardware_capability& hw) const { if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) { - // plus<> operator — fully specialized dispatch if (operation_t == op_kind_t::plus || operation_t == op_kind_t::min || operation_t == op_kind_t::max) { - // accum_size=1 (int8, uint8, bool) if (accum_size == 1) { auto [i, t] = scale_mem_bound(bi100_plus_accum1_o4::threads, bi100_plus_accum1_o4::items, accum_size); @@ -241,7 +234,6 @@ struct policy_selector { BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - // accum_size=2 (int16, float16, bfloat16) if (accum_size == 2) { auto [i, t] = scale_mem_bound(bi100_plus_accum2_o4::threads, bi100_plus_accum2_o4::items, accum_size); @@ -249,7 +241,6 @@ struct policy_selector { BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - // accum_size=4 (float32, int32) if (accum_size == 4) { int vec = (offset_size <= 4) ? bi100_plus_float32_o4::vec : bi100_plus_float32_o8::vec; @@ -258,7 +249,6 @@ struct policy_selector { ReducePassPolicy rp{t, i, vec, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - // accum_size=8 (float64, int64) if (accum_size == 8) { if (accum_t == type_t::float64) { int vec = (offset_size <= 4) ? bi100_plus_float64_o4::vec @@ -268,7 +258,6 @@ struct policy_selector { ReducePassPolicy rp{t, i, vec, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - // int64 and other 8-byte types int vec = (offset_size <= 4) ? bi100_plus_int64_o4::vec : bi100_plus_int64_o8::vec; auto [i, t] = scale_mem_bound(bi100_plus_int64_o4::threads, @@ -276,7 +265,6 @@ struct policy_selector { ReducePassPolicy rp{t, i, vec, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - // accum_size=16 (int128, complex) if (accum_size == 16) { auto [i, t] = scale_mem_bound(bi100_plus_accum16_o4::threads, bi100_plus_accum16_o4::items, accum_size); @@ -286,7 +274,6 @@ struct policy_selector { } } } - // Fallback: SM60-equivalent with SM=16 tile optimization auto [i, t] = scale_mem_bound(bi100_default::threads, bi100_default::items, accum_size); ReducePassPolicy rp{t, i, bi100_default::vec, @@ -294,7 +281,6 @@ struct policy_selector { return {rp, rp}; } - // --- Main entry point --- constexpr ReducePolicy operator()(const hardware_capability& hw) const { if (determinism == determinism_t::gpu_to_gpu) return get_deterministic(hw);