diff --git a/muh/include/muh/tuning/tuning_reduce.cuh b/muh/include/muh/tuning/tuning_reduce.cuh index ecc1c7a2..55f040e5 100644 --- a/muh/include/muh/tuning/tuning_reduce.cuh +++ b/muh/include/muh/tuning/tuning_reduce.cuh @@ -1,10 +1,25 @@ // muh/include/muh/tuning/tuning_reduce.cuh — BI-V100 reduce tuning // // Mirrors: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_reduce.cuh -// Pattern: policy_selector functor, dispatches on muh::hardware_capability -// // vllm impact: Attention score reduction in multi-head attention // Competition weight: Output TPS × 16.796 (highest priority) +// +// DERIVATION MODEL (not copy-paste from SM100): +// +// BI-V100 vs SM100 (B200): +// SMEM: 48KB vs 48KB (default) — same +// L2: 6MB vs 50MB — 8.3x smaller +// BW: 900 GB/s vs 8000 GB/s — 8.9x lower +// SM count: 50 vs 148 — 3x fewer +// BW/SM: 18 GB/s vs 54 GB/s — 3x lower (≈ A100 level) +// +// Constraint: tile_size = threads * items * accum_size <= SMEM (48KB) +// +// SM100 reduce float64 uses threads=640, items=16 → tile = 81920 bytes. +// 81920 > 49152 (BI-V100 SMEM). This would CRASH on BI-V100. +// Similarly int64 uses threads=512, items=15 → tile = 61440 > 49152. +// +// Fix: derive threads/items from SMEM constraint, not copy from SM100. #pragma once @@ -13,7 +28,6 @@ namespace muh::tuning::reduce { -/// Policy for a single reduction pass (mirrors cub::ReducePassPolicy) struct ReducePassPolicy { int threads_per_block; int items_per_thread; @@ -22,94 +36,85 @@ struct ReducePassPolicy { CacheLoadModifier load_modifier; }; -/// Full reduction policy (mirrors cub::ReducePolicy) struct ReducePolicy { ReducePassPolicy multi_tile; ReducePassPolicy single_tile; }; -/// Determinism modes (mirrors cuda::execution::determinism::__determinism_t) enum class determinism_t { - run_to_run, // default: WARP_REDUCTIONS + LOAD_LDG - gpu_to_gpu, // deterministic: RAKING + LOAD_DEFAULT - not_guaranteed, // nondeterministic: WARP_REDUCTIONS_NONDETERMINISTIC + run_to_run, + gpu_to_gpu, + not_guaranteed, }; // ============================================================ -// BI-V100 tuning values — initialized from CCCL SM100 benchmarks -// Status: PENDING BI-V100 BENCHMARK (values will change) +// BI-V100 tuning values — DERIVED from hardware constraints // -// CCCL SM100 benchmark results (from tuning_reduce.cuh): -// float32+o4: ipt=16, tpb=512, ipv=2 → 1.061x geo, 1.167x max -// float64+o4: ipt=16, tpb=640, ipv=1 → 1.018x geo, 1.057x max -// int64+o4: ipt=15, tpb=512, ipv=2 → 1.020x geo, 1.058x max -// int64+o8: ipt=15, tpb=512, ipv=1 → 1.019x geo, 1.057x max -// -// CCCL SM90 deterministic benchmark results: -// float32: ipt=13, tpb=224 → 1.107x geo, 1.317x max -// float64: ipt=11, tpb=128 → 1.232x geo, 1.582x max +// Key constraint: tile_bytes = threads * items * accum_size <= 48KB +// SM100 values that violate this are WRONG for BI-V100. // ============================================================ -// --- Non-deterministic (default) tunings --- - struct bi100_float32_plus_o4 { - // BI-V100: TBD — SM100 ref: ipt_16.tpb_512.ipv_2 1.061 1.000 1.065 1.167 + // accum_size=4, tile = 512*16*4 = 32768 ≤ 49152 ✓ + // SM100 ref: ipt_16.tpb_512.ipv_2 1.061 1.000 1.065 1.167 + // Derivation: SMEM OK, threads=512 for occupancy on 50 SMs. Keep. static constexpr int items = 16; static constexpr int threads = 512; static constexpr int items_per_vec_load = 2; }; struct bi100_float64_plus_o4 { - // BI-V100: TBD — SM100 ref: ipt_16.tpb_640.ipv_1 1.018 1.000 1.016 1.057 - static constexpr int items = 16; - static constexpr int threads = 640; + // SM100: threads=640, items=16 → tile = 640*16*8 = 81920 > 49152 ✗ OVERFLOW + // Derivation: max items at 512 threads = 49152/(512*8) = 12 + // SM90 used threads=256, items=16 → tile = 32768 (conservative) + // Choose: threads=512, items=12 → tile = 49152 (max utilization) + static constexpr int items = 12; + static constexpr int threads = 512; static constexpr int items_per_vec_load = 1; }; struct bi100_int64_plus_o4 { - // BI-V100: TBD — SM100 ref: ipt_15.tpb_512.ipv_2 1.020 1.000 1.018 1.058 - static constexpr int items = 15; - static constexpr int threads = 512; + // SM100: threads=512, items=15 → tile = 512*15*8 = 61440 > 49152 ✗ OVERFLOW + // Derivation: max items at 384 threads = 49152/(384*8) = 16 + // Choose: threads=384, items=16 → tile = 49152 (max utilization) + static constexpr int items = 16; + static constexpr int threads = 384; static constexpr int items_per_vec_load = 2; }; struct bi100_int64_plus_o8 { - // BI-V100: TBD — SM100 ref: ipt_15.tpb_512.ipv_1 1.019 1.000 1.017 1.057 - static constexpr int items = 15; - static constexpr int threads = 512; + // SM100: threads=512, items=15 → same overflow + // Derivation: same as o4 but vec=1 (8-byte offset reduces vectorization) + static constexpr int items = 16; + static constexpr int threads = 384; static constexpr int items_per_vec_load = 1; }; -// --- Deterministic tunings (BLOCK_REDUCE_RAKING) --- -// CCCL uses these when determinism == gpu_to_gpu -// vec_size is forced to 1 for deterministic reduction - +// Deterministic tunings: BLOCK_REDUCE_RAKING, vec_size=1 struct bi100_det_float32 { - // BI-V100: TBD — SM90 ref: ipt_13.tpb_224 1.107 1.010 1.097 1.317 + // SM90 ref: ipt_13.tpb_224 1.107 1.010 1.097 1.317 + // tile = 224*13*4 = 11648 ≤ 49152 ✓ (safe, same as SM90) static constexpr int items = 13; static constexpr int threads = 224; }; struct bi100_det_float64 { - // BI-V100: TBD — SM86 ref: ipt_11.tpb_128 1.232 1.002 1.245 1.582 + // SM86 ref: ipt_11.tpb_128 1.232 1.002 1.245 1.582 + // tile = 128*11*8 = 11264 ≤ 49152 ✓ static constexpr int items = 11; static constexpr int threads = 128; }; -/// Fallback for types without specific tuning struct bi100_default { + // SM60-equivalent fallback: tile = 256*16*accum_size + // At accum_size=8: 256*16*8 = 32768 ≤ 49152 ✓ static constexpr int items = 16; static constexpr int threads = 256; static constexpr int items_per_vec_load = 4; }; // ============================================================ -// policy_selector -// -// Dispatch logic mirrors CCCL's exactly: -// 1. if determinism == gpu_to_gpu → get_deterministic_tuning() -// 2. else → get_two_phase_tuning() -// 3. if determinism == not_guaranteed → override reduce_algorithm +// policy_selector — three determinism modes matching CCCL // ============================================================ struct policy_selector { @@ -119,8 +124,6 @@ struct policy_selector { int accum_size; determinism_t determinism = determinism_t::run_to_run; - /// Deterministic reduction: BLOCK_REDUCE_RAKING, vec_size=1, LOAD_DEFAULT - /// Matches CCCL get_deterministic_tuning() constexpr ReducePolicy get_deterministic(const hardware_capability& hw) const { if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) { if (accum_t == type_t::float32) { @@ -136,74 +139,56 @@ struct policy_selector { return {rp, rp}; } } - - // Fallback deterministic auto [t, i] = scale_mem_bound(256, 16, accum_size); ReducePassPolicy rp{t, i, 1, BLOCK_REDUCE_RAKING, LOAD_DEFAULT}; return {rp, rp}; } - /// Standard two-phase reduction: BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG - /// Matches CCCL get_two_phase_tuning() constexpr ReducePolicy get_two_phase(const hardware_capability& hw) const { if (operation_t == op_kind_t::plus && hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) { if (accum_t == type_t::float32 && offset_size == 4 && accum_size == 4) { - auto [t, i] = scale_mem_bound( - bi100_float32_plus_o4::threads, - bi100_float32_plus_o4::items, accum_size); + auto [t, i] = scale_mem_bound(bi100_float32_plus_o4::threads, + bi100_float32_plus_o4::items, accum_size); ReducePassPolicy rp{t, i, bi100_float32_plus_o4::items_per_vec_load, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - if (accum_t == type_t::float64 && offset_size == 4 && accum_size == 8) { - auto [t, i] = scale_mem_bound( - bi100_float64_plus_o4::threads, - bi100_float64_plus_o4::items, accum_size); + auto [t, i] = scale_mem_bound(bi100_float64_plus_o4::threads, + bi100_float64_plus_o4::items, accum_size); ReducePassPolicy rp{t, i, bi100_float64_plus_o4::items_per_vec_load, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - if (offset_size == 4 && accum_size == 8) { - auto [t, i] = scale_mem_bound( - bi100_int64_plus_o4::threads, - bi100_int64_plus_o4::items, accum_size); + auto [t, i] = scale_mem_bound(bi100_int64_plus_o4::threads, + bi100_int64_plus_o4::items, accum_size); ReducePassPolicy rp{t, i, bi100_int64_plus_o4::items_per_vec_load, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - if (offset_size == 8 && accum_size == 8) { - auto [t, i] = scale_mem_bound( - bi100_int64_plus_o8::threads, - bi100_int64_plus_o8::items, accum_size); + auto [t, i] = scale_mem_bound(bi100_int64_plus_o8::threads, + bi100_int64_plus_o8::items, accum_size); ReducePassPolicy rp{t, i, bi100_int64_plus_o8::items_per_vec_load, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } } - - // Fallback: SM60-equivalent conservative policy - auto [t, i] = scale_mem_bound( - bi100_default::threads, bi100_default::items, accum_size); + auto [t, i] = scale_mem_bound(bi100_default::threads, bi100_default::items, accum_size); ReducePassPolicy rp{t, i, bi100_default::items_per_vec_load, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; return {rp, rp}; } - /// Main dispatch — mirrors CCCL's operator()(compute_capability) constexpr ReducePolicy operator()(const hardware_capability& hw) const { - if (determinism == determinism_t::gpu_to_gpu) { + if (determinism == determinism_t::gpu_to_gpu) return get_deterministic(hw); - } - auto policy = get_two_phase(hw); - if (determinism == determinism_t::not_guaranteed) { + if (determinism == determinism_t::not_guaranteed) policy.multi_tile.reduce_algorithm = BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC; - } return policy; } }; diff --git a/muh/include/muh/tuning/tuning_scan.cuh b/muh/include/muh/tuning/tuning_scan.cuh index e68c525c..41732aac 100644 --- a/muh/include/muh/tuning/tuning_scan.cuh +++ b/muh/include/muh/tuning/tuning_scan.cuh @@ -5,6 +5,13 @@ // // vllm impact: Prefix scan in paged attention block table lookup // Competition weight: Input TPS × 2.799 +// +// DERIVATION (not copy-paste from SM100): +// - SMEM constraint: tile = threads * items * value_size <= 48KB +// SM100 8B tunings (416*23*8=76544, 320*22*8=56320) OVERFLOW on BI-V100 +// - Delay parameters: SM100 L2=50MB, BI-V100 L2=6MB (8.3x smaller) +// Smaller L2 → faster coherence → shorter delays +// Heuristic: ns *= 0.5, l2w *= 0.6 (to be refined by benchmark) #pragma once @@ -68,7 +75,7 @@ struct bi100_lookback_1B_o4 { static constexpr int threads = 512; static constexpr int items = 18; static constexpr LookbackDelayPolicy delay = { - LookbackDelayAlgorithm::exponential_backon, 768, 820}; + LookbackDelayAlgorithm::exponential_backon, 384, 492}; static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_WARP_TRANSPOSE; static constexpr BlockStoreAlgorithm store_algo = BLOCK_STORE_WARP_TRANSPOSE; static constexpr CacheLoadModifier load_mod = LOAD_DEFAULT; @@ -79,7 +86,7 @@ struct bi100_lookback_2B_o4 { static constexpr int threads = 512; static constexpr int items = 13; static constexpr LookbackDelayPolicy delay = { - LookbackDelayAlgorithm::exponential_backon, 1384, 720}; + LookbackDelayAlgorithm::exponential_backon, 692, 432}; static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_WARP_TRANSPOSE; static constexpr BlockStoreAlgorithm store_algo = BLOCK_STORE_WARP_TRANSPOSE; static constexpr CacheLoadModifier load_mod = LOAD_DEFAULT; @@ -90,7 +97,7 @@ struct bi100_lookback_4B_o4 { static constexpr int threads = 384; static constexpr int items = 22; static constexpr LookbackDelayPolicy delay = { - LookbackDelayAlgorithm::exponential_backon_jitter, 1904, 830}; + LookbackDelayAlgorithm::exponential_backon_jitter, 952, 498}; static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_WARP_TRANSPOSE; static constexpr BlockStoreAlgorithm store_algo = BLOCK_STORE_WARP_TRANSPOSE; static constexpr CacheLoadModifier load_mod = LOAD_DEFAULT; @@ -101,29 +108,31 @@ struct bi100_lookback_4B_o8 { static constexpr int threads = 416; static constexpr int items = 19; static constexpr LookbackDelayPolicy delay = { - LookbackDelayAlgorithm::exponential_backon, 956, 550}; + LookbackDelayAlgorithm::exponential_backon, 478, 330}; static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_WARP_TRANSPOSE; static constexpr BlockStoreAlgorithm store_algo = BLOCK_STORE_WARP_TRANSPOSE; static constexpr CacheLoadModifier load_mod = LOAD_CA; }; struct bi100_lookback_8B_o4 { - // SM100 ref: ipt_23.tpb_416.ns_772.dcid_5.l2w_710 → 1.089x + // SM100 ref: ipt_23.tpb_416 → tile=76544 > 49152 SMEM OVERFLOW + // Derived: items = 49152/(416*8) = 14. Delay halved (L2 6MB vs 50MB). static constexpr int threads = 416; - static constexpr int items = 23; + static constexpr int items = 14; static constexpr LookbackDelayPolicy delay = { - LookbackDelayAlgorithm::exponential_backon_jitter_window, 772, 710}; + LookbackDelayAlgorithm::exponential_backon_jitter_window, 386, 426}; static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_WARP_TRANSPOSE; static constexpr BlockStoreAlgorithm store_algo = BLOCK_STORE_WARP_TRANSPOSE; static constexpr CacheLoadModifier load_mod = LOAD_DEFAULT; }; struct bi100_lookback_8B_o8 { - // SM100 ref: ipt_22.tpb_320.ns_328.dcid_2.l2w_965 → 1.080x + // SM100 ref: ipt_22.tpb_320 → tile=56320 > 49152 SMEM OVERFLOW + // Derived: items = 49152/(320*8) = 19. Delay: ns*0.5, l2w*0.6. static constexpr int threads = 320; - static constexpr int items = 22; + static constexpr int items = 19; static constexpr LookbackDelayPolicy delay = { - LookbackDelayAlgorithm::exponential_backoff, 328, 965}; + LookbackDelayAlgorithm::exponential_backoff, 164, 579}; static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_WARP_TRANSPOSE; static constexpr BlockStoreAlgorithm store_algo = BLOCK_STORE_WARP_TRANSPOSE; static constexpr CacheLoadModifier load_mod = LOAD_DEFAULT; diff --git a/muh/include/muh/tuning/tuning_transform.cuh b/muh/include/muh/tuning/tuning_transform.cuh index 6ebd0041..4664969f 100644 --- a/muh/include/muh/tuning/tuning_transform.cuh +++ b/muh/include/muh/tuning/tuning_transform.cuh @@ -92,9 +92,10 @@ struct policy_selector { if (items_for_vec < 1) items_for_vec = 1; // items_for_latency: enough items to hide memory latency - // CCCL uses cc_to_min_bytes_in_flight(cc) which is ~48KB for SM90+ - // For BI-V100: estimate 48KB in flight, 256 threads - int bytes_in_flight = 48 * 1024; + // CCCL cc_to_min_bytes_in_flight: B200=64KB, H100=48KB, A100=16KB, V100=12KB + // BI-V100 per-SM BW = 900/50 = 18 GB/s ≈ A100 (2000/108 = 18.5 GB/s) + // → Use 16KB (A100-level), not 48-64KB + int bytes_in_flight = 16 * 1024; int items_for_latency = bytes_in_flight / (256 * min_elem_size); if (items_for_latency < 1) items_for_latency = 1; diff --git a/muh/test/compile_test.cpp b/muh/test/compile_test.cpp index 51bad6a5..8570e9c5 100644 --- a/muh/test/compile_test.cpp +++ b/muh/test/compile_test.cpp @@ -106,6 +106,17 @@ int main() { CHECK_EQ(p.lookback.threads_per_block, 384, "scan.f32.lookback.threads=384"); CHECK_EQ(p.lookback.items_per_thread, 22, "scan.f32.lookback.items=22"); CHECK_NONZERO(p.lookahead.reduce_and_scan_warps, "scan.f32.lookahead.warps"); + + // 8-byte scan: was SMEM overflow with SM100 values (416*23*8=76544 > 49152) + auto ps8 = policy_selector{ + .input_value_size = 8, .accum_size = 8, .offset_size = 4, + .input_type = type_t::int64, .accum_type = type_t::int64, + .operation_t = op_kind_t::plus, .is_primitive_accum = true, + }; + auto p8 = ps8(hw); + CHECK_EQ(p8.lookback.items_per_thread, 14, "scan.8B.items=14(derived)"); + CHECK_TRUE(p8.lookback.threads_per_block * p8.lookback.items_per_thread * 8 <= 49152, + "scan.8B.tile_fits_48KB_smem"); } // --- batch_memcpy: two-tier ---