dylanyunlong/project_6
1
0
Fork 0
Code Issues Pull Requests Actions Projects Releases Wiki Activity

[MUH] Add C++/CUDA tuning headers — the real muh, not Python wrappers

Browse Source 
The core of muh is now C++ headers that mirror CCCL's tuning architecture:

muh/include/muh/
├── hardware.cuh              — hardware_capability descriptor (replaces cuda::compute_capability)
├── muh.cuh                   — top-level include + scoring formula
└── tuning/
    ├── common.cuh            — shared types, compatible with CCCL's common.cuh
    ├── tuning_reduce.cuh     — P0: attention reduction (5 type specializations)
    ├── tuning_topk.cuh       — P0: sampling top-k/top-p (2B/4B key specializations)
    ├── tuning_scan.cuh       — P0: prefix scan (6 lookback + 6 lookahead specializations)
    ├── tuning_transform.cuh  — P1: activation elementwise (SiLU/GELU/RMSNorm)
    ├── tuning_batch_memcpy.cuh — P1: KV cache block copy
    └── tuning_for.cuh        — P2: RoPE position encoding

Architecture:
  - Each tuning header has a policy_selector struct with operator()(hardware_capability)
  - Dispatches on muh::hardware_capability instead of cuda::compute_capability
  - bi100_* structs hold per-type tuning values (initialized from CCCL SM100 reference)
  - When CCCL headers are available, re-exports their enum types
  - When standalone, provides compatible enum definitions

Python files (extract.py, parse.py, gen_yaml.py, gen_patch.py) remain as tooling.
The C++ headers are what actually gets compiled into the vllm binary.
This commit is contained in:
dylanyunlon
2026-07-30 14:01:07 +00:00
parent 9b21a13119
commit 5f880bb279
9 changed files with 1029 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

62
muh/include/muh/hardware.cuh Normal file
View File

@@ -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

65
muh/include/muh/muh.cuh Normal file
View File

@@ -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 <muh/muh.cuh>
//
// 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

139
muh/include/muh/tuning/common.cuh Normal file
View File

@@ -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(<cub/config.cuh>)
#include <cub/block/block_load.cuh>
#include <cub/block/block_store.cuh>
#include <cub/block/block_reduce.cuh>
#include <cub/block/block_scan.cuh>
#include <cub/device/dispatch/tuning/common.cuh>
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(<cub/config.cuh>)
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

63
muh/include/muh/tuning/tuning_batch_memcpy.cuh Normal file
View File

@@ -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

51
muh/include/muh/tuning/tuning_for.cuh Normal file
View File

@@ -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

152
muh/include/muh/tuning/tuning_reduce.cuh Normal file
View File

@@ -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 <geo> <min> <avg> <max>
// 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

293
muh/include/muh/tuning/tuning_scan.cuh Normal file
View File

@@ -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

105
muh/include/muh/tuning/tuning_topk.cuh Normal file
View File

@@ -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

99
muh/include/muh/tuning/tuning_transform.cuh Normal file
View File

@@ -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