diff --git a/baseline.muh b/baseline.muh index c19500b6..3885bcc3 100644 --- a/baseline.muh +++ b/baseline.muh @@ -1,24 +1,14 @@ -# baseline.muh — Competition reference configuration -# Corresponds to: dev.modelhub.org.cn EngineX-Iluvatar/enginex-vllm-bi100-qwen36 +# baseline.muh — Competition vllm launch configuration # -# This is the starting point. All tuning values are pending BI-V100 benchmarks. -# Child .muh files use 'extends: baseline.muh' to override specific algorithms. - -# --- Hardware description --- -hardware: - name: Iluvatar-BI-V100-50c-200G - gpu_count: 4 - # These need to be confirmed on actual hardware: - warp_size: 32 - max_threads_per_block: 1024 - max_shared_memory_per_block: 49152 - max_registers_per_thread: 255 - l2_cache_size_bytes: 6291456 - memory_bandwidth_gbps: 900 - compute_capability: iluvatar_bi100 +# This file stores ONLY the vllm server launch config. +# Kernel tuning values live in muh/include/muh/tuning/tuning_*.cuh +# as constexpr structs — NOT here. +# +# Pipeline: +# muh/tuning/*.cuh (bi100_* values) → gen_patch.py → vllm kernel patches +# baseline.muh (vllm config) → gen_yaml.py → computility-run.yaml # --- vllm launch configuration --- -# Maps directly to computility-run.yaml command vllm: model_path: /model served_model_name: llm @@ -37,86 +27,7 @@ vllm: reasoning_parser: qwen3 enable_prefix_caching: true -# --- Concurrency --- concurrency: 1 -# --- Environment --- env: VLLM_ENGINE_ITERATION_TIMEOUT_S: 3600 - -# --- Tuning overrides (per CCCL algorithm) --- -# Each key corresponds to a tuning_*.cuh schema in muh/schema/ -# Values are TBD until we run benchmarks on BI-V100 -# -# Priority order (by competition score impact): -# 1. reduce — attention reduction (Output TPS × 16.796) -# 2. topk — sampling top-k/top-p (Output TPS × 16.796) -# 3. scan — prefix scan in paged attention -# 4. transform — activation kernels (SiLU, GELU) -# 5. batch_memcpy — KV cache management (Cache TPS × 0.56) -# 6. for — RoPE position encoding - -tuning: - reduce: - _priority: P0 - _vllm_impact: attention_reduction - _score_weight: Output TPS × 16.796 - # CCCL SM90 reference: threads=128, items=24, vec_size=4 - # CCCL SM100 reference: threads varies by accum_size - threads_per_block: null - items_per_thread: null - vec_size: null - - topk: - _priority: P0 - _vllm_impact: sampling_decode - _score_weight: Output TPS × 16.796 - # CCCL reference: threads=512, items=4, bits_per_pass=11 - threads_per_block: null - items_per_thread: null - bits_per_pass: null - - scan: - _priority: P0 - _vllm_impact: paged_attention_prefix_scan - _score_weight: Input TPS × 2.799 - # CCCL SM90 lookback: threads=128, items=24, delay=fixed(688, 1140) for float32 - # CCCL SM100 lookback: threads=384, items=22, delay=exponential_backon(1904, 830) - # CCCL SM100 lookahead: warps=4, items=80-1, lookahead_items=3 - threads_per_block: null - items_per_thread: null - load_algorithm: null - store_algorithm: null - scan_algorithm: null - - transform: - _priority: P1 - _vllm_impact: activation_elementwise - _score_weight: Output TPS × 16.796 - threads_per_block: null - items_per_thread: null - - batch_memcpy: - _priority: P1 - _vllm_impact: kv_cache_copy - _score_weight: Cache TPS × 0.56 - threads_per_block: null - - for: - _priority: P2 - _vllm_impact: rope_position_encoding - threads_per_block: null - items_per_thread: null - - radix_sort: - _priority: P2 - _vllm_impact: beam_search_token_sort - threads_per_block: null - items_per_thread: null - radix_bits: null - - merge: - _priority: P2 - _vllm_impact: sequence_merging - threads_per_block: null - items_per_thread: null diff --git a/muh/gen_patch.py b/muh/gen_patch.py index 253276ef..ef335634 100644 --- a/muh/gen_patch.py +++ b/muh/gen_patch.py @@ -1,244 +1,228 @@ #!/usr/bin/env python3 -"""muh/gen_patch.py — Generate vllm kernel patches from .muh tuning configuration +"""muh/gen_patch.py — Generate vllm kernel patches from C++ tuning headers -Given a .muh file with tuning overrides for BI-V100, generates unified diff -patches that can be applied to the vllm source tree to inject optimized -kernel parameters. +Reads muh/include/muh/tuning/tuning_*.cuh, extracts bi100_* struct values, +and generates unified diff patches for the vllm source tree. -The key insight: vllm's CUDA kernels (attention, sampling, layernorm) have -hardcoded launch configs. This script generates patches that replace those -hardcodes with values tuned for Iluvatar BI-V100 via CCCL benchmark data. +The previous version read from .muh YAML files. This version reads directly +from C++ headers — single source of truth, no YAML middleman. Usage: - python3 muh/gen_patch.py baseline.muh [-o patches/] [--vllm-root /path/to/vllm] + python3 muh/gen_patch.py [--header-dir muh/include/muh/tuning] [-o patches/] """ +import re import os import sys +import glob import argparse from datetime import datetime -sys.path.insert(0, os.path.dirname(__file__)) -from parse import load_muh + +def extract_bi100_structs(filepath): + """Extract all bi100_* struct constexpr values from a C++ header. + + Returns list of (struct_name, {field: value, ...}) tuples. + """ + with open(filepath, 'r') as f: + content = f.read() + + structs = [] + # Split on struct definitions + # Pattern: struct bi100_xxx { ... }; + pattern = re.compile( + r'struct\s+(bi100_\w+)\s*\{(.*?)\};', + re.DOTALL + ) + + for m in pattern.finditer(content): + name = m.group(1) + body = m.group(2) + fields = {} + + # Extract: static constexpr int threads = 512; + for fm in re.finditer( + r'static\s+constexpr\s+int\s+(\w+)\s*=\s*(\d+)', + body + ): + fields[fm.group(1)] = int(fm.group(2)) + + # Extract: static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_DIRECT; + for fm in re.finditer( + r'static\s+constexpr\s+\w+\s+(\w+)\s*=\s*(\w+)', + body + ): + if fm.group(1) not in fields: # don't overwrite int extractions + fields[fm.group(1)] = fm.group(2) + + # Extract LookbackDelayPolicy: {LookbackDelayAlgorithm::xxx, N, M} + delay_m = re.search( + r'LookbackDelayPolicy\s+\w+\s*=\s*\{\s*' + r'LookbackDelayAlgorithm::(\w+)\s*,\s*(\d+)\s*,\s*(\d+)\s*\}', + body + ) + if delay_m: + fields['delay_algo'] = delay_m.group(1) + fields['delay_ns'] = int(delay_m.group(2)) + fields['delay_l2w'] = int(delay_m.group(3)) + + if fields: + structs.append((name, fields)) + + return structs -# --- Kernel location mapping --- -# Maps CCCL algorithm names to vllm source files and the specific -# constants/defines that control kernel launch parameters. +def algo_from_filename(filepath): + """tuning_reduce.cuh → reduce""" + base = os.path.basename(filepath) + return base.replace('tuning_', '').replace('.cuh', '') -VLLM_KERNEL_MAP = { - "reduce": { - "description": "Attention score reduction in multi-head attention", - "files": [ - "csrc/attention/attention_kernels.cu", - "csrc/attention/paged_attention_v2.cu", - ], - "params": { - "threads_per_block": { - "pattern": "NUM_THREADS", - "default": 128, - "locations": ["#define NUM_THREADS 128"], - }, - "items_per_thread": { - "pattern": "NUM_ITEMS_PER_THREAD", - "default": 8, - }, - "vec_size": { - "pattern": "VEC_SIZE", - "default": 4, - }, - }, - }, - "topk": { - "description": "Top-k / top-p sampling in decode stage", - "files": [ - "csrc/sampling/sampling_kernels.cu", - ], - "params": { - "threads_per_block": { - "pattern": "SAMPLING_BLOCK_SIZE", - "default": 256, - }, - "bits_per_pass": { - "pattern": "RADIX_BITS", - "default": 8, - }, - }, - }, - "scan": { - "description": "Prefix scan in paged attention block table lookup", - "files": [ - "csrc/attention/paged_attention_v1.cu", - ], - "params": { - "threads_per_block": { - "pattern": "SCAN_BLOCK_SIZE", - "default": 128, - }, - }, - }, - "transform": { - "description": "Elementwise activation kernels (SiLU, GELU, RMSNorm)", - "files": [ - "csrc/activation_kernels.cu", - "csrc/layernorm_kernels.cu", - ], - "params": { - "threads_per_block": { - "pattern": "ACTIVATION_BLOCK_SIZE", - "default": 512, - }, - }, - }, - "batch_memcpy": { - "description": "KV cache block copy between GPU memory regions", - "files": [ - "csrc/cache_kernels.cu", - ], - "params": { - "threads_per_block": { - "pattern": "COPY_BLOCK_SIZE", - "default": 256, - }, - }, - }, - "for": { - "description": "Elementwise for-each kernels (position embeddings, rope)", - "files": [ - "csrc/pos_encoding_kernels.cu", - ], - "params": { - "threads_per_block": { - "pattern": "ROPE_BLOCK_SIZE", - "default": 512, - }, - }, - }, + +# --- vllm kernel mapping --- +# Maps (algorithm, struct_field) → (vllm_file, define/variable, context) +# This must be updated when we have access to actual vllm-bi100 source tree. +# For now, these are the known injection points from enginex-vllm-bi100-qwen36. + +VLLM_INJECTION_POINTS = { + ('reduce', 'threads'): [ + ('csrc/attention/attention_kernels.cu', 'NUM_THREADS'), + ('csrc/attention/paged_attention_v2.cu', 'NUM_THREADS'), + ], + ('reduce', 'items'): [ + ('csrc/attention/attention_kernels.cu', 'NUM_ITEMS_PER_THREAD'), + ], + ('reduce', 'items_per_vec_load'): [ + ('csrc/attention/attention_kernels.cu', 'VEC_SIZE'), + ], + ('topk', 'threads'): [ + ('csrc/sampling/sampling_kernels.cu', 'SAMPLING_BLOCK_SIZE'), + ], + ('topk', 'bits_per_pass'): [ + ('csrc/sampling/sampling_kernels.cu', 'RADIX_BITS'), + ], + ('scan', 'threads'): [ + ('csrc/attention/paged_attention_v1.cu', 'SCAN_BLOCK_SIZE'), + ], + ('transform', 'threads'): [ + ('csrc/activation_kernels.cu', 'ACTIVATION_BLOCK_SIZE'), + ('csrc/layernorm_kernels.cu', 'LAYERNORM_BLOCK_SIZE'), + ], + ('batch_memcpy', 'threads'): [ + ('csrc/cache_kernels.cu', 'COPY_BLOCK_SIZE'), + ], + ('for', 'threads'): [ + ('csrc/pos_encoding_kernels.cu', 'ROPE_BLOCK_SIZE'), + ], } -def generate_define_patch(algo, param_name, old_value, new_value, define_name, filepath): - """Generate a unified diff snippet for a #define change.""" - lines = [] - lines.append(f"--- a/{filepath}") - lines.append(f"+++ b/{filepath}") - lines.append(f"@@ -1,1 +1,1 @@") - lines.append(f"-#define {define_name} {old_value}") - lines.append(f"+#define {define_name} {new_value} // muh: tuned for BI-V100 ({algo}.{param_name})") - return "\n".join(lines) - - -def generate_patches(config, vllm_root=None): - """Generate all patches from tuning config.""" - tuning = config.get("tuning", {}) +def generate_patches(header_dir): + """Read all tuning headers, extract bi100 values, generate patches.""" patches = [] summary = [] - for algo, algo_params in tuning.items(): - if not isinstance(algo_params, dict): + headers = sorted(glob.glob(os.path.join(header_dir, 'tuning_*.cuh'))) + if not headers: + print(f"ERROR: No tuning_*.cuh found in {header_dir}", file=sys.stderr) + return [], [] + + for hpath in headers: + algo = algo_from_filename(hpath) + structs = extract_bi100_structs(hpath) + + if not structs: + summary.append(f"SKIP {algo}: no bi100_* structs found") continue - mapping = VLLM_KERNEL_MAP.get(algo) - if mapping is None: - summary.append(f"SKIP {algo}: no vllm kernel mapping defined") - continue + # Use the first non-default struct as the primary tuning + # (default is fallback; prefer the type-specific ones) + primary = None + for name, fields in structs: + if 'default' not in name: + primary = (name, fields) + break + if primary is None: + primary = structs[0] - for param_name, new_value in algo_params.items(): - if param_name.startswith("_"): - continue - if new_value is None: + pname, pfields = primary + summary.append(f"READ {algo}: {pname} → {pfields}") + + for field_name, value in pfields.items(): + key = (algo, field_name) + if key not in VLLM_INJECTION_POINTS: continue - param_spec = mapping.get("params", {}).get(param_name) - if param_spec is None: - continue - - old_value = param_spec.get("default") - define_name = param_spec.get("pattern", param_name.upper()) - - for filepath in mapping.get("files", []): - patch = generate_define_patch( - algo, param_name, old_value, new_value, define_name, filepath + for vllm_file, define_name in VLLM_INJECTION_POINTS[key]: + patch_text = ( + f"--- a/{vllm_file}\n" + f"+++ b/{vllm_file}\n" + f"@@ muh tuning injection @@\n" + f"-// {define_name}: default\n" + f"+#define {define_name} {value} " + f"// muh: from {pname}.{field_name} (tuning_{algo}.cuh)\n" ) patches.append({ - "algo": algo, - "param": param_name, - "file": filepath, - "old": old_value, - "new": new_value, - "diff": patch, + 'algo': algo, + 'struct': pname, + 'field': field_name, + 'value': value, + 'vllm_file': vllm_file, + 'define': define_name, + 'diff': patch_text, }) summary.append( - f"PATCH {filepath}: {define_name} {old_value} → {new_value} " - f"(from {algo}.{param_name})" + f" PATCH {vllm_file}: {define_name} = {value} " + f"(from {pname}.{field_name})" ) return patches, summary def write_patches(patches, out_dir): - """Write patches to individual .patch files.""" + """Write combined patch file.""" os.makedirs(out_dir, exist_ok=True) - # Combined patch - combined_path = os.path.join(out_dir, "muh_bi100_tuning.patch") - with open(combined_path, 'w') as f: + combined = os.path.join(out_dir, 'muh_bi100_tuning.patch') + with open(combined, 'w') as f: f.write(f"# muh kernel tuning patch for Iluvatar BI-V100\n") f.write(f"# Generated: {datetime.now().isoformat()}\n") - f.write(f"# Algorithms patched: {len(set(p['algo'] for p in patches))}\n") - f.write(f"# Total changes: {len(patches)}\n\n") + f.write(f"# Source: muh/include/muh/tuning/tuning_*.cuh bi100_* structs\n") + f.write(f"# Patches: {len(patches)}\n\n") for p in patches: - f.write(p["diff"]) - f.write("\n\n") + f.write(p['diff']) + f.write('\n') - # Per-algorithm patches - by_algo = {} - for p in patches: - by_algo.setdefault(p["algo"], []).append(p) - - for algo, algo_patches in by_algo.items(): - algo_path = os.path.join(out_dir, f"{algo}.patch") - with open(algo_path, 'w') as f: - f.write(f"# muh tuning patch: {algo} for BI-V100\n\n") - for p in algo_patches: - f.write(p["diff"]) - f.write("\n\n") - - return combined_path + return combined def main(): - parser = argparse.ArgumentParser(description="Generate vllm kernel patches from .muh") - parser.add_argument("muh_file", help="Path to .muh file") - parser.add_argument("-o", "--output-dir", default="patches", - help="Output directory for patches (default: patches)") - parser.add_argument("--vllm-root", default=None, - help="Path to vllm source tree (for verification)") - parser.add_argument("--dry-run", action="store_true", - help="Print patches to stdout instead of writing files") - args = parser.parse_args() + p = argparse.ArgumentParser(description='Generate vllm patches from muh C++ headers') + p.add_argument('--header-dir', default='muh/include/muh/tuning', + help='Directory containing tuning_*.cuh headers') + p.add_argument('-o', '--output-dir', default='patches', + help='Output directory for patches') + p.add_argument('--dry-run', action='store_true', + help='Print to stdout instead of writing') + args = p.parse_args() - config = load_muh(args.muh_file) - patches, summary = generate_patches(config, args.vllm_root) + patches, summary = generate_patches(args.header_dir) - print(f"muh gen_patch: {len(patches)} patches from {args.muh_file}\n") + print(f"muh gen_patch: scanned {args.header_dir}\n") for s in summary: print(f" {s}") if not patches: - print("\nNo patches generated. Add tuning overrides to your .muh file.") + print("\nNo patches generated.") return if args.dry_run: - print("\n--- Patches ---\n") + print(f"\n--- {len(patches)} patches ---\n") for p in patches: - print(p["diff"]) - print() + print(p['diff']) else: combined = write_patches(patches, args.output_dir) - print(f"\nWritten to {args.output_dir}/") - print(f"Combined: {combined}") + print(f"\nWritten: {combined}") -if __name__ == "__main__": +if __name__ == '__main__': main() diff --git a/muh/include/CMakeLists.txt b/muh/include/CMakeLists.txt new file mode 100644 index 00000000..1eeafd8e --- /dev/null +++ b/muh/include/CMakeLists.txt @@ -0,0 +1,40 @@ +cmake_minimum_required(VERSION 3.18) +project(muh LANGUAGES CXX) + +# muh is a header-only library +add_library(muh INTERFACE) +target_include_directories(muh INTERFACE ${CMAKE_CURRENT_SOURCE_DIR}) +target_compile_features(muh INTERFACE cxx_std_17) + +# If CCCL is available, link it +if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/../../cccl_upstream/cub/cub/config.cuh") + target_include_directories(muh INTERFACE + ${CMAKE_CURRENT_SOURCE_DIR}/../../cccl_upstream/cub + ${CMAKE_CURRENT_SOURCE_DIR}/../../cccl_upstream/thrust + ${CMAKE_CURRENT_SOURCE_DIR}/../../cccl_upstream/libcudacxx/include + ) + target_compile_definitions(muh INTERFACE MUH_HAS_CCCL=1) +endif() + +# Compile test — verifies all headers parse without errors +# This is a host-only test (no GPU needed) +option(MUH_BUILD_TESTS "Build muh compile tests" ON) + +if(MUH_BUILD_TESTS) + add_executable(muh_compile_test + ${CMAKE_CURRENT_SOURCE_DIR}/../test/compile_test.cpp + ) + target_link_libraries(muh_compile_test PRIVATE muh) + + # If we have a CUDA compiler, also test .cu compilation + include(CheckLanguage) + check_language(CUDA) + if(CMAKE_CUDA_COMPILER) + enable_language(CUDA) + add_executable(muh_cuda_compile_test + ${CMAKE_CURRENT_SOURCE_DIR}/../test/cuda_compile_test.cu + ) + target_link_libraries(muh_cuda_compile_test PRIVATE muh) + set_target_properties(muh_cuda_compile_test PROPERTIES CUDA_STANDARD 17) + endif() +endif() diff --git a/muh/include/muh/tuning/tuning_scan.cuh b/muh/include/muh/tuning/tuning_scan.cuh index fca3a537..e68c525c 100644 --- a/muh/include/muh/tuning/tuning_scan.cuh +++ b/muh/include/muh/tuning/tuning_scan.cuh @@ -200,10 +200,8 @@ struct policy_selector { /// 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 (hw.at_least(hardware_capability::vendor_t::iluvatar, 100) + && 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, @@ -222,6 +220,7 @@ struct policy_selector { 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}; + default: break; } } else if (offset_size == 8) { switch (input_value_size) { @@ -233,11 +232,12 @@ struct policy_selector { 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}; + default: break; } } } - fallback: + // 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, diff --git a/muh/test/compile_test.cpp b/muh/test/compile_test.cpp new file mode 100644 index 00000000..3549da31 --- /dev/null +++ b/muh/test/compile_test.cpp @@ -0,0 +1,140 @@ +// 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 + +#include "muh/muh.cuh" + +#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) + +#define CHECK_TRUE(expr, name) \ + do { \ + if (!(expr)) { \ + std::fprintf(stderr, "FAIL: %s\n", name); \ + failures++; \ + } else { \ + passes++; \ + } \ + } while(0) + +int main() { + 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"); + 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 --- + { + 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, + }; + 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]"); + } + + // --- Test topk policy_selector --- + { + 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]"); + } + + // --- Test scan policy_selector --- + { + using namespace muh::tuning::scan; + auto ps = 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, + }; + 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"); + } + + // --- Test transform policy_selector --- + { + using namespace muh::tuning::transform; + auto ps = policy_selector{ + .min_elem_size = 2, + .max_elem_size = 2, + .num_inputs = 1, + }; + auto policy = ps(hw); + CHECK_NONZERO(policy.bulk.threads_per_block, + "transform.bulk.threads_per_block"); + } + + // --- 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 --- + { + 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"); + } + + // --- Report --- + std::printf("\nmuh compile test: %d passed, %d failed\n", passes, failures); + return failures > 0 ? 1 : 0; +} diff --git a/muh/test/cuda_compile_test.cu b/muh/test/cuda_compile_test.cu new file mode 100644 index 00000000..876b4a2d --- /dev/null +++ b/muh/test/cuda_compile_test.cu @@ -0,0 +1,41 @@ +// muh/test/cuda_compile_test.cu — CUDA compilation test +// +// Verifies muh headers compile under nvcc/clang CUDA mode. +// Does NOT require GPU execution — just compilation. + +#include "muh/muh.cuh" + +__global__ void dummy_kernel() { + // Instantiate policy selectors in device code to verify + // all constexpr paths compile on the device side + auto hw = muh::hardware_capability::bi_v100(); + + // Reduce + auto rp = 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); + (void)rp; + + // Topk + auto tp = muh::tuning::topk::policy_selector{.key_size = 2}(hw); + (void)tp; +} + +int main() { + // Host-side test (same as compile_test.cpp core) + auto hw = muh::target_hw; + + 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); + + printf("CUDA compile test passed: reduce.threads=%d\n", + reduce_policy.multi_tile.threads_per_block); + return 0; +}