From c7a63bc2c8eb512421d9d84d16af90a58862b25e Mon Sep 17 00:00:00 2001 From: Claude Date: Thu, 30 Jul 2026 14:37:38 +0000 Subject: [PATCH] =?UTF-8?q?[MUH]=20Fix=207=20structural=20discrepancies=20?= =?UTF-8?q?vs=20CCCL=20=E2=80=94=20read=20source,=20not=20grep?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Fixes found by reading all 17 muh files + 6 CCCL counterpart policy_selectors as full source code input: 1. topk: BLOCK_LOAD_DIRECT → BLOCK_LOAD_VECTORIZE (CCCL SM90+ uses VECTORIZE). bits_per_pass was wrong (muh: ks<=4→9, CCCL: ks>=2→11). items now computed dynamically (4*4/key_size) not hardcoded. 2. reduce: added determinism dispatch — three modes matching CCCL: gpu_to_gpu (BLOCK_REDUCE_RAKING, vec_size=1, LOAD_DEFAULT), run_to_run (WARP_REDUCTIONS, LOAD_LDG, default), not_guaranteed (WARP_REDUCTIONS_NONDETERMINISTIC). Added bi100_det_float32 and bi100_det_float64 tuning structs with SM90 benchmark reference values. 3. batch_memcpy: flat single-tier → SmallBuffer+LargeBuffer two-tier matching CCCL structure (128 threads small, 256 threads large, warp_threshold=128, block_threshold=8192). 4. transform: single BulkPolicy → three-policy structure (VectorizedPolicy + AsyncCopyPolicy + PrefetchPolicy) matching CCCL. items_per_thread computed from bytes_in_flight / (threads * elem_size). 5. compile_test: 17 checks → 33 checks. Now verifies exact values: reduce determinism modes, topk VECTORIZE + bits=11, batch_memcpy two-tier thresholds, transform three-policy structure. 6. gen_patch: added fallback extraction for inline policy_selector values (topk now generates SAMPLING_BLOCK_SIZE patch). 7. MUH_PROJECT_CHECKPOINT.md: 'PRD设计阶段还没有代码' → actual status. 7 files changed, 413 insertions, 265 deletions. --- MUH_PROJECT_CHECKPOINT.md | 8 +- muh/gen_patch.py | 37 +++- .../muh/tuning/tuning_batch_memcpy.cuh | 87 +++++---- muh/include/muh/tuning/tuning_reduce.cuh | 117 +++++++++--- muh/include/muh/tuning/tuning_topk.cuh | 118 ++++++------ muh/include/muh/tuning/tuning_transform.cuh | 133 +++++++------ muh/test/compile_test.cpp | 174 +++++++++--------- 7 files changed, 411 insertions(+), 263 deletions(-) diff --git a/MUH_PROJECT_CHECKPOINT.md b/MUH_PROJECT_CHECKPOINT.md index c39b2df8..3995abc8 100644 --- a/MUH_PROJECT_CHECKPOINT.md +++ b/MUH_PROJECT_CHECKPOINT.md @@ -45,7 +45,7 @@ muh 是我们设计的 **tuning DSL（领域特定语言）**，用于： **为什么需要它**: CCCL 有 27 个 tuning_*.cuh 文件（17000+ 行），每个算法都有针对不同 NVIDIA SM 架构的特化参数。天垓100 不是 NVIDIA GPU，不能直接用这些参数，但 tuning 的维度（block size、warp 策略、shared memory 用量、prefetch 策略）是通用的。muh 让迁移过程变成"改配置 + 跑 benchmark"而不是"手改 kernel + 祈祷"。 -**muh 的状态**: PRD 设计阶段，还没有代码。 +**muh 的状态**: v0.3 — 6个算法的C++ tuning headers已就绪(reduce/scan/topk/transform/batch_memcpy/for)，compile_test 33项通过，gen_patch.py从C++ headers提取bi100值生成vllm patches。参数值从CCCL SM100复制，等BI-V100实测替换。 ## 四、已完成的工作 @@ -93,9 +93,9 @@ muh 是我们设计的 **tuning DSL（领域特定语言）**，用于： ## 五、还没做的（下一步） -1. **muh 语言 PRD 设计** — 定义 muh 的 schema、语法、codegen target、参数空间 -2. **muh PRD items 写入 project 6** — 作为真实 Issue，带 label 和测试用例 -3. **从 CCCL tuning_*.cuh 提取参数空间** — 建立"CCCL tuning 维度 → muh 配置项"的映射 +1. ~~muh 语言 PRD 设计~~ ✅ Done — muh是C++ header-only lib，不是独立语言 +2. ~~从 CCCL tuning_*.cuh 提取参数空间~~ ✅ Done — 6个算法的policy_selector已实现 +3. **在BI-V100上跑benchmark** — 用实测数据替换bi100_*中的SM100复制值 4. **获取 enginex-vllm-bi100-qwen36 的实际代码** — 需要在 Phanthy Cloud 开发环境里操作 5. **设计 muh → vllm kernel 的 codegen 管道** 6. **实际在天垓100 上跑 benchmark** diff --git a/muh/gen_patch.py b/muh/gen_patch.py index ef335634..3dc222bb 100644 --- a/muh/gen_patch.py +++ b/muh/gen_patch.py @@ -116,6 +116,36 @@ VLLM_INJECTION_POINTS = { } +def extract_hardcoded_values(filepath): + """Fallback: extract key values from policy_selector return statements. + + For algorithms where bi100_* structs don't exist (values computed inline). + Extracts the threads_per_block from the first return in the iluvatar branch. + """ + with open(filepath, 'r') as f: + content = f.read() + + algo = algo_from_filename(filepath) + + # Find iluvatar branch + iluvatar_match = re.search( + r'hw\.at_least\(.*iluvatar.*?\)\s*\{(.*?)(?=\n\s{2,4}\})', + content, re.DOTALL + ) + if not iluvatar_match: + return [] + + branch = iluvatar_match.group(1) + + # Find return {N, ...} — first integer is typically threads_per_block + return_match = re.search(r'return\s*\{(\d+)', branch) + if not return_match: + return [] + + threads = int(return_match.group(1)) + return [('__inline__', {'threads': threads})] + + def generate_patches(header_dir): """Read all tuning headers, extract bi100 values, generate patches.""" patches = [] @@ -131,8 +161,11 @@ def generate_patches(header_dir): structs = extract_bi100_structs(hpath) if not structs: - summary.append(f"SKIP {algo}: no bi100_* structs found") - continue + # Fallback: try extracting inline values from policy_selector + structs = extract_hardcoded_values(hpath) + if not structs: + summary.append(f"SKIP {algo}: no bi100_* structs and no inline values found") + continue # Use the first non-default struct as the primary tuning # (default is fallback; prefer the type-specific ones) diff --git a/muh/include/muh/tuning/tuning_batch_memcpy.cuh b/muh/include/muh/tuning/tuning_batch_memcpy.cuh index a1449556..ab476aed 100644 --- a/muh/include/muh/tuning/tuning_batch_memcpy.cuh +++ b/muh/include/muh/tuning/tuning_batch_memcpy.cuh @@ -4,6 +4,12 @@ // // vllm impact: KV cache block copy between GPU memory regions // Competition weight: Cache TPS × 0.56 +// +// CCCL structure: two-tier (SmallBuffer handled by single block, +// LargeBuffer by multi-block collaboration). Thresholds: +// warp_level: 128 bytes +// block_level: 8 KiB +// muh must replicate this structure, not flatten it. #pragma once @@ -12,51 +18,68 @@ namespace muh::tuning::batch_memcpy { -/// Batch memcpy policy -struct BatchMemcpyPolicy { +/// Small buffer policy: single thread block handles many small buffers +struct SmallBufferPolicy { int threads_per_block; + int buffers_per_thread; + int bytes_per_thread; + bool prefer_pow2_bits; + int block_level_tile_size; + int warp_level_threshold; + int block_level_threshold; LookbackDelayPolicy buffer_lookback_delay; LookbackDelayPolicy block_lookback_delay; }; -// ============================================================ -// BI-V100 tuning values -// -// CCCL reference from tuning_batch_memcpy.cuh: -// Default: threads=256 -// Buffer lookback delay and block lookback delay are separate — -// they control the two decoupled lookback scans that coordinate -// the batch copy across thread blocks. -// -// For vllm: KV cache copies are typically large contiguous blocks -// (head_dim × num_layers × sizeof(half)), so high throughput matters -// more than latency. -// ============================================================ +/// Large buffer policy: multiple blocks collaborate on one large buffer +struct LargeBufferPolicy { + int threads_per_block; + int bytes_per_thread; +}; -struct bi100_default { - static constexpr int threads = 256; - static constexpr LookbackDelayPolicy buffer_delay = { - LookbackDelayAlgorithm::fixed_delay, 350, 450}; - static constexpr LookbackDelayPolicy block_delay = { - LookbackDelayAlgorithm::fixed_delay, 350, 450}; +/// Full batch memcpy policy +struct BatchMemcpyPolicy { + SmallBufferPolicy small_buffer; + LargeBufferPolicy large_buffer; }; // ============================================================ -// policy_selector +// BI-V100 tuning values — from CCCL SM70+ defaults +// +// CCCL policy_selector (all architectures): +// small: 128 threads, 4 bufs/thread, 8 bytes/thread +// prefer_pow2_bits = (cc < 7.0) +// warp_threshold = 128, block_threshold = 8192 +// delays = default_delay_constructor_policy(true) +// large: 256 threads, 32 bytes/thread +// +// For BI-V100: start with CCCL defaults. +// The delay policy is arch-sensitive (CCCL uses +// default_delay_constructor_policy which picks fixed_delay for +// primitive types). We use fixed_delay as starting point. // ============================================================ struct policy_selector { - constexpr BatchMemcpyPolicy operator()(const hardware_capability& hw) const { - if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) { - return {bi100_default::threads, - bi100_default::buffer_delay, - bi100_default::block_delay}; - } - // Fallback - return {256, - {LookbackDelayAlgorithm::fixed_delay, 350, 450}, - {LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + constexpr BatchMemcpyPolicy operator()(const hardware_capability& hw) const { + // BI-V100: assume >= SM70 equivalent (no prefer_pow2_bits) + bool prefer_pow2 = false; + + LargeBufferPolicy large{256, 32}; + + SmallBufferPolicy small{ + /* threads_per_block = */ 128, + /* buffers_per_thread = */ 4, + /* bytes_per_thread = */ 8, + /* prefer_pow2_bits = */ prefer_pow2, + /* block_level_tile_size = */ large.threads_per_block * large.bytes_per_thread, + /* warp_level_threshold = */ 128, + /* block_level_threshold = */ 8 * 1024, + /* buffer_lookback_delay = */ {LookbackDelayAlgorithm::fixed_delay, 350, 450}, + /* block_lookback_delay = */ {LookbackDelayAlgorithm::fixed_delay, 350, 450}, + }; + + return {small, large}; } }; diff --git a/muh/include/muh/tuning/tuning_reduce.cuh b/muh/include/muh/tuning/tuning_reduce.cuh index da6d3a53..ecc1c7a2 100644 --- a/muh/include/muh/tuning/tuning_reduce.cuh +++ b/muh/include/muh/tuning/tuning_reduce.cuh @@ -28,54 +28,75 @@ struct ReducePolicy { 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 +}; + // ============================================================ -// BI-V100 tuning values -// Status: PENDING BENCHMARK +// BI-V100 tuning values — initialized from CCCL SM100 benchmarks +// Status: PENDING BI-V100 BENCHMARK (values will change) // -// These are initialized from CCCL SM90/SM100 reference values. -// Must be replaced with actual BI-V100 benchmark results. +// 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 SM100 reference (for comparison): -// float32, offset_4: ipt=16, tpb=512, ipv=2 → 1.061x speedup -// float64, offset_4: ipt=16, tpb=640, ipv=1 → 1.018x speedup -// int64, offset_4: ipt=15, tpb=512, ipv=2 → 1.020x speedup -// int64, offset_8: ipt=15, tpb=512, ipv=1 → 1.019x speedup +// 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 // ============================================================ -/// BI-V100 tuning for sum(float32), 4-byte offset +// --- Non-deterministic (default) tunings --- + struct bi100_float32_plus_o4 { - // Benchmark annotation format: ipt_N.tpb_M.ipv_K - // BI-V100: TBD — using SM100 reference as starting point + // BI-V100: TBD — SM100 ref: ipt_16.tpb_512.ipv_2 1.061 1.000 1.065 1.167 static constexpr int items = 16; static constexpr int threads = 512; static constexpr int items_per_vec_load = 2; }; -/// BI-V100 tuning for sum(float64), 4-byte offset struct bi100_float64_plus_o4 { - // BI-V100: TBD + // 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; static constexpr int items_per_vec_load = 1; }; -/// BI-V100 tuning for sum(int64), 4-byte offset struct bi100_int64_plus_o4 { - // BI-V100: TBD + // 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; static constexpr int items_per_vec_load = 2; }; -/// BI-V100 tuning for sum(int64), 8-byte offset struct bi100_int64_plus_o8 { - // BI-V100: TBD + // 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; static constexpr int items_per_vec_load = 1; }; -/// BI-V100 default fallback +// --- Deterministic tunings (BLOCK_REDUCE_RAKING) --- +// CCCL uses these when determinism == gpu_to_gpu +// vec_size is forced to 1 for deterministic reduction + +struct bi100_det_float32 { + // BI-V100: TBD — SM90 ref: ipt_13.tpb_224 1.107 1.010 1.097 1.317 + 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 + static constexpr int items = 11; + static constexpr int threads = 128; +}; + +/// Fallback for types without specific tuning struct bi100_default { static constexpr int items = 16; static constexpr int threads = 256; @@ -83,12 +104,12 @@ struct bi100_default { }; // ============================================================ -// policy_selector — the core dispatch functor -// Follows CCCL's exact pattern: -// 1. Accept hardware description -// 2. Match type/op/size to a tuning struct -// 3. Apply mem-bound scaling -// 4. Return ReducePolicy +// 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 // ============================================================ struct policy_selector { @@ -96,10 +117,35 @@ struct policy_selector { op_kind_t operation_t; int offset_size; int accum_size; + determinism_t determinism = determinism_t::run_to_run; - /// Dispatch for BI-V100 - constexpr ReducePolicy operator()(const hardware_capability& hw) const { - // Only tuned for sum currently (matching CCCL's approach) + /// 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) { + auto [t, i] = scale_mem_bound(bi100_det_float32::threads, + bi100_det_float32::items, accum_size); + ReducePassPolicy rp{t, i, 1, BLOCK_REDUCE_RAKING, LOAD_DEFAULT}; + return {rp, rp}; + } + if (accum_t == type_t::float64) { + auto [t, i] = scale_mem_bound(bi100_det_float64::threads, + bi100_det_float64::items, accum_size); + ReducePassPolicy rp{t, i, 1, BLOCK_REDUCE_RAKING, LOAD_DEFAULT}; + 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)) { @@ -140,13 +186,26 @@ struct policy_selector { } } - // Fallback: conservative policy + // Fallback: SM60-equivalent conservative policy 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) { + return get_deterministic(hw); + } + + auto policy = get_two_phase(hw); + if (determinism == determinism_t::not_guaranteed) { + policy.multi_tile.reduce_algorithm = BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC; + } + return policy; + } }; } // namespace muh::tuning::reduce diff --git a/muh/include/muh/tuning/tuning_topk.cuh b/muh/include/muh/tuning/tuning_topk.cuh index 1ebebb3d..a186de43 100644 --- a/muh/include/muh/tuning/tuning_topk.cuh +++ b/muh/include/muh/tuning/tuning_topk.cuh @@ -4,6 +4,18 @@ // // vllm impact: Top-k / top-p sampling in decode stage // Competition weight: Output TPS × 16.796 (highest priority, tied with reduce) +// +// CCCL SM90+ policy_selector (the ground truth): +// bits_per_pass = calc_bits_per_pass(key_size): +// key_size 1 → 8 +// key_size 2 → 8 (but note: CCCL default returns 11 for 2/4/8, +// only key_size=1 returns 8. See switch below.) +// key_size 4 → 11 +// key_size 8 → 11 +// items_per_thread = max(1, 4 * 4 / key_size) // 16 bytes per thread +// threads_per_block = 512 +// load_algorithm = BLOCK_LOAD_VECTORIZE (NOT BLOCK_LOAD_DIRECT) +// scan_algorithm = BLOCK_SCAN_WARP_SCANS #pragma once @@ -12,7 +24,7 @@ namespace muh::tuning::topk { -/// Top-k policy (mirrors cub's topk_policy) +/// Top-k policy (mirrors cub::detail::topk::topk_policy) struct TopkPolicy { int threads_per_block; int items_per_thread; @@ -22,82 +34,68 @@ struct TopkPolicy { }; // ============================================================ -// BI-V100 tuning values +// bits_per_pass calculation — must match CCCL exactly // -// CCCL reference from tuning_topk.cuh policy_selector: -// The CCCL topk uses radix-based selection with configurable bits_per_pass. -// Key insight: bits_per_pass trades off #passes vs per-pass work. -// For small key_size (1-2B): 4-6 bits optimal -// For large key_size (4-8B): 8-11 bits optimal +// CCCL source (tuning_topk.cuh): +// case 1: default: return 8; +// case 2: case 4: case 8: return 11; // -// Default: threads=512, items=nominal_4b*4/key_size, bits=calc_bits_per_pass(key_size) -// calc_bits_per_pass: key_size<=2 → 8, key_size<=4 → 9, key_size<=8 → 10, else → 11 +// Previous muh version had a wrong mapping: +// key_size<=2 → 8, key_size<=4 → 9, key_size<=8 → 10 +// This was WRONG. CCCL's actual function returns 11 for 2/4/8. // ============================================================ -/// BI-V100 topk for 2-byte keys (float16/bfloat16 — most relevant for LLM logits) -struct bi100_topk_2B { - // LLM decode: logits are typically fp16/bf16, so key_size=2 - // CCCL default for 2B: bits_per_pass=8, items=4*4/2=8 - static constexpr int threads = 512; - static constexpr int items = 8; // nominal_4b_items=4, scaled: 4*4/2=8 - static constexpr int bits_per_pass = 8; - static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_DIRECT; - static constexpr BlockScanAlgorithm scan_algo = BLOCK_SCAN_WARP_SCANS; -}; - -/// BI-V100 topk for 4-byte keys (float32) -struct bi100_topk_4B { - static constexpr int threads = 512; - static constexpr int items = 4; // nominal_4b_items=4, scaled: 4*4/4=4 - static constexpr int bits_per_pass = 9; - static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_DIRECT; - static constexpr BlockScanAlgorithm scan_algo = BLOCK_SCAN_WARP_SCANS; -}; - -/// BI-V100 default fallback -struct bi100_topk_default { - static constexpr int threads = 256; - static constexpr int items = 4; - static constexpr int bits_per_pass = 8; - static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_DIRECT; - static constexpr BlockScanAlgorithm scan_algo = BLOCK_SCAN_WARP_SCANS; -}; +constexpr int calc_bits_per_pass(int key_size) { + switch (key_size) { + case 1: + default: + return 8; + case 2: + case 4: + case 8: + return 11; + } +} // ============================================================ // policy_selector +// +// Matches CCCL's SM90+ path exactly: +// threads = 512 +// items = max(1, nominal_4b(4) * 4 / key_size) +// load = BLOCK_LOAD_VECTORIZE +// scan = BLOCK_SCAN_WARP_SCANS +// bits = calc_bits_per_pass(key_size) +// +// BI-V100 values: using SM100 as starting point. +// Once benchmarked on BI-V100, items/threads may diverge. // ============================================================ struct policy_selector { int key_size; - static constexpr int calc_bits_per_pass(int ks) { - if (ks <= 2) return 8; - if (ks <= 4) return 9; - if (ks <= 8) return 10; - return 11; - } - constexpr TopkPolicy operator()(const hardware_capability& hw) const { if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) { - switch (key_size) { - case 1: - case 2: - return {bi100_topk_2B::threads, bi100_topk_2B::items, - bi100_topk_2B::load_algo, bi100_topk_2B::scan_algo, - bi100_topk_2B::bits_per_pass}; - case 4: - return {bi100_topk_4B::threads, bi100_topk_4B::items, - bi100_topk_4B::load_algo, bi100_topk_4B::scan_algo, - bi100_topk_4B::bits_per_pass}; - } + // SM90+ path from CCCL: 16 bytes per thread + constexpr int nominal_4b_items = 4; + int items = nominal_4b_items * 4 / key_size; + if (items < 1) items = 1; + + return {512, items, + BLOCK_LOAD_VECTORIZE, // CCCL uses VECTORIZE, not DIRECT + BLOCK_SCAN_WARP_SCANS, + calc_bits_per_pass(key_size)}; } - // Fallback - int items = (4 * 4) / key_size; + // Fallback: older arch path + constexpr int nominal_4b_items = 4; + int items = nominal_4b_items * 4 / key_size; if (items < 1) items = 1; - if (items > 4) items = 4; - return {bi100_topk_default::threads, items, - bi100_topk_default::load_algo, bi100_topk_default::scan_algo, + if (items > nominal_4b_items) items = nominal_4b_items; + + return {512, items, + BLOCK_LOAD_VECTORIZE, + BLOCK_SCAN_WARP_SCANS, calc_bits_per_pass(key_size)}; } }; diff --git a/muh/include/muh/tuning/tuning_transform.cuh b/muh/include/muh/tuning/tuning_transform.cuh index f08d0970..6ebd0041 100644 --- a/muh/include/muh/tuning/tuning_transform.cuh +++ b/muh/include/muh/tuning/tuning_transform.cuh @@ -4,6 +4,14 @@ // // vllm impact: Activation functions (SiLU, GELU), RMSNorm, residual add // Competition weight: Output TPS × 16.796 +// +// CCCL structure: three transform policy types selected by iterator properties: +// 1. TransformVectorizedPolicy — contiguous + trivially_relocatable inputs +// 2. TransformAsyncCopyPolicy — SM90+ bulk copy (cp.async.bulk) +// 3. TransformPrefetchPolicy — fallback when stable_address needed +// +// For vllm: activations are contiguous dense fp16/bf16 tensors. +// → primarily hits the vectorized path. #pragma once @@ -12,86 +20,103 @@ namespace muh::tuning::transform { -/// Bulk transform policy (for contiguous input/output) -struct BulkPolicy { +/// Vectorized transform: coalesced loads via vector types +struct VectorizedPolicy { int threads_per_block; int items_per_thread; - int vec_size; // vectorization width + int vec_size; }; -/// No-input policy (for fill/memset operations) +/// Async copy transform: SM90+ bulk shared memory copy +struct AsyncCopyPolicy { + int threads_per_block; + int min_items_per_thread; + int store_vec_size; +}; + +/// Prefetch transform: prefetch-based fallback +struct PrefetchPolicy { + int threads_per_block; +}; + +/// Fill policy (no input, e.g. memset) struct FillPolicy { int threads_per_block; - int items_per_thread_no_input; + int items_per_thread; }; -/// Full transform policy +/// Full transform policy — selected at dispatch time based on iterator properties struct TransformPolicy { - BulkPolicy bulk; + VectorizedPolicy vectorized; + AsyncCopyPolicy async_copy; + PrefetchPolicy prefetch; FillPolicy fill; }; // ============================================================ -// BI-V100 tuning values +// BI-V100 tuning // -// CCCL reference from tuning_transform.cuh: -// Default bulk: threads=256, items=auto, vec_size=auto -// The transform kernel is bandwidth-bound for large elementwise ops. -// Key insight: vec_size should match the hardware's natural vector width. -// For NVIDIA: vec_size is typically 4 (128-bit loads). -// For BI-V100: TBD, likely also 4 or 8. +// CCCL SM90+ vectorized path: +// threads = 256 (SM90) or 128 (SM100) +// items = computed from: max(items_for_vec, items_for_latency) +// items_for_vec = ceil(vec_bytes / min_elem_size) +// items_for_latency = (min_bytes_in_flight) / (threads * elem_size) +// vec_size = auto (power-of-2 aligned to hardware vector width) // -// items_per_thread is computed as: -// items_for_vec = ceil(vector_bytes / min_elem_size) -// items_for_latency = (latency * bandwidth) / (threads * elem_size) -// items = max(items_for_vec, items_for_latency) -// ============================================================ - -struct bi100_bulk_default { - static constexpr int threads = 256; - static constexpr int items = 8; // conservative starting point - static constexpr int vec_size = 4; // 128-bit vector loads -}; - -struct bi100_fill_default { - static constexpr int threads = 256; - static constexpr int items = 2; -}; - -// ============================================================ -// policy_selector +// CCCL SM90+ async_copy path: +// threads = 256 (SM90) or 128 (SM100) +// min_items_per_thread = computed from SMEM capacity +// store_vec_size = auto_ublkcp_store_vec_size(output.value_type_size) +// +// CCCL prefetch: +// threads = 256 +// +// For BI-V100: start with SM100 values (128 threads for bulk/async, +// 256 for prefetch). vec_size = 4 (128-bit loads, standard for most GPUs). // ============================================================ struct policy_selector { - int min_elem_size; // minimum element size across all inputs - int max_elem_size; // maximum element size - int num_inputs; // number of input iterators + int min_elem_size; + int max_elem_size; + int num_inputs; + bool all_contiguous; + bool all_trivially_relocatable; + bool requires_stable_address; constexpr TransformPolicy operator()(const hardware_capability& hw) const { - if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) { - // For BI-V100: start with CCCL defaults, tune via benchmark - int vec_size = bi100_bulk_default::vec_size; + // Compute vectorization params + int vec_size = 4; // 128-bit, matches CCCL default - // Compute items_per_thread based on vectorization - int items_for_vec = (vec_size * 4) / min_elem_size; // 4 = sizeof(int) - if (items_for_vec < 1) items_for_vec = 1; + // items_for_vec: how many items fit in one vector load + int items_for_vec = (vec_size * 4) / min_elem_size; // 4 = sizeof(int) + if (items_for_vec < 1) items_for_vec = 1; - // Ensure items is a multiple of vec_size - int items = items_for_vec; - if (items % vec_size != 0) { - items = ((items / vec_size) + 1) * vec_size; - } + // 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; + int items_for_latency = bytes_in_flight / (256 * min_elem_size); + if (items_for_latency < 1) items_for_latency = 1; - return { - {bi100_bulk_default::threads, items, vec_size}, - {bi100_fill_default::threads, bi100_fill_default::items} - }; + int bulk_items = items_for_vec > items_for_latency ? items_for_vec : items_for_latency; + + // Ensure items is a multiple of vec_size for aligned access + if (bulk_items % vec_size != 0) { + bulk_items = ((bulk_items / vec_size) + 1) * vec_size; } - // Fallback + // BI-V100: 128 threads for bulk (SM100-like), 256 for prefetch + int bulk_threads = hw.at_least(hardware_capability::vendor_t::iluvatar, 100) ? 128 : 256; + return { - {bi100_bulk_default::threads, bi100_bulk_default::items, bi100_bulk_default::vec_size}, - {bi100_fill_default::threads, bi100_fill_default::items} + // vectorized + {bulk_threads, bulk_items, vec_size}, + // async_copy (BI-V100 may not support cp.async.bulk — conservative) + {bulk_threads, 4, vec_size}, + // prefetch + {256}, + // fill + {256, 2}, }; } }; diff --git a/muh/test/compile_test.cpp b/muh/test/compile_test.cpp index 3549da31..51bad6a5 100644 --- a/muh/test/compile_test.cpp +++ b/muh/test/compile_test.cpp @@ -1,10 +1,5 @@ // muh/test/compile_test.cpp — Compile-time verification of muh tuning headers // -// This test does NOT require a GPU. It verifies: -// 1. All headers parse without errors -// 2. All policy_selector functors instantiate and return valid policies -// 3. All bi100_* struct values are non-zero (not forgotten placeholders) -// // Build: g++ -std=c++17 -I muh/include muh/test/compile_test.cpp -o muh_test // Run: ./muh_test @@ -13,125 +8,140 @@ #include #include -// Helper: verify a value is non-zero (catches forgotten TBD placeholders) #define CHECK_NONZERO(expr, name) \ - do { \ - auto _v = (expr); \ - if (_v == 0) { \ - std::fprintf(stderr, "FAIL: %s == 0 (placeholder not filled)\n", name); \ - failures++; \ - } else { \ - passes++; \ - } \ - } while(0) + do { auto _v = (expr); if (_v == 0) { std::fprintf(stderr, "FAIL: %s == 0\n", name); failures++; } else { passes++; } } while(0) + +#define CHECK_EQ(expr, expected, name) \ + do { auto _v = (expr); if (_v != (expected)) { std::fprintf(stderr, "FAIL: %s == %d, expected %d\n", name, (int)_v, (int)(expected)); failures++; } else { passes++; } } while(0) #define CHECK_TRUE(expr, name) \ - do { \ - if (!(expr)) { \ - std::fprintf(stderr, "FAIL: %s\n", name); \ - failures++; \ - } else { \ - passes++; \ - } \ - } while(0) + do { if (!(expr)) { std::fprintf(stderr, "FAIL: %s\n", name); failures++; } else { passes++; } } while(0) int main() { + using namespace muh::tuning; // bring enum values into scope int passes = 0; int failures = 0; auto hw = muh::target_hw; - // --- Verify hardware descriptor --- - CHECK_TRUE(hw.vendor == muh::hardware_capability::vendor_t::iluvatar, - "target_hw.vendor == iluvatar"); + // --- Hardware descriptor --- + CHECK_TRUE(hw.vendor == muh::hardware_capability::vendor_t::iluvatar, "target_hw.vendor"); CHECK_NONZERO(hw.warp_size, "target_hw.warp_size"); - CHECK_NONZERO(hw.max_threads_per_block, "target_hw.max_threads_per_block"); - // --- Test reduce policy_selector --- + // --- reduce: default (run_to_run) --- { using namespace muh::tuning::reduce; auto ps = policy_selector{ .accum_t = muh::tuning::type_t::float32, .operation_t = muh::tuning::op_kind_t::plus, - .offset_size = 4, - .accum_size = 4, + .offset_size = 4, .accum_size = 4, }; - auto policy = ps(hw); - CHECK_NONZERO(policy.multi_tile.threads_per_block, - "reduce.float32.threads_per_block"); - CHECK_NONZERO(policy.multi_tile.items_per_thread, - "reduce.float32.items_per_thread"); - CHECK_NONZERO(policy.multi_tile.vec_size, - "reduce.float32.vec_size"); - - // Verify known bi100 value matches - CHECK_TRUE(policy.multi_tile.threads_per_block > 0 && - policy.multi_tile.threads_per_block <= 1024, - "reduce.threads_per_block in [1, 1024]"); + auto p = ps(hw); + CHECK_EQ(p.multi_tile.threads_per_block, 512, "reduce.f32.threads"); + CHECK_EQ(p.multi_tile.vec_size, 2, "reduce.f32.vec_size"); + CHECK_EQ(p.multi_tile.reduce_algorithm, BLOCK_REDUCE_WARP_REDUCTIONS, "reduce.f32.algo"); } - // --- Test topk policy_selector --- + // --- reduce: deterministic (gpu_to_gpu) --- + { + using namespace muh::tuning::reduce; + auto ps = policy_selector{ + .accum_t = muh::tuning::type_t::float32, + .operation_t = muh::tuning::op_kind_t::plus, + .offset_size = 4, .accum_size = 4, + .determinism = determinism_t::gpu_to_gpu, + }; + auto p = ps(hw); + CHECK_EQ(p.multi_tile.reduce_algorithm, BLOCK_REDUCE_RAKING, "reduce.det.algo=RAKING"); + CHECK_EQ(p.multi_tile.vec_size, 1, "reduce.det.vec_size=1"); + CHECK_EQ(p.multi_tile.load_modifier, LOAD_DEFAULT, "reduce.det.load=DEFAULT"); + } + + // --- reduce: nondeterministic --- + { + using namespace muh::tuning::reduce; + auto ps = policy_selector{ + .accum_t = muh::tuning::type_t::float32, + .operation_t = muh::tuning::op_kind_t::plus, + .offset_size = 4, .accum_size = 4, + .determinism = determinism_t::not_guaranteed, + }; + auto p = ps(hw); + CHECK_EQ(p.multi_tile.reduce_algorithm, BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC, + "reduce.nondet.algo=NONDETERMINISTIC"); + } + + // --- topk: verify VECTORIZE and correct bits_per_pass --- { using namespace muh::tuning::topk; - auto ps = policy_selector{.key_size = 2}; - auto policy = ps(hw); - CHECK_NONZERO(policy.threads_per_block, "topk.2B.threads_per_block"); - CHECK_NONZERO(policy.items_per_thread, "topk.2B.items_per_thread"); - CHECK_NONZERO(policy.bits_per_pass, "topk.2B.bits_per_pass"); - CHECK_TRUE(policy.bits_per_pass >= 4 && policy.bits_per_pass <= 11, - "topk.bits_per_pass in [4, 11]"); + // 2-byte keys (fp16 logits — LLM hot path) + auto p2 = policy_selector{.key_size = 2}(hw); + CHECK_EQ(p2.load_algorithm, BLOCK_LOAD_VECTORIZE, "topk.2B.load=VECTORIZE"); + CHECK_EQ(p2.bits_per_pass, 11, "topk.2B.bits=11"); // CCCL: case 2 → 11 + CHECK_EQ(p2.items_per_thread, 8, "topk.2B.items=8"); // 4*4/2=8 + CHECK_EQ(p2.threads_per_block, 512, "topk.2B.threads=512"); + + // 4-byte keys + auto p4 = policy_selector{.key_size = 4}(hw); + CHECK_EQ(p4.bits_per_pass, 11, "topk.4B.bits=11"); + CHECK_EQ(p4.items_per_thread, 4, "topk.4B.items=4"); // 4*4/4=4 + + // 1-byte keys + auto p1 = policy_selector{.key_size = 1}(hw); + CHECK_EQ(p1.bits_per_pass, 8, "topk.1B.bits=8"); + CHECK_EQ(p1.items_per_thread, 16, "topk.1B.items=16"); // 4*4/1=16 } - // --- Test scan policy_selector --- + // --- scan: lookback + lookahead --- { using namespace muh::tuning::scan; auto ps = policy_selector{ - .input_value_size = 4, - .accum_size = 4, - .offset_size = 4, + .input_value_size = 4, .accum_size = 4, .offset_size = 4, .input_type = muh::tuning::type_t::float32, .accum_type = muh::tuning::type_t::float32, .operation_t = muh::tuning::op_kind_t::plus, .is_primitive_accum = true, }; - auto policy = ps(hw); - CHECK_NONZERO(policy.lookback.threads_per_block, - "scan.float32.lookback.threads_per_block"); - CHECK_NONZERO(policy.lookback.items_per_thread, - "scan.float32.lookback.items_per_thread"); + auto p = ps(hw); + 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"); } - // --- Test transform policy_selector --- + // --- batch_memcpy: two-tier --- + { + using namespace muh::tuning::batch_memcpy; + auto p = policy_selector{}(hw); + CHECK_EQ(p.small_buffer.threads_per_block, 128, "batch_memcpy.small.threads=128"); + CHECK_EQ(p.small_buffer.buffers_per_thread, 4, "batch_memcpy.small.bufs=4"); + CHECK_EQ(p.small_buffer.warp_level_threshold, 128, "batch_memcpy.small.warp_thresh=128"); + CHECK_EQ(p.small_buffer.block_level_threshold, 8192, "batch_memcpy.small.block_thresh=8192"); + CHECK_EQ(p.large_buffer.threads_per_block, 256, "batch_memcpy.large.threads=256"); + CHECK_EQ(p.large_buffer.bytes_per_thread, 32, "batch_memcpy.large.bytes=32"); + } + + // --- transform: three-policy --- { using namespace muh::tuning::transform; auto ps = policy_selector{ - .min_elem_size = 2, - .max_elem_size = 2, - .num_inputs = 1, + .min_elem_size = 2, .max_elem_size = 2, .num_inputs = 1, + .all_contiguous = true, .all_trivially_relocatable = true, + .requires_stable_address = false, }; - auto policy = ps(hw); - CHECK_NONZERO(policy.bulk.threads_per_block, - "transform.bulk.threads_per_block"); + auto p = ps(hw); + CHECK_NONZERO(p.vectorized.threads_per_block, "transform.vectorized.threads"); + CHECK_NONZERO(p.vectorized.vec_size, "transform.vectorized.vec_size"); + CHECK_NONZERO(p.async_copy.threads_per_block, "transform.async_copy.threads"); + CHECK_EQ(p.prefetch.threads_per_block, 256, "transform.prefetch.threads=256"); + CHECK_EQ(p.fill.threads_per_block, 256, "transform.fill.threads=256"); } - // --- Test batch_memcpy policy_selector --- - { - using namespace muh::tuning::batch_memcpy; - auto ps = policy_selector{}; - auto policy = ps(hw); - CHECK_NONZERO(policy.threads_per_block, - "batch_memcpy.threads_per_block"); - } - - // --- Test for_each policy_selector --- + // --- for_each --- { using namespace muh::tuning::for_each; - auto ps = policy_selector{}; - auto policy = ps(hw); - CHECK_NONZERO(policy.threads_per_block, - "for_each.threads_per_block"); - CHECK_NONZERO(policy.items_per_thread, - "for_each.items_per_thread"); + auto p = policy_selector{}(hw); + CHECK_EQ(p.threads_per_block, 256, "for.threads=256"); + CHECK_EQ(p.items_per_thread, 4, "for.items=4"); } // --- Report ---