12ad7a3190
每个文件都是直接 cat 读完整 CCCL 源码后理解全部参数语义，
然后用大模型生成 BI-V100 适配版本。不使用 grep/sed/批量脚本。
segmented_sort.cuh (46→189 lines):
- 三层策略完整移植: large(RadixSort), medium(SubWarpMergeSort 16T), small(SubWarpMergeSort 2-8T)
- SM86 tuning: radix_bits=key>1?6:4, scale_reg_bound(256,23)
- BI-V100 SMEM cap for all three tiers
merge_sort.cuh (43→83 lines):
- SM50{256,11} SM52{512,15} SM60+{256,17} 三代参数
- nominal_4b_items_to_items scaling + unroll flag
merge.cuh (55→89 lines):
- SM52/SM60/SM80/SM90/SM100 五代参数
- bulk_copy=false (BI-V100 无 cp.async.bulk)
adjacent_difference.cuh (46→77 lines):
- nominal_8b_items_to_items(7) scaling
- may_alias → LOAD_CA vs LOAD_LDG
batch_memcpy.cuh (86→95 lines):
- small{128T,4buf,8B} + large{256T,32B} 双策略
- prefer_pow2_bits=false (SM70+)
find.cuh (32→39 lines):
- scale_mem_bound(128,16) + vec_size=4
find_bound_sorted_values.cuh (33→47 lines):
- SM80+: {512, N4B(15)} / SM60+: {256} / SM50: LOAD_LDG
40 lines
1.1 KiB
Plaintext
40 lines
1.1 KiB
Plaintext
// muh/include/muh/tuning/tuning_find.cuh — BI-V100
|
|
//
|
|
// Mirrors: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_find.cuh
|
|
// CCCL source: 90 lines. Single policy for all CC:
|
|
// {scale_mem_bound(128, 16, input_size), vec_size=4, LOAD_LDG}
|
|
//
|
|
// vllm relevance: DeviceFind::FindIf for locating EOS tokens, stop sequences,
|
|
// and special token positions in output sequences.
|
|
//
|
|
// No SMEM usage (vectorized global memory scan), BI-V100 safe.
|
|
|
|
#pragma once
|
|
|
|
#include "muh/hardware.cuh"
|
|
#include "muh/tuning/common.cuh"
|
|
|
|
namespace muh::tuning::find {
|
|
|
|
struct FindIfPolicy {
|
|
int threads_per_block;
|
|
int items_per_thread;
|
|
int vec_size;
|
|
CacheLoadModifier load_modifier;
|
|
};
|
|
|
|
// CCCL policy_selector (all CC):
|
|
// scale_mem_bound(128, 16, input_type_size) → threads, items
|
|
// vec_size = 4, load_modifier = LOAD_LDG
|
|
|
|
struct policy_selector {
|
|
int input_type_size;
|
|
|
|
constexpr FindIfPolicy operator()(const hardware_capability& hw) const {
|
|
auto [items, threads] = scale_mem_bound(128, 16, input_type_size);
|
|
return FindIfPolicy{threads, items, 4, LOAD_LDG};
|
|
}
|
|
};
|
|
|
|
} // namespace muh::tuning::find
|