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
96 lines
3.3 KiB
Plaintext
96 lines
3.3 KiB
Plaintext
// muh/include/muh/tuning/tuning_batch_memcpy.cuh — BI-V100
|
|
//
|
|
// Mirrors: cccl_upstream/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh
|
|
// CCCL source: 227 lines. DeviceMemcpy (batched copy) tuning for KV cache block copy.
|
|
//
|
|
// CCCL has a two-tier policy:
|
|
// small_buffer: {128 threads, 4 buffers/thread, 8 bytes/thread, prefer_pow2_bits,
|
|
// block_tile=8192, warp_threshold=128, block_threshold=8192,
|
|
// buffer_delay=default, block_delay=default}
|
|
// large_buffer: {256 threads, 32 bytes/thread}
|
|
//
|
|
// vllm relevance: KV cache block copy in paged attention — when sequences are
|
|
// forked (beam search) or compacted, batches of small/medium KV cache blocks
|
|
// need to be copied efficiently. This is the [muh] batch_memcpy kernel.
|
|
//
|
|
// SMEM: small buffer kernel uses shared memory for:
|
|
// - buffer metadata: buffers_per_tile * (src_ptr + dst_ptr + size) = tile * 24B
|
|
// - byte staging: threads * bytes_per_thread = 128 * 8 = 1024B
|
|
// - prefix scan: small (scan of buffer counts)
|
|
// Total ≈ 4KB for default config → well within 48KB.
|
|
|
|
#pragma once
|
|
|
|
#include "muh/hardware.cuh"
|
|
#include "muh/tuning/common.cuh"
|
|
|
|
namespace muh::tuning::batch_memcpy {
|
|
|
|
// ============================================================================
|
|
// Policy structs — matching CCCL exactly
|
|
// ============================================================================
|
|
|
|
struct BatchedCopySmallBufferPolicy {
|
|
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;
|
|
};
|
|
|
|
struct BatchedCopyLargeBufferPolicy {
|
|
int threads_per_block;
|
|
int bytes_per_thread;
|
|
};
|
|
|
|
struct BatchedCopyLookbackPolicy {
|
|
BatchedCopySmallBufferPolicy small_buffer;
|
|
BatchedCopyLargeBufferPolicy large_buffer;
|
|
};
|
|
|
|
enum class BatchedCopyAlgorithm { lookback };
|
|
|
|
struct BatchedCopyPolicy {
|
|
BatchedCopyAlgorithm algorithm;
|
|
BatchedCopyLookbackPolicy lookback;
|
|
};
|
|
|
|
// ============================================================================
|
|
// policy_selector — matches CCCL exactly
|
|
//
|
|
// CCCL uses the same policy for all CC, only prefer_pow2_bits differs:
|
|
// - SM < 7.0: prefer_pow2_bits = true
|
|
// - SM >= 7.0: prefer_pow2_bits = false
|
|
//
|
|
// BI-V100: equivalent to SM70+ → prefer_pow2_bits = false
|
|
// ============================================================================
|
|
|
|
struct policy_selector {
|
|
constexpr BatchedCopyPolicy operator()(const hardware_capability& /*hw*/) const {
|
|
auto large = BatchedCopyLargeBufferPolicy{256, 32};
|
|
|
|
auto small = BatchedCopySmallBufferPolicy{
|
|
128, // threads_per_block
|
|
4, // buffers_per_thread
|
|
8, // bytes_per_thread
|
|
false, // prefer_pow2_bits (SM70+ = false)
|
|
large.threads_per_block * large.bytes_per_thread, // block_level_tile_size = 8192
|
|
128, // warp_level_threshold
|
|
8 * 1024, // block_level_threshold = 8192
|
|
default_lookback_delay(4), // buffer offset delay (BufferOffsetT = int32)
|
|
default_lookback_delay(4), // block offset delay
|
|
};
|
|
|
|
return BatchedCopyPolicy{
|
|
BatchedCopyAlgorithm::lookback,
|
|
BatchedCopyLookbackPolicy{small, large}
|
|
};
|
|
}
|
|
};
|
|
|
|
} // namespace muh::tuning::batch_memcpy
|