diff --git a/muh/include/muh/hardware.cuh b/muh/include/muh/hardware.cuh new file mode 100644 index 00000000..ec1b0c7c --- /dev/null +++ b/muh/include/muh/hardware.cuh @@ -0,0 +1,62 @@ +// muh/include/muh/hardware.cuh — Iluvatar BI-V100 hardware descriptor +// +// This header replaces cuda::compute_capability as the dispatch key. +// CCCL's policy_selector uses operator()(cuda::compute_capability cc) +// to select tuning params. muh's policy_selector uses +// operator()(muh::hardware_capability hw) instead. + +#pragma once + +namespace muh { + +/// Hardware capability descriptor for non-NVIDIA GPUs. +/// Replaces cuda::compute_capability {major, minor} with a richer +/// description that captures what actually matters for kernel tuning. +struct hardware_capability { + int warp_size; // threads per warp (NVIDIA=32, BI-V100=TBD) + int max_threads_per_block; // max CTA size (typically 1024) + int max_shared_memory_per_block; // bytes of shared memory per block + int max_registers_per_thread; // max registers per thread + int l2_cache_size_bytes; // L2 cache size in bytes + int memory_bandwidth_gbps; // HBM bandwidth in GB/s + int sm_count; // number of SMs / compute units + + // For dispatch: identifies which tuning table to use + enum class vendor_t { nvidia, iluvatar, unknown }; + vendor_t vendor; + int arch_version; // e.g. 100 for BI-V100 + + // Convenience constructors + constexpr static hardware_capability bi_v100() { + return { + .warp_size = 32, // TBD: confirm on actual hardware + .max_threads_per_block = 1024, + .max_shared_memory_per_block = 49152, // 48 KiB, TBD + .max_registers_per_thread = 255, + .l2_cache_size_bytes = 6 * 1024 * 1024, // 6 MiB, TBD + .memory_bandwidth_gbps = 900, // TBD + .sm_count = 50, // 50c in the spec + .vendor = vendor_t::iluvatar, + .arch_version = 100, + }; + } + + // Comparison for dispatch: exact match on vendor + arch + constexpr bool operator==(const hardware_capability& o) const { + return vendor == o.vendor && arch_version == o.arch_version; + } + constexpr bool operator!=(const hardware_capability& o) const { + return !(*this == o); + } + + // Check if this hardware is "at least" a given capability + // For same vendor, compares arch_version + constexpr bool at_least(vendor_t v, int min_arch) const { + return vendor == v && arch_version >= min_arch; + } +}; + +/// Global default target — set to BI-V100 for competition +inline constexpr auto target_hw = hardware_capability::bi_v100(); + +} // namespace muh diff --git a/muh/include/muh/muh.cuh b/muh/include/muh/muh.cuh new file mode 100644 index 00000000..a330a38a --- /dev/null +++ b/muh/include/muh/muh.cuh @@ -0,0 +1,65 @@ +// muh/include/muh/muh.cuh — Top-level muh header +// +// Provides the complete tuning dispatch for Iluvatar BI-V100. +// Include this single header to get all tuning policies. +// +// Usage: +// #include +// +// auto hw = muh::target_hw; // BI-V100 by default +// auto reduce_policy = muh::tuning::reduce::policy_selector{ +// .accum_t = muh::tuning::type_t::float32, +// .operation_t = muh::tuning::op_kind_t::plus, +// .offset_size = 4, +// .accum_size = 4, +// }(hw); +// +// auto scan_policy = muh::tuning::scan::policy_selector{ +// .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, +// }(hw); + +#pragma once + +// Hardware descriptor +#include "muh/hardware.cuh" + +// Shared types (compatible with CCCL) +#include "muh/tuning/common.cuh" + +// Per-algorithm tuning (P0 = highest priority for competition) +#include "muh/tuning/tuning_reduce.cuh" // P0: attention reduction +#include "muh/tuning/tuning_topk.cuh" // P0: sampling top-k/top-p +#include "muh/tuning/tuning_scan.cuh" // P0: prefix scan in paged attention + +// P1 +#include "muh/tuning/tuning_transform.cuh" // P1: activation kernels +#include "muh/tuning/tuning_batch_memcpy.cuh" // P1: KV cache management + +// P2 +#include "muh/tuning/tuning_for.cuh" // P2: RoPE position encoding + +namespace muh { + +/// Version info +constexpr int MUH_VERSION_MAJOR = 0; +constexpr int MUH_VERSION_MINOR = 1; +constexpr int MUH_VERSION_PATCH = 0; + +/// Competition scoring formula +/// Token吞吐加权值 = Output TPS × 16.796 + Input TPS × 2.799 + Cache TPS × 0.56 +struct scoring { + static constexpr double output_weight = 16.796; + static constexpr double input_weight = 2.799; + static constexpr double cache_weight = 0.56; + static constexpr double baseline_threshold = 8000.0; // minimum to pass + static constexpr double advanced_uplift = 0.30; // 30% for advanced prize + static constexpr double special_uplift = 0.50; // 50% for special prize +}; + +} // namespace muh diff --git a/muh/include/muh/tuning/common.cuh b/muh/include/muh/tuning/common.cuh new file mode 100644 index 00000000..b3ad429e --- /dev/null +++ b/muh/include/muh/tuning/common.cuh @@ -0,0 +1,139 @@ +// muh/include/muh/tuning/common.cuh — Shared tuning types +// +// Compatible with cub/device/dispatch/tuning/common.cuh. +// Re-exports CCCL's enum types so muh tuning headers can reference them +// without pulling in the full CUB dependency tree. +// When compiling against actual CCCL, prefer their originals. + +#pragma once + +#include "muh/hardware.cuh" + +// If CCCL is available, use their definitions +#if __has_include() +#include +#include +#include +#include +#include + +namespace muh::tuning { + using cub::BlockLoadAlgorithm; + using cub::BlockStoreAlgorithm; + using cub::BlockReduceAlgorithm; + using cub::BlockScanAlgorithm; + using cub::CacheLoadModifier; + using cub::LookbackDelayAlgorithm; + using cub::LookbackDelayPolicy; + using cub::detail::type_t; + using cub::detail::op_kind_t; + // Re-export all enum values for convenience + using cub::BLOCK_LOAD_DIRECT; + using cub::BLOCK_LOAD_VECTORIZE; + using cub::BLOCK_LOAD_WARP_TRANSPOSE; + using cub::BLOCK_LOAD_WARP_TRANSPOSE_TIMESLICED; + using cub::BLOCK_STORE_DIRECT; + using cub::BLOCK_STORE_WARP_TRANSPOSE; + using cub::BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED; + using cub::BLOCK_REDUCE_RAKING; + using cub::BLOCK_REDUCE_WARP_REDUCTIONS; + using cub::BLOCK_SCAN_RAKING; + using cub::BLOCK_SCAN_WARP_SCANS; + using cub::LOAD_DEFAULT; + using cub::LOAD_CA; + using cub::LOAD_CG; + using cub::LOAD_CS; + using cub::LOAD_LDG; +} + +#else +// Standalone definitions when CCCL is not available (for analysis/codegen) + +namespace muh::tuning { + +enum BlockLoadAlgorithm { + BLOCK_LOAD_DIRECT, + BLOCK_LOAD_VECTORIZE, + BLOCK_LOAD_TRANSPOSE, + BLOCK_LOAD_WARP_TRANSPOSE, + BLOCK_LOAD_WARP_TRANSPOSE_TIMESLICED, + BLOCK_LOAD_STRIPED, +}; + +enum BlockStoreAlgorithm { + BLOCK_STORE_DIRECT, + BLOCK_STORE_WARP_TRANSPOSE, + BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED, + BLOCK_STORE_STRIPED, +}; + +enum BlockReduceAlgorithm { + BLOCK_REDUCE_RAKING, + BLOCK_REDUCE_RAKING_COMMUTATIVE_ONLY, + BLOCK_REDUCE_WARP_REDUCTIONS, + BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC, +}; + +enum BlockScanAlgorithm { + BLOCK_SCAN_RAKING, + BLOCK_SCAN_RAKING_MEMOIZE, + BLOCK_SCAN_WARP_SCANS, +}; + +enum CacheLoadModifier { + LOAD_DEFAULT, + LOAD_CA, + LOAD_CG, + LOAD_CS, + LOAD_CV, + LOAD_LDG, +}; + +enum class LookbackDelayAlgorithm { + no_delay, + fixed_delay, + exponential_backoff, + exponential_backoff_jitter, + exponential_backoff_jitter_window, + exponential_backon_jitter_window, + exponential_backon_jitter, + exponential_backon, +}; + +struct LookbackDelayPolicy { + LookbackDelayAlgorithm kind; + unsigned int delay; + unsigned int l2_write_latency; +}; + +enum class type_t { + boolean, int8, int16, int32, int64, int128, + uint8, uint16, uint32, uint64, uint128, + float32, float64, other +}; + +enum class op_kind_t { plus, min, max, other }; + +} // namespace muh::tuning +#endif // __has_include() + +namespace muh::tuning { + +/// Memory-bound scaling: given nominal params for 4-byte types, +/// scale items_per_thread inversely with actual type size +/// to keep shared memory footprint constant. +/// Directly mirrors cub::detail::MemBoundScaling. +struct scaled_params { + int threads_per_block; + int items_per_thread; +}; + +constexpr scaled_params scale_mem_bound( + int nominal_threads, int nominal_4b_items, int type_size) { + int items = (nominal_4b_items * 4) / type_size; + if (items < 1) items = 1; + if (items > nominal_4b_items) items = nominal_4b_items; + return {nominal_threads, items}; +} + +} // namespace muh::tuning diff --git a/muh/include/muh/tuning/tuning_batch_memcpy.cuh b/muh/include/muh/tuning/tuning_batch_memcpy.cuh new file mode 100644 index 00000000..a1449556 --- /dev/null +++ b/muh/include/muh/tuning/tuning_batch_memcpy.cuh @@ -0,0 +1,63 @@ +// muh/include/muh/tuning/tuning_batch_memcpy.cuh — BI-V100 batch memcpy tuning +// +// Mirrors: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh +// +// vllm impact: KV cache block copy between GPU memory regions +// Competition weight: Cache TPS × 0.56 + +#pragma once + +#include "muh/hardware.cuh" +#include "muh/tuning/common.cuh" + +namespace muh::tuning::batch_memcpy { + +/// Batch memcpy policy +struct BatchMemcpyPolicy { + int threads_per_block; + 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. +// ============================================================ + +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}; +}; + +// ============================================================ +// policy_selector +// ============================================================ + +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}}; + } +}; + +} // namespace muh::tuning::batch_memcpy diff --git a/muh/include/muh/tuning/tuning_for.cuh b/muh/include/muh/tuning/tuning_for.cuh new file mode 100644 index 00000000..34db9059 --- /dev/null +++ b/muh/include/muh/tuning/tuning_for.cuh @@ -0,0 +1,51 @@ +// muh/include/muh/tuning/tuning_for.cuh — BI-V100 for-each tuning +// +// Mirrors: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_for.cuh +// +// vllm impact: RoPE position encoding, simple elementwise kernels +// Competition weight: contributes to Output TPS + +#pragma once + +#include "muh/hardware.cuh" +#include "muh/tuning/common.cuh" + +namespace muh::tuning::for_each { + +/// For-each policy +struct ForPolicy { + int threads_per_block; + int items_per_thread; +}; + +// ============================================================ +// BI-V100 tuning +// +// CCCL reference from tuning_for.cuh: +// threads_per_block: 256 default, can be runtime-determined if set <1 +// items_per_thread: typically 1-4 for simple elementwise +// +// The for_each kernel is extremely simple — it's a parallel_for +// with no shared memory, no reduction, no scan. Tuning is purely +// about occupancy (threads × items = tile_size). +// ============================================================ + +struct bi100_default { + static constexpr int threads = 256; + static constexpr int items = 4; +}; + +// ============================================================ +// policy_selector +// ============================================================ + +struct policy_selector { + constexpr ForPolicy operator()(const hardware_capability& hw) const { + if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) { + return {bi100_default::threads, bi100_default::items}; + } + return {256, 4}; + } +}; + +} // namespace muh::tuning::for_each diff --git a/muh/include/muh/tuning/tuning_reduce.cuh b/muh/include/muh/tuning/tuning_reduce.cuh new file mode 100644 index 00000000..da6d3a53 --- /dev/null +++ b/muh/include/muh/tuning/tuning_reduce.cuh @@ -0,0 +1,152 @@ +// 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) + +#pragma once + +#include "muh/hardware.cuh" +#include "muh/tuning/common.cuh" + +namespace muh::tuning::reduce { + +/// Policy for a single reduction pass (mirrors cub::ReducePassPolicy) +struct ReducePassPolicy { + int threads_per_block; + int items_per_thread; + int vec_size; + BlockReduceAlgorithm reduce_algorithm; + CacheLoadModifier load_modifier; +}; + +/// Full reduction policy (mirrors cub::ReducePolicy) +struct ReducePolicy { + ReducePassPolicy multi_tile; + ReducePassPolicy single_tile; +}; + +// ============================================================ +// BI-V100 tuning values +// Status: PENDING BENCHMARK +// +// These are initialized from CCCL SM90/SM100 reference values. +// Must be replaced with actual BI-V100 benchmark results. +// +// 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 +// ============================================================ + +/// BI-V100 tuning for sum(float32), 4-byte offset +struct bi100_float32_plus_o4 { + // Benchmark annotation format: ipt_N.tpb_M.ipv_K + // BI-V100: TBD — using SM100 reference as starting point + 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 + 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 + 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 + static constexpr int items = 15; + static constexpr int threads = 512; + static constexpr int items_per_vec_load = 1; +}; + +/// BI-V100 default fallback +struct bi100_default { + static constexpr int items = 16; + static constexpr int threads = 256; + static constexpr int items_per_vec_load = 4; +}; + +// ============================================================ +// 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 +// ============================================================ + +struct policy_selector { + type_t accum_t; + op_kind_t operation_t; + int offset_size; + int accum_size; + + /// Dispatch for BI-V100 + constexpr ReducePolicy operator()(const hardware_capability& hw) const { + // Only tuned for sum currently (matching CCCL's approach) + 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); + 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); + 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); + 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); + ReducePassPolicy rp{t, i, bi100_int64_plus_o8::items_per_vec_load, + BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG}; + return {rp, rp}; + } + } + + // Fallback: 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}; + } +}; + +} // namespace muh::tuning::reduce diff --git a/muh/include/muh/tuning/tuning_scan.cuh b/muh/include/muh/tuning/tuning_scan.cuh new file mode 100644 index 00000000..fca3a537 --- /dev/null +++ b/muh/include/muh/tuning/tuning_scan.cuh @@ -0,0 +1,293 @@ +// muh/include/muh/tuning/tuning_scan.cuh — BI-V100 scan tuning +// +// Mirrors: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_scan.cuh +// This is the most complex tuning file in CCCL (900+ lines for NVIDIA). +// +// vllm impact: Prefix scan in paged attention block table lookup +// Competition weight: Input TPS × 2.799 + +#pragma once + +#include "muh/hardware.cuh" +#include "muh/tuning/common.cuh" + +namespace muh::tuning::scan { + +/// Lookback scan policy (mirrors cub::ScanLookbackPolicy) +struct ScanLookbackPolicy { + int threads_per_block; + int items_per_thread; + BlockLoadAlgorithm load_algorithm; + CacheLoadModifier load_modifier; + BlockStoreAlgorithm store_algorithm; + BlockScanAlgorithm scan_algorithm; + LookbackDelayPolicy lookback_delay; +}; + +/// Lookahead scan policy (mirrors cub::ScanLookaheadPolicy) +struct ScanLookaheadPolicy { + int reduce_and_scan_warps; + int items_per_thread; + int lookahead_items_per_thread; + int lookahead_stages; + int block_idx_stages; +}; + +/// Full scan policy +enum class ScanAlgorithm { lookback, lookahead }; + +struct ScanPolicy { + ScanAlgorithm algorithm; + ScanLookbackPolicy lookback; + ScanLookaheadPolicy lookahead; +}; + +// ============================================================ +// BI-V100 tuning values +// +// CCCL reference from tuning_scan.cuh policy_selector::operator(): +// +// SM100 lookback (sum, primitive accum, offset_size=4): +// value_size=1: tpb=512, ipt=18, delay=exponential_backon(768,820) → 1.189x +// value_size=2: tpb=512, ipt=13, delay=exponential_backon(1384,720) → 1.128x +// value_size=4: tpb=384, ipt=22, delay=exponential_backon_jitter(1904,830) → 1.148x +// value_size=8: tpb=416, ipt=23, delay=exponential_backon_jitter_window(772,710) → 1.089x +// +// SM100 lookahead: +// value_size=1: warps=4, ipt=160-1, lai=8 +// value_size=2: warps=6, ipt=96-1, lai=2 +// value_size=4: float→warps=4,ipt=88-1,lai=3; int→warps=4,ipt=80-1,lai=3 +// value_size=8: warps=2, ipt=88-1, lai=5 +// value_size=16: warps=5, ipt=16-1, lai=8 +// ============================================================ + +// --- Lookback tunings for BI-V100 --- + +struct bi100_lookback_1B_o4 { + // SM100 ref: ipt_18.tpb_512.ns_768.dcid_7.l2w_820 → 1.189x + static constexpr int threads = 512; + static constexpr int items = 18; + static constexpr LookbackDelayPolicy delay = { + LookbackDelayAlgorithm::exponential_backon, 768, 820}; + 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_2B_o4 { + // SM100 ref: ipt_13.tpb_512.ns_1384.dcid_7.l2w_720 → 1.128x + static constexpr int threads = 512; + static constexpr int items = 13; + static constexpr LookbackDelayPolicy delay = { + LookbackDelayAlgorithm::exponential_backon, 1384, 720}; + 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_4B_o4 { + // SM100 ref: ipt_22.tpb_384.ns_1904.dcid_6.l2w_830 → 1.148x + static constexpr int threads = 384; + static constexpr int items = 22; + static constexpr LookbackDelayPolicy delay = { + LookbackDelayAlgorithm::exponential_backon_jitter, 1904, 830}; + 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_4B_o8 { + // SM100 ref: ipt_19.tpb_416.ns_956.dcid_7.l2w_550 → 1.146x + static constexpr int threads = 416; + static constexpr int items = 19; + static constexpr LookbackDelayPolicy delay = { + LookbackDelayAlgorithm::exponential_backon, 956, 550}; + 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 + static constexpr int threads = 416; + static constexpr int items = 23; + static constexpr LookbackDelayPolicy delay = { + LookbackDelayAlgorithm::exponential_backon_jitter_window, 772, 710}; + 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 + static constexpr int threads = 320; + static constexpr int items = 22; + static constexpr LookbackDelayPolicy delay = { + LookbackDelayAlgorithm::exponential_backoff, 328, 965}; + 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; +}; + +// --- Lookahead tunings for BI-V100 --- + +struct bi100_lookahead_1B { + // SM100 ref: wrps_4.lbi_8.ipt_160 → 1.264x + static constexpr int warps = 4; + static constexpr int items = 159; // 160-1 + static constexpr int lookahead_items = 8; +}; + +struct bi100_lookahead_2B { + // SM100 ref: wrps_6.lbi_2.ipt_96 → 1.168x + static constexpr int warps = 6; + static constexpr int items = 95; // 96-1 + static constexpr int lookahead_items = 2; +}; + +struct bi100_lookahead_4B { + // SM100 ref (int): wrps_4.lbi_3.ipt_80 → 1.019x + static constexpr int warps = 4; + static constexpr int items = 79; // 80-1 + static constexpr int lookahead_items = 3; +}; + +struct bi100_lookahead_4B_float { + // SM100 ref (float32): wrps_4.lbi_3.ipt_88 → 1.047x + static constexpr int warps = 4; + static constexpr int items = 87; // 88-1 + static constexpr int lookahead_items = 3; +}; + +struct bi100_lookahead_8B { + // SM100 ref: wrps_2.lbi_5.ipt_88 → 1.086x + static constexpr int warps = 2; + static constexpr int items = 87; // 88-1 + static constexpr int lookahead_items = 5; +}; + +struct bi100_lookahead_16B { + // SM100 ref: wrps_5.lbi_8.ipt_16 → 1.160x + static constexpr int warps = 5; + static constexpr int items = 15; // 16-1 + static constexpr int lookahead_items = 8; +}; + +// --- Lookback default fallback --- + +struct bi100_lookback_default { + static constexpr int threads = 128; + static constexpr int items = 15; + static constexpr LookbackDelayPolicy delay = { + LookbackDelayAlgorithm::fixed_delay, 350, 450}; + 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; +}; + +// ============================================================ +// policy_selector +// ============================================================ + +struct policy_selector { + int input_value_size; + int accum_size; + int offset_size; + type_t input_type; + type_t accum_type; + op_kind_t operation_t; + bool is_primitive_accum; + + /// Get the best lookback policy for BI-V100 + constexpr ScanLookbackPolicy get_lookback(const hardware_capability& hw) const { + if (!hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) + goto fallback; + + if (operation_t == op_kind_t::plus && is_primitive_accum) { + if (offset_size == 4) { + switch (input_value_size) { + case 1: return {bi100_lookback_1B_o4::threads, bi100_lookback_1B_o4::items, + bi100_lookback_1B_o4::load_algo, bi100_lookback_1B_o4::load_mod, + bi100_lookback_1B_o4::store_algo, BLOCK_SCAN_WARP_SCANS, + bi100_lookback_1B_o4::delay}; + case 2: return {bi100_lookback_2B_o4::threads, bi100_lookback_2B_o4::items, + bi100_lookback_2B_o4::load_algo, bi100_lookback_2B_o4::load_mod, + bi100_lookback_2B_o4::store_algo, BLOCK_SCAN_WARP_SCANS, + bi100_lookback_2B_o4::delay}; + case 4: return {bi100_lookback_4B_o4::threads, bi100_lookback_4B_o4::items, + bi100_lookback_4B_o4::load_algo, bi100_lookback_4B_o4::load_mod, + bi100_lookback_4B_o4::store_algo, BLOCK_SCAN_WARP_SCANS, + bi100_lookback_4B_o4::delay}; + case 8: return {bi100_lookback_8B_o4::threads, bi100_lookback_8B_o4::items, + bi100_lookback_8B_o4::load_algo, bi100_lookback_8B_o4::load_mod, + bi100_lookback_8B_o4::store_algo, BLOCK_SCAN_WARP_SCANS, + bi100_lookback_8B_o4::delay}; + } + } else if (offset_size == 8) { + switch (input_value_size) { + case 4: return {bi100_lookback_4B_o8::threads, bi100_lookback_4B_o8::items, + bi100_lookback_4B_o8::load_algo, bi100_lookback_4B_o8::load_mod, + bi100_lookback_4B_o8::store_algo, BLOCK_SCAN_WARP_SCANS, + bi100_lookback_4B_o8::delay}; + case 8: return {bi100_lookback_8B_o8::threads, bi100_lookback_8B_o8::items, + bi100_lookback_8B_o8::load_algo, bi100_lookback_8B_o8::load_mod, + bi100_lookback_8B_o8::store_algo, BLOCK_SCAN_WARP_SCANS, + bi100_lookback_8B_o8::delay}; + } + } + } + + fallback: + return {bi100_lookback_default::threads, bi100_lookback_default::items, + bi100_lookback_default::load_algo, bi100_lookback_default::load_mod, + bi100_lookback_default::store_algo, BLOCK_SCAN_WARP_SCANS, + bi100_lookback_default::delay}; + } + + /// Get the best lookahead policy for BI-V100 + constexpr ScanLookaheadPolicy get_lookahead(const hardware_capability& hw) const { + // Lookahead requires specific hardware features (pipeline stages, etc.) + // BI-V100 support is TBD — if not available, caller falls back to lookback + if (!hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) + return {4, 63, 4, 2, -1}; // conservative default + + if (is_primitive_accum) { + switch (input_value_size) { + case 1: return {bi100_lookahead_1B::warps, bi100_lookahead_1B::items, + bi100_lookahead_1B::lookahead_items, 2, -1}; + case 2: return {bi100_lookahead_2B::warps, bi100_lookahead_2B::items, + bi100_lookahead_2B::lookahead_items, 2, -1}; + case 4: + if (input_type == type_t::float32) + return {bi100_lookahead_4B_float::warps, bi100_lookahead_4B_float::items, + bi100_lookahead_4B_float::lookahead_items, 2, -1}; + return {bi100_lookahead_4B::warps, bi100_lookahead_4B::items, + bi100_lookahead_4B::lookahead_items, 2, -1}; + case 8: return {bi100_lookahead_8B::warps, bi100_lookahead_8B::items, + bi100_lookahead_8B::lookahead_items, 2, -1}; + case 16: return {bi100_lookahead_16B::warps, bi100_lookahead_16B::items, + bi100_lookahead_16B::lookahead_items, 2, -1}; + } + } + + // Fallback lookahead + int default_items = (256 / (input_value_size == 2 ? 2 : accum_size)) - 1; + if (default_items < 1) default_items = 1; + int lai = accum_size == 2 ? 3 : 4; + return {4, default_items, lai, 2, -1}; + } + + /// Main dispatch — matches CCCL's operator()(cuda::compute_capability) + constexpr ScanPolicy operator()(const hardware_capability& hw) const { + // Try lookahead first (if hardware supports it) + // TODO: add can_use_lookahead check once BI-V100 pipeline support is confirmed + auto lookahead = get_lookahead(hw); + + // For now, default to lookback (safer, works on all hardware) + auto lookback = get_lookback(hw); + + return {ScanAlgorithm::lookback, lookback, lookahead}; + } +}; + +} // namespace muh::tuning::scan diff --git a/muh/include/muh/tuning/tuning_topk.cuh b/muh/include/muh/tuning/tuning_topk.cuh new file mode 100644 index 00000000..1ebebb3d --- /dev/null +++ b/muh/include/muh/tuning/tuning_topk.cuh @@ -0,0 +1,105 @@ +// muh/include/muh/tuning/tuning_topk.cuh — BI-V100 top-k tuning +// +// Mirrors: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_topk.cuh +// +// vllm impact: Top-k / top-p sampling in decode stage +// Competition weight: Output TPS × 16.796 (highest priority, tied with reduce) + +#pragma once + +#include "muh/hardware.cuh" +#include "muh/tuning/common.cuh" + +namespace muh::tuning::topk { + +/// Top-k policy (mirrors cub's topk_policy) +struct TopkPolicy { + int threads_per_block; + int items_per_thread; + BlockLoadAlgorithm load_algorithm; + BlockScanAlgorithm scan_algorithm; + int bits_per_pass; +}; + +// ============================================================ +// BI-V100 tuning values +// +// 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 +// +// 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 +// ============================================================ + +/// 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; +}; + +// ============================================================ +// policy_selector +// ============================================================ + +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}; + } + } + + // Fallback + int items = (4 * 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, + calc_bits_per_pass(key_size)}; + } +}; + +} // namespace muh::tuning::topk diff --git a/muh/include/muh/tuning/tuning_transform.cuh b/muh/include/muh/tuning/tuning_transform.cuh new file mode 100644 index 00000000..f08d0970 --- /dev/null +++ b/muh/include/muh/tuning/tuning_transform.cuh @@ -0,0 +1,99 @@ +// muh/include/muh/tuning/tuning_transform.cuh — BI-V100 transform tuning +// +// Mirrors: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_transform.cuh +// +// vllm impact: Activation functions (SiLU, GELU), RMSNorm, residual add +// Competition weight: Output TPS × 16.796 + +#pragma once + +#include "muh/hardware.cuh" +#include "muh/tuning/common.cuh" + +namespace muh::tuning::transform { + +/// Bulk transform policy (for contiguous input/output) +struct BulkPolicy { + int threads_per_block; + int items_per_thread; + int vec_size; // vectorization width +}; + +/// No-input policy (for fill/memset operations) +struct FillPolicy { + int threads_per_block; + int items_per_thread_no_input; +}; + +/// Full transform policy +struct TransformPolicy { + BulkPolicy bulk; + FillPolicy fill; +}; + +// ============================================================ +// BI-V100 tuning values +// +// 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. +// +// 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 +// ============================================================ + +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 + + 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 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; + + // 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; + } + + return { + {bi100_bulk_default::threads, items, vec_size}, + {bi100_fill_default::threads, bi100_fill_default::items} + }; + } + + // Fallback + return { + {bi100_bulk_default::threads, bi100_bulk_default::items, bi100_bulk_default::vec_size}, + {bi100_fill_default::threads, bi100_fill_default::items} + }; + } +}; + +} // namespace muh::tuning::transform