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

[MUH] Fix 7 structural discrepancies vs CCCL — read source, not grep

Browse Source 
Fixes found by reading all 17 muh files + 6 CCCL counterpart
policy_selectors as full source code input:

1. topk: BLOCK_LOAD_DIRECT → BLOCK_LOAD_VECTORIZE (CCCL SM90+ uses
   VECTORIZE). bits_per_pass was wrong (muh: ks<=4→9, CCCL: ks>=2→11).
   items now computed dynamically (4*4/key_size) not hardcoded.

2. reduce: added determinism dispatch — three modes matching CCCL:
   gpu_to_gpu (BLOCK_REDUCE_RAKING, vec_size=1, LOAD_DEFAULT),
   run_to_run (WARP_REDUCTIONS, LOAD_LDG, default),
   not_guaranteed (WARP_REDUCTIONS_NONDETERMINISTIC).
   Added bi100_det_float32 and bi100_det_float64 tuning structs
   with SM90 benchmark reference values.

3. batch_memcpy: flat single-tier → SmallBuffer+LargeBuffer two-tier
   matching CCCL structure (128 threads small, 256 threads large,
   warp_threshold=128, block_threshold=8192).

4. transform: single BulkPolicy → three-policy structure
   (VectorizedPolicy + AsyncCopyPolicy + PrefetchPolicy) matching CCCL.
   items_per_thread computed from bytes_in_flight / (threads * elem_size).

5. compile_test: 17 checks → 33 checks. Now verifies exact values:
   reduce determinism modes, topk VECTORIZE + bits=11, batch_memcpy
   two-tier thresholds, transform three-policy structure.

6. gen_patch: added fallback extraction for inline policy_selector
   values (topk now generates SAMPLING_BLOCK_SIZE patch).

7. MUH_PROJECT_CHECKPOINT.md: 'PRD设计阶段还没有代码' → actual status.

7 files changed, 413 insertions, 265 deletions.
This commit is contained in:
Claude
2026-07-30 14:37:38 +00:00
parent e02134a3ce
commit c7a63bc2c8
7 changed files with 411 additions and 263 deletions
Download Patch File Download Diff File Expand all files Collapse all files

87
muh/include/muh/tuning/tuning_batch_memcpy.cuh
View File

@@ -4,6 +4,12 @@
//
// vllm impact: KV cache block copy between GPU memory regions
// Competition weight: Cache TPS × 0.56
//
// CCCL structure: two-tier (SmallBuffer handled by single block,
// LargeBuffer by multi-block collaboration). Thresholds:
// warp_level: 128 bytes
// block_level: 8 KiB
// muh must replicate this structure, not flatten it.
#pragma once
@@ -12,51 +18,68 @@
namespace muh::tuning::batch_memcpy {
/// Batch memcpy policy
struct BatchMemcpyPolicy {
/// Small buffer policy: single thread block handles many small buffers
struct SmallBufferPolicy {
int threads_per_block;
int buffers_per_thread;
int bytes_per_thread;
bool prefer_pow2_bits;
int block_level_tile_size;
int warp_level_threshold;
int block_level_threshold;
LookbackDelayPolicy buffer_lookback_delay;
LookbackDelayPolicy block_lookback_delay;
};
// ============================================================
// BI-V100 tuning values
//
// CCCL reference from tuning_batch_memcpy.cuh:
// Default: threads=256
// Buffer lookback delay and block lookback delay are separate —
// they control the two decoupled lookback scans that coordinate
// the batch copy across thread blocks.
//
// For vllm: KV cache copies are typically large contiguous blocks
// (head_dim × num_layers × sizeof(half)), so high throughput matters
// more than latency.
// ============================================================
/// Large buffer policy: multiple blocks collaborate on one large buffer
struct LargeBufferPolicy {
int threads_per_block;
int bytes_per_thread;
};
struct bi100_default {
static constexpr int threads = 256;
static constexpr LookbackDelayPolicy buffer_delay = {
LookbackDelayAlgorithm::fixed_delay, 350, 450};
static constexpr LookbackDelayPolicy block_delay = {
LookbackDelayAlgorithm::fixed_delay, 350, 450};
/// Full batch memcpy policy
struct BatchMemcpyPolicy {
SmallBufferPolicy small_buffer;
LargeBufferPolicy large_buffer;
};
// ============================================================
// policy_selector
// BI-V100 tuning values — from CCCL SM70+ defaults
//
// CCCL policy_selector (all architectures):
// small: 128 threads, 4 bufs/thread, 8 bytes/thread
// prefer_pow2_bits = (cc < 7.0)
// warp_threshold = 128, block_threshold = 8192
// delays = default_delay_constructor_policy(true)
// large: 256 threads, 32 bytes/thread
//
// For BI-V100: start with CCCL defaults.
// The delay policy is arch-sensitive (CCCL uses
// default_delay_constructor_policy which picks fixed_delay for
// primitive types). We use fixed_delay as starting point.
// ============================================================
struct policy_selector {
constexpr BatchMemcpyPolicy operator()(const hardware_capability& hw) const {
if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) {
return {bi100_default::threads,
bi100_default::buffer_delay,
bi100_default::block_delay};
}
// Fallback
return {256,
{LookbackDelayAlgorithm::fixed_delay, 350, 450},
{LookbackDelayAlgorithm::fixed_delay, 350, 450}};
constexpr BatchMemcpyPolicy operator()(const hardware_capability& hw) const {
// BI-V100: assume >= SM70 equivalent (no prefer_pow2_bits)
bool prefer_pow2 = false;
LargeBufferPolicy large{256, 32};
SmallBufferPolicy small{
/* threads_per_block = */ 128,
/* buffers_per_thread = */ 4,
/* bytes_per_thread = */ 8,
/* prefer_pow2_bits = */ prefer_pow2,
/* block_level_tile_size = */ large.threads_per_block * large.bytes_per_thread,
/* warp_level_threshold = */ 128,
/* block_level_threshold = */ 8 * 1024,
/* buffer_lookback_delay = */ {LookbackDelayAlgorithm::fixed_delay, 350, 450},
/* block_lookback_delay = */ {LookbackDelayAlgorithm::fixed_delay, 350, 450},
};
return {small, large};
}
};

117
muh/include/muh/tuning/tuning_reduce.cuh
View File

@@ -28,54 +28,75 @@ struct ReducePolicy {
ReducePassPolicy single_tile;
};
/// Determinism modes (mirrors cuda::execution::determinism::__determinism_t)
enum class determinism_t {
run_to_run, // default: WARP_REDUCTIONS + LOAD_LDG
gpu_to_gpu, // deterministic: RAKING + LOAD_DEFAULT
not_guaranteed, // nondeterministic: WARP_REDUCTIONS_NONDETERMINISTIC
};
// ============================================================
// BI-V100 tuning values
// Status: PENDING BENCHMARK
// BI-V100 tuning values — initialized from CCCL SM100 benchmarks
// Status: PENDING BI-V100 BENCHMARK (values will change)
//
// These are initialized from CCCL SM90/SM100 reference values.
// Must be replaced with actual BI-V100 benchmark results.
// CCCL SM100 benchmark results (from tuning_reduce.cuh):
// float32+o4: ipt=16, tpb=512, ipv=2 → 1.061x geo, 1.167x max
// float64+o4: ipt=16, tpb=640, ipv=1 → 1.018x geo, 1.057x max
// int64+o4: ipt=15, tpb=512, ipv=2 → 1.020x geo, 1.058x max
// int64+o8: ipt=15, tpb=512, ipv=1 → 1.019x geo, 1.057x max
//
// CCCL SM100 reference (for comparison):
// float32, offset_4: ipt=16, tpb=512, ipv=2 → 1.061x speedup
// float64, offset_4: ipt=16, tpb=640, ipv=1 → 1.018x speedup
// int64, offset_4: ipt=15, tpb=512, ipv=2 → 1.020x speedup
// int64, offset_8: ipt=15, tpb=512, ipv=1 → 1.019x speedup
// CCCL SM90 deterministic benchmark results:
// float32: ipt=13, tpb=224 → 1.107x geo, 1.317x max
// float64: ipt=11, tpb=128 → 1.232x geo, 1.582x max
// ============================================================
/// BI-V100 tuning for sum(float32), 4-byte offset
// --- Non-deterministic (default) tunings ---
struct bi100_float32_plus_o4 {
// Benchmark annotation format: ipt_N.tpb_M.ipv_K <geo> <min> <avg> <max>
// BI-V100: TBD — using SM100 reference as starting point
// BI-V100: TBD — SM100 ref: ipt_16.tpb_512.ipv_2 1.061 1.000 1.065 1.167
static constexpr int items = 16;
static constexpr int threads = 512;
static constexpr int items_per_vec_load = 2;
};
/// BI-V100 tuning for sum(float64), 4-byte offset
struct bi100_float64_plus_o4 {
// BI-V100: TBD
// BI-V100: TBD — SM100 ref: ipt_16.tpb_640.ipv_1 1.018 1.000 1.016 1.057
static constexpr int items = 16;
static constexpr int threads = 640;
static constexpr int items_per_vec_load = 1;
};
/// BI-V100 tuning for sum(int64), 4-byte offset
struct bi100_int64_plus_o4 {
// BI-V100: TBD
// BI-V100: TBD — SM100 ref: ipt_15.tpb_512.ipv_2 1.020 1.000 1.018 1.058
static constexpr int items = 15;
static constexpr int threads = 512;
static constexpr int items_per_vec_load = 2;
};
/// BI-V100 tuning for sum(int64), 8-byte offset
struct bi100_int64_plus_o8 {
// BI-V100: TBD
// BI-V100: TBD — SM100 ref: ipt_15.tpb_512.ipv_1 1.019 1.000 1.017 1.057
static constexpr int items = 15;
static constexpr int threads = 512;
static constexpr int items_per_vec_load = 1;
};
/// BI-V100 default fallback
// --- Deterministic tunings (BLOCK_REDUCE_RAKING) ---
// CCCL uses these when determinism == gpu_to_gpu
// vec_size is forced to 1 for deterministic reduction
struct bi100_det_float32 {
// BI-V100: TBD — SM90 ref: ipt_13.tpb_224 1.107 1.010 1.097 1.317
static constexpr int items = 13;
static constexpr int threads = 224;
};
struct bi100_det_float64 {
// BI-V100: TBD — SM86 ref: ipt_11.tpb_128 1.232 1.002 1.245 1.582
static constexpr int items = 11;
static constexpr int threads = 128;
};
/// Fallback for types without specific tuning
struct bi100_default {
static constexpr int items = 16;
static constexpr int threads = 256;
@@ -83,12 +104,12 @@ struct bi100_default {
};
// ============================================================
// policy_selector — the core dispatch functor
// Follows CCCL's exact pattern:
// 1. Accept hardware description
// 2. Match type/op/size to a tuning struct
// 3. Apply mem-bound scaling
// 4. Return ReducePolicy
// policy_selector
//
// Dispatch logic mirrors CCCL's exactly:
// 1. if determinism == gpu_to_gpu → get_deterministic_tuning()
// 2. else → get_two_phase_tuning()
// 3. if determinism == not_guaranteed → override reduce_algorithm
// ============================================================
struct policy_selector {
@@ -96,10 +117,35 @@ struct policy_selector {
op_kind_t operation_t;
int offset_size;
int accum_size;
determinism_t determinism = determinism_t::run_to_run;
/// Dispatch for BI-V100
constexpr ReducePolicy operator()(const hardware_capability& hw) const {
// Only tuned for sum currently (matching CCCL's approach)
/// Deterministic reduction: BLOCK_REDUCE_RAKING, vec_size=1, LOAD_DEFAULT
/// Matches CCCL get_deterministic_tuning()
constexpr ReducePolicy get_deterministic(const hardware_capability& hw) const {
if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) {
if (accum_t == type_t::float32) {
auto [t, i] = scale_mem_bound(bi100_det_float32::threads,
bi100_det_float32::items, accum_size);
ReducePassPolicy rp{t, i, 1, BLOCK_REDUCE_RAKING, LOAD_DEFAULT};
return {rp, rp};
}
if (accum_t == type_t::float64) {
auto [t, i] = scale_mem_bound(bi100_det_float64::threads,
bi100_det_float64::items, accum_size);
ReducePassPolicy rp{t, i, 1, BLOCK_REDUCE_RAKING, LOAD_DEFAULT};
return {rp, rp};
}
}
// Fallback deterministic
auto [t, i] = scale_mem_bound(256, 16, accum_size);
ReducePassPolicy rp{t, i, 1, BLOCK_REDUCE_RAKING, LOAD_DEFAULT};
return {rp, rp};
}
/// Standard two-phase reduction: BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG
/// Matches CCCL get_two_phase_tuning()
constexpr ReducePolicy get_two_phase(const hardware_capability& hw) const {
if (operation_t == op_kind_t::plus &&
hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) {
@@ -140,13 +186,26 @@ struct policy_selector {
}
}
// Fallback: conservative policy
// Fallback: SM60-equivalent conservative policy
auto [t, i] = scale_mem_bound(
bi100_default::threads, bi100_default::items, accum_size);
ReducePassPolicy rp{t, i, bi100_default::items_per_vec_load,
BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG};
return {rp, rp};
}
/// Main dispatch — mirrors CCCL's operator()(compute_capability)
constexpr ReducePolicy operator()(const hardware_capability& hw) const {
if (determinism == determinism_t::gpu_to_gpu) {
return get_deterministic(hw);
}
auto policy = get_two_phase(hw);
if (determinism == determinism_t::not_guaranteed) {
policy.multi_tile.reduce_algorithm = BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC;
}
return policy;
}
};
} // namespace muh::tuning::reduce

118
muh/include/muh/tuning/tuning_topk.cuh
View File

@@ -4,6 +4,18 @@
//
// vllm impact: Top-k / top-p sampling in decode stage
// Competition weight: Output TPS × 16.796 (highest priority, tied with reduce)
//
// CCCL SM90+ policy_selector (the ground truth):
// bits_per_pass = calc_bits_per_pass(key_size):
// key_size 1 → 8
// key_size 2 → 8 (but note: CCCL default returns 11 for 2/4/8,
// only key_size=1 returns 8. See switch below.)
// key_size 4 → 11
// key_size 8 → 11
// items_per_thread = max(1, 4 * 4 / key_size) // 16 bytes per thread
// threads_per_block = 512
// load_algorithm = BLOCK_LOAD_VECTORIZE (NOT BLOCK_LOAD_DIRECT)
// scan_algorithm = BLOCK_SCAN_WARP_SCANS
#pragma once
@@ -12,7 +24,7 @@
namespace muh::tuning::topk {
/// Top-k policy (mirrors cub's topk_policy)
/// Top-k policy (mirrors cub::detail::topk::topk_policy)
struct TopkPolicy {
int threads_per_block;
int items_per_thread;
@@ -22,82 +34,68 @@ struct TopkPolicy {
};
// ============================================================
// BI-V100 tuning values
// bits_per_pass calculation — must match CCCL exactly
//
// CCCL reference from tuning_topk.cuh policy_selector:
// The CCCL topk uses radix-based selection with configurable bits_per_pass.
// Key insight: bits_per_pass trades off #passes vs per-pass work.
// For small key_size (1-2B): 4-6 bits optimal
// For large key_size (4-8B): 8-11 bits optimal
// CCCL source (tuning_topk.cuh):
// case 1: default: return 8;
// case 2: case 4: case 8: return 11;
//
// Default: threads=512, items=nominal_4b*4/key_size, bits=calc_bits_per_pass(key_size)
// calc_bits_per_pass: key_size<=2 → 8, key_size<=4 → 9, key_size<=8 → 10, else → 11
// Previous muh version had a wrong mapping:
// key_size<=2 → 8, key_size<=4 → 9, key_size<=8 → 10
// This was WRONG. CCCL's actual function returns 11 for 2/4/8.
// ============================================================
/// BI-V100 topk for 2-byte keys (float16/bfloat16 — most relevant for LLM logits)
struct bi100_topk_2B {
// LLM decode: logits are typically fp16/bf16, so key_size=2
// CCCL default for 2B: bits_per_pass=8, items=4*4/2=8
static constexpr int threads = 512;
static constexpr int items = 8; // nominal_4b_items=4, scaled: 4*4/2=8
static constexpr int bits_per_pass = 8;
static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_DIRECT;
static constexpr BlockScanAlgorithm scan_algo = BLOCK_SCAN_WARP_SCANS;
};
/// BI-V100 topk for 4-byte keys (float32)
struct bi100_topk_4B {
static constexpr int threads = 512;
static constexpr int items = 4; // nominal_4b_items=4, scaled: 4*4/4=4
static constexpr int bits_per_pass = 9;
static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_DIRECT;
static constexpr BlockScanAlgorithm scan_algo = BLOCK_SCAN_WARP_SCANS;
};
/// BI-V100 default fallback
struct bi100_topk_default {
static constexpr int threads = 256;
static constexpr int items = 4;
static constexpr int bits_per_pass = 8;
static constexpr BlockLoadAlgorithm load_algo = BLOCK_LOAD_DIRECT;
static constexpr BlockScanAlgorithm scan_algo = BLOCK_SCAN_WARP_SCANS;
};
constexpr int calc_bits_per_pass(int key_size) {
switch (key_size) {
case 1:
default:
return 8;
case 2:
case 4:
case 8:
return 11;
}
}
// ============================================================
// policy_selector
//
// Matches CCCL's SM90+ path exactly:
// threads = 512
// items = max(1, nominal_4b(4) * 4 / key_size)
// load = BLOCK_LOAD_VECTORIZE
// scan = BLOCK_SCAN_WARP_SCANS
// bits = calc_bits_per_pass(key_size)
//
// BI-V100 values: using SM100 as starting point.
// Once benchmarked on BI-V100, items/threads may diverge.
// ============================================================
struct policy_selector {
int key_size;
static constexpr int calc_bits_per_pass(int ks) {
if (ks <= 2) return 8;
if (ks <= 4) return 9;
if (ks <= 8) return 10;
return 11;
}
constexpr TopkPolicy operator()(const hardware_capability& hw) const {
if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) {
switch (key_size) {
case 1:
case 2:
return {bi100_topk_2B::threads, bi100_topk_2B::items,
bi100_topk_2B::load_algo, bi100_topk_2B::scan_algo,
bi100_topk_2B::bits_per_pass};
case 4:
return {bi100_topk_4B::threads, bi100_topk_4B::items,
bi100_topk_4B::load_algo, bi100_topk_4B::scan_algo,
bi100_topk_4B::bits_per_pass};
}
// SM90+ path from CCCL: 16 bytes per thread
constexpr int nominal_4b_items = 4;
int items = nominal_4b_items * 4 / key_size;
if (items < 1) items = 1;
return {512, items,
BLOCK_LOAD_VECTORIZE, // CCCL uses VECTORIZE, not DIRECT
BLOCK_SCAN_WARP_SCANS,
calc_bits_per_pass(key_size)};
}
// Fallback
int items = (4 * 4) / key_size;
// Fallback: older arch path
constexpr int nominal_4b_items = 4;
int items = nominal_4b_items * 4 / key_size;
if (items < 1) items = 1;
if (items > 4) items = 4;
return {bi100_topk_default::threads, items,
bi100_topk_default::load_algo, bi100_topk_default::scan_algo,
if (items > nominal_4b_items) items = nominal_4b_items;
return {512, items,
BLOCK_LOAD_VECTORIZE,
BLOCK_SCAN_WARP_SCANS,
calc_bits_per_pass(key_size)};
}
};

133
muh/include/muh/tuning/tuning_transform.cuh
View File

@@ -4,6 +4,14 @@
//
// vllm impact: Activation functions (SiLU, GELU), RMSNorm, residual add
// Competition weight: Output TPS × 16.796
//
// CCCL structure: three transform policy types selected by iterator properties:
// 1. TransformVectorizedPolicy — contiguous + trivially_relocatable inputs
// 2. TransformAsyncCopyPolicy — SM90+ bulk copy (cp.async.bulk)
// 3. TransformPrefetchPolicy — fallback when stable_address needed
//
// For vllm: activations are contiguous dense fp16/bf16 tensors.
// → primarily hits the vectorized path.
#pragma once
@@ -12,86 +20,103 @@
namespace muh::tuning::transform {
/// Bulk transform policy (for contiguous input/output)
struct BulkPolicy {
/// Vectorized transform: coalesced loads via vector types
struct VectorizedPolicy {
int threads_per_block;
int items_per_thread;
int vec_size; // vectorization width
int vec_size;
};
/// No-input policy (for fill/memset operations)
/// Async copy transform: SM90+ bulk shared memory copy
struct AsyncCopyPolicy {
int threads_per_block;
int min_items_per_thread;
int store_vec_size;
};
/// Prefetch transform: prefetch-based fallback
struct PrefetchPolicy {
int threads_per_block;
};
/// Fill policy (no input, e.g. memset)
struct FillPolicy {
int threads_per_block;
int items_per_thread_no_input;
int items_per_thread;
};
/// Full transform policy
/// Full transform policy — selected at dispatch time based on iterator properties
struct TransformPolicy {
BulkPolicy bulk;
VectorizedPolicy vectorized;
AsyncCopyPolicy async_copy;
PrefetchPolicy prefetch;
FillPolicy fill;
};
// ============================================================
// BI-V100 tuning values
// BI-V100 tuning
//
// CCCL reference from tuning_transform.cuh:
// Default bulk: threads=256, items=auto, vec_size=auto
// The transform kernel is bandwidth-bound for large elementwise ops.
// Key insight: vec_size should match the hardware's natural vector width.
// For NVIDIA: vec_size is typically 4 (128-bit loads).
// For BI-V100: TBD, likely also 4 or 8.
// CCCL SM90+ vectorized path:
// threads = 256 (SM90) or 128 (SM100)
// items = computed from: max(items_for_vec, items_for_latency)
// items_for_vec = ceil(vec_bytes / min_elem_size)
// items_for_latency = (min_bytes_in_flight) / (threads * elem_size)
// vec_size = auto (power-of-2 aligned to hardware vector width)
//
// items_per_thread is computed as:
// items_for_vec = ceil(vector_bytes / min_elem_size)
// items_for_latency = (latency * bandwidth) / (threads * elem_size)
// items = max(items_for_vec, items_for_latency)
// ============================================================
struct bi100_bulk_default {
static constexpr int threads = 256;
static constexpr int items = 8; // conservative starting point
static constexpr int vec_size = 4; // 128-bit vector loads
};
struct bi100_fill_default {
static constexpr int threads = 256;
static constexpr int items = 2;
};
// ============================================================
// policy_selector
// CCCL SM90+ async_copy path:
// threads = 256 (SM90) or 128 (SM100)
// min_items_per_thread = computed from SMEM capacity
// store_vec_size = auto_ublkcp_store_vec_size(output.value_type_size)
//
// CCCL prefetch:
// threads = 256
//
// For BI-V100: start with SM100 values (128 threads for bulk/async,
// 256 for prefetch). vec_size = 4 (128-bit loads, standard for most GPUs).
// ============================================================
struct policy_selector {
int min_elem_size; // minimum element size across all inputs
int max_elem_size; // maximum element size
int num_inputs; // number of input iterators
int min_elem_size;
int max_elem_size;
int num_inputs;
bool all_contiguous;
bool all_trivially_relocatable;
bool requires_stable_address;
constexpr TransformPolicy operator()(const hardware_capability& hw) const {
if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) {
// For BI-V100: start with CCCL defaults, tune via benchmark
int vec_size = bi100_bulk_default::vec_size;
// Compute vectorization params
int vec_size = 4; // 128-bit, matches CCCL default
// Compute items_per_thread based on vectorization
int items_for_vec = (vec_size * 4) / min_elem_size; // 4 = sizeof(int)
if (items_for_vec < 1) items_for_vec = 1;
// items_for_vec: how many items fit in one vector load
int items_for_vec = (vec_size * 4) / min_elem_size; // 4 = sizeof(int)
if (items_for_vec < 1) items_for_vec = 1;
// Ensure items is a multiple of vec_size
int items = items_for_vec;
if (items % vec_size != 0) {
items = ((items / vec_size) + 1) * vec_size;
}
// items_for_latency: enough items to hide memory latency
// CCCL uses cc_to_min_bytes_in_flight(cc) which is ~48KB for SM90+
// For BI-V100: estimate 48KB in flight, 256 threads
int bytes_in_flight = 48 * 1024;
int items_for_latency = bytes_in_flight / (256 * min_elem_size);
if (items_for_latency < 1) items_for_latency = 1;
return {
{bi100_bulk_default::threads, items, vec_size},
{bi100_fill_default::threads, bi100_fill_default::items}
};
int bulk_items = items_for_vec > items_for_latency ? items_for_vec : items_for_latency;
// Ensure items is a multiple of vec_size for aligned access
if (bulk_items % vec_size != 0) {
bulk_items = ((bulk_items / vec_size) + 1) * vec_size;
}
// Fallback
// BI-V100: 128 threads for bulk (SM100-like), 256 for prefetch
int bulk_threads = hw.at_least(hardware_capability::vendor_t::iluvatar, 100) ? 128 : 256;
return {
{bi100_bulk_default::threads, bi100_bulk_default::items, bi100_bulk_default::vec_size},
{bi100_fill_default::threads, bi100_fill_default::items}
// vectorized
{bulk_threads, bulk_items, vec_size},
// async_copy (BI-V100 may not support cp.async.bulk — conservative)
{bulk_threads, 4, vec_size},
// prefetch
{256},
// fill
{256, 2},
};
}
};