feat: optimize dispatch, KV cache, RPC lifecycle, and ROCm sampling - #1447
Conversation
| } | ||
|
|
||
| const auto decoder_batch_size = params.sequence_lengths.size(0); | ||
| if (decoder_batch_size > 0 && params.no_repeat_ngram_size.has_value()) { |
There was a problem hiding this comment.
[P2] Generic fallback 中 do_sample=false 的 logits 恢复会一并撤销 no_repeat_ngram 与 repetition_penalty
在 generic PyTorch fallback 的 sampleGreedy 中,先 clone original_logits(当 do_sample 不全为 true 时),随后依次在 params.logits 上原地施加 temperature、repetition_penalty、no_repeat_ngram(index_put_ 置 -inf),最后执行 params.logits.copy_(torch::where(do_sample_device, params.logits, original_logits)),将 do_sample=false 的行整体还原为原始 logits。这导致这些行不仅跳过 temperature,也一并撤销了 repetition_penalty 和 no_repeat_ngram 的 -inf 禁止项。CUDA/ROCm 路径中 no_repeat_ngram 在 top_k=1 快速路径(argmax/贪心)之前生效,因此贪心解码会尊重 ngram 禁止;fallback 在此处语义分叉,若 CPU 构建同时启用 no_repeat_ngram 与 do_sample=false,禁止 token 会被重新采样/选中。
建议:将 do_sample=false 的还原范围缩小到仅 temperature(或把 penalty/ngram 的施加移到 restore 之后),保证 no_repeat_ngram 与 repetition_penalty 对 do_sample=false 的行仍然生效。
评审版本:e249c6d5a1d6
| // active device is not the one which mapped the host allocation. Stage every array on | ||
| // the logits device before issuing batched sampling launches. | ||
| if (!all_top_k_one && !all_top_k_no_limit) { | ||
| std::transform(top_k_ptr, top_k_ptr + batch_size, top_k_ptr, [](auto t) { return t <= 0 ? 1 << 30 : t; }); |
There was a problem hiding this comment.
[P2] ROCm 采样原地改写调用方 pinned top_k/top_p 缓冲区,跨层破坏只读契约
本补丁在 ROCm sampleGreedy 中新增 std::transform(top_k_ptr, ..., [](auto t){ return t <= 0 ? 1 << 30 : t; })(CudaSampleOp.cc:751 附近),把 top_k<=0 的行原地改写为 int32 值 1073741824(0x40000000),随后 auto top_k_t = params.top_k.to(device).contiguous() 携带该值传给采样与重归一化内核。在混合分支(非 all_top_k_one 且非 all_top_k_no_limit)中新增调用 top_k_renorm_probs(probs_t, sampling_probs_t, top_k_t, 0, ...)(约 810-820 行)。同文件 api.cc 中 top_k_sampling_from_probs 的转换已由 -static_cast<float*>(maybe_top_k_arr->data_ptr()) 改为 +static_cast<int*>(...),而 top_k_renorm_probs 的转换在本补丁中未改动(仍为 float*)。因此无限制行的 int32 位型 0x40000000 被重解释为 float 2.0f,重归一化输出会把该行错误限制为 top-2,进而污染 cum_log_probs 与 output_all_probs。触发条件:混合批次(部分行 top_k>0、部分行 top_k<=0)且请求 cum_log_probs 或 output_all_probs。旧代码走 torch::multinomial 纯 PyTorch 路径,不经过该 renorm 内核,故为本轮新引入。
建议:不要原地改写调用方输入:对 top_k/top_p 先 clone(或仅在 .to(device) 之后的 GPU 副本上做 transform),保持 params.top_k/params.top_p 只读;或在 Sampler 侧把 top_k/top_p 也纳入持久化缓冲并显式重写,避免哨兵残留。
评审版本:8c39857cd0ee
| at::RecordFunctionGuard record_function_guard(record_functions); | ||
| c10::InferenceMode inference_guard(true); | ||
| cuda_graph::GraphStreamGuard stream_guard(dispatch_stream); | ||
| dispatchSingleStream(stream, |
There was a problem hiding this comment.
[P2] 并行分发时所有 worker 线程共用同一个 CUDA stream 并发入队
dispatch() 在主线程捕获 dispatch_stream = cuda_graph::graphGetCurrentStream()(line 185),随后每个 worker 任务内执行 cuda_graph::GraphStreamGuard stream_guard(dispatch_stream)(line 209)把线程局部当前流设为同一个 dispatch_stream,并在 dispatchSingleStream 中通过 at::cuda::getCurrentCUDAStream().stream() 对同一流执行 cudaSoftmaxInplace 及 .cpu() 设备到主机拷贝。多线程向同一 CUDA stream 并发 enqueue 的先后顺序是不确定的;且每个 worker 的 .cpu() 同步会等待该流上所有已入队操作(含其它 worker 的 kernel),使 GPU 侧工作被串行化,削弱并行分发的收益。若该 dispatch 恰在 CUDA graph capture 阶段(非 warm_up 的惰性 capture)触发,从非 capture 线程向 capture stream 入队属于未定义行为。
建议:为每个 worker 分配独立的 stream(例如从 stream pool 取用),或至少不在 worker 内复用同一 dispatch_stream 做 device-host 同步;若确需共享,应评估 capture 阶段是否可能启用 thread_pool 并显式规避。
评审版本:e249c6d5a1d6
| static constexpr int64_t kStopStreamWaitTimeoutMs = 2000; | ||
|
|
||
| protected: | ||
| void cancelStreamOnTeardown() noexcept; |
There was a problem hiding this comment.
[P2] cancelStreamOnTeardown() 声明为 noexcept 却调用可抛出的虚函数与 reportError,析构路径可能触发 std::terminate
GenerateContext.h:91 声明 void cancelStreamOnTeardown() noexcept;。其实现(GenerateContext.cc:162)在析构链 ~GenerateContext() -> stopStream() -> cancelStreamOnTeardown() 中调用虚函数 isRequestCancelled() 以及 stream_->reportError(...)。同 PR 的测试 PrefillBatchRpcServerTest.cc 中 ThrowOnceContext 重写了 isRequestCancelled() 并 throw std::runtime_error("test fetch failure"),证明该虚函数可以被派生类改写为抛出。一旦在 teardown 路径(析构函数本身隐式 noexcept)中抛出,noexcept 会立即调用 std::terminate 而非让调用方有机会兜底。
建议:teardown 函数不应依赖 noexcept 来“吞”异常——noexcept 在抛异常时是 terminate 而非忽略。建议将 cancelStreamOnTeardown 的 noexcept 去掉并改用 try/catch(...) 包裹体,或在调用 reportError/isRequestCancelled 前保证它们不抛(例如捕获并记录日志后继续),避免析构阶段因派生类重写或 reportError 内部异常导致进程崩溃。
评审版本:e249c6d5a1d6
| for (int i = 0; i < old_batch_size; ++i) { | ||
| const auto& blocks = kv_cache_resource->blocks(i, 0); | ||
| if (batch_fork_count[i] == 0 && !blocks.empty()) { | ||
| reference_updates.push_back({blocks.back(), 1, 0}); |
There was a problem hiding this comment.
[P2] updateKVBlock 对 fork_count==0 的 beam 只释放尾部 block,前缀 block 引用在 DeviceBlockPool 中永久泄漏
改动前对 fork_count==0 的 beam 释放其全部 block(- 行 full_kv_cache_group_->unreference(kv_cache_resource->blocks(old_batch_idx, 0))),改动后 SingleTypeKVCacheAllocator.cc:552-553 仅对 blocks.back()(尾部 block)生成 {block,1,0} 递减;前缀 block 既不在此循环、也不在第二个循环(563-565 行只处理 forks>1 的 beam 前缀 {blocks[j],1,forks})中产生任何更新。消费端 DeviceBlockPool.cc 按 refcounts_[block] = refcounts_[block] - old_count + new_count 执行,故前缀 block 的旧引用永不被减。触发条件:beam search 中某序列长度>1(>1 个 block)的 beam 被完全剪枝(fork_count==0)。(a) 若其前缀被某 fork_count>1 的 beam 共享,该 fork 分支的 {blocks[j],1,forks} 净 +forks-1,而缺失去掉的 -1,导致净多 1 个引用(如 blocks[0] 由 2 变 3);(b) 若前缀为私有,引用计数永不为 0,block 永不回收。两种情形都会在长期 beam search 中抬高引用计数并耗尽 KV cache 容量。同补丁新增测试 ReusesDroppedPrivatePrefixAndTailAtFullCapacity 期望被丢弃的私有前缀 blocks[2] 被复用为 replacement,与该实现相矛盾(仅 blocks[3] 尾部被释放,blocks[2] 泄漏导致 required_free_blocks 不足、updateKVBlock 返回 false),进一步印证该缺陷。
建议:对 fork_count==0 的 beam 遍历其全部 block 生成 {block,1,0}(而非仅 blocks.back()),使前缀与尾部 block 的引用都被释放;并补充一个『淘汰 beam 长度>1 且前缀私有』的单测验证 freeBlocksNum 恢复。
评审版本:e249c6d5a1d6
| // Each row already owns an independent generator state. Mixing the current | ||
| // batch row into the Philox subsequence makes a request's stream change when | ||
| // continuous batching reorders it or changes its co-batch size. | ||
| hiprand_init(philox_seed[bx], 0, philox_offset[bx], &state); |
There was a problem hiding this comment.
[P2] ROCm Philox 子序列改为 0 后随机流唯一性完全依赖 seed/offset,默认 generator 路径与预留空间收缩叠加退化
kernel.cuh:394/531/664 三处采样内核将 - hiprand_init(philox_seed[bx], bx, philox_offset[bx], &state); 改为 + hiprand_init(philox_seed[bx], 0, philox_offset[bx], &state);,去掉 batch 行号 bx 作为 Philox 子序列,随机流唯一性从此完全依赖 (seed, offset)。消费方 CudaSampleOp.cc ROCm 分支同时把 - auto [sd, ofst] = get_seed_and_offset(batch_size * 32, params.generator[i].defined() ? ... : std::nullopt); 改为 + auto [seed, offset] = get_seed_and_offset(max_sampling_rounds, generator);(max_sampling_rounds=32,:786-788),把每行预留的 Philox 计数空间从 batch_size*32 收缩到 32。而 api.cc:25-32 的 get_seed_and_offset 对 undefined generator 走 at::get_generator_or_default<CUDAGeneratorImpl>(generator, getDefaultCUDAGenerator()),即多行共享进程级默认 generator:seed 固定、行间唯一性只剩全局 offset 自增(每次 +32)这一个维度。因果链:子序列去 bx + 默认 generator 固定 seed → 行间唯一性退化为单 offset 维度;且每行预留空间缩到 32,而 kernel.cuh 的 do-while 采样循环(:414 起)只有 while (low < high) 收敛条件、无 round < 32 硬上限,病态分布一旦超过 32 轮会越过 32 的预留边界与相邻行 offset 重叠,叠加子序列=0 后产生相同随机流、token 输出相关。当前生产调用方(每行独立 generator,RocmSamplerOpTest.cc 用 arange 唯一 seed)均未触发,属契约收紧后的潜在退变。
建议:在 kernel 内对 round 加硬上限(如 >= max_sampling_rounds 即 break/置失败),或在 get_seed_and_offset 对 nullopt 分支为每行派生唯一 seed(而非依赖全局 offset),并在测试中补一条『多行共享默认 generator』的用例验证无相关性。
评审版本:e249c6d5a1d6
| ) | ||
| for stop_word_id in stop_word_ids: | ||
| if stop_word_id and np.array_equal(tokens[-len(stop_word_id) :], stop_word_id): | ||
| if stop_word_id and tokens[-len(stop_word_id) :] == stop_word_id: |
There was a problem hiding this comment.
[P3] stop_word_id 比较从 np.array_equal 改为列表 == 对非 list 类型更脆弱
改动将 np.array_equal(tokens[-len(stop_word_id):], stop_word_id) 改为 tokens[-len(stop_word_id):] == stop_word_id(补丁 - 行 if stop_word_id and np.array_equal(tokens[-len(stop_word_id) :], stop_word_id):)。当 stop_word_id 为 numpy 数组时,list == np.array 返回逐元素布尔数组,if 判断会抛 "truth value of an array is ambiguous";当为 tuple 时 list == tuple 恒为 False。类型注解虽为 List[List[int]],但原实现对这些类型更宽容。
建议:若确认调用方只传 list 可保持不变;否则建议保留 np.array_equal 或显式转换为 list 后再比较。
评审版本:e249c6d5a1d6
| @@ -281,7 +283,7 @@ DeviceSamplingFromProb(uint32_t | |||
|
|
|||
There was a problem hiding this comment.
[P2] DeviceSamplingFromProb 新增 greater_than_u[j] 条件可能跳过合法的逆 CDF 采样位置
改动前为 - if (greater_than_u_diff[j]) {,改动后为 + if (greater_than_u_diff[j] && greater_than_u[j]) {。greater_than_u_diff[j] 是 CDF 跨越 u 的检测(inclusive_cdf[j]>u 且 inclusive_cdf[j-1]<=u),而 greater_than_u[j] 从命名看是 value[j]>u。在尖峰分布下,CDF 跨越 u 的位置其增量 value[j] 可能很小(例如 probs=[0.9,0.05,0.05]、u=0.92,跨越在 j=1 但 value[1]=0.05<=u),此时新增的 && greater_than_u[j] 为 false,导致该位置不被 atomicMin 记录,sampled_id 保持 d、last_valid_id=-1,最终合法行被误判为失败(output=-1、success=false)。现有测试(如 CombinedSamplingMatchesReturnedDistribution 用均匀 logits)未覆盖这种尖峰分布场景。
建议:确认 greater_than_u 的确切语义;若它表示 value>u,则应移除该附加条件,仅保留 CDF 跨越检测 greater_than_u_diff[j],并补充尖峰分布(低概率 token 被采样)的测试用例。
评审版本:e249c6d5a1d6
| bool DETERMINISTIC> | ||
| __device__ __forceinline__ bool | ||
| SampleProbabilityRound(const float* probs, | ||
| uint32_t row_idx, |
There was a problem hiding this comment.
[P3] SampleProbabilityRound 硬编码 const float 与 kernel 的 DType 模板不一致*
新增的 SampleProbabilityRound 函数签名参数为 const float* probs,而调用它的 TopKSamplingFromProbKernel/TopPSamplingFromProbKernel/TopKTopPSamplingFromProbKernel 均以 DType* probs 作为入参(kernel 模板参数 DType)。当前 api.cc 仅以 float 实例化,故可编译通过;但若未来以半精度等非 float 类型实例化该 kernel,DType* 无法隐式转换为 const float*,将导致编译失败。
建议:将 SampleProbabilityRound 的 probs 参数改为模板类型 DType(与 kernel 一致),或在函数内沿用 cast_load 的 DType 语义。
评审版本:e249c6d5a1d6
| // Clear the full probability row only on failure; never gather with the -1 sentinel. | ||
| __global__ void | ||
| FinalizeSamplingProbKernel(float* probs, const int* samples, const bool* success, float* log_probs, uint32_t d) { | ||
| const uint32_t row = blockIdx.x; |
There was a problem hiding this comment.
[P2] finalize_sampling_probs 缺少 indices 参数,与采样内核的 indices 重排契约不一致,beam 重排时 log_prob 会按 batch 位次读错行
FinalizeSamplingProbKernel 中 log_probs[row] = logf(probs[static_cast<size_t>(row) * d + samples[row]]) 仅在 success[row]==true 时执行。但 success 由采样内核在采样成功时置 true,samples[row] 此时应为合法 token;然而 finalize_sampling_probs 的 api.cc 校验只检查 samples 的 dtype/shape,并不校验其取值。若调用方传入的 samples 含 -1(例如 all_top_k_one 之外的路径在 sampling 失败但 success 未正确同步时),或 samples 与 success 来自不同 producer,healthy 分支会以负数索引越界读 probs。
建议:给 finalize_sampling_probs(及 FinalizeSamplingProbKernel)增加与采样内核一致的 std::optional<torch::Tensor> maybe_indices 参数,在 indices 非空时用 row_idx = indices[bx] 索引 probs/samples/log_probs;或在 api.cc 中 TORCH_CHECK 断言 finalize 的调用场景 indices 必须为空,避免将来 beam 重排时静默读错行。
评审版本:a9bd8a974369
| static_cast<int*>(output.data_ptr()), | ||
| maybe_indices.has_value() ? static_cast<int*>(maybe_indices->data_ptr()) : nullptr, | ||
| has_top_k_arr ? static_cast<float*>(maybe_top_k_arr->data_ptr()) : nullptr, | ||
| has_top_k_arr ? static_cast<int*>(maybe_top_k_arr->data_ptr()) : nullptr, |
There was a problem hiding this comment.
[P3] top_k_arr 强转为 int 却缺少 int32 类型校验*
top_k_sampling_from_probs 中 maybe_top_k_arr 仅做了 CHECK_INPUT / CHECK_DEVICE(api.cc:117-119),随后直接 static_cast<int*>(maybe_top_k_arr->data_ptr())(api.cc:132)。本次改动将 top_k_arr 从 float* 改回 int*(补丁 - 行:has_top_k_arr ? static_cast<float*>(maybe_top_k_arr->data_ptr()) : nullptr),但对比同函数对 success 的处理(success->scalar_type() == torch::kBool 显式 TORCH_CHECK),top_k_arr 缺少 dtype 校验。若存在调用方误传 float/非 int32 tensor,会被静默按 int32 位模式读取,产生错误的 top_k 值。当前补丁内调用方(CudaSampleOp.cc、RocmSamplerOpTest.cc)均传 int32,故仅为防御性缺口。
建议:在 maybe_top_k_arr.has_value() 分支内增加 TORCH_CHECK(maybe_top_k_arr->scalar_type() == torch::kInt32, ...),与 success 的类型校验保持一致,避免未来调用方误传类型导致静默错误。
评审版本:e249c6d5a1d6
| decrementRefCounterImpl<false>(block_indices); | ||
| } | ||
|
|
||
| std::vector<int> decrementRefCounterWithFreeInfo(const std::vector<int>& block_indices) { |
There was a problem hiding this comment.
[P3] 新增 BlockRefCounter::decrementRefCounterWithFreeInfo 是死代码,KV cache 原子引用更新实际走 DeviceBlockPool::tryReplaceRequestReferences
跨文件因果链:① BlockRefCounter.h:43 新增 std::vector<int> decrementRefCounterWithFreeInfo(...),cache/BUILD:400 新增 name="block_ref_counter" cc_library,utils/test/BUILD:23 新增 block_ref_counter_test;② 但本 PR 同时重写的 SingleTypeKVCacheAllocator.cc::updateKVBlock(补丁中旧实现删除了 full_kv_cache_group_->unreference/reference 逐块计数逻辑)改为直接调用 block_pool_->tryReplaceRequestReferences(reference_updates, ...),其内部用 DeviceBlockPool 自有的 refcounts_(unordered_map)计数,从未调用 BlockRefCounter。grep 全仓:decrementRefCounterWithFreeInfo 仅出现在 BlockRefCounter.h 定义处与 BlockRefCounterTest.cc 测试处,无任何生产 .cc 调用;block_ref_counter 目标仅被 block_ref_counter_test 依赖。即新增的'带释放信息递减'API 与其 BUILD target 是为旧逐块计数方案准备的,但最终方案切换后未清理,成为与片 1 的 StlUtil/vectorRemoveIf 同类的死代码(且连同 BlockRefCounter 类一起失去生产使用者)。
建议:删除未使用的 decrementRefCounterWithFreeInfo 及其测试,或将其真正接入 DeviceBlockPool 的引用更新路径;避免留下与生产机制脱节的死代码与误导性测试。
评审版本:a9bd8a974369
PR #1447 评审:未发现阻塞项,等待人工审核
无阻塞项非阻塞发现
自动代码评审 · head |
rtp-llm-review-bot
left a comment
There was a problem hiding this comment.
lgtm ready to ci
评审版本:e249c6d5a1d69c767860bebd7d4baa8e1195d6bd
e249c6d to
8c39857
Compare
| #include <numeric> | ||
| #include <set> | ||
| #include <stdexcept> | ||
| #include "torch/csrc/autograd/profiler_kineto.h" |
There was a problem hiding this comment.
[P3] 测试依赖 torch 内部头文件与内部 API,随版本升级脆弱
新增 #include "torch/csrc/autograd/profiler_kineto.h",并直接使用 torch::profiler::impl::ProfilerConfig、torch::profiler::impl::ActivityType、at::RecordFunction::currentThreadId() 等内部/非稳定 API(torch/csrc 路径不属于公开 API)。第 126 行 EXPECT_EQ(event.startThreadId() == at::RecordFunction::currentThreadId(), GetParam() == 0) 依赖内部线程 ID 表示一致性,torch 升级后可能编译失败或断言误判。
建议:改用公开的 torch::profiler 接口,或避免依赖内部线程 ID 断言(例如只断言 update 事件数量与串/并行分支)。
评审版本:8c39857cd0ee
|
|
||
| NormalOutputDispatcher dispatcher({}, 2); | ||
| ASSERT_NE(dispatcher.thread_pool_, nullptr); | ||
| StreamGroups stream_groups({failing_stream, delayed_stream}); |
There was a problem hiding this comment.
[P2] testParallelDispatchWaitsForAllWorkersBeforePropagatingException 遗漏 sampler_output.success,依赖未定义张量的脆弱处理
该测试构造 MergedOutput 时只设置 sampler_output.token_ids,未设置 sampler_output.success;同文件其余所有新增/改动的 dispatch 测试(testParallelDispatchMultipleStreams、testMixedPromptLengthsAndBeamExpansion、testDispatchPreservesCpuProfilingAndAsyncDisableGuard)均显式设置 success=torch::tensor({true,...},kBool)。dispatch() 中 success_cpu = merge_outputs.sampler_output.success.cpu() 对未定义张量返回未定义张量,随后 collectStreamSamplerError 在 updateOutput 抛异常之前访问 success_cpu;若该函数未检查 defined()(或 SamplerOutput 默认构造 success 为空张量导致越界),则测试在验证异常传播前即崩溃,属于未定义行为依赖。
建议:显式设置 merge_outputs.sampler_output.success = torch::tensor({true, true}, torch::kBool),与同文件其他 dispatch 测试保持一致,避免依赖未定义/空张量的隐式处理。
评审版本:8c39857cd0ee
| params.runtime_config.output_dispatcher_worker_count = 2; | ||
| params.py_model = py::none(); | ||
| // Exercise the executor constructor without starting TP collectives. | ||
| NormalExecutor executor(params, nullptr); |
There was a problem hiding this comment.
[P3] 为验证一行条件而构造完整 NormalExecutor,引入全局静态副作用
testDispatchWorkersOnlyCreatedOnOutputRank 在循环中两次构造完整 NormalExecutor executor(params, nullptr)(rank 0/1),仅为验证 tp_rank_ > 0 ? 0 : worker_count。NormalExecutor 构造函数(NormalExecutor.cc)会创建 CUDA stream 池 dispatch_runner_(cuda_graph::graphGetStreamFromPool(true)) 并调用静态 LogitsProcessorFactory::init(...)(LogitsProcessorFactory.h 中为 static,会重初始化静态 grammar backend),引入全局状态与资源分配,可能造成测试间顺序依赖。
建议:直接构造 NormalBatchStreamProcessor 或 NormalOutputDispatcher 断言 worker 数量,避免拉起完整 executor 及其全局副作用。
评审版本:8c39857cd0ee
| } // namespace | ||
|
|
||
| NormalOutputDispatcher::NormalOutputDispatcher(std::vector<int64_t> output_vocab_ids, int async_worker_count): | ||
| output_vocab_ids_(std::move(output_vocab_ids)) { |
There was a problem hiding this comment.
[P3] NormalOutputDispatcher 构造函数无条件打 INFO 日志,串行分发路径每次构造都会产生日志噪音
构造函数新增分支 else { RTP_LLM_LOG_INFO("output dispatcher worker count is 0; using serial dispatch"); }(NormalOutputDispatcher.cc 约 50 行)。该路径在 worker_count=0(默认值)时每次构造都会执行,而 NormalOutputDispatcher 会随 NormalBatchStreamProcessor/NormalExecutor 多次构造(包括 warm_up、propose、tp_rank>0 等场景),属于非 per-forward 但高频的构造期日志。
建议:降为 DEBUG 级别,或仅在显式配置 worker_count 时打印一次。
评审版本:8c39857cd0ee
rtp-llm-review-bot
left a comment
There was a problem hiding this comment.
lgtm ready to ci
评审版本:8c39857cd0eeb1809befbd2ae614fb2aa7e63f23
8c39857 to
b222887
Compare
| kv_indptr = self.fmha_params.decode_page_indptr_d | ||
| last_page_len = self.fmha_params.paged_kv_last_page_len_d | ||
| active_batch_size = last_page_len.numel() | ||
| batch_size = qo_indptr.numel() - 1 |
There was a problem hiding this comment.
[P2] _plan_prefill_wrapper 把 qo_indptr.numel()-1 当作 capture batch size,非 copy-params 路径下 shrink 会误抛 "captured batch size"
第 452 行 batch_size = qo_indptr.numel() - 1 在 CUDA graph 分支中被当作 capture batch size 参与校验(第 464 行 graph_kv_indptr.numel() != batch_size + 1 与 graph_last_page_len.numel() != batch_size)。但该假设只在 prefill_cuda_graph_copy_params is not None 时成立:prepare 中该分支把 qo_indptr 固定为 self.qo_indptr(numel=capture+1);而 prefill_cuda_graph_copy_params is None 时(第 390-392 行)qo_indptr = attn_inputs.cu_seqlens_device[: active+1] 是 active 输入(numel=active+1)。此时 graph_kv_indptr.numel() 恒为 capture+1,batch_size+1 为 active+1,shrink(active<capture)时二者不等,会误判为 "captured batch size" 抛 ValueError,而非进入第 479 行的 shrink/pad 逻辑。已核实生产 capture 流程 cuda_graph_runner.cc:1707-1713 始终设置 prefill_cuda_graph_copy_params,故该路径当前可能不可达,但 _plan_prefill_wrapper 未强制这一不变量,属潜在缺陷。
建议:用 graph buffer 自身推导 capture batch size(如 graph_kv_indptr.numel() - 1 或 graph_last_page_len.numel()),不要依赖 qo_indptr.numel()-1;或让非 copy-params 路径也复用固定 capture 的 qo_indptr 缓冲。
评审版本:b222887bfbc5
| # FlashInfer derives its batch size from qo_indptr. Inactive | ||
| # query slots must have empty KV ranges; passing shortened KV | ||
| # tensors makes the native planner read beyond their bounds. | ||
| graph_last_page_len[active_batch_size:].zero_() |
There was a problem hiding this comment.
[P3] 对 captured CUDA graph 输入缓冲区的异步写入未与 replay stream 显式同步
第 479 行 graph_last_page_len[active_batch_size:].zero_() 与第 480-484 行 graph_kv_indptr[active_batch_size+1:].copy_(..., non_blocking=True) 在 prepare() 阶段向 captured graph 的输入缓冲区(_paged_kv_indptr_buf/_paged_kv_last_page_len_buf)发起异步设备端写入,enqueue 到当前 stream。这些缓冲在 forward() 的 graph replay 中会被读取;若 replay 运行在不同 stream,或跨请求 prepare/forward 流水线重叠,则元数据写入与 replay 读取之间无顺序保证("若存在")。
建议:确保这些元数据写入与 graph replay 使用同一 stream,或在 replay 前对相关 stream 做显式同步。
评审版本:b222887bfbc5
Keep capture-sized KV metadata views at stable addresses and pad inactive slots before FlashInfer planning. Restore native tensor shapes before capacity checks so replay can grow after a smaller batch. Add deterministic metadata checks and host/device shrink-and-grow coverage. Native CUDA build passes; all 27 replay tests pass on H20 under Compute Sanitizer with zero errors. The adjacent eager suite has numerical failures that also occur with the baseline operator.
b222887 to
a9bd8a9
Compare
| auto top_p_t = params.top_p.to(params.logits.device(), /*non_blocking=*/true).contiguous(); | ||
|
|
||
| // 6. Sample. Kernels write the per-row status directly, without a comparison launch. | ||
| auto success = |
There was a problem hiding this comment.
[P2] ROCm 采样原地改写调用方 pinned top_k/top_p 缓冲区,跨层破坏只读契约
跨文件因果链(生产者→聚合者→消费者):
(1) 生产者 CudaSampleOp.cc:757-758(ROCm 分支)auto success = all_top_k_one ? torch::Tensor() : torch::empty(...)——当全部行 top_k==1 时 success 是未定义张量;随后 :760 走 torch::argmax(probs_t, -1),对含 NaN/Inf/负值的 logits 行 argmax 静默返回索引 0,既不写 -1 哨兵也不写 success=false;:874 return {success} 把未定义 success 返回。对照改动前 - return GreedyOutput{};(本就无 success),本 PR 新增的 + return {success}; 在 all_top_k_one 分支反而退化为未定义。
(2) 聚合者 Sampler.cc:252-255 if (greedy_output.success.defined()) { success.copy_(greedy_output.success); } else { success.fill_(true); }——未定义 success 被无条件 fill_(true),把「无失败信息」转成「全部成功」。
(3) 消费者 NormalOutputDispatcher.cc:96-100(collectStreamSamplerError)仅在 !success[batch_idx_in+i] 时 set_first_error("sampler generate token id failed"),因 success 全 true,采样失败永不被上报(:418 的 output_vocab 分支同理)。
后果:本 PR 为 ROCm 采样新增的失败检测(kernel.cuh 写 -1 哨兵 + success=false,RocmSamplerInvalidProbTest 验证 NaN/Inf/负值行被拒绝)在 top_k=1 路径被整条链静默绕过——请求带 NaN logits(FP16 溢出或 logits processor 产 NaN)且设置 cum_log_probs/output_all_probs 时,得到 token 0 与 NaN cum_log_probs(:867 sampling_probs_t.gather(...).log() 对 NaN 行取 log 得 NaN),且下游无任何报错,核心价值路径被破坏。
建议:不要原地改写调用方输入:对 top_k/top_p 先 clone(或仅在 .to(device) 之后的 GPU 副本上做 transform),保持 params.top_k/params.top_p 只读;或在 Sampler 侧把 top_k/top_p 也纳入持久化缓冲并显式重写,避免哨兵残留。
评审版本:a9bd8a974369
| uintptr_t stream = 0); | ||
| uintptr_t stream = 0, | ||
| std::optional<torch::Tensor> success = std::nullopt); | ||
|
|
There was a problem hiding this comment.
[P3] sampling.h 文件末尾缺少换行符
三个采样函数(top_k/top_p/top_k_top_p_sampling_from_probs)的 success 均为 std::optionaltorch::Tensor success = std::nullopt(可选,向后兼容),而新增的 finalize_sampling_probs 把 success 声明为必填 torch::Tensor success(非 optional)。采样内核无论是否传入 success 都会在失败行写入 output[bx]=-1 哨兵(kernel.cuh 中 if(tx==0){output[bx]=-1; if(success!=nullptr) success[bx]=false;}),但清理失败行 probs 的 finalize 却强制要求 success。若存在调用方沿用旧接口不传 success,将既拿不到 success 区分失败行,也无法调用 finalize 清理失败行残留的 NaN/负值 probs,仅能靠 output<0 识别 -1 哨兵。当前唯一调用方 CudaSampleOp.cc 恒传 success(success.defined() 才调用 finalize),故无实际触发路径。
建议:在文件末尾补上换行符,避免某些编译器/工具链告警。
评审版本:a9bd8a974369
| } | ||
| if (async_worker_count > 0) { | ||
| thread_pool_ = std::make_unique<autil::LockFreeThreadPool>( | ||
| async_worker_count, 2 * async_worker_count, nullptr, "OutputDispatcher"); |
There was a problem hiding this comment.
[P3] NormalOutputDispatcher 构造函数无条件打 INFO 日志,串行分发路径每次构造都会产生日志噪音
跨文件因果链:① engine_group_args.py _non_negative_int 仅 if parsed < 0 拒绝负值,不设上界;② ConfigModules.h int output_dispatcher_worker_count = 0(32 位,经 pybind def_readwrite 绑定);③ NormalExecutor.cc:171 const int async_worker_count = ... params.runtime_config.output_dispatcher_worker_count 原样透传;④ NormalOutputDispatcher.cc:56 thread_pool_ = std::make_unique<autil::LockFreeThreadPool>(async_worker_count, 2 * async_worker_count, ...) 中 2 * async_worker_count 为 int*int 有符号运算。后果:用户传 --output_dispatcher_worker_count 1500000000(介于 INT_MAX/2 与 INT_MAX 之间)时,argparse 校验通过、pybind 赋值不抛 OverflowError(值仍在 int32 范围内),但 2 * 1500000000 = 3000000000 溢出 int32 触发未定义行为,回绕为负的队列长度传给 LockFreeThreadPool,导致线程池被错误配置或崩溃。该值避开了分片评审提到的「>INT_MAX 时 pybind 抛 OverflowError」边界,是更隐蔽的二次溢出。
建议:降为 DEBUG 级别,或仅在显式配置 worker_count 时打印一次。
评审版本:a9bd8a974369
| graph_last_page_len[active_batch_size:].zero_() | ||
| graph_kv_indptr[active_batch_size + 1 :].copy_( | ||
| kv_indptr[active_batch_size : active_batch_size + 1].expand( | ||
| batch_size - active_batch_size |
There was a problem hiding this comment.
[P2] 对 captured CUDA graph 输入缓冲区的异步写入未与 replay stream 显式同步
_plan_prefill_wrapper 中 batch_size = qo_indptr.numel()-1(query 数),active_batch_size = last_page_len.numel()(KV 序列数)。plan 接口同时接收 page_indice_d 作为 query→KV 映射,说明支持多个 query 共享同一 KV 序列;此时 KV 序列数可小于 query 数。新增的 elif active_batch_size != batch_size: raise ValueError("Paged prefill query and KV batch sizes must match") 在非 cuda-graph 路径会对此类合法场景误抛异常(旧实现直接透传 decode_page_indptr_d,无此限制)。若该 prefill 路径确实存在 prefix 共享调用方,则引入回归。
建议:确保这些元数据写入与 graph replay 使用同一 stream,或在 replay 前对相关 stream 做显式同步。
评审版本:a9bd8a974369
| batch_resource[batch_idx] = std::move(resource); | ||
| } | ||
|
|
||
| void swap(BatchKVCacheResource& other) noexcept { |
There was a problem hiding this comment.
[P2] swap 无条件交换 cache_keys_initialized_,成功 updateKVBlock 后标志被误置回 false,与失败路径保留标志的语义不一致
BatchKVCacheResource.h 新增 void swap(...) { batch_resource.swap(...); std::swap(cache_keys_initialized_, other.cache_keys_initialized_); }(221-224 行)。SingleTypeKVCacheAllocator.cc 的 updateKVBlock 中 staged_resource 仅 resetBatchSize(new_batch_size) 后逐 fork setCacheKeys,从未调用 markCacheKeysInitialized,故 staged 侧 cache_keys_initialized_ 恒为默认 false;成功路径 kv_cache_resource->swap(staged_resource) 后 kv_cache_resource 的标志被置回 false,尽管 fork/retained 的 cache keys 均已保留。失败路径(tryReplaceRequestReferences 返回 false 提前 return、无 swap)标志保持不变——测试 UpdateKVBlockReservationFailureLeavesResourceUnchanged 显式断言失败后 EXPECT_TRUE(resource->cacheKeysInitialized()),证明该标志应跨成功更新保留。而消费者 KVCacheManager.cc:410 const bool keys_already_initialized = malloc_info.batch_kv_cache_resource->cacheKeysInitialized(); 依赖此标志区分『首次 malloc 的重试』与『真正的首次尝试』(405 行注释),标志被错误清 false 会触发 initCacheKeys 重算与 prefill-cache-hit 指标重复上报。
建议:swap 不应交换 cache_keys_initialized_(或 updateKVBlock 成功提交后显式 kv_cache_resource->markCacheKeysInitialized()),使成功/失败路径对标志的语义一致:cache keys 保留则标志保留。
评审版本:a9bd8a974369
| { | ||
| std::lock_guard<std::mutex> lock(mutex_); | ||
| checkInitializedNoLock(); | ||
| for (size_t i = 0; i < changes.size(); ++i) { |
There was a problem hiding this comment.
[P3] tree_hold 上限为 1 使 request_refs 在多种 tree 引用并存时被高估,弱化引用计数校验
tryReplaceRequestReferences 中 const uint32_t tree_hold = treeRefCountNoLock(change.block) > 0 ? 1 : 0; 将 tree 引用数截断为 0/1,随后 request_refs = refcounts_[change.block] - tree_hold。若同一 block 同时持有多种 BlockTreeRefType(如 LOAD 与 CACHE 各 1),treeRefCountNoLock 返回 2,tree_hold 仍为 1,request_refs 比真实请求引用数多 1。这使 RTP_LLM_CHECK_WITH_INFO(request_refs >= change.old_count) 校验过松,可能掩盖释放超过实际持有的请求引用数的错误(reclaimable 判定因额外要求 tree_hold == 0 而不受影响,但 sanity check 被弱化)。
建议:将 tree_hold 改为直接使用 treeRefCountNoLock 的返回值(uint32_t),或单独以 treeRefCountNoLock()>0 判断是否存在 tree 引用,避免截断导致 request_refs 高估。
评审版本:a9bd8a974369
| namespace rtp_llm { | ||
|
|
||
| template<typename T, typename Pred> | ||
| void vectorRemoveIf(std::vector<T>& vec, Pred&& pred) { |
There was a problem hiding this comment.
[P3] 新增 StlUtil.h/vectorRemoveIf 及 stl_utils BUILD target 为无调用方的死代码
本补丁新增 rtp_llm/cpp/utils/StlUtil.h(定义模板 vectorRemoveIf)并在 rtp_llm/cpp/utils/BUILD 新增 name="stl_utils" 的 cc_library。全仓 grep 显示 vectorRemoveIf 仅出现在 StlUtil.h 定义处,stl_utils 仅出现在自身 BUILD 定义处,无任何调用方,也无其他 BUILD target 依赖 //rtp_llm/cpp/utils:stl_utils。
建议:删除 StlUtil.h 与 stl_utils target;若为后续改动预留,应推迟到有实际消费者时再引入。
评审版本:a9bd8a974369
| } | ||
|
|
||
| void PrefillGenerateContext::setStream(const std::shared_ptr<GenerateStream>& stream) { | ||
| if (stream_ && stream_ != stream) { |
There was a problem hiding this comment.
[P3] PrefillGenerateContext::setStream 与基类 GenerateContext::setStream 重复实现 retry 取消逻辑(DRY),且出队调用语义不一致
本补丁在 GenerateContext::setStream 新增了 if (stream_ && stream_ != stream) { stopStreamForRetry(); }(GenerateContext.cc 中 stopStreamForRetry 内部用 meta->dequeue(request_id, stream_)),同时在 PrefillGenerateContext::setStream 内联复制了几乎相同的 if (stream_ && stream_ != stream) { reportError(CANCELLED, "cancel abandoned retry attempt"); dequeueStreamFromRuntimeMeta(); }。两处对旧 stream 的取消语义重复,且出队方式不同(基类用 request_id 调 meta->dequeue,Prefill 用 dequeueStreamFromRuntimeMeta),后续若只改一处会导致 retry 时旧 stream 未被正确出队/取消。
建议:抽取一个可重载的出队钩子(如 protected virtual dequeueStream()),让基类 stopStreamForRetry 复用同一段 retry 取消+出队逻辑,消除两处重复。
评审版本:a9bd8a974369
|
|
||
| DeferredPrefillContext::~DeferredPrefillContext() { | ||
| finishLogicalTrace(); | ||
| // Batch contexts span multiple handlers. The last owner marks completion |
There was a problem hiding this comment.
[P3] DeferredPrefillContext 仅在非 ACTIVE 终态标记 RPC 完成,ACTIVE 销毁会触发基类 GenerateContext 析构的 ERROR 日志误报
GenerateContext.cc 析构新增 if (!rpc_handling_completed_) { RTP_LLM_LOG_ERROR("GenerateContext destroyed before RPC handling completed", ...); },而 Decode/Local/Prefill 三个 handler 通过 autil::ScopeGuard(std::uncaught_exceptions() 相等才 markRpcHandlingCompleted())为普通路径打标。但 PrefillBatchRpcServer.cc 的 DeferredPrefillContext::~DeferredPrefillContext() 采用不同判定:if (context && context->terminalCause() != PrefillTerminalCause::ACTIVE) { if (!logical_status.ok()) context->error_status = logical_status; context->markRpcHandlingCompleted(); }。两条打标路径的契约不一致:batch 路径在 terminalCause 仍为 ACTIVE(如服务关闭/异常路径在终态发布前销毁 deferred context)时不调用 markRpcHandlingCompleted,随后基类析构无条件输出一条 ERROR 级『destroyed before RPC handling completed』日志,造成关闭期日志噪音并误导排障。
建议:对齐两条路径的完成判定:DeferredPrefillContext 析构在 ACTIVE 状态下也应显式标记 RPC 处理已完成(或在基类析构中区分『已发布终态』与『异常提前销毁』两种场景,避免对正常关闭路径报 ERROR)。
评审版本:a9bd8a974369
|
|
||
| def _non_negative_int(value: str) -> int: | ||
| try: | ||
| parsed = int(value) |
There was a problem hiding this comment.
[P2] output_dispatcher_worker_count 无上界,pybind int 边界与 2*async_worker_count 二次溢出构成跨文件崩溃链
跨文件因果链:① 生产者 engine_group_args.py:8 parsed = int(value) 接受任意精度整数,仅 engine_group_args.py:13 if parsed < 0 校验非负,无上界;② 该值经 engine_group_args.py:52 bind_to=(runtime_config, "output_dispatcher_worker_count") 绑定到 ConfigInit.cc:1563 .def_readwrite("output_dispatcher_worker_count", &RuntimeConfig::output_dispatcher_worker_count),而目标成员 ConfigModules.h:492 声明为 int output_dispatcher_worker_count = 0(32 位有符号);③ 消费者 NormalOutputDispatcher.cc:56-57 thread_pool_ = std::make_unique<autil::LockFreeThreadPool>(async_worker_count, 2 * async_worker_count, nullptr, ...) 中 2 * async_worker_count 为 int 算术。触发路径:--output_dispatcher_worker_count 1073741824(2^30)能通过 pybind 的 int 赋值(< INT_MAX=2^31-1),随后 2 * 1073741824 = 2147483648 有符号溢出(UB,实际回绕为负),负值作为 LockFreeThreadPool 队列大小传入,导致崩溃或超大分配;即便用更“合理”的 2e9 也会在 2* 处溢出。注意 async_worker_count 仅在 tp_rank==0 且非 warm_up 且非 propose 时取该配置值(NormalExecutor.cc),但生产可关闭 warm_up,路径真实可达。
建议:在 _non_negative_int 增加合理上界(如 < 1024,或至少 < INT_MAX/2);同时在 NormalOutputDispatcher 构造中对 async_worker_count 做 2 * async_worker_count 溢出检查(改用 int64_t 或 RTP_LLM_CHECK(async_worker_count <= std::numeric_limits<int>::max()/2))。
评审版本:a9bd8a974369
a9bd8a9 to
e82a106
Compare
| ASSERT_EQ(beam->nextBatchSize(), 2); | ||
|
|
||
| const auto logits = torch::tensor({0.f, 1.f, 2.f, 3.f, 4.f, 5.f, 6.f, 7.f, 8.f, 9.f, 9.f, 7.f, 5.f, 3.f, 1.f, | ||
| 0.f, 2.f, 4.f, 6.f, 8.f, 2.f, 0.f, 4.f, 1.f, 6.f, 3.f, 8.f, 5.f, 9.f, 7.f}) |
There was a problem hiding this comment.
[P2] testMixedPromptLengthsAndBeamExpansion 中 first 流的 token_ids 行与期望输出不一致,新 token 4 被放在第 3 列而非 seqLength 列
token_ids 张量为 torch::tensor({0,9,9,4,1,0,2,9,1,0,3,9,2,2,1,5}).reshape({4,4}),第 0 行为 [0,9,9,4];first 流由 make_stream({0},false) 构造,输入长度为 1,按本文件其他测试(如 testParallelDispatchMultipleStreams 中 stream1 输入 {0} 且 token_ids 行 [0,1] 得到 completeTokenIdsVec=={0,1})确立的约定,新 token 应位于第 seqLength==1 列(值 9),但测试断言 EXPECT_EQ(first->completeTokenIdsVec(0), (std::vector{0,4})) 期望新 token 为 4,而 4 位于第 3 列。beam 流(L=2)与 last 流(L=3)的新 token 分别位于第 2、3 列,与约定一致,唯独 first 行错位。若 dispatch 按 seqLength 列读取,该测试将失败或掩盖 offset 读取错误。
建议:将 token_ids 第 0 行改为 {0,4,9,9}(或等价地让新 token 4 位于第 1 列),使 first 流的 token 布局与 beam/last 流及既有测试约定一致。
评审版本:e82a10668cff
| } | ||
|
|
||
| TEST_F(NormalBatchStreamProcessorTest, testLoss) { | ||
| TEST_P(OutputDispatchTest, testLoss) { |
There was a problem hiding this comment.
[P3] 参数化的 testLoss/testSoftmaxProbs 在并行模式下未设置 sampler_output.success,成功/失败路径未被覆盖
本轮将 testLoss(第1230行)与 testSoftmaxProbs(第491行)从 TEST_F 改为 TEST_P(OutputDispatchTest),INSTANTIATE_TEST_SUITE_P 用 ::testing::Values(0,2) 使二者新增 worker_count=2 的并行分发路径。但这两个用例只设置了 sampler_output.token_ids / cum_log_probs,未设置 sampler_output.success(全文件 grep success 仅命中 102/578/647/813/850/881/1180 行,不含这两个用例)。NormalOutputDispatcher::dispatch 中 success_cpu = copyToPinnedCpuAsync(sampler_output.success, ...) 对未定义张量直接原样返回,collectStreamSamplerError 里 if (success_cpu.defined()) 分支被跳过。因此并行模式下这两个用例的采样失败->报错链路从未被真正执行,参数化新增的并行路径对错误收集逻辑零覆盖,是台账 P2(异常用例遗漏 success)在参数化改动下引入的新变体。
建议:为 testLoss 与 testSoftmaxProbs 补充 sampler_output.success = torch::tensor({true,...}, torch::kBool)(或按批次行数填充),使并行分发的错误收集路径在两种 worker_count 下都被实际覆盖。
评审版本:e82a10668cff
|
|
||
| tpi::ProfilerConfig config(tpi::ProfilerState::KINETO, /*report_input_shapes=*/false); | ||
| const std::set<tpi::ActivityType> activities{tpi::ActivityType::CPU}; | ||
| tap::prepareProfiler(config, activities); |
There was a problem hiding this comment.
[P3] testDispatchPreservesCpuProfilingAndAsyncDisableGuard 混用已废弃的 torch::autograd::profiler 遗留 API 与 KINETO 新 API
该测试在第 106-107 行调用 tap::prepareProfiler(config, activities) 与 tap::enableProfiler(config, activities)(torch::autograd::profiler 命名空间下的遗留 profiler API,PyTorch 1.8+ 已弃用并在部分版本移除),却同时构造 tpi::ProfilerConfig(torch::profiler::impl 的 KINETO 新配置类型),并在第 128 行依赖 at::RecordFunction::currentThreadId() 静态内部方法。这是对台账 [P3](torch 内部头文件与内部 API)的一个新变体:这里不只是 include 内部头文件,而是实际调用了已废弃的遗留函数,跨 torch 版本升级时该测试会编译失败或事件名/线程 ID 语义不一致。
建议:改用 torch::profiler 的 record_function / ProfilerConfig 新 API 重写该测试,或将遗留 API 调用集中到一处并加版本兼容层,避免 prepareProfiler/enableProfiler 与 KINETO 配置的跨命名空间混用。
评审版本:e82a10668cff
Add a configurable output dispatcher worker pool with serial dispatch by default. Propagate inference and device stream context, drain workers before propagating exceptions, and join dispatch work when stopping the engine. Wire the worker count through runtime configuration, Python bindings, and server arguments, with serial/parallel dispatch and shutdown regression coverage.
e82a106 to
360abdf
Compare
| "//rtp_llm/models_py/bindings/core:exec_ops_test_lib", | ||
| "//rtp_llm/models_py/bindings/cuda/ops:cuda_impl", | ||
| ], | ||
| }) + cuda13_torch_link_deps, |
There was a problem hiding this comment.
[P2] 新增缓存更新测试硬编码 CUDA GPU 执行属性,与 ROCm 依赖选择不一致
新增 cc_test 的 deps 使用 select({"@//:using_rocm": ["//rtp_llm/models_py/bindings/core:exec_ctx_rocm"], "//conditions:default": [...]}) 表明该测试需在 ROCm 与 CUDA 两平台编译运行,但 exec_properties 无条件写死 {"gpu": "A10"}(NVIDIA GPU)。同 PR 中 ROCm 专用测试使用 exec_properties = {'gpu':'MI308X-ROCM7', 'gpu_count':'2'} 并带 tags=["rocm"]。在 ROCm 构建/CI 中该目标会以 ROCm 依赖编译,却被调度到不存在的 A10 资源,导致测试无法运行或调度失败。
建议:按平台选择执行属性或使用 device_test_envs()/device_impl_target() 统一处理,例如 select({"@//:using_rocm": {"gpu": "MI308X-ROCM7", "gpu_count": "2"}, "//conditions:default": {"gpu": "A10"}}),并补充相应 tags。
评审版本:360abdf475f8
Summary
SingleTypeKVCacheAllocator.Tests and coverage