88db0ed89c
tuning_reduce.cuh (201→311 lines): - Add accum_size=1/2/16 branches (int8, bfloat16, int128) - Add min/max op dispatch (same params as plus for BI-V100) - SM=16 tile maximization: det_float32 tile 11648→49152 (23%→100% SMEM) - SM=16 tile maximization: det_float64 tile 11264→49152 (23%→100% SMEM) - Add float32_o8, int64_o4/o8 variants with vec_size dispatch - Increase float32 items 16→24 (32768→49152, fill SMEM for fewer CTAs) tuning_scan.cuh: - Fix 1B tile from 9216→16384 (19%→33% SMEM, scan needs 2x buffer) - Fix 2B tile from 13312→24576 (27%→100% SMEM with double buffer) - Fix 8B_o4 tile: threads 416→384 for warp alignment, items 14→16 - Update header comments with confirmed SM=16 hardware profile - Document lookback delay heuristic for L2=6MB tuning_transform.cuh (128→168 lines): - CRITICAL: bytes_in_flight 16KB→32KB (was based on 900/50=18 GB/s, actual is 900/16=56 GB/s — 3× error) - Add full PrefetchPolicy struct matching CCCL upstream - Add AsyncCopyPolicy with BI-V100 fallback (no cp.async support) - Document CCCL cc_to_min_bytes_in_flight reference values - Add vec_size calculation from element size (16-byte vector loads) - Cap items_per_thread at 32 to prevent register pressure hardware.cuh: - Add SMEM 48KB vs 32KB disambiguation note
63 lines
2.5 KiB
Plaintext
63 lines
2.5 KiB
Plaintext
// 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 vs _custom_ops.py's 32KB claim
|
|
.max_registers_per_thread = 255,
|
|
.l2_cache_size_bytes = 6 * 1024 * 1024, // 6 MiB, TBD
|
|
.memory_bandwidth_gbps = 900, // Confirmed: 1200MHz mem clock // TBD
|
|
.sm_count = 16, // CONFIRMED: ixsmi shows 16 SMs per BI-V100 // 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
|