已开启
feat: 对齐CPU训练一致性 #168
lan_xin创建于 9 天前
feat: 对齐CPU训练一致性 #168
已开启
lan_xin创建于 9 天前
lan_xin成员
9 天前

问题/功能描述

修改方案描述

开发自检

============================================================
  NPU 索引准备: NpuIndexIVFPQ (nlist=262144, M=32, nbits=8)
============================================================
检测到 16 个 NPU 设备
NPU 资源获取成功
  [distributed] deviceList = [0, 1, 2, 3, 4, 5, 6, 7]
  [distributed] cp.max_points_per_centroid=20
创建索引完成 (0.00s)
  [distributed] deviceList(actual) = [0, 1, 2, 3, 4, 5, 6, 7]
NPU 训练中...
训练完成 (87.72s)
添加向量 (200000000 条)...
添加完成, ntotal=200000000 (637.04s)
nprobe 设置为 32
  [cache] NPU 索引无法直接序列化, 转出 CPU 索引再保存: ivfpq_cache/ivfpq_32_262144_200000000_cea2b34c09904a30.index
  [cache] 已保存 CPU 索引: ivfpq_cache/ivfpq_32_262144_200000000_cea2b34c09904a30.index
NPU 搜索完成 (0.8662s)
NPU 搜索有效 label: 640/640

============================================================
  NPU -> CPU copyTo roundtrip 检查
============================================================
copyTo 完成 (23.88s), ntotal=200000000
copyTo 后 CPU search 完成: valid_labels=640/640, metric=IP
  [cache] 已加载 CPU 索引: ivfpq_cache/ivfpq_32_262144_200000000_cea2b34c09904a30.index (ntotal=200000000)
CPU 搜索完成 (0.0258s)
likedislike
合并受阻
AtlasAccount
AtlasAccount成员
9 天前 评论:

compile

likedislike
atomgit-bot
atomgit-bot
9 天前 评论:

变更摘要

该 PR 以「对齐 CPU 训练一致性」为目标,将 NPU 侧 IVFPQ 的 K-means 训练与赋值流程向 CPU(faiss::Clustering)行为靠拢:为 K-means 引入 nredo 多次重跑并按目标函数择优、使赋值算子同时支持 L2 与内积(IP)度量、将残差编码回退到 CPU 路径以规避已知算子缺陷,并顺带修正算子 kernel 的同步与偏移问题,保证 NPU 训练结果与 CPU 一致。

主要改动

  • K-means 多轮重跑择优:在 IVFPQ::trainKMeansOnNputrainKMeansOnNpuDistributed 中新增 nredo 参数,循环执行训练并用新增的 computeKMeansObjective 按目标值(L2 取最小、IP 取最大)选出最优质心,trainImpl_ 返回值由 void 改为 float objective,对齐 faiss::Clustering 的 redo 语义;同时 NpuIndexIVFPQ 改为透传 pq.cp.nredopq.cp.sphericalmetric_type 等参数。
  • 赋值度量扩展:训练与赋值链路新增 assignmentMetricfaiss::METRIC_L2 / METRIC_INNER_PRODUCT)贯穿 runKMeans_runTrainAssignOp_ 等;assignCentroidOnDevice 的 IP 路径不再构造零质心范数缓冲,改由算子内部处理。
  • ascendc_ivfpq_coarse_assign 算子支持度量类型:算子定义新增必填 metric_type 属性,tiling 校验其取值(0/1)并写入 tilingData.metricType;kernel 按 metricType 分支,IP 用 -q.c 的 argmin 形式(避免范数累积与相消),L2 保持 |q-c|^2 计算。
  • 残差编码回退 CPUIVFPQ::encodeResidualsOnDevice 直接返回 false,因 AscendcIvfpqResidualEncode 算子存在已知 task-step bug(vectorCoreNum_ = usedCoreNum_ * 2 跳过一半任务,导致 dim=1024 残差召回率降至约 2%),改为走 CPU fallback 路径。
  • 一致性细节修正normL2_ 改用 FAISS 的 faiss::fvec_renorm_L2 实现;updateCentroids_ 的计数改为 float(新增 centroidCounts 字段);kernel 的 CopyLabels 增加 PIPE_SPIPE_MTE3 的同步并修正 tail tile 输出偏移(改用 formerTileLength 索引);相关测试将 kDim 由 128 改为 1024,测试数据由均匀分布 np.random.random 改为标准正态 np.random.standard_normal
likedislike
不准确?
atomgit-bot
atomgit-bot
9 天前 评论:

代码审查

✅ 未发现问题

likedislike
不准确?
Xxiangjie10成员
9 天前 添加了label:pr-audit-failed
xiangjie10成员
9 天前 评论:
🔍 PR 规范审计未通过,以下项目需要修正:
  • PR 未关联里程碑或 Issue

请修正后重新提交,或联系仓库管理员。

likedislike
ascend-robotascend-robot成员
9 天前 添加了label:stat/needs-squash
AtlasAccountAtlasAccount成员
9 天前 添加了label:ci-pipeline-running
ascend-robotascend-robot成员
9 天前 添加了label:ascend-cla/yes
ascend-robot
ascend-robot成员
9 天前 评论:

Thanks for your pull-request.
The full list of commands accepted by me can be found at here.
You can get sig-info at here.
You can self-configure the PR merge rules for this repository. For more details, please refer to Here.


PR Approval Progress

⚠️ This PR does not yet meet the following requirements:lgtm (requires ≥ 2 person(s) per module)、approve (requires ≥ 1 person(s) per module)

Module Approval Details

module lgtm status approve status
repo-Ascend/faiss ❌ (0/2)(You can also ask: 陈治涵, 朱伯扬, 易浩, 翁旭涛, 蒋立) ❌ (0/1)(You can also ask: 于金正, 易浩, 蒋立, 王永俊, 向杰)

💡 Tip:

  • Committer can comment /approve or /lgtm
  • Commenting /approve implies both code review (lgtm) and intent to merge (approve)

CLA Signature Pass

gcw_npKwvWT2, thanks for your pull request. All authors of the commits have signed the CLA. 👍

likedislike
AtlasAccountAtlasAccount成员
9 天前 删除了label:ci-pipeline-running
AtlasAccountAtlasAccount成员
9 天前 添加了label:ci-pipeline-passed
AtlasAccount
AtlasAccount成员
9 天前 评论:
流水线 PR-pipeline_faiss#562 [ commitID:b05a52ac ] 已完成
>>>代码风格自动修复执行成功(无修复内容)
阶段 任务名 状态 详情
编译构建 PR-build-SAST-check ✅ COMPLETED >>>
Build_arm ✅ COMPLETED >>>
恶意代码检查 Antipoison ✅ COMPLETED >>>
编码安全与规范检查 pre-commit ✅ COMPLETED >>>
开源片段检查 SCA ✅ COMPLETED >>>
开发者测试 PreSmoke ✅ COMPLETED >>>
流水线 PR-pipeline_faiss ✅ COMPLETED >>>
此流水线已支持下列评论快捷指令,仅PR创建者和白名单成员评论有效
  • compile : 运行流水线
  • retry : 重试流水线所有失败子任务
  • retry <任务名> : 仅重试指定失败子任务
  • stop : 停止流水线
likedislike
lan_xin成员9 天前进行代码检视2
faiss/npu/impl/IVFPQ.cpp
已过期
@@ -168,0 +170,4 @@
170+ const std::vector<float>& data,
171+ const std::vector<float>& centroids,
172+ int64_t metricType) {
173+ float objective = 0.0f;
lan_xin9 天前评论:

严重程度: 建议
问题: computeKMeansObjectiveobjectivefloat,对 totalSize * dim 个分量累加,大训练集下精度不足。
原因: float 仅 24 位尾数。以 totalSize=50000、dim=1024 为例,累加约 5×10⁷ 次,低位会丢失,导致 nredo 各轮 objective 比较不可靠,可能选不到真正最优的 centroids,削弱 nredo 重做机制的意义。
怎么改: 用 double 累加 objective,仅在比较/返回时转 float:

double objective = 0.0;
for (int i = 0; i < totalSize; ++i) {
    ...
    double value = 0.0;
    for (int j = 0; j < dim; ++j) {
        if (metricType == faiss::METRIC_INNER_PRODUCT) {
            value += (double)x[j] * c[j];
        } else {
            const double diff = (double)x[j] - c[j];
            value += diff * diff;
        }
    }
    objective += value;
}
return static_cast<float>(objective);
likedislike
System
系统消息系统
8 天前 评论:

changed this line on 614624cd view diff detail

lan_xin成员9 天前进行代码检视2
faiss/npu/impl/IVFPQ.cpp
已过期
@@ -899,0 +1000,4 @@
1000+ // (vectorCoreNum_ = usedCoreNum_ * 2 skips half of the (vectorTile,
1001+ // subspace) tasks), collapsing dim=1024 residual recall to ~2%. The
1002+ // operator path below is kept but disabled until the bug is fixed.
1003+ return false;
lan_xin9 天前评论:

严重程度: 提示
问题: encodeResidualsOnDevice 开头直接 return false;,后续约 140 行算子路径代码全部不可达。
原因: 为临时禁用存在 task-step bug 的 AscendcIvfpqResidualEncode 算子,但保留 dead code 增加维护负担,且编译器可能不告警,后续维护者易误读为有效逻辑。
怎么改: 用宏或配置开关显式控制,或删除不可达分支并在注释中记录禁用原因与恢复条件:

// 算子存在 task-step bug,暂时走 CPU fallback。
// TODO: bug 修复后恢复算子路径与 dimPerSubQuantizer_ % 8 判定。
return false;
likedislike
System
系统消息系统
8 天前 评论:

changed this line on 614624cd view diff detail

lan_xin成员9 天前进行代码检视2
faiss/npu/impl/IVFPQ.cpp
已过期
@@ -737,2 +790,2 @@
737- for (auto& worker : workers) {
738- worker.get();
790+ CentroidUpdateWorkspace_ centroidWorkspace;
791+ float objective = bestObjective;
lan_xin9 天前评论:

严重程度: 提示
问题: float objective = bestObjective; 初始值在 iter 循环内立即被 computeKMeansObjective 覆盖,初始值无实际作用却易被误读为"以历史最优为基准累加"。
原因: 变量初始化值选择不当,与 nredo 选优语义无关联但形式上产生歧义。
怎么改: 改为中性初值,明确语义:

float objective = 0.0f;
for (int iter = 0; iter < niter; ++iter) {
    ...
    objective = computeKMeansObjective(...);
    ...
}
likedislike
System
系统消息系统
8 天前 评论:

changed this line on 614624cd view diff detail

lan_xin成员9 天前进行代码检视2
faiss/npu/impl/IVFPQ.cpp
已过期
@@ -566,0 +597,4 @@
597+ x,
598+ trainDataHost,
599+ centroidsHost,
600+ actualSeed + redo * 15486557LL,
lan_xin9 天前评论:

严重程度: 提示
问题: actualSeed + redo * 15486557LL 中 15486557 为魔法数字,未定义命名常量也未注释。
原因: 该质数用作 redo 间种子步长以保证各轮初始化互异,但裸数字缺乏自解释性,后续维护者难以判断其用途与取值依据。
怎么改: 提取为命名常量并注释:

// 质数步长,确保 nredo 各轮 initTraining_ 使用互异种子
constexpr int64_t kRedoSeedStride = 15486557LL;
...
initTraining_(..., actualSeed + redo * kRedoSeedStride, spherical);
likedislike
System
系统消息系统
8 天前 评论:

changed this line on 614624cd view diff detail

lan_xin成员9 天前进行代码检视2
faiss/npu/ops/op_impl/ascendc_ivfpq_coarse_assign/op_kernel/ascendc_ivfpq_coarse_assign.h
已过期
@@ -271,3 +263,1 @@
271- codesDoubleLocal,
272- distResultLocal[i * this->codesLength],
273- this->codesLength);
263+ if (this->metricType == 0) {
lan_xin9 天前评论:

严重程度: 提示
问题: Kernel 内 if (this->metricType == 0) 与 op_host tiling 校验 *metricType != 0 && *metricType != 1 均使用裸数字 0/1,未用命名常量。
原因: 0/1 隐式对应 faiss::METRIC_INNER_PRODUCT(0) 与 faiss::METRIC_L2(1),但 Kernel 侧无法直接引用 faiss 枚举。裸数字可读性差,且若上游枚举值变更会静默出错。
怎么改: 在 Kernel 头文件定义本地命名常量并注释映射关系:

// 与 faiss::METRIC_INNER_PRODUCT(0) / METRIC_L2(1) 对应
constexpr int32_t kMetricInnerProduct = 0;
constexpr int32_t kMetricL2 = 1;
...
if (this->metricType == kMetricInnerProduct) { ... }
likedislike
System
系统消息系统
8 天前 评论:

changed this line on 614624cd view diff detail

lan_xin成员9 天前进行代码检视1
faiss/npu/NpuIndexIVFPQ.cpp
@@ -735,0 +730,4 @@
730+ pq.cp.nredo,
731+ pq.cp.seed,
732+ pq.cp.spherical,
733+ faiss::METRIC_L2,
lan_xin9 天前评论:

严重程度: 提示
问题: PQ 子量化器训练 indexTrainImpl_(..., faiss::METRIC_L2, ...) 硬编码 L2 度量,且无注释说明,与 coarse quantizer 使用 metric_type 的做法不一致。
原因: PQ 残差编码始终使用 L2 是 faiss 的标准做法(残差范数最小化),但缺少注释使该设计意图不明显,后续维护者可能误改为跟随索引 metric_type。
怎么改: 补充注释说明设计依据:

// PQ 残差量化始终使用 L2:残差编码以 ||x - c||² 最小化为目标,
// 与索引的 coarse metric 无关(faiss CPU 路径同此约定)。
indexTrainImpl_(..., pq.cp.spherical, faiss::METRIC_L2, trainedCentroids[m]);
likedislike
lan_xin成员9 天前进行代码检视2
faiss/npu/impl/IVFPQ.cpp
已过期
@@ -1294,30 +1425,30 @@
12941425 if (useBucketedUpdate) {
12951426 // Use fixed logical chunks instead of OpenMP thread ids so the bucket
12961427 // layout remains valid even when OpenMP changes the active team size.
@@ -1314,13 +1445,14 @@ void IVFPQ::updateCentroids_(
13141445 count += workspace.threadOffsets
13151446 [static_cast<size_t>(rank) * nlist + ci];
13161447 }
1317- hassign[ci] = count;
1448+ workspace.clusterCounts[ci] = count;
1449+ hassign[ci] = static_cast<float>(count);
13181450 }
13191451 
13201452 workspace.clusterOffsets[0] = 0;
13211453 for (int ci = 0; ci < nlist; ++ci) {
13221454 workspace.clusterOffsets[ci + 1] =
1323- workspace.clusterOffsets[ci] + hassign[ci];
1455+ workspace.clusterOffsets[ci] + workspace.clusterCounts[ci];
13241456 }
13251457 
13261458#pragma omp parallel for
@@ -1380,7 +1512,7 @@ void IVFPQ::updateCentroids_(
13801512 0,
13811513 static_cast<size_t>(nlist) * static_cast<size_t>(dim) *
13821514 sizeof(float));
1383- hassign.assign(nlist, 0);
1515+ hassign.assign(nlist, 0.0f);
13841516#pragma omp parallel
13851517 {
13861518 const int nt = omp_get_num_threads();
@@ -1408,7 +1540,7 @@ void IVFPQ::updateCentroids_(
14081540 if (hassign[ci] == 0) {
lan_xin9 天前评论:

严重程度: 提示
问题: if (hassign[ci] == 0)hassign 已由 std::vector<int> 改为 std::vector<float>,对 float 与 0 直接 == 比较风格不佳。
原因: 虽当前赋值均为整数转 0.0f,比较精确成立,但 float == 比较在风格检查中通常告警,且若未来引入浮点累加会埋下隐患。
怎么改: 使用范围比较更稳健:

if (hassign[ci] <= 0.0f) {
    continue;
}
likedislike
System
系统消息系统
8 天前 评论:

changed this line on 614624cd view diff detail

Llan_xin成员
8 天前 推送  1 个提交:614624cd-fix code review
ascend-robot
ascend-robot成员
8 天前 评论:

compile

likedislike
ascend-robotascend-robot成员
8 天前 删除了label:ci-pipeline-passed
ascend-robotascend-robot成员
8 天前 删除了label:ascend-cla/yes
ascend-robotascend-robot成员
8 天前 添加了label:ascend-cla/yes
ascend-robotascend-robot成员
8 天前 添加了label:ci-pipeline-running
ascend-robotascend-robot成员
8 天前 删除了label:ci-pipeline-running
ascend-robotascend-robot成员
8 天前 添加了label:ci-pipeline-failed
ascend-robot
ascend-robot成员
8 天前 评论:
流水线 PR-pipeline_faiss#571 [ commitID:614624cd ] 运行失败
>>>代码风格自动修复执行成功(无修复内容)
阶段 任务名 状态 详情
编译构建 PR-build-SAST-check ✅ COMPLETED >>>
Build_arm ✅ COMPLETED >>>
恶意代码检查 Antipoison ✅ COMPLETED >>>
编码安全与规范检查 pre-commit ✅ COMPLETED >>>
开源片段检查 SCA ✅ COMPLETED >>>
开发者测试 PreSmoke ✅ COMPLETED >>>
UT ❌ FAILED >>>
流水线 PR-pipeline_faiss ❌ FAILED >>>
此流水线已支持下列评论快捷指令,仅PR创建者和白名单成员评论有效
  • compile : 运行流水线
  • retry : 重试流水线所有失败子任务
  • retry <任务名> : 仅重试指定失败子任务
  • stop : 停止流水线
likedislike
Llan_xin成员
7 天前 修改了pull request 的描述
yihao1234成员6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -168,0 +187,4 @@
187+ }
188+ objective += value;
189+ }
190+ return static_cast<float>(objective);
yihao12346 天前评论:

严重程度: 建议

问题: computeKMeansObjective 内部使用 double 累加目标函数值,但返回类型为 float,大规模数据集下可能溢出导致 redo 比较失效

原因: 该函数遍历 totalSize 条向量计算 K-means 目标函数,内部用 double 累加避免了精度问题,但最终 static_cast(objective) 将结果截断为 float。当训练数据量大(如 PR 描述中的 2 亿向量 x 1024 维)且向量值较大时,累加后的 double 值可能超过 float 的最大值(约 3.4e38),导致返回 +inf。在 redo 循环中,如果所有 redo 的 objective 都是 +inf,比较 objective < bestObjective(L2)永远为 false,bestCentroids 将保持为空,最终 centroidsOut = std::move(bestCentroids) 会产生空向量,导致后续使用崩溃。

怎么改:
建议将 computeKMeansObjective 的返回类型改为 double,同时将 trainImpl_ 和 redo 循环中的 bestObjective/objective 变量也改为 double。或者,在 redo 循环中添加对 bestCentroids 为空的防护:如果所有 redo 都未产生更好的结果(例如因溢出),保留最后一次的 centroidsHost 作为兜底。

likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -1074,0 +1181,4 @@
1181+ true,
1182+ assignmentMetric);
1183+ 
1184+ objective = computeKMeansObjective(
yihao12346 天前评论:

严重程度: 建议

问题: computeKMeansObjective 在 trainImpl_ 和 trainKMeansOnNpuDistributed 的每次 K-means 迭代中都被调用,但只有最后一次迭代的目标值用于 redo 间的比较

原因: computeKMeansObjective 的复杂度为 O(totalSize * dim),对于大规模训练数据(数百万向量 x 数百维度),这是一次昂贵的全量 CPU 计算。在 niter 次迭代中,每次迭代后都调用该函数计算目标值,但 trainImpl_ 只返回最后一次的 objective,trainKMeansOnNpuDistributed 中也只取最后一次迭代后的 objective 进行 redo 比较。前 niter-1 次的计算结果被丢弃,造成不必要的性能开销。

怎么改:
在 trainImpl_ 中,仅在最后一次迭代(iter == niter - 1)时调用 computeKMeansObjective,而非每次迭代都调用。在 trainKMeansOnNpuDistributed 中同样可以仅在最后一次迭代后计算 objective。

likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -899,0 +999,4 @@
999+ // behaves like dim=128 (where dsub%8!=0 already used CPU). The
1000+ // AscendcIvfpqResidualEncode operator has a known task-step bug
1001+ // (vectorCoreNum_ = usedCoreNum_ * 2 skips half of the (vectorTile,
1002+ // subspace) tasks), collapsing dim=1024 residual recall to ~2%. The
yihao12346 天前评论:

严重程度: 提示

问题: encodeResidualsOnDevice 中新增的注释末尾不完整,以 "The" 结束后直接接代码

原因: 注释最后一句为 'collapsing dim=1024 residual recall to ~2%. The',以 'The' 结尾后直接跟 if 语句,似乎是一句未写完的话。读者无法确定 'The' 后面想表达什么内容,可能是 'The following check...' 或其他说明。这影响了注释的可读性和可维护性。

怎么改:
补全注释最后一句,例如 'The following condition forces CPU fallback for all dimensions affected by this bug.' 或直接删除未完成的 'The'。

likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -567,0 +621,4 @@
621+ const bool isBetter = assignmentMetric == faiss::METRIC_INNER_PRODUCT
622+ ? objective > bestObjective
623+ : objective < bestObjective;
624+ if (isBetter) {
yihao12346 天前评论:

严重程度: 建议

问题: redo 循环结束后直接 std::move(bestCentroids) 到 centroidsOut,未处理 bestCentroids 可能为空的情况

原因: 当 nredo=1 时,第一次 redo 的 objective 会与初始值(HUGE_VALF 或 -HUGE_VALF)比较。正常情况下第一次结果总是 'better',bestCentroids 会被赋值。但如果 objective 因浮点溢出变为 +inf(L2)或 -inf(IP),比较结果为 false,bestCentroids 保持为空,std::move 后 centroidsOut 为空向量,后续代码访问空向量的数据会导致未定义行为或崩溃。trainKMeansOnNpuDistributed 中存在同样的问题。

怎么改:
在 redo 循环后添加防护:if (bestCentroids.empty()) { bestCentroids = centroidsHost; } 确保即使比较逻辑异常,也有兜底结果。或在循环中保证至少第一次 redo 的结果总是被采纳。

likedislike
duliqiang6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -896,1 +1000,4 @@
1000+ // AscendcIvfpqResidualEncode operator has a known task-step bug
1001+ // (vectorCoreNum_ = usedCoreNum_ * 2 skips half of the (vectorTile,
1002+ // subspace) tasks), collapsing dim=1024 residual recall to ~2%. The
8961003 if (dimPerSubQuantizer_ % 8 != 0) {
duliqiang6 天前评论:

严重程度: 严重

问题: encodeResidualsOnDevice 注释声称「强制走 CPU fallback 以规避 dim=1024 task-step bug」,但实际仍只在 dimPerSubQuantizer_ % 8 != 0return false;注释还以残缺的 The 结尾。

原因: 典型配置 dim=1024, M=32dsub=32% 8 == 0,会继续进入下方 AscendcIvfpqResidualEncode 算子路径。注释描述的防护并未生效,且与「对齐 CPU 训练一致性 / 修复残差召回约 2%」的目标相矛盾。

怎么改:

// AscendcIvfpqResidualEncode 存在 task-step bug(vectorCoreNum_ = usedCoreNum_ * 2),
// 会导致 dim=1024 残差召回率降至约 2%。bug 修复前一律走 CPU fallback。
// TODO: 算子修复后恢复 NPU 路径与 dsub%8 判定。
return false;

或保留算子路径,但删除/改写误导性注释,并补全末尾残缺句子。

likedislike
duliqiang6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -567,0 +626,4 @@
626+ bestCentroids = centroidsHost;
627+ }
628+ }
629+ centroidsOut = std::move(bestCentroids);
duliqiang6 天前评论:

严重程度: 建议

问题: nredo 循环结束后直接 centroidsOut = std::move(bestCentroids);,未处理 bestCentroids 仍为空的情况(单卡与 distributed 两处同构)。

原因:objective 因 float 截断/溢出变成 ±inf,或出现 NaN 时,isBetter 比较失败,bestCentroids 保持默认空向量;后续访问会 UB/崩溃。与 computeKMeansObjective 返回 float 的风险叠加。

怎么改:

if (bestCentroids.empty()) {
    bestCentroids = centroidsHost;  // 至少兜底最后一轮
}
centroidsOut = std::move(bestCentroids);

或保证第一次 redo 无条件采纳;同时建议 objective 全程用 doubletrainKMeansOnNpuDistributed 中对称处理。

likedislike
duliqiang6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -168,0 +187,4 @@
187+ }
188+ objective += value;
189+ }
190+ return static_cast<float>(objective);
duliqiang6 天前评论:

严重程度: 建议

问题: computeKMeansObjective 内部用 double 累加,但 return static_cast<float>(objective),返回类型仍是 float

原因: 大规模训练(如 PR 自检中的 2e8×1024)下,目标值可能超出 float 范围变成 ±inf,使 nredo 择优失效,并与「空 bestCentroids」风险耦合。精度层面 float 累加比较也不如 double 稳。

怎么改: 将返回类型及 trainImpl_ / redo 循环中的 bestObjective/objective 改为 double;若必须对外保持 float,至少在比较前检测 !std::isfinite(objective) 并跳过/兜底。

likedislike
duliqiang6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -1074,0 +1181,4 @@
1181+ true,
1182+ assignmentMetric);
1183+ 
1184+ objective = computeKMeansObjective(
duliqiang6 天前评论:

严重程度: 建议

问题: trainImpl_(以及 distributed 路径)在每次 K-means iter 后都调用 computeKMeansObjective,但 redo 间比较只用最后一次返回值。

原因: 复杂度为 O(totalSize * dim) 的全量 CPU 扫表;前 niter-1 次结果被丢弃,大规模训练时白白放大耗时。

怎么改:

if (iter == niter - 1) {
    objective = computeKMeansObjective(...);
}

distributed 路径同样仅在最后一轮 iter 后计算。

likedislike
duliqiang6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -168,0 +190,4 @@
190+ return static_cast<float>(objective);
191+}
192+ 
193+// Force exact CPU labels after the NPU training assignment. This isolates
duliqiang6 天前评论:

问题: computeKMeansObjective 结束后紧跟一段「Force exact CPU labels after the NPU training assignment…」注释,却挂在 verifyTrainingCoreCount 上,语义不匹配。

likedislike
duliqiang6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -1050,3 +1158,4 @@
10501158 int totalSize = (int)(trainDataHost.size() / dim);
10511159 std::vector<int32_t> totalAssigns(totalSize);
10521160 CentroidUpdateWorkspace_ centroidWorkspace;
1161+ float objective = assignmentMetric == faiss::METRIC_INNER_PRODUCT
duliqiang6 天前评论:

问题: trainImpl_float objective = ±HUGE_VALF; 的初值在循环内第一次 computeKMeansObjective 时即被覆盖,语义上像「带历史最优基准」,实际无用。

likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/NpuIndexIVFPQ.cpp
@@ -165,7 +165,7 @@
165165 index->metric_arg,
166166 index->nlist,
167167 normalizeIVFPQConfig(config)),
168168 pq(index->pq),
yihao12346 天前评论:

严重程度: 建议

问题: 常量 IVF_PQ_REDO_SEED_STRIDE = 15486557LL 是一个魔法数字,没有注释解释其取值依据。该值用于在 K-means 多次重做(nredo)时为每次重做生成不同的随机种子(actualSeed + redo * IVF_PQ_REDO_SEED_STRIDE),但为什么选择 15486557 而非其他值未做说明。

原因: 在随机数生成中,种子间隔值的选择会影响不同 redo 之间随机序列的独立性。如果间隔值不够大或者与内部随机数生成器的周期存在公约数,可能导致不同 redo 的随机序列存在相关性,降低 K-means 多次重做寻找最优聚类的效果。缺乏注释使得后续维护者无法判断此值是否经过选择验证,也不清楚是否可以随意修改。此外该常量在分布式训练路径(trainKMeansOnNpuDistributed)中也使用相同的间隔值,两处使用需要保持一致。

怎么改:
建议添加注释说明取值依据,例如:

// 使用质数作为种子间隔,确保不同 redo 之间的随机种子有足够的差异。
// 15486557 是一个大于常见训练规模的质数,避免与 K-means 内部采样数量产生周期重叠。
constexpr int64_t IVF_PQ_REDO_SEED_STRIDE = 15486557LL;

如果该值是随意选取的,考虑使用更标准的做法,如直接对 actualSeed 做哈希或使用 std::mt19937_64 重新生成种子。

likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/NpuIndexIVFPQ.cpp
@@ -688,8 +688,7 @@ void NpuIndexIVFPQ::trainPQCodeBook_(
688688 std::vector<float> sampled;
689689 const float* trainData = x;
690690 if (trainCount < n) {
691- sampled = sampleTrainData(
692- x, n, this->d, trainCount, ivfpqConfig_.cp.seed);
691+ sampled = sampleTrainData(x, n, this->d, trainCount, pq.cp.seed);
693692 trainData = sampled.data();
694693 }
695694 
@@ -708,7 +707,7 @@ void NpuIndexIVFPQ::trainPQCodeBook_(
708707 }
709708 
710709 std::vector<std::vector<float>> trainedCentroids(pq.M);
711- const int pqNiter = ivfpqConfig_.cp.niter;
710+ const int pqNiter = pq.cp.niter;
712711 const size_t workerCount =
713712 std::min(devices.size(), static_cast<size_t>(pq.M));
714713 std::vector<std::future<void>> workers;
@@ -728,7 +727,10 @@ void NpuIndexIVFPQ::trainPQCodeBook_(
728727 devices[worker],
729728 trainingResources[worker].get(),
730729 pqNiter,
yihao12346 天前评论:

严重程度: 建议

问题: 在 trainPQCodeBook_ 中调用 indexTrainImpl_ 时,assignmentMetric 参数被硬编码为 faiss::METRIC_L2,而粗量化器训练(train 方法中)使用的是 metric_type。PQ 子量化器训练始终使用 L2 距离进行分配,即使整体索引使用内积度量(metric_type == METRIC_INNER_PRODUCT)。

原因: 如果索引配置为内积度量(如 spherical=true 时),粗量化器训练使用 metric_type(可能是 INNER_PRODUCT)进行分配,但 PQ 子量化器训练却固定使用 L2。这种不一致可能是设计意图(PQ 编码残差时通常使用 L2),但也可能是遗漏。在原始代码中,trainPQCodeBook_ 使用 ivfpqConfig_.cp.spherical 控制 NPU 训练行为,但没有传递 metric type。新代码引入了 assignmentMetric 参数但只在 PQ 训练中固定为 L2,这与粗量化器的 metric_type 传递方式不一致,可能导致内积索引的 PQ 编码质量次优。

怎么改:
如果 PQ 训练确实应始终使用 L2(这是 faiss CPU 实现中的常见做法,因为 PQ 编码的是残差向量),建议在代码中添加注释说明:

// PQ 子量化器始终使用 L2 距离进行聚类分配,因为 PQ 编码的是残差向量,
// 残差空间中的距离度量与原始空间的度量类型无关。
faiss::METRIC_L2,

如果 PQ 训练应跟随索引的 metric_type,则应将 metric_type 传递给 trainPQCodeBook_ 而非硬编码 L2。

likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -566,1 +588,3 @@
566- centroidsOut.assign(centroidsHost.begin(), centroidsHost.end());
588+ std::vector<float> bestCentroids;
589+ float bestObjective = assignmentMetric == faiss::METRIC_INNER_PRODUCT
590+ ? -HUGE_VALF
yihao12346 天前评论:

严重程度: 建议

问题: trainImpl_ 函数的返回类型从 void 改为 float(返回 K-means 目标函数值),但 trainKMeansOnNpuDistributed 中的分布式训练路径也调用了 trainImpl_,其返回值的使用方式需要验证。在分布式路径中,多个 worker 各自训练子片段,直接比较单个 worker 的 objective 值可能不反映全局聚类质量。

原因: 在单机路径(trainKMeansOnNpu)中,trainImpl_ 返回的 objective 值用于比较不同 redo 的聚类质量,选择最优质心。但在分布式路径(trainKMeansOnNpuDistributed)中,每个 worker 只处理数据的一个分片,trainImpl_ 返回的是局部 objective。如果直接使用局部 objective 进行 redo 间的比较,可能选择的是局部最优而非全局最优。需要确认分布式路径是否聚合了所有 worker 的 objective 值。

怎么改:
建议检查分布式训练路径中 objective 的聚合逻辑。如果尚未实现聚合,添加类似以下代码:

// 在分布式路径中聚合各 worker 的 objective
float totalObjective = 0.0;
for (auto& f : workerFutures) {
    totalObjective += f.get();
}
// 使用 totalObjective 进行 redo 间比较

如果分布式路径目前不支持 nredo(nredo 始终为 1),则应添加断言或注释说明此限制。

likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -168,0 +186,4 @@
186+ }
187+ }
188+ objective += value;
189+ }
yihao12346 天前评论:

严重程度: 建议

问题: computeKMeansObjective 内部使用 double 累积目标值,但返回类型为 float。对于大规模训练集(如 PR 描述中提到的 200M 向量场景,即使经过子采样),目标值累加可能超出 float 精度范围,导致不同 redo 之间的比较不可靠。

原因: K-means 的 redo 选择依赖于 objective 的相对大小比较。当训练数据量较大或维度较高时,sum over all vectors 的 objective 值可能很大,float 的 23 位尾数(约 7 位有效数字)可能无法区分相近的 objective 值,使得 redo 选择退化为随机。内部已经用 double 累积说明作者意识到了精度需求,但返回时的 float 截断浪费了这份精度。

怎么改:
将返回类型改为 double,同时修改调用方 trainImpl_ 和 trainKMeansOnNpuDistributed 中 bestObjective 的类型为 double:

static double computeKMeansObjective(...)
{
    ...
    return objective;  // 直接返回 double,无需 cast
}
likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -898,30 +1005,30 @@
8981005 }
@@ -1028,7 +1135,7 @@ bool IVFPQ::encodeResidualsOnDevice(
10281135}
10291136 
10301137// Training loop: runKMeans_ + updateCentroids_ + normL2_ per iteration.
1031-void IVFPQ::trainImpl_(
1138+float IVFPQ::trainImpl_(
10321139 NpuResources* resources,
10331140 int n,
10341141 const float* x,
@@ -1043,13 +1150,17 @@ void IVFPQ::trainImpl_(
10431150 std::vector<float>& centroidsHost,
10441151 std::vector<float>& trainDataHost,
10451152 int niter,
1046- bool spherical) {
1153+ bool spherical,
1154+ int64_t assignmentMetric) {
10471155 (void)n;
10481156 (void)x;
10491157 
10501158 int totalSize = (int)(trainDataHost.size() / dim);
10511159 std::vector<int32_t> totalAssigns(totalSize);
10521160 CentroidUpdateWorkspace_ centroidWorkspace;
1161+ float objective = assignmentMetric == faiss::METRIC_INNER_PRODUCT
1162+ ? -HUGE_VALF
1163+ : HUGE_VALF;
10531164 for (int iter = 0; iter < niter; iter++) {
10541165 updateCentroidsToDevice_(
10551166 nlist, dim, deviceDim, centroidsHost, centroidsDev);
@@ -1067,7 +1178,16 @@ void IVFPQ::trainImpl_(
10671178 centroidsDev,
yihao12346 天前评论:

严重程度: 建议

问题: computeKMeansObjective 在 trainImpl_ 的每次 K-means 迭代中都被调用,但只有最后一次迭代的返回值被使用(用于 redo 间的最优选择)。同样的模式也出现在 trainKMeansOnNpuDistributed 中。

原因: computeKMeansObjective 遍历全部训练向量 × 全部维度,复杂度为 O(totalSize × dim)。在典型配置下(如 niter=25, nredo=3),该函数会被调用 75 次,但仅最后一次的结果有效。前 74 次计算完全浪费 CPU 时间,且在 NPU 训练流程中这段 CPU 计算会阻塞流水线,降低整体训练吞吐。

怎么改:
将 objective 计算移到迭代循环之后,或仅在最后一次迭代时计算:

for (int iter = 0; iter < niter; iter++) {
    // ... runKMeans_, updateCentroids_ ...
    if (iter == niter - 1) {
        objective = computeKMeansObjective(
                dim, totalSize, totalAssigns,
                trainDataHost, centroidsHost, assignmentMetric);
    }
}

同理修改 trainKMeansOnNpuDistributed 中的对应逻辑。

likedislike
yihao1234成员6 天前进行代码检视1
faiss/npu/impl/IVFPQ.cpp
@@ -1195,30 +1322,30 @@
11951322 stream,
11961323 workspaceDev);
11971324 }
@@ -1225,6 +1352,7 @@ void IVFPQ::runTrainAssignOp_(
12251352 int32_t* labelsDev,
12261353 int nlist,
12271354 int dim,
1355+ int64_t assignmentMetric,
12281356 aclrtStream stream,
12291357 DeviceVector<uint8_t>& workspaceDev) {
12301358 (void)resources;
@@ -1247,6 +1375,7 @@ void IVFPQ::runTrainAssignOp_(
12471375 queryTensor.get(),
12481376 centroidsTensor.get(),
12491377 centroidsSqrTensor.get(),
1378+ assignmentMetric,
12501379 labelsTensor.get(),
12511380 &workspaceSize,
12521381 &executor));
@@ -1280,6 +1409,7 @@ void IVFPQ::updateCentroids_(
12801409 static_cast<size_t>(maxThreads) * nlist);
12811410 workspace.clusterOffsets.resize(static_cast<size_t>(nlist) + 1);
12821411 workspace.clusterCounts.resize(nlist);
1412+ workspace.centroidCounts.resize(nlist);
12831413 workspace.bucketedTrainData.resize(
12841414 static_cast<size_t>(totalSize) * dim);
12851415 } catch (const std::bad_alloc&) {
@@ -1288,9 +1418,9 @@ void IVFPQ::updateCentroids_(
12881418 }
12891419 }
12901420 
yihao12346 天前评论:

严重程度: 提示

问题: updateCentroids_ 中 hassign(聚类计数)的类型从 std::vector 改为 std::vector,同时新增了 workspace.centroidCounts(float)与原有 workspace.clusterCounts(int)并存。

原因: 聚类计数本质上是整数值,使用 float 存储虽然对于典型 nlist 范围内精确无误,但语义上不够清晰。同时 clusterCounts(int)和 centroidCounts(float)并存容易让后续维护者困惑两者的关系和用途。

怎么改:
如果 float 的引入是为了后续支持加权聚类(非整数权重),建议添加注释说明设计意图。如果仅是为了统一类型,考虑直接将 clusterCounts 也改为 float 并合并为同一字段,减少冗余。

likedislike