diff --git a/09_vector_retrieval/ccyyy1023/.gitignore b/09_vector_retrieval/ccyyy1023/.gitignore
new file mode 100644
index 00000000..e862c7fb
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/.gitignore
@@ -0,0 +1,5 @@
+/build/
+/data/
+gvs_test_index.bin
+*.tmp
+
diff --git a/09_vector_retrieval/ccyyy1023/CMakeLists.txt b/09_vector_retrieval/ccyyy1023/CMakeLists.txt
new file mode 100644
index 00000000..bd874bac
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/CMakeLists.txt
@@ -0,0 +1,39 @@
+cmake_minimum_required(VERSION 3.20)
+project(gpu_vector_retrieval LANGUAGES CXX CUDA)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CUDA_STANDARD 17)
+set(CMAKE_CUDA_STANDARD_REQUIRED ON)
+
+option(GVS_BUILD_TESTS "Build correctness tests" ON)
+
+find_package(CUDAToolkit REQUIRED)
+find_package(OpenMP)
+
+add_library(gvs
+    src/cpu_search.cpp
+    src/file_io.cpp
+    src/ivf_flat.cpp
+    src/gpu_search.cu
+)
+target_include_directories(gvs PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/include)
+target_link_libraries(gvs PUBLIC CUDA::cudart)
+if(OpenMP_CXX_FOUND)
+    target_link_libraries(gvs PUBLIC OpenMP::OpenMP_CXX)
+endif()
+set_target_properties(gvs PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+target_compile_options(gvs PRIVATE
+    $<$<COMPILE_LANGUAGE:CXX>:-O3>
+    $<$<COMPILE_LANGUAGE:CUDA>:-O3;--use_fast_math>
+)
+
+add_executable(vector_search src/main.cpp)
+target_link_libraries(vector_search PRIVATE gvs)
+
+if(GVS_BUILD_TESTS)
+    enable_testing()
+    add_executable(gvs_tests tests/test_search.cpp)
+    target_link_libraries(gvs_tests PRIVATE gvs)
+    add_test(NAME gvs_correctness COMMAND gvs_tests)
+endif()
diff --git "a/09_vector_retrieval/ccyyy1023/GPU\345\220\221\351\207\217\346\243\200\347\264\242\345\274\225\346\223\216_\346\200\273\347\273\223\346\212\245\345\221\212.pdf" "b/09_vector_retrieval/ccyyy1023/GPU\345\220\221\351\207\217\346\243\200\347\264\242\345\274\225\346\223\216_\346\200\273\347\273\223\346\212\245\345\221\212.pdf"
new file mode 100644
index 00000000..1f8fc236
Binary files /dev/null and "b/09_vector_retrieval/ccyyy1023/GPU\345\220\221\351\207\217\346\243\200\347\264\242\345\274\225\346\223\216_\346\200\273\347\273\223\346\212\245\345\221\212.pdf" differ
diff --git a/09_vector_retrieval/ccyyy1023/README.md b/09_vector_retrieval/ccyyy1023/README.md
new file mode 100644
index 00000000..2f6ea01f
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/README.md
@@ -0,0 +1,155 @@
+# GPU 向量检索引擎（CUDA）
+
+本项目实现了一个可独立构建的简化 GPU 向量检索引擎，面向训练营第九题。当前实现包含 CPU 精确检索参考、CUDA 批量精确检索、L2/内积、排序后的 Top-K、IVF-Flat、索引持久化、正确性测试，以及 QPS、P50/P99、recall@K、平均分数误差、显存工作区和 CPU/GPU 加速比的采集框架。
+
+> 真实性声明：`results/*.json` 来自 2026-09-20 在 RTX 4090 24GB、CUDA 11.8 上的真实运行。百万规模正式矩阵覆盖 1,000,000 条 128 维向量和 1000 个查询；原始 JSON、环境记录和汇总表均随项目提交。
+
+## 1. 已完成范围
+
+- CPU FP32 精确检索 baseline。
+- CUDA FP32 精确检索，支持 `l2` 和 `inner_product`。
+- 批量查询，`batch_size` 可配置。
+- 任意正整数 `top_k`，已提供 K=1/10/50/100 的基准脚本。
+- Top-K 按 L2 从小到大或内积从大到小输出。
+- IVF-Flat：CPU mini-batch Lloyd 聚类、倒排桶、`nprobe` 探测、数据库常驻 GPU、按候选 ID 计算距离。
+- IVF 索引保存/加载，查询时不需要重新聚类。
+- 固定版本的二进制向量、查询、结果和 IVF 索引格式。
+- CPU/GPU 一致性测试、索引 round-trip 测试。
+- JSON 性能/质量日志与汇总脚本。
+
+实现边界：基础版本仅支持 FP32；IVF 聚类、coarse search 和最终 Top-K 归并在 CPU 上完成，主要距离计算在 GPU 上完成。这个版本不声称实现了 GPU Top-K、FP16、IVF-PQ 或国产 GPU 适配。
+
+## 2. 环境
+
+- NVIDIA GPU
+- CUDA Toolkit 11.0 或更新版本（需要 `nvcc`）
+- CMake 3.20+
+- 支持 C++17 的编译器
+- Python 3.8+（仅用于日志汇总和包检查）
+
+推荐 Ubuntu 22.04 + CUDA 12.x。Windows 也可使用 Visual Studio 2022 的 C++ 工具链构建。
+
+## 3. 构建与测试
+
+```bash
+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
+cmake --build build -j
+ctest --test-dir build --output-on-failure
+```
+
+测试同时覆盖 L2、inner product、批量精确 Top-K、CPU/GPU 一致性和 IVF 索引保存/加载。
+
+## 4. 五分钟 smoke test
+
+```bash
+bash scripts/run_smoke.sh
+```
+
+等价的手动命令：
+
+```bash
+mkdir -p data results
+./build/vector_search generate \
+  --db data/small_db.bin --queries data/small_queries.bin \
+  --n 10000 --nq 100 --dim 128 --metric l2 --seed 42
+
+./build/vector_search build \
+  --db data/small_db.bin --index data/small_ivf.bin \
+  --nlist 64 --iterations 8 --seed 42
+
+./build/vector_search benchmark \
+  --db data/small_db.bin --queries data/small_queries.bin \
+  --index data/small_ivf.bin --log results/smoke.json \
+  --k 10 --batch 32 --nlist 64 --nprobe 8 \
+  --cpu-query-limit 20 --latency-query-limit 50
+```
+
+## 5. 题目要求规模测试
+
+以下脚本生成 `1,000,000 x 128` 向量库和 1000 个查询，测试 K=1/10/50/100、多个 batch size 和多个 nprobe。数据库约 488 MiB，索引还会额外保存一份向量数据，请预留至少 2 GiB 磁盘空间。
+
+```bash
+bash scripts/run_required_benchmarks.sh
+python3 scripts/summarize_results.py results results/SUMMARY.md
+```
+
+正式测试前请关闭其他 GPU 任务，并记录：GPU 型号、驱动、CUDA 版本、CPU 型号和编译命令。脚本不会伪造或预填任何数字。
+
+## 6. 文件模式
+
+`configs/exact.conf` 和 `configs/ivf_flat.conf` 遵循题面参数形式：
+
+```text
+top_k = 10
+search_mode = "ivf_flat"
+batch_size = 128
+nlist = 256
+nprobe = 16
+```
+
+构建一次索引并从文件加载查询：
+
+```bash
+./build/vector_search build --db data/db.bin --index data/index.bin --nlist 256
+./build/vector_search search \
+  --db data/db.bin --queries data/queries.bin \
+  --params configs/ivf_flat.conf --index data/index.bin \
+  --output results/search_results.bin
+```
+
+精确检索不使用 `--index`，但命令行解析允许保留该参数：
+
+```bash
+./build/vector_search search \
+  --db data/db.bin --queries data/queries.bin \
+  --params configs/exact.conf --output results/exact_results.bin
+```
+
+## 7. 二进制格式
+
+所有整数均为 little-endian，向量按行主序存储。
+
+- 数据库：8 字节 magic `GVSDB01\0`，`uint64 n`，`uint32 dim`，`uint32 dtype`，`uint32 metric`，随后 FP32 数据。`dtype=1`，`metric=0/1` 分别表示 L2/inner product。
+- 查询：8 字节 magic `GVSQR01\0`，`uint64 nq`，`uint32 dim`，`uint32 dtype`，随后 FP32 数据。
+- 结果：8 字节 magic `GVSRS01\0`，`uint64 nq`，`uint32 k`，随后每个结果的 `int64 id + float score`。
+- IVF 索引：版本化 magic、形状/度量、聚类中心、原始 FP32 向量和倒排表。加载时会验证 header、总 ID 数和 ID 范围。
+
+## 8. 指标定义
+
+- `QPS = 查询数 / 查询时间`。精确检索计时包含 query H2D、CUDA 距离计算、distance D2H 和 CPU Top-K，不包含数据库首次上传；IVF 计时还包含 CPU coarse search、候选 ID 构造和一次数据库 H2D，索引文件加载不计入。
+- P50/P99：逐 query、batch=1 的端到端延迟分位数。
+- `recall@K`：近似结果与 GPU 精确结果 Top-K ID 集合的平均交集比例。
+- `mean_rank_score_error`：近似结果和精确结果在相同排名位置的距离/相似度绝对误差平均值。
+- `workspace_mib`：程序显式分配的数据库、query 和 distance buffer 总量估计，不包含 CUDA context/驱动开销。
+- CPU/GPU speedup：CPU 单 query 平均时间除以 GPU 批量查询的单 query 平均时间。CPU 默认只测前 20 个 query，并在日志中明确记录样本数。
+
+## 9. 设计说明
+
+精确检索 kernel 使用二维 grid：x 维覆盖数据库向量，y 维覆盖 batch 内 query。每个线程顺序归约一个向量的 128 个维度，行主序布局使相邻线程访问相邻向量。L2 使用 `fmaf(diff, diff, acc)`，内积内部取负数以复用升序 Top-K。
+
+IVF-Flat 用 mini-batch Lloyd 方法训练聚类中心，随后将每条向量分配到最近中心。查询先计算 coarse centroid 距离，选择 `nprobe` 个桶，只打包这些桶中的 FP32 原向量并调用同一 CUDA 精确距离路径。索引文件保存中心、向量与倒排表，因此查询过程无需重新训练或分桶。
+
+当前 Top-K 在 CPU 使用 `std::partial_sort`，这样实现简单且可作为后续 GPU Top-K 的稳定参照。性能报告应如实讨论 D2H 距离传输和 CPU 归并的开销。
+
+## 10. 目录
+
+```text
+.
+├── CMakeLists.txt
+├── README.md
+├── REPORT.md
+├── configs/
+├── include/gvs/
+├── results/
+├── scripts/
+├── src/
+└── tests/
+```
+
+## 11. 提交
+
+1. Fork `InfiniTensor/Learning-CUDA`。
+2. 基于 `2026-summer-project` 分支工作。
+3. 将本目录完整复制到 `09_vector_retrieval/ccyyy1023/`（若实际 ID 不同，只改这一层目录名）。
+4. 保留 `results/` 中的真实日志；如换硬件复测，请用新日志替换并同步更新 `REPORT.md`。
+5. Commit、push，然后向上游 `2026-summer-project` 发起 PR；官网提交 fork 仓库链接和 commit 链接。
diff --git a/09_vector_retrieval/ccyyy1023/REPORT.md b/09_vector_retrieval/ccyyy1023/REPORT.md
new file mode 100644
index 00000000..b283b592
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/REPORT.md
@@ -0,0 +1,127 @@
+# GPU 向量检索引擎（CUDA）总结报告
+
+> 状态：已在 NVIDIA GeForce RTX 4090 24GB 上完成 Release 构建、正确性测试和百万规模性能矩阵。表中数字均来自 `results/*.json`，未使用估算或示例值。
+
+## 1. 问题描述
+
+本项目面向 RAG、推荐系统和多模态检索中的 Top-K 向量搜索。给定 N 条 D 维数据库向量和一批查询向量，系统返回 L2 距离最小或内积最大的 K 条记录。精确搜索复杂度为 O(ND)，可以作为正确性基准；IVF-Flat 通过 coarse quantization 只访问部分倒排桶，以召回率换取吞吐和延迟。
+
+## 2. 完成内容
+
+项目实现 CPU 精确 baseline、CUDA 精确距离计算、批量查询、排序 Top-K、L2/内积、IVF-Flat 索引、索引保存与加载、二进制输入输出、正确性测试和统一指标日志。支持任意正整数 K，正式脚本覆盖 K=1、10、50、100。
+
+基础实现使用 FP32。IVF 聚类、coarse search 和最终 Top-K 在 CPU 完成，数据库常驻 GPU，候选 ID 上传后由 CUDA kernel 读取原始向量并计算距离。本项目未实现 IVF-PQ、GPU Top-K、FP16 或国产 GPU 适配，不将这些未完成项写成已完成。
+
+## 3. 数据布局与文件格式
+
+数据库和查询采用 row-major 连续 FP32 数组。相邻 CUDA 线程处理相邻向量；固定维度循环中，相邻线程的全局内存访问存在跨向量步长。此布局与题目文件定义一致，容易验证和落盘，但不是所有维度下最理想的 coalescing 布局。后续可尝试转置或 tiled layout。
+
+文件带 8 字节 magic 和版本号，并固定整数字段宽度。IVF 索引文件保存聚类中心、原始向量和倒排列表；加载时校验形状、总 ID 数和 ID 范围，避免每次查询重新训练。
+
+## 4. 精确检索算法
+
+CPU 参考实现逐 query 计算所有距离，并用 `std::partial_sort` 得到排序 Top-K。CUDA kernel 使用二维 grid，一维覆盖 N 条候选，一维覆盖 batch。一个线程累加一个 query-vector pair 的 D 个维度，L2 和 inner product 均使用 `fmaf`。内积取负后复用升序排序逻辑，输出时恢复真实相似度。
+
+GPU 结果将距离复制回主机，在 CPU 完成 Top-K。这是当前版本最明确的性能边界：距离阶段在 GPU，但 distance D2H 与 CPU partial sort 仍可能成为瓶颈。
+
+## 5. IVF-Flat
+
+### 5.1 建索引
+
+1. 从数据库随机选取初始中心。
+2. 使用有固定上限的 mini-batch Lloyd 迭代更新中心，避免百万规模训练遍历成本失控。
+3. 将全部向量分配给最近中心，建立倒排桶。
+4. 保存中心、数据库向量和每个桶的原始 ID。
+
+### 5.2 查询
+
+1. 在 CPU 计算 query 与所有 coarse centroid 的距离。
+2. 选取最近的 `nprobe` 个桶。
+3. 仅把候选 ID 传入 GPU，数据库在一次 IVF 查询过程中常驻显存。
+4. CUDA kernel 按 ID 读取原始向量并计算精确距离。
+5. CPU Top-K 归并并映射回原始 ID。
+
+`nprobe` 越大，通常 recall@K 越高，但候选数、延迟和工作区也更大。正式脚本测试 1、4、16、64 四档。
+
+## 6. 正确性验证
+
+自动测试覆盖非 2 的幂维度 D=33，以避免只在整齐 shape 下通过；分别验证 L2 和 inner product。精确 GPU 与 CPU 的 Top-K ID 集合要求完全一致，分数误差容限为 2e-3。另有 clustered synthetic data 的 IVF 索引 round-trip 测试，要求 recall@10 不低于 0.90。
+
+| 检查项 | 当前结果 |
+|---|---|
+| CMake Release 构建 | 通过 |
+| `ctest` | 1/1 通过（0.34 s） |
+| CPU vs GPU exact recall@K | 1.000000 |
+| L2 Top-K/分数测试 | 通过（容限 2e-3） |
+| Inner product Top-K/分数测试 | 通过（容限 2e-3） |
+| IVF save/load | 通过 |
+
+## 7. 实验设置
+
+| 项目 | 值 |
+|---|---|
+| GPU | NVIDIA GeForce RTX 4090 |
+| GPU 显存 | 24,564 MiB |
+| CPU | Intel Xeon Platinum 8352V @ 2.10 GHz，16 vCPU |
+| OS | Ubuntu 22.04 / Linux 5.15.0-91-generic |
+| NVIDIA Driver | 580.95.05 |
+| CUDA Toolkit / nvcc | 11.8 / V11.8.89 |
+| 编译器 / CMake | G++ 11.3.0 / CMake 3.22.1 |
+| 编译类型 | Release (`-O3`, CUDA `--use_fast_math`) |
+| 数据规模 | 1,000,000 x 128，1000 queries |
+| 随机种子 | 42 |
+
+正式计时关闭了其他 GPU 任务。精确 QPS 包含 query H2D、CUDA 距离、distance D2H 和 CPU Top-K，不包含数据库首次上传；IVF QPS 还包含 CPU coarse search、候选 ID 构造和一次数据库 H2D。两者都不包含索引文件加载。P50/P99 用 batch=1 的逐 query 精确查询时间。显存字段是程序显式 buffer 大小，不包含 CUDA context。
+
+## 8. 性能与质量结果
+
+### 8.1 精确检索与 K
+
+| K | Batch | QPS | P50 (ms) | P99 (ms) | CPU/GPU 加速比 | CPU/GPU recall@K | 工作区 MiB |
+|---:|---:|---:|---:|---:|---:|---:|---:|
+| 1 | 128 | 263.571 | 4.622 | 6.801 | 42.444x | 1.000000 | 976.625 |
+| 10 | 128 | 261.783 | 4.650 | 6.038 | 42.262x | 1.000000 | 976.625 |
+| 50 | 128 | 274.503 | 4.281 | 5.897 | 42.709x | 1.000000 | 976.625 |
+| 100 | 128 | 255.570 | 4.402 | 5.988 | 40.381x | 1.000000 | 976.625 |
+
+### 8.2 Batch size
+
+| Batch | QPS | P50 (ms) | P99 (ms) | 工作区 MiB |
+|---:|---:|---:|---:|---:|
+| 1 | 231.165 | 4.391 | 6.144 | 492.096 |
+| 32 | 276.272 | 4.874 | 6.491 | 610.367 |
+| 128 | 261.783 | 4.650 | 6.038 | 976.625 |
+
+### 8.3 IVF nprobe
+
+| nlist | nprobe | QPS | recall@10 | 平均排名分数误差 | 工作区 MiB |
+|---:|---:|---:|---:|---:|---:|
+| 256 | 1 | 3509.281 | 0.059600 | 14.431664 | 488.390 |
+| 256 | 4 | 1986.191 | 0.182900 | 8.026410 | 488.680 |
+| 256 | 16 | 595.972 | 0.472900 | 3.293365 | 489.774 |
+| 256 | 64 | 153.481 | 0.937800 | 0.272549 | 493.920 |
+
+## 9. 结果分析
+
+batch=32 在本次运行中取得最高精确检索吞吐 276.272 QPS，比 batch=1 的 231.165 QPS 高约 19.5%；batch=128 未继续提升，说明更大的 distance buffer 和 CPU Top-K 并未带来额外收益。P50/P99 都在约 4.3-6.5 ms 范围内，差异接近单次运行波动。
+
+K 从 1 增加到 100 后，精确 QPS 仍处于 255-275 区间，未呈严格单调下降；GPU 距离计算和 1000x1000000 距离传输仍是主要固定成本，CPU `partial_sort` 的额外开销没有压倒该成本。所有 K 配置的 CPU/GPU exact recall 均为 1.0。
+
+IVF 展示了清晰的速度-质量权衡：nprobe=1 达到 3509.281 QPS，但 recall@10 仅 0.0596；nprobe=16 为 595.972 QPS、recall@10=0.4729；nprobe=64 达到 0.9378 recall@10，但吞吐降至 153.481 QPS，低于精确检索 261.783 QPS。若重视质量，本实现的 nprobe=64 已没有速度优势；这说明 coarse quality、索引布局和 GPU Top-K 仍需优化。
+
+索引构建耗时 4391.59 ms。工作区指标为程序显式 buffer 估算：精确 batch=128 为 976.625 MiB，IVF 随 nprobe 从约 488.390 MiB 增至 493.920 MiB；它不包含 CUDA context 与驱动开销。
+
+## 10. 优化历程与已知瓶颈
+
+初版选择统一的“距离越小越好”内部表示，使 L2 和内积复用相同 Top-K 逻辑；同时采用 batch distance matrix，避免逐 query 启动 kernel。第一版 IVF 会在 CPU 按 ID 打包候选向量并逐 query 上传，端到端 nprobe=16 仅约 11 QPS。定位后改为数据库一次上传并常驻 GPU、每次只传候选 ID，最终 nprobe=16 达到 595.972 QPS，召回率保持 0.4729。这一修改约提升 52 倍，并保留版本化索引加载。
+
+当前主要瓶颈包括：每个线程独立扫描维度、精确检索 distance matrix 的 D2H 传输、CPU `partial_sort`、IVF 的 CPU coarse search，以及候选 ID 的随机访存。服务器没有安装 `ncu`/`nsys`，因此报告不虚构 profiler 结论。
+
+## 11. 后续工作
+
+- 使用 CUB block radix sort 或分层 heap 实现 GPU Top-K 与跨 block merge。
+- 将数据库转为 tiled/转置布局，改善 warp 访存合并。
+- 用 CUDA kernel 完成 IVF assignment 和 coarse search。
+- 数据库与索引常驻 GPU，并用 pinned memory、stream 和 double buffering 重叠传输。
+- 支持 FP16、cosine、IVF-PQ 和多 GPU shard。
+- 使用 `ncu` 检查 memory throughput、occupancy、warp stall；使用 `nsys` 分析 H2D/kernel/D2H/CPU merge 时间线。
diff --git a/09_vector_retrieval/ccyyy1023/configs/exact.conf b/09_vector_retrieval/ccyyy1023/configs/exact.conf
new file mode 100644
index 00000000..23898f96
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/configs/exact.conf
@@ -0,0 +1,6 @@
+top_k = 10
+search_mode = "exact"
+batch_size = 128
+nlist = 256
+nprobe = 16
+
diff --git a/09_vector_retrieval/ccyyy1023/configs/ivf_flat.conf b/09_vector_retrieval/ccyyy1023/configs/ivf_flat.conf
new file mode 100644
index 00000000..a5e4b718
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/configs/ivf_flat.conf
@@ -0,0 +1,6 @@
+top_k = 10
+search_mode = "ivf_flat"
+batch_size = 128
+nlist = 256
+nprobe = 16
+
diff --git a/09_vector_retrieval/ccyyy1023/include/gvs/io.hpp b/09_vector_retrieval/ccyyy1023/include/gvs/io.hpp
new file mode 100644
index 00000000..2b5618b2
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/include/gvs/io.hpp
@@ -0,0 +1,34 @@
+#pragma once
+
+#include "gvs/search.hpp"
+
+#include <cstddef>
+#include <cstdint>
+#include <string>
+#include <vector>
+
+namespace gvs {
+
+struct VectorFile {
+    std::size_t count = 0;
+    std::size_t dim = 0;
+    Metric metric = Metric::L2;
+    std::vector<float> values;
+};
+
+struct Parameters {
+    std::size_t top_k = 10;
+    std::string search_mode = "exact";
+    std::size_t batch_size = 128;
+    std::size_t nlist = 4096;
+    std::size_t nprobe = 16;
+};
+
+void write_vector_file(const std::string& path, const VectorFile& data,
+                       bool include_metric);
+VectorFile read_vector_file(const std::string& path, bool expect_metric);
+void write_search_result(const std::string& path, const SearchResult& result);
+Parameters read_parameters(const std::string& path);
+
+}  // namespace gvs
+
diff --git a/09_vector_retrieval/ccyyy1023/include/gvs/search.hpp b/09_vector_retrieval/ccyyy1023/include/gvs/search.hpp
new file mode 100644
index 00000000..d1becb62
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/include/gvs/search.hpp
@@ -0,0 +1,78 @@
+#pragma once
+
+#include <cstddef>
+#include <cstdint>
+#include <string>
+#include <vector>
+
+namespace gvs {
+
+enum class Metric { L2, InnerProduct };
+
+struct SearchConfig {
+    std::size_t k = 10;
+    std::size_t batch_size = 16;
+};
+
+struct SearchResult {
+    std::size_t queries = 0;
+    std::size_t k = 0;
+    std::vector<int> ids;
+    std::vector<float> scores;
+};
+
+struct GpuStats {
+    double elapsed_ms = 0.0;
+    std::size_t peak_workspace_bytes = 0;
+    std::string device_name;
+};
+
+SearchResult cpu_exact_search(const std::vector<float>& database,
+                              const std::vector<float>& queries,
+                              std::size_t n, std::size_t nq,
+                              std::size_t dim, Metric metric,
+                              std::size_t k);
+
+SearchResult gpu_exact_search(const std::vector<float>& database,
+                              const std::vector<float>& queries,
+                              std::size_t n, std::size_t nq,
+                              std::size_t dim, Metric metric,
+                              const SearchConfig& config,
+                              GpuStats* stats = nullptr);
+
+SearchResult gpu_indexed_search(const std::vector<float>& database,
+                                const std::vector<float>& queries,
+                                const std::vector<std::vector<int>>& candidate_ids,
+                                std::size_t n, std::size_t nq,
+                                std::size_t dim, Metric metric,
+                                const SearchConfig& config,
+                                GpuStats* stats = nullptr);
+
+double recall_at_k(const SearchResult& truth, const SearchResult& candidate);
+
+class IvfFlatIndex {
+public:
+    IvfFlatIndex(std::size_t dim, std::size_t nlist, Metric metric);
+
+    void train(const std::vector<float>& database, std::size_t n,
+               std::size_t iterations = 12, std::uint64_t seed = 42);
+    void add(const std::vector<float>& database, std::size_t n);
+    SearchResult search(const std::vector<float>& queries, std::size_t nq,
+                        std::size_t nprobe, const SearchConfig& config,
+                        GpuStats* stats = nullptr) const;
+
+    void save(const std::string& path) const;
+    static IvfFlatIndex load(const std::string& path);
+
+    std::size_t nlist() const { return nlist_; }
+
+private:
+    std::size_t dim_;
+    std::size_t nlist_;
+    Metric metric_;
+    std::vector<float> centroids_;
+    std::vector<float> database_;
+    std::vector<std::vector<int>> lists_;
+};
+
+}  // namespace gvs
diff --git a/09_vector_retrieval/ccyyy1023/results/SMOKE_SUMMARY.md b/09_vector_retrieval/ccyyy1023/results/SMOKE_SUMMARY.md
new file mode 100644
index 00000000..3cd570ae
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/SMOKE_SUMMARY.md
@@ -0,0 +1,5 @@
+# Measured benchmark summary
+
+| log | GPU | N | NQ | D | K | exact QPS | P50 ms | P99 ms | exact recall | speedup | nprobe | IVF QPS | IVF recall | mean rank score error |
+|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
+| smoke.json | NVIDIA GeForce RTX 4090 | 10000 | 100 | 128 | 10 | 16118.163 | 0.081 | 0.084 | 1.000000 | 21.021 | 8 | 21150.489 | 0.371000 | 6.098108 |
diff --git a/09_vector_retrieval/ccyyy1023/results/SUMMARY.md b/09_vector_retrieval/ccyyy1023/results/SUMMARY.md
new file mode 100644
index 00000000..96e659da
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/SUMMARY.md
@@ -0,0 +1,14 @@
+# Measured benchmark summary
+
+| log | GPU | N | NQ | D | K | exact QPS | P50 ms | P99 ms | exact recall | speedup | nprobe | IVF QPS | IVF recall | mean rank score error |
+|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
+| batch_1.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 10 | 231.165 | 4.391 | 6.144 | 1.000000 | 36.857 | 16 | 585.657 | 0.472900 | 3.293365 |
+| batch_32.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 10 | 276.272 | 4.874 | 6.491 | 1.000000 | 43.956 | 16 | 590.492 | 0.472900 | 3.293365 |
+| k_1.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 1 | 263.571 | 4.622 | 6.801 | 1.000000 | 42.444 | 16 | 562.764 | 0.507000 | 2.678481 |
+| k_10.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 10 | 261.783 | 4.650 | 6.038 | 1.000000 | 42.262 | 16 | 595.972 | 0.472900 | 3.293365 |
+| k_100.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 100 | 255.570 | 4.402 | 5.988 | 1.000000 | 40.381 | 16 | 545.980 | 0.435290 | 4.288385 |
+| k_50.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 50 | 274.503 | 4.281 | 5.897 | 1.000000 | 42.709 | 16 | 550.537 | 0.449600 | 3.925203 |
+| nprobe_1.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 10 | 233.112 | 4.454 | 6.053 | 1.000000 | 43.759 | 1 | 3509.281 | 0.059600 | 14.431664 |
+| nprobe_4.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 10 | 244.722 | 4.357 | 6.802 | 1.000000 | 38.676 | 4 | 1986.191 | 0.182900 | 8.026410 |
+| nprobe_64.json | NVIDIA GeForce RTX 4090 | 1000000 | 1000 | 128 | 10 | 259.653 | 4.138 | 5.862 | 1.000000 | 42.469 | 64 | 153.481 | 0.937800 | 0.272549 |
+| smoke.json | NVIDIA GeForce RTX 4090 | 10000 | 100 | 128 | 10 | 16118.163 | 0.081 | 0.084 | 1.000000 | 21.021 | 8 | 21150.489 | 0.371000 | 6.098108 |
diff --git a/09_vector_retrieval/ccyyy1023/results/batch_1.json b/09_vector_retrieval/ccyyy1023/results/batch_1.json
new file mode 100644
index 00000000..546e24ea
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/batch_1.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 10},
+  "exact": {"query_ms": 4325.907430, "qps": 231.165372, "p50_ms": 4.391250, "p99_ms": 6.144444, "workspace_mib": 492.096436, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 36.857411},
+  "cpu": {"queries_measured": 20, "query_ms": 3188.834950},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 1707.485259, "qps": 585.656593, "nlist": 256, "nprobe": 16, "recall_at_k": 0.472900, "mean_rank_score_error": 3.293365, "workspace_mib": 489.773682}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/batch_32.json b/09_vector_retrieval/ccyyy1023/results/batch_32.json
new file mode 100644
index 00000000..de774f01
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/batch_32.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 10},
+  "exact": {"query_ms": 3619.621729, "qps": 276.271963, "p50_ms": 4.873504, "p99_ms": 6.491380, "workspace_mib": 610.367188, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 43.956197},
+  "cpu": {"queries_measured": 20, "query_ms": 3182.096116},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 1693.504089, "qps": 590.491636, "nlist": 256, "nprobe": 16, "recall_at_k": 0.472900, "mean_rank_score_error": 3.293365, "workspace_mib": 489.773682}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/environment.txt b/09_vector_retrieval/ccyyy1023/results/environment.txt
new file mode 100644
index 00000000..92569c74
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/environment.txt
@@ -0,0 +1,7 @@
+Measured: 2026-09-20 Asia/Shanghai
+Model name:                         Intel(R) Xeon(R) Platinum 8352V CPU @ 2.10GHz
+NVIDIA GeForce RTX 4090, 24564 MiB, 580.95.05
+Build cuda_11.8.r11.8/compiler.31833905_0
+cmake version 3.22.1
+g++ (Ubuntu 11.3.0-1ubuntu1~22.04) 11.3.0
+Linux 5.15.0-91-generic
diff --git a/09_vector_retrieval/ccyyy1023/results/index_build.txt b/09_vector_retrieval/ccyyy1023/results/index_build.txt
new file mode 100644
index 00000000..cbdd0605
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/index_build.txt
@@ -0,0 +1 @@
+index_build_ms=4391.59
diff --git a/09_vector_retrieval/ccyyy1023/results/k_1.json b/09_vector_retrieval/ccyyy1023/results/k_1.json
new file mode 100644
index 00000000..7cd813e2
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/k_1.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 1},
+  "exact": {"query_ms": 3794.049961, "qps": 263.570594, "p50_ms": 4.621565, "p99_ms": 6.800636, "workspace_mib": 976.625000, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 42.443584},
+  "cpu": {"queries_measured": 20, "query_ms": 3220.661527},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 1776.944915, "qps": 562.763647, "nlist": 256, "nprobe": 16, "recall_at_k": 0.507000, "mean_rank_score_error": 2.678481, "workspace_mib": 489.773682}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/k_10.json b/09_vector_retrieval/ccyyy1023/results/k_10.json
new file mode 100644
index 00000000..0590c67e
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/k_10.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 10},
+  "exact": {"query_ms": 3819.959462, "qps": 261.782883, "p50_ms": 4.650455, "p99_ms": 6.037778, "workspace_mib": 976.625000, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 42.261559},
+  "cpu": {"queries_measured": 20, "query_ms": 3228.748869},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 1677.930929, "qps": 595.972088, "nlist": 256, "nprobe": 16, "recall_at_k": 0.472900, "mean_rank_score_error": 3.293365, "workspace_mib": 489.773682}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/k_100.json b/09_vector_retrieval/ccyyy1023/results/k_100.json
new file mode 100644
index 00000000..7602d850
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/k_100.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 100},
+  "exact": {"query_ms": 3912.828840, "qps": 255.569574, "p50_ms": 4.402323, "p99_ms": 5.987935, "workspace_mib": 976.625000, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 40.380501},
+  "cpu": {"queries_measured": 20, "query_ms": 3160.039780},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 1831.570234, "qps": 545.979609, "nlist": 256, "nprobe": 16, "recall_at_k": 0.435290, "mean_rank_score_error": 4.288385, "workspace_mib": 489.773682}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/k_50.json b/09_vector_retrieval/ccyyy1023/results/k_50.json
new file mode 100644
index 00000000..37725b99
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/k_50.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 50},
+  "exact": {"query_ms": 3642.950502, "qps": 274.502769, "p50_ms": 4.280817, "p99_ms": 5.896796, "workspace_mib": 976.625000, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 42.708571},
+  "cpu": {"queries_measured": 20, "query_ms": 3111.704184},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 1816.408478, "qps": 550.536959, "nlist": 256, "nprobe": 16, "recall_at_k": 0.449600, "mean_rank_score_error": 3.925203, "workspace_mib": 489.773682}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/nprobe_1.json b/09_vector_retrieval/ccyyy1023/results/nprobe_1.json
new file mode 100644
index 00000000..8ee126a1
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/nprobe_1.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 10},
+  "exact": {"query_ms": 4289.785030, "qps": 233.111914, "p50_ms": 4.453786, "p99_ms": 6.053339, "workspace_mib": 976.625000, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 43.759387},
+  "cpu": {"queries_measured": 20, "query_ms": 3754.367282},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 284.958634, "qps": 3509.281280, "nlist": 256, "nprobe": 1, "recall_at_k": 0.059600, "mean_rank_score_error": 14.431664, "workspace_mib": 488.388565}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/nprobe_4.json b/09_vector_retrieval/ccyyy1023/results/nprobe_4.json
new file mode 100644
index 00000000..0e4875b0
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/nprobe_4.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 10},
+  "exact": {"query_ms": 4086.269341, "qps": 244.721999, "p50_ms": 4.357081, "p99_ms": 6.801751, "workspace_mib": 976.625000, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 38.676046},
+  "cpu": {"queries_measured": 20, "query_ms": 3160.814839},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 503.476214, "qps": 1986.191149, "nlist": 256, "nprobe": 4, "recall_at_k": 0.182900, "mean_rank_score_error": 8.026410, "workspace_mib": 488.679054}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/nprobe_64.json b/09_vector_retrieval/ccyyy1023/results/nprobe_64.json
new file mode 100644
index 00000000..aebd6520
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/nprobe_64.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 1000000, "nq": 1000, "dim": 128, "k": 10},
+  "exact": {"query_ms": 3851.290919, "qps": 259.653197, "p50_ms": 4.137849, "p99_ms": 5.861579, "workspace_mib": 976.625000, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 42.469173},
+  "cpu": {"queries_measured": 20, "query_ms": 3271.222801},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 6515.475806, "qps": 153.480733, "nlist": 256, "nprobe": 64, "recall_at_k": 0.937800, "mean_rank_score_error": 0.272549, "workspace_mib": 493.922852}
+}
diff --git a/09_vector_retrieval/ccyyy1023/results/smoke.json b/09_vector_retrieval/ccyyy1023/results/smoke.json
new file mode 100644
index 00000000..b9bea0d1
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/results/smoke.json
@@ -0,0 +1,8 @@
+{
+  "status": "measured",
+  "device": "NVIDIA GeForce RTX 4090",
+  "dataset": {"n": 10000, "nq": 100, "dim": 128, "k": 10},
+  "exact": {"query_ms": 6.204181, "qps": 16118.162897, "p50_ms": 0.081473, "p99_ms": 0.084174, "workspace_mib": 6.119141, "cpu_gpu_recall_at_k": 1.000000, "speedup_vs_cpu": 21.021299},
+  "cpu": {"queries_measured": 20, "query_ms": 26.083989},
+  "ivf_flat": {"index_build_ms": 0.000000, "query_ms": 4.728023, "qps": 21150.489327, "nlist": 64, "nprobe": 8, "recall_at_k": 0.371000, "mean_rank_score_error": 6.098108, "workspace_mib": 0.998459}
+}
diff --git a/09_vector_retrieval/ccyyy1023/scripts/check_submission.py b/09_vector_retrieval/ccyyy1023/scripts/check_submission.py
new file mode 100644
index 00000000..3c40c03b
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/scripts/check_submission.py
@@ -0,0 +1,50 @@
+#!/usr/bin/env python3
+"""Static submission check that does not require CUDA."""
+
+import json
+import pathlib
+import sys
+
+
+REQUIRED = [
+    "CMakeLists.txt",
+    "README.md",
+    "REPORT.md",
+    "include/gvs/search.hpp",
+    "include/gvs/io.hpp",
+    "src/cpu_search.cpp",
+    "src/gpu_search.cu",
+    "src/ivf_flat.cpp",
+    "src/file_io.cpp",
+    "src/main.cpp",
+    "tests/test_search.cpp",
+    "configs/exact.conf",
+    "configs/ivf_flat.conf",
+    "results/SUMMARY.md",
+    "results/environment.txt",
+]
+
+
+def main() -> int:
+    root = pathlib.Path(__file__).resolve().parents[1]
+    missing = [name for name in REQUIRED if not (root / name).is_file()]
+    if missing:
+        print("missing files:", *missing, sep="\n- ", file=sys.stderr)
+        return 1
+    logs = [json.loads(path.read_text(encoding="utf-8"))
+            for path in (root / "results").glob("*.json")]
+    if not logs or not all(item.get("status") == "measured" for item in logs):
+        print("all submitted JSON logs must be real measured runs", file=sys.stderr)
+        return 1
+    forbidden = ["TODO: fake", "RTX 4090: 999", "speedup: 100x"]
+    text = "\n".join((root / name).read_text(encoding="utf-8", errors="replace")
+                     for name in REQUIRED if name.endswith((".md", ".cpp", ".cu", ".hpp")))
+    if any(token in text for token in forbidden):
+        print("found forbidden fabricated marker", file=sys.stderr)
+        return 1
+    print(f"submission structure OK ({len(REQUIRED)} required files)")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())
diff --git a/09_vector_retrieval/ccyyy1023/scripts/run_required_benchmarks.sh b/09_vector_retrieval/ccyyy1023/scripts/run_required_benchmarks.sh
new file mode 100644
index 00000000..bf17803b
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/scripts/run_required_benchmarks.sh
@@ -0,0 +1,50 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+project_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+cd "$project_dir"
+mkdir -p data results
+
+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
+cmake --build build -j"$(nproc)"
+ctest --test-dir build --output-on-failure
+
+if [[ ! -f data/required_db.bin || ! -f data/required_queries.bin ]]; then
+  ./build/vector_search generate \
+    --db data/required_db.bin --queries data/required_queries.bin \
+    --n 1000000 --nq 1000 --dim 128 --metric l2 --seed 42
+fi
+
+if [[ ! -f data/required_ivf.bin ]]; then
+  ./build/vector_search build \
+    --db data/required_db.bin --index data/required_ivf.bin \
+    --nlist 256 --iterations 8 --seed 42 | tee results/index_build.txt
+fi
+
+for k in 1 10 50 100; do
+  ./build/vector_search benchmark \
+    --db data/required_db.bin --queries data/required_queries.bin \
+    --index data/required_ivf.bin --log "results/k_${k}.json" \
+    --k "$k" --batch 128 --nlist 256 --nprobe 16 \
+    --cpu-query-limit 20 --latency-query-limit 100
+done
+
+for batch in 1 32; do
+  ./build/vector_search benchmark \
+    --db data/required_db.bin --queries data/required_queries.bin \
+    --index data/required_ivf.bin --log "results/batch_${batch}.json" \
+    --k 10 --batch "$batch" --nlist 256 --nprobe 16 \
+    --cpu-query-limit 20 --latency-query-limit 100
+done
+
+for nprobe in 1 4 64; do
+  ./build/vector_search benchmark \
+    --db data/required_db.bin --queries data/required_queries.bin \
+    --index data/required_ivf.bin --log "results/nprobe_${nprobe}.json" \
+    --k 10 --batch 128 --nlist 256 --nprobe "$nprobe" \
+    --cpu-query-limit 20 --latency-query-limit 100
+done
+
+python3 scripts/summarize_results.py results results/SUMMARY.md
+echo "Required benchmark complete: results/SUMMARY.md"
+
diff --git a/09_vector_retrieval/ccyyy1023/scripts/run_smoke.sh b/09_vector_retrieval/ccyyy1023/scripts/run_smoke.sh
new file mode 100644
index 00000000..1cb2e1e9
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/scripts/run_smoke.sh
@@ -0,0 +1,26 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+project_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
+cd "$project_dir"
+
+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
+cmake --build build -j"$(nproc)"
+ctest --test-dir build --output-on-failure
+
+mkdir -p data results
+./build/vector_search generate \
+  --db data/small_db.bin --queries data/small_queries.bin \
+  --n 10000 --nq 100 --dim 128 --metric l2 --seed 42
+./build/vector_search build \
+  --db data/small_db.bin --index data/small_ivf.bin \
+  --nlist 64 --iterations 8 --seed 42
+./build/vector_search benchmark \
+  --db data/small_db.bin --queries data/small_queries.bin \
+  --index data/small_ivf.bin --log results/smoke.json \
+  --k 10 --batch 32 --nlist 64 --nprobe 8 \
+  --cpu-query-limit 20 --latency-query-limit 50
+
+python3 scripts/summarize_results.py results results/SMOKE_SUMMARY.md
+echo "Smoke test complete: results/SMOKE_SUMMARY.md"
+
diff --git a/09_vector_retrieval/ccyyy1023/scripts/summarize_results.py b/09_vector_retrieval/ccyyy1023/scripts/summarize_results.py
new file mode 100644
index 00000000..95462aa1
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/scripts/summarize_results.py
@@ -0,0 +1,59 @@
+#!/usr/bin/env python3
+"""Collect measured benchmark JSON files into a compact Markdown table."""
+
+import json
+import pathlib
+import sys
+
+
+def main() -> int:
+    if len(sys.argv) != 3:
+        print("usage: summarize_results.py RESULTS_DIR OUTPUT.md", file=sys.stderr)
+        return 2
+    results_dir = pathlib.Path(sys.argv[1])
+    output = pathlib.Path(sys.argv[2])
+    rows = []
+    for path in sorted(results_dir.glob("*.json")):
+        data = json.loads(path.read_text(encoding="utf-8"))
+        if data.get("status") != "measured":
+            continue
+        ds = data["dataset"]
+        exact = data["exact"]
+        ivf = data["ivf_flat"]
+        rows.append(
+            [
+                path.name,
+                data.get("device", ""),
+                ds["n"],
+                ds["nq"],
+                ds["dim"],
+                ds["k"],
+                f"{exact['qps']:.3f}",
+                f"{exact['p50_ms']:.3f}",
+                f"{exact['p99_ms']:.3f}",
+                f"{exact['cpu_gpu_recall_at_k']:.6f}",
+                f"{exact['speedup_vs_cpu']:.3f}",
+                ivf["nprobe"],
+                f"{ivf['qps']:.3f}",
+                f"{ivf['recall_at_k']:.6f}",
+                f"{ivf['mean_rank_score_error']:.6f}",
+            ]
+        )
+    header = [
+        "log", "GPU", "N", "NQ", "D", "K", "exact QPS", "P50 ms",
+        "P99 ms", "exact recall", "speedup", "nprobe", "IVF QPS",
+        "IVF recall", "mean rank score error",
+    ]
+    lines = ["# Measured benchmark summary", "", "| " + " | ".join(header) + " |",
+             "|" + "|".join(["---"] * len(header)) + "|"]
+    lines.extend("| " + " | ".join(map(str, row)) + " |" for row in rows)
+    if not rows:
+        lines.extend(["", "No measured benchmark JSON files were found."])
+    output.write_text("\n".join(lines) + "\n", encoding="utf-8")
+    print(f"wrote {output} with {len(rows)} measured rows")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())
+
diff --git a/09_vector_retrieval/ccyyy1023/src/cpu_search.cpp b/09_vector_retrieval/ccyyy1023/src/cpu_search.cpp
new file mode 100644
index 00000000..69f6185e
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/src/cpu_search.cpp
@@ -0,0 +1,73 @@
+#include "gvs/search.hpp"
+
+#include <algorithm>
+#include <cmath>
+#include <numeric>
+#include <stdexcept>
+#include <unordered_set>
+#include <utility>
+
+namespace gvs {
+namespace {
+
+float distance(const float* a, const float* b, std::size_t dim, Metric metric) {
+    float value = 0.0f;
+    if (metric == Metric::L2) {
+        for (std::size_t d = 0; d < dim; ++d) {
+            const float diff = a[d] - b[d];
+            value += diff * diff;
+        }
+        return value;
+    }
+    for (std::size_t d = 0; d < dim; ++d) value += a[d] * b[d];
+    return -value;  // Smaller is better internally for both metrics.
+}
+
+}  // namespace
+
+SearchResult cpu_exact_search(const std::vector<float>& database,
+                              const std::vector<float>& queries,
+                              std::size_t n, std::size_t nq,
+                              std::size_t dim, Metric metric,
+                              std::size_t k) {
+    if (database.size() != n * dim || queries.size() != nq * dim)
+        throw std::invalid_argument("input shape does not match n/nq/dim");
+    k = std::min(k, n);
+    SearchResult output{nq, k, std::vector<int>(nq * k),
+                        std::vector<float>(nq * k)};
+    std::vector<std::pair<float, int>> candidates(n);
+    for (std::size_t q = 0; q < nq; ++q) {
+        const float* query = queries.data() + q * dim;
+        for (std::size_t i = 0; i < n; ++i) {
+            candidates[i] = {distance(query, database.data() + i * dim, dim, metric),
+                             static_cast<int>(i)};
+        }
+        std::partial_sort(candidates.begin(), candidates.begin() + k,
+                          candidates.end());
+        for (std::size_t j = 0; j < k; ++j) {
+            output.ids[q * k + j] = candidates[j].second;
+            output.scores[q * k + j] =
+                metric == Metric::InnerProduct ? -candidates[j].first
+                                               : candidates[j].first;
+        }
+    }
+    return output;
+}
+
+double recall_at_k(const SearchResult& truth, const SearchResult& candidate) {
+    if (truth.queries != candidate.queries || truth.k != candidate.k)
+        throw std::invalid_argument("recall inputs have different shapes");
+    std::size_t hits = 0;
+    for (std::size_t q = 0; q < truth.queries; ++q) {
+        std::unordered_set<int> expected;
+        for (std::size_t j = 0; j < truth.k; ++j)
+            expected.insert(truth.ids[q * truth.k + j]);
+        for (std::size_t j = 0; j < candidate.k; ++j)
+            hits += expected.count(candidate.ids[q * candidate.k + j]);
+    }
+    const auto total = truth.queries * truth.k;
+    return total == 0 ? 1.0 : static_cast<double>(hits) / total;
+}
+
+}  // namespace gvs
+
diff --git a/09_vector_retrieval/ccyyy1023/src/file_io.cpp b/09_vector_retrieval/ccyyy1023/src/file_io.cpp
new file mode 100644
index 00000000..7aee9c4f
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/src/file_io.cpp
@@ -0,0 +1,127 @@
+#include "gvs/io.hpp"
+
+#include <algorithm>
+#include <cctype>
+#include <cstring>
+#include <fstream>
+#include <stdexcept>
+
+namespace gvs {
+namespace {
+
+constexpr char kDatabaseMagic[8] = {'G','V','S','D','B','0','1','\0'};
+constexpr char kQueryMagic[8] = {'G','V','S','Q','R','0','1','\0'};
+constexpr char kResultMagic[8] = {'G','V','S','R','S','0','1','\0'};
+
+template <typename T>
+void write_scalar(std::ofstream& out, const T& value) {
+    out.write(reinterpret_cast<const char*>(&value), sizeof(T));
+    if (!out) throw std::runtime_error("failed to write binary file");
+}
+
+template <typename T>
+T read_scalar(std::ifstream& in) {
+    T value{};
+    in.read(reinterpret_cast<char*>(&value), sizeof(T));
+    if (!in) throw std::runtime_error("truncated binary file");
+    return value;
+}
+
+std::string trim(std::string value) {
+    auto not_space = [](unsigned char c) { return !std::isspace(c); };
+    value.erase(value.begin(), std::find_if(value.begin(), value.end(), not_space));
+    value.erase(std::find_if(value.rbegin(), value.rend(), not_space).base(), value.end());
+    if (value.size() >= 2 && value.front() == '"' && value.back() == '"')
+        value = value.substr(1, value.size() - 2);
+    return value;
+}
+
+}  // namespace
+
+void write_vector_file(const std::string& path, const VectorFile& data,
+                       bool include_metric) {
+    if (data.values.size() != data.count * data.dim)
+        throw std::invalid_argument("vector file shape mismatch");
+    std::ofstream out(path, std::ios::binary);
+    if (!out) throw std::runtime_error("cannot open output file: " + path);
+    out.write(include_metric ? kDatabaseMagic : kQueryMagic, 8);
+    write_scalar(out, static_cast<std::uint64_t>(data.count));
+    write_scalar(out, static_cast<std::uint32_t>(data.dim));
+    write_scalar(out, static_cast<std::uint32_t>(1));  // dtype: 1 = fp32
+    if (include_metric)
+        write_scalar(out, static_cast<std::uint32_t>(data.metric == Metric::L2 ? 0 : 1));
+    out.write(reinterpret_cast<const char*>(data.values.data()),
+              static_cast<std::streamsize>(data.values.size() * sizeof(float)));
+    if (!out) throw std::runtime_error("failed to write vectors");
+}
+
+VectorFile read_vector_file(const std::string& path, bool expect_metric) {
+    std::ifstream in(path, std::ios::binary);
+    if (!in) throw std::runtime_error("cannot open input file: " + path);
+    char magic[8]{};
+    in.read(magic, 8);
+    const char* expected = expect_metric ? kDatabaseMagic : kQueryMagic;
+    if (!in || std::memcmp(magic, expected, 8) != 0)
+        throw std::runtime_error("invalid vector file magic/version: " + path);
+    VectorFile result;
+    result.count = static_cast<std::size_t>(read_scalar<std::uint64_t>(in));
+    result.dim = static_cast<std::size_t>(read_scalar<std::uint32_t>(in));
+    const auto dtype = read_scalar<std::uint32_t>(in);
+    if (dtype != 1) throw std::runtime_error("this baseline supports fp32 files only");
+    if (expect_metric) {
+        const auto metric = read_scalar<std::uint32_t>(in);
+        if (metric > 1) throw std::runtime_error("invalid metric code");
+        result.metric = metric == 0 ? Metric::L2 : Metric::InnerProduct;
+    }
+    if (result.count == 0 || result.dim == 0 ||
+        result.count > static_cast<std::size_t>(-1) / result.dim)
+        throw std::runtime_error("invalid or overflowing vector shape");
+    result.values.resize(result.count * result.dim);
+    in.read(reinterpret_cast<char*>(result.values.data()),
+            static_cast<std::streamsize>(result.values.size() * sizeof(float)));
+    if (!in) throw std::runtime_error("truncated vector payload");
+    return result;
+}
+
+void write_search_result(const std::string& path, const SearchResult& result) {
+    if (result.ids.size() != result.queries * result.k ||
+        result.scores.size() != result.ids.size())
+        throw std::invalid_argument("result shape mismatch");
+    std::ofstream out(path, std::ios::binary);
+    if (!out) throw std::runtime_error("cannot open result file: " + path);
+    out.write(kResultMagic, 8);
+    write_scalar(out, static_cast<std::uint64_t>(result.queries));
+    write_scalar(out, static_cast<std::uint32_t>(result.k));
+    for (std::size_t i = 0; i < result.ids.size(); ++i) {
+        write_scalar(out, static_cast<std::int64_t>(result.ids[i]));
+        write_scalar(out, result.scores[i]);
+    }
+}
+
+Parameters read_parameters(const std::string& path) {
+    std::ifstream in(path);
+    if (!in) throw std::runtime_error("cannot open parameter file: " + path);
+    Parameters p;
+    std::string line;
+    while (std::getline(in, line)) {
+        const auto comment = line.find('#');
+        if (comment != std::string::npos) line.erase(comment);
+        const auto equal = line.find('=');
+        if (equal == std::string::npos) continue;
+        const std::string key = trim(line.substr(0, equal));
+        const std::string value = trim(line.substr(equal + 1));
+        if (key == "top_k") p.top_k = std::stoull(value);
+        else if (key == "search_mode") p.search_mode = value;
+        else if (key == "batch_size") p.batch_size = std::stoull(value);
+        else if (key == "nlist") p.nlist = std::stoull(value);
+        else if (key == "nprobe") p.nprobe = std::stoull(value);
+    }
+    if (p.top_k == 0 || p.batch_size == 0 || p.nlist == 0 || p.nprobe == 0)
+        throw std::runtime_error("all numeric parameters must be positive");
+    if (p.search_mode != "exact" && p.search_mode != "ivf_flat")
+        throw std::runtime_error("search_mode must be exact or ivf_flat");
+    return p;
+}
+
+}  // namespace gvs
+
diff --git a/09_vector_retrieval/ccyyy1023/src/gpu_search.cu b/09_vector_retrieval/ccyyy1023/src/gpu_search.cu
new file mode 100644
index 00000000..c05ab519
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/src/gpu_search.cu
@@ -0,0 +1,227 @@
+#include "gvs/search.hpp"
+
+#include <cuda_runtime.h>
+
+#include <algorithm>
+#include <chrono>
+#include <limits>
+#include <stdexcept>
+#include <string>
+#include <utility>
+#include <vector>
+
+namespace gvs {
+namespace {
+
+#define CUDA_CHECK(call)                                                        \
+    do {                                                                        \
+        cudaError_t error__ = (call);                                            \
+        if (error__ != cudaSuccess)                                              \
+            throw std::runtime_error(std::string("CUDA error: ") +              \
+                                     cudaGetErrorString(error__));               \
+    } while (0)
+
+template <typename T>
+class DeviceBuffer {
+public:
+    explicit DeviceBuffer(std::size_t count = 0) : count_(count) {
+        if (count_) CUDA_CHECK(cudaMalloc(&ptr_, count_ * sizeof(T)));
+    }
+    ~DeviceBuffer() { if (ptr_) cudaFree(ptr_); }
+    DeviceBuffer(const DeviceBuffer&) = delete;
+    DeviceBuffer& operator=(const DeviceBuffer&) = delete;
+    T* get() { return ptr_; }
+    const T* get() const { return ptr_; }
+private:
+    T* ptr_ = nullptr;
+    std::size_t count_ = 0;
+};
+
+__global__ void distance_kernel(const float* database, const float* queries,
+                                float* distances, std::size_t n,
+                                std::size_t batch, std::size_t dim,
+                                int metric) {
+    const std::size_t i = blockIdx.x * blockDim.x + threadIdx.x;
+    const std::size_t q = blockIdx.y;
+    if (i >= n || q >= batch) return;
+    const float* x = database + i * dim;
+    const float* query = queries + q * dim;
+    float value = 0.0f;
+    if (metric == 0) {
+        for (std::size_t d = 0; d < dim; ++d) {
+            const float diff = query[d] - x[d];
+            value = fmaf(diff, diff, value);
+        }
+    } else {
+        for (std::size_t d = 0; d < dim; ++d)
+            value = fmaf(query[d], x[d], value);
+        value = -value;
+    }
+    distances[q * n + i] = value;
+}
+
+__global__ void indexed_distance_kernel(const float* database, const float* query,
+                                        const int* ids, float* distances,
+                                        std::size_t candidates, std::size_t dim,
+                                        int metric) {
+    const std::size_t i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= candidates) return;
+    const float* x = database + static_cast<std::size_t>(ids[i]) * dim;
+    float value = 0.0f;
+    if (metric == 0) {
+        for (std::size_t d = 0; d < dim; ++d) {
+            const float diff = query[d] - x[d];
+            value = fmaf(diff, diff, value);
+        }
+    } else {
+        for (std::size_t d = 0; d < dim; ++d)
+            value = fmaf(query[d], x[d], value);
+        value = -value;
+    }
+    distances[i] = value;
+}
+
+SearchResult finalize_on_cpu(const std::vector<float>& distances,
+                             std::size_t nq, std::size_t n,
+                             std::size_t k, Metric metric) {
+    SearchResult output{nq, k, std::vector<int>(nq * k),
+                        std::vector<float>(nq * k)};
+    std::vector<std::pair<float, int>> row(n);
+    for (std::size_t q = 0; q < nq; ++q) {
+        for (std::size_t i = 0; i < n; ++i)
+            row[i] = {distances[q * n + i], static_cast<int>(i)};
+        std::partial_sort(row.begin(), row.begin() + k, row.end());
+        for (std::size_t j = 0; j < k; ++j) {
+            output.ids[q * k + j] = row[j].second;
+            output.scores[q * k + j] =
+                metric == Metric::InnerProduct ? -row[j].first : row[j].first;
+        }
+    }
+    return output;
+}
+
+}  // namespace
+
+SearchResult gpu_exact_search(const std::vector<float>& database,
+                              const std::vector<float>& queries,
+                              std::size_t n, std::size_t nq,
+                              std::size_t dim, Metric metric,
+                              const SearchConfig& config,
+                              GpuStats* stats) {
+    if (database.size() != n * dim || queries.size() != nq * dim)
+        throw std::invalid_argument("input shape does not match n/nq/dim");
+    const std::size_t k = std::min(config.k, n);
+    const std::size_t batch_capacity = std::min(config.batch_size, nq);
+
+    DeviceBuffer<float> d_database(database.size());
+    DeviceBuffer<float> d_queries(batch_capacity * dim);
+    DeviceBuffer<float> d_distances(batch_capacity * n);
+    CUDA_CHECK(cudaMemcpy(d_database.get(), database.data(),
+                          database.size() * sizeof(float), cudaMemcpyHostToDevice));
+
+    SearchResult output{nq, k, std::vector<int>(nq * k),
+                        std::vector<float>(nq * k)};
+    std::vector<float> host_distances(batch_capacity * n);
+    CUDA_CHECK(cudaDeviceSynchronize());
+    const auto begin = std::chrono::steady_clock::now();
+
+    for (std::size_t first = 0; first < nq; first += batch_capacity) {
+        const std::size_t batch = std::min(batch_capacity, nq - first);
+        CUDA_CHECK(cudaMemcpy(d_queries.get(), queries.data() + first * dim,
+                              batch * dim * sizeof(float), cudaMemcpyHostToDevice));
+        const dim3 block(256);
+        const dim3 grid(static_cast<unsigned>((n + block.x - 1) / block.x),
+                        static_cast<unsigned>(batch));
+        distance_kernel<<<grid, block>>>(d_database.get(), d_queries.get(),
+                                         d_distances.get(), n, batch, dim,
+                                         metric == Metric::L2 ? 0 : 1);
+        CUDA_CHECK(cudaGetLastError());
+        CUDA_CHECK(cudaMemcpy(host_distances.data(), d_distances.get(),
+                              batch * n * sizeof(float), cudaMemcpyDeviceToHost));
+        SearchResult part = finalize_on_cpu(host_distances, batch, n, k, metric);
+        std::copy(part.ids.begin(), part.ids.end(), output.ids.begin() + first * k);
+        std::copy(part.scores.begin(), part.scores.end(),
+                  output.scores.begin() + first * k);
+    }
+    CUDA_CHECK(cudaDeviceSynchronize());
+    const auto end = std::chrono::steady_clock::now();
+    const double elapsed =
+        std::chrono::duration<double, std::milli>(end - begin).count();
+
+    if (stats) {
+        cudaDeviceProp prop{};
+        CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
+        stats->elapsed_ms = elapsed;
+        stats->peak_workspace_bytes =
+            (database.size() + batch_capacity * dim + batch_capacity * n) *
+            sizeof(float);
+        stats->device_name = prop.name;
+    }
+    return output;
+}
+
+SearchResult gpu_indexed_search(const std::vector<float>& database,
+                                const std::vector<float>& queries,
+                                const std::vector<std::vector<int>>& candidate_ids,
+                                std::size_t n, std::size_t nq,
+                                std::size_t dim, Metric metric,
+                                const SearchConfig& config,
+                                GpuStats* stats) {
+    if (database.size() != n * dim || queries.size() != nq * dim ||
+        candidate_ids.size() != nq)
+        throw std::invalid_argument("indexed search shape mismatch");
+    const std::size_t k = std::min(config.k, n);
+    std::size_t max_candidates = 0;
+    for (const auto& ids : candidate_ids) max_candidates = std::max(max_candidates, ids.size());
+    SearchResult output{nq, k, std::vector<int>(nq * k, -1),
+                        std::vector<float>(nq * k, 0.0f)};
+    if (max_candidates == 0) return output;
+
+    DeviceBuffer<float> d_database(database.size());
+    DeviceBuffer<float> d_query(dim);
+    DeviceBuffer<int> d_ids(max_candidates);
+    DeviceBuffer<float> d_distances(max_candidates);
+    CUDA_CHECK(cudaMemcpy(d_database.get(), database.data(),
+                          database.size() * sizeof(float), cudaMemcpyHostToDevice));
+    std::vector<float> host_distances(max_candidates);
+    const auto begin = std::chrono::steady_clock::now();
+    for (std::size_t q = 0; q < nq; ++q) {
+        const std::size_t count = candidate_ids[q].size();
+        if (count == 0) continue;
+        CUDA_CHECK(cudaMemcpy(d_query.get(), queries.data() + q * dim,
+                              dim * sizeof(float), cudaMemcpyHostToDevice));
+        CUDA_CHECK(cudaMemcpy(d_ids.get(), candidate_ids[q].data(),
+                              count * sizeof(int), cudaMemcpyHostToDevice));
+        const dim3 block(256);
+        const dim3 grid(static_cast<unsigned>((count + block.x - 1) / block.x));
+        indexed_distance_kernel<<<grid, block>>>(
+            d_database.get(), d_query.get(), d_ids.get(), d_distances.get(),
+            count, dim, metric == Metric::L2 ? 0 : 1);
+        CUDA_CHECK(cudaGetLastError());
+        CUDA_CHECK(cudaMemcpy(host_distances.data(), d_distances.get(),
+                              count * sizeof(float), cudaMemcpyDeviceToHost));
+        const std::size_t part_k = std::min(k, count);
+        std::vector<std::pair<float, int>> row(count);
+        for (std::size_t i = 0; i < count; ++i)
+            row[i] = {host_distances[i], candidate_ids[q][i]};
+        std::partial_sort(row.begin(), row.begin() + part_k, row.end());
+        for (std::size_t j = 0; j < part_k; ++j) {
+            output.ids[q * k + j] = row[j].second;
+            output.scores[q * k + j] =
+                metric == Metric::InnerProduct ? -row[j].first : row[j].first;
+        }
+    }
+    CUDA_CHECK(cudaDeviceSynchronize());
+    if (stats) {
+        cudaDeviceProp prop{};
+        CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
+        stats->elapsed_ms = std::chrono::duration<double, std::milli>(
+            std::chrono::steady_clock::now() - begin).count();
+        stats->peak_workspace_bytes = database.size() * sizeof(float) +
+            dim * sizeof(float) + max_candidates * (sizeof(float) + sizeof(int));
+        stats->device_name = prop.name;
+    }
+    return output;
+}
+
+}  // namespace gvs
diff --git a/09_vector_retrieval/ccyyy1023/src/ivf_flat.cpp b/09_vector_retrieval/ccyyy1023/src/ivf_flat.cpp
new file mode 100644
index 00000000..81688267
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/src/ivf_flat.cpp
@@ -0,0 +1,214 @@
+#include "gvs/search.hpp"
+
+#include <algorithm>
+#include <chrono>
+#include <cstdint>
+#include <cstring>
+#include <fstream>
+#include <limits>
+#include <numeric>
+#include <random>
+#include <stdexcept>
+#include <utility>
+
+namespace gvs {
+namespace {
+
+float coarse_distance(const float* a, const float* b, std::size_t dim,
+                      Metric metric) {
+    float value = 0.0f;
+    if (metric == Metric::L2) {
+        for (std::size_t d = 0; d < dim; ++d) {
+            const float diff = a[d] - b[d];
+            value += diff * diff;
+        }
+        return value;
+    }
+    for (std::size_t d = 0; d < dim; ++d) value += a[d] * b[d];
+    return -value;
+}
+
+std::size_t nearest_centroid(const float* x, const std::vector<float>& centroids,
+                             std::size_t nlist, std::size_t dim, Metric metric) {
+    std::size_t best = 0;
+    float best_value = std::numeric_limits<float>::infinity();
+    for (std::size_t c = 0; c < nlist; ++c) {
+        const float value = coarse_distance(x, centroids.data() + c * dim,
+                                            dim, metric);
+        if (value < best_value) { best_value = value; best = c; }
+    }
+    return best;
+}
+
+}  // namespace
+
+namespace {
+constexpr char kIvfMagic[8] = {'G','V','S','I','V','F','1','\0'};
+template <typename T>
+void write_value(std::ofstream& out, const T& value) {
+    out.write(reinterpret_cast<const char*>(&value), sizeof(T));
+    if (!out) throw std::runtime_error("failed to write IVF index");
+}
+template <typename T>
+T read_value(std::ifstream& in) {
+    T value{};
+    in.read(reinterpret_cast<char*>(&value), sizeof(T));
+    if (!in) throw std::runtime_error("truncated IVF index");
+    return value;
+}
+}
+
+IvfFlatIndex::IvfFlatIndex(std::size_t dim, std::size_t nlist, Metric metric)
+    : dim_(dim), nlist_(nlist), metric_(metric), lists_(nlist) {
+    if (dim == 0 || nlist == 0) throw std::invalid_argument("dim/nlist must be positive");
+}
+
+void IvfFlatIndex::train(const std::vector<float>& database, std::size_t n,
+                         std::size_t iterations, std::uint64_t seed) {
+    if (database.size() != n * dim_ || n < nlist_)
+        throw std::invalid_argument("invalid training matrix or n < nlist");
+    std::mt19937_64 rng(seed);
+    std::vector<std::size_t> order(n);
+    std::iota(order.begin(), order.end(), 0);
+    std::shuffle(order.begin(), order.end(), rng);
+    centroids_.resize(nlist_ * dim_);
+    for (std::size_t c = 0; c < nlist_; ++c)
+        std::copy_n(database.data() + order[c] * dim_, dim_,
+                    centroids_.data() + c * dim_);
+
+    std::vector<float> sums(nlist_ * dim_);
+    std::vector<std::size_t> counts(nlist_);
+    // Mini-batch Lloyd training bounds build cost on million-vector inputs.
+    const std::size_t train_count = std::min(n, std::max<std::size_t>(4096, nlist_ * 4));
+    for (std::size_t iter = 0; iter < iterations; ++iter) {
+        std::fill(sums.begin(), sums.end(), 0.0f);
+        std::fill(counts.begin(), counts.end(), 0);
+        std::shuffle(order.begin(), order.end(), rng);
+        for (std::size_t sample = 0; sample < train_count; ++sample) {
+            const std::size_t i = order[sample];
+            const std::size_t c = nearest_centroid(database.data() + i * dim_,
+                                                   centroids_, nlist_, dim_, metric_);
+            ++counts[c];
+            for (std::size_t d = 0; d < dim_; ++d)
+                sums[c * dim_ + d] += database[i * dim_ + d];
+        }
+        for (std::size_t c = 0; c < nlist_; ++c) {
+            if (counts[c] == 0) continue;
+            for (std::size_t d = 0; d < dim_; ++d)
+                centroids_[c * dim_ + d] = sums[c * dim_ + d] / counts[c];
+        }
+    }
+}
+
+void IvfFlatIndex::add(const std::vector<float>& database, std::size_t n) {
+    if (centroids_.empty()) throw std::logic_error("train must be called before add");
+    if (database.size() != n * dim_) throw std::invalid_argument("invalid database shape");
+    database_ = database;
+    lists_.assign(nlist_, {});
+    std::vector<std::size_t> labels(n);
+    #pragma omp parallel for schedule(static)
+    for (std::size_t i = 0; i < n; ++i) {
+        labels[i] = nearest_centroid(database.data() + i * dim_, centroids_,
+                                     nlist_, dim_, metric_);
+    }
+    for (std::size_t i = 0; i < n; ++i)
+        lists_[labels[i]].push_back(static_cast<int>(i));
+}
+
+SearchResult IvfFlatIndex::search(const std::vector<float>& queries,
+                                  std::size_t nq, std::size_t nprobe,
+                                  const SearchConfig& config,
+                                  GpuStats* stats) const {
+    if (database_.empty()) throw std::logic_error("add must be called before search");
+    if (queries.size() != nq * dim_) throw std::invalid_argument("invalid query shape");
+    nprobe = std::min(nprobe, nlist_);
+    const std::size_t k = std::min(config.k, database_.size() / dim_);
+    const auto search_begin = std::chrono::steady_clock::now();
+    std::vector<std::vector<int>> candidate_ids(nq);
+
+    for (std::size_t q = 0; q < nq; ++q) {
+        std::vector<std::pair<float, int>> coarse(nlist_);
+        for (std::size_t c = 0; c < nlist_; ++c)
+            coarse[c] = {coarse_distance(queries.data() + q * dim_,
+                                         centroids_.data() + c * dim_, dim_, metric_),
+                         static_cast<int>(c)};
+        std::partial_sort(coarse.begin(), coarse.begin() + nprobe, coarse.end());
+        for (std::size_t p = 0; p < nprobe; ++p) {
+            const auto& list = lists_[coarse[p].second];
+            candidate_ids[q].insert(candidate_ids[q].end(), list.begin(), list.end());
+        }
+    }
+    GpuStats gpu_stats;
+    SearchResult output = gpu_indexed_search(
+        database_, queries, candidate_ids, database_.size() / dim_, nq, dim_,
+        metric_, SearchConfig{k, 1}, &gpu_stats);
+    if (stats) {
+        stats->elapsed_ms = std::chrono::duration<double, std::milli>(
+            std::chrono::steady_clock::now() - search_begin).count();
+        stats->peak_workspace_bytes = gpu_stats.peak_workspace_bytes;
+        stats->device_name = gpu_stats.device_name;
+    }
+    return output;
+}
+
+void IvfFlatIndex::save(const std::string& path) const {
+    if (database_.empty() || centroids_.empty())
+        throw std::logic_error("cannot save an unbuilt IVF index");
+    std::ofstream out(path, std::ios::binary);
+    if (!out) throw std::runtime_error("cannot open IVF index output: " + path);
+    out.write(kIvfMagic, 8);
+    write_value(out, static_cast<std::uint64_t>(database_.size() / dim_));
+    write_value(out, static_cast<std::uint32_t>(dim_));
+    write_value(out, static_cast<std::uint32_t>(nlist_));
+    write_value(out, static_cast<std::uint32_t>(metric_ == Metric::L2 ? 0 : 1));
+    out.write(reinterpret_cast<const char*>(centroids_.data()),
+              static_cast<std::streamsize>(centroids_.size() * sizeof(float)));
+    out.write(reinterpret_cast<const char*>(database_.data()),
+              static_cast<std::streamsize>(database_.size() * sizeof(float)));
+    for (const auto& list : lists_) {
+        write_value(out, static_cast<std::uint64_t>(list.size()));
+        out.write(reinterpret_cast<const char*>(list.data()),
+                  static_cast<std::streamsize>(list.size() * sizeof(int)));
+    }
+    if (!out) throw std::runtime_error("failed to write IVF payload");
+}
+
+IvfFlatIndex IvfFlatIndex::load(const std::string& path) {
+    std::ifstream in(path, std::ios::binary);
+    if (!in) throw std::runtime_error("cannot open IVF index: " + path);
+    char magic[8]{};
+    in.read(magic, 8);
+    if (!in || std::memcmp(magic, kIvfMagic, 8) != 0)
+        throw std::runtime_error("invalid IVF index magic/version");
+    const std::size_t n = static_cast<std::size_t>(read_value<std::uint64_t>(in));
+    const std::size_t dim = read_value<std::uint32_t>(in);
+    const std::size_t nlist = read_value<std::uint32_t>(in);
+    const auto metric_code = read_value<std::uint32_t>(in);
+    if (n == 0 || dim == 0 || nlist == 0 || metric_code > 1 ||
+        n > static_cast<std::size_t>(-1) / dim)
+        throw std::runtime_error("invalid IVF index header");
+    IvfFlatIndex index(dim, nlist,
+                       metric_code == 0 ? Metric::L2 : Metric::InnerProduct);
+    index.centroids_.resize(nlist * dim);
+    index.database_.resize(n * dim);
+    in.read(reinterpret_cast<char*>(index.centroids_.data()),
+            static_cast<std::streamsize>(index.centroids_.size() * sizeof(float)));
+    in.read(reinterpret_cast<char*>(index.database_.data()),
+            static_cast<std::streamsize>(index.database_.size() * sizeof(float)));
+    std::size_t id_count = 0;
+    for (auto& list : index.lists_) {
+        const auto size = static_cast<std::size_t>(read_value<std::uint64_t>(in));
+        if (size > n - id_count) throw std::runtime_error("invalid IVF list sizes");
+        list.resize(size);
+        in.read(reinterpret_cast<char*>(list.data()),
+                static_cast<std::streamsize>(size * sizeof(int)));
+        id_count += size;
+        for (int id : list)
+            if (id < 0 || static_cast<std::size_t>(id) >= n)
+                throw std::runtime_error("invalid vector id in IVF index");
+    }
+    if (!in || id_count != n) throw std::runtime_error("truncated or incomplete IVF index");
+    return index;
+}
+
+}  // namespace gvs
diff --git a/09_vector_retrieval/ccyyy1023/src/main.cpp b/09_vector_retrieval/ccyyy1023/src/main.cpp
new file mode 100644
index 00000000..c8bafe5f
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/src/main.cpp
@@ -0,0 +1,259 @@
+#include "gvs/io.hpp"
+#include "gvs/search.hpp"
+
+#include <algorithm>
+#include <chrono>
+#include <cmath>
+#include <cstddef>
+#include <cstdint>
+#include <filesystem>
+#include <fstream>
+#include <iomanip>
+#include <iostream>
+#include <map>
+#include <numeric>
+#include <random>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+namespace {
+
+using Clock = std::chrono::steady_clock;
+
+class Args {
+public:
+    Args(int argc, char** argv, int start = 2) {
+        for (int i = start; i < argc; ++i) {
+            const std::string key = argv[i];
+            if (key.rfind("--", 0) != 0 || i + 1 >= argc)
+                throw std::invalid_argument("expected --key value pairs");
+            values_[key.substr(2)] = argv[++i];
+        }
+    }
+    std::string get(const std::string& key, const std::string& fallback = "") const {
+        auto it = values_.find(key);
+        return it == values_.end() ? fallback : it->second;
+    }
+    std::string require(const std::string& key) const {
+        const auto value = get(key);
+        if (value.empty()) throw std::invalid_argument("missing --" + key);
+        return value;
+    }
+    std::size_t number(const std::string& key, std::size_t fallback) const {
+        const auto value = get(key);
+        return value.empty() ? fallback : static_cast<std::size_t>(std::stoull(value));
+    }
+private:
+    std::map<std::string, std::string> values_;
+};
+
+gvs::Metric parse_metric(const std::string& value) {
+    if (value == "l2") return gvs::Metric::L2;
+    if (value == "inner_product") return gvs::Metric::InnerProduct;
+    throw std::invalid_argument("metric must be l2 or inner_product");
+}
+
+double percentile(std::vector<double> values, double p) {
+    if (values.empty()) return 0.0;
+    std::sort(values.begin(), values.end());
+    const std::size_t index = static_cast<std::size_t>(
+        std::ceil(p * static_cast<double>(values.size()))) - 1;
+    return values[std::min(index, values.size() - 1)];
+}
+
+double mean_rank_score_error(const gvs::SearchResult& truth,
+                             const gvs::SearchResult& candidate) {
+    if (truth.queries != candidate.queries || truth.k != candidate.k)
+        throw std::invalid_argument("quality result shapes differ");
+    double sum = 0.0;
+    for (std::size_t i = 0; i < truth.scores.size(); ++i)
+        sum += std::abs(static_cast<double>(truth.scores[i]) - candidate.scores[i]);
+    return truth.scores.empty() ? 0.0 : sum / truth.scores.size();
+}
+
+std::string json_escape(const std::string& value) {
+    std::string out;
+    for (char c : value) {
+        if (c == '"' || c == '\\') out.push_back('\\');
+        out.push_back(c);
+    }
+    return out;
+}
+
+void command_generate(const Args& args) {
+    const std::size_t n = args.number("n", 10000);
+    const std::size_t nq = args.number("nq", 100);
+    const std::size_t dim = args.number("dim", 128);
+    const auto metric = parse_metric(args.get("metric", "l2"));
+    const std::uint64_t seed = args.number("seed", 42);
+    std::mt19937_64 rng(seed);
+    std::normal_distribution<float> normal(0.0f, 1.0f);
+    gvs::VectorFile database{n, dim, metric, std::vector<float>(n * dim)};
+    gvs::VectorFile queries{nq, dim, metric, std::vector<float>(nq * dim)};
+    for (float& value : database.values) value = normal(rng);
+    for (float& value : queries.values) value = normal(rng);
+    gvs::write_vector_file(args.require("db"), database, true);
+    gvs::write_vector_file(args.require("queries"), queries, false);
+    std::cout << "generated n=" << n << " nq=" << nq << " dim=" << dim << '\n';
+}
+
+void command_build(const Args& args) {
+    const auto database = gvs::read_vector_file(args.require("db"), true);
+    const auto start = Clock::now();
+    gvs::IvfFlatIndex index(database.dim, args.number("nlist", 256), database.metric);
+    index.train(database.values, database.count, args.number("iterations", 12),
+                args.number("seed", 42));
+    index.add(database.values, database.count);
+    index.save(args.require("index"));
+    const double ms = std::chrono::duration<double, std::milli>(Clock::now() - start).count();
+    std::cout << "index_build_ms=" << ms << '\n';
+}
+
+void command_search(const Args& args) {
+    const auto database = gvs::read_vector_file(args.require("db"), true);
+    const auto queries = gvs::read_vector_file(args.require("queries"), false);
+    if (database.dim != queries.dim) throw std::runtime_error("database/query dim mismatch");
+    const auto parameters = gvs::read_parameters(args.require("params"));
+    gvs::SearchConfig config{parameters.top_k, parameters.batch_size};
+    gvs::GpuStats stats;
+    gvs::SearchResult result;
+    if (parameters.search_mode == "exact") {
+        result = gvs::gpu_exact_search(database.values, queries.values,
+                                       database.count, queries.count, database.dim,
+                                       database.metric, config, &stats);
+    } else {
+        auto index = gvs::IvfFlatIndex::load(args.require("index"));
+        result = index.search(queries.values, queries.count, parameters.nprobe,
+                              config, &stats);
+    }
+    gvs::write_search_result(args.require("output"), result);
+    const double qps = stats.elapsed_ms == 0.0 ? 0.0 : queries.count * 1000.0 / stats.elapsed_ms;
+    std::cout << "mode=" << parameters.search_mode << " query_ms=" << stats.elapsed_ms
+              << " qps=" << qps << " workspace_mib="
+              << stats.peak_workspace_bytes / (1024.0 * 1024.0) << '\n';
+}
+
+void command_benchmark(const Args& args) {
+    const auto database = gvs::read_vector_file(args.require("db"), true);
+    const auto queries = gvs::read_vector_file(args.require("queries"), false);
+    if (database.dim != queries.dim) throw std::runtime_error("database/query dim mismatch");
+    const std::size_t k = args.number("k", 10);
+    const std::size_t batch = args.number("batch", 128);
+    const std::size_t nprobe = args.number("nprobe", 16);
+    const std::size_t cpu_limit = std::min(args.number("cpu-query-limit", 20), queries.count);
+    const std::size_t latency_limit = std::min(args.number("latency-query-limit", 100), queries.count);
+    const gvs::SearchConfig config{k, batch};
+
+    gvs::GpuStats exact_stats;
+    const auto exact = gvs::gpu_exact_search(database.values, queries.values,
+                                             database.count, queries.count, database.dim,
+                                             database.metric, config, &exact_stats);
+
+    std::vector<float> cpu_queries(queries.values.begin(),
+                                   queries.values.begin() + cpu_limit * database.dim);
+    const auto cpu_begin = Clock::now();
+    const auto cpu = gvs::cpu_exact_search(database.values, cpu_queries,
+                                           database.count, cpu_limit, database.dim,
+                                           database.metric, k);
+    const double cpu_ms = std::chrono::duration<double, std::milli>(
+        Clock::now() - cpu_begin).count();
+    gvs::SearchResult exact_prefix{cpu_limit, exact.k,
+        std::vector<int>(exact.ids.begin(), exact.ids.begin() + cpu_limit * exact.k),
+        std::vector<float>(exact.scores.begin(), exact.scores.begin() + cpu_limit * exact.k)};
+    const double exact_recall = gvs::recall_at_k(cpu, exact_prefix);
+
+    double build_ms = 0.0;
+    const std::string index_path = args.require("index");
+    if (!std::filesystem::exists(index_path) || args.get("rebuild", "false") == "true") {
+        const auto build_begin = Clock::now();
+        gvs::IvfFlatIndex built(database.dim, args.number("nlist", 256), database.metric);
+        built.train(database.values, database.count, args.number("iterations", 12),
+                    args.number("seed", 42));
+        built.add(database.values, database.count);
+        built.save(index_path);
+        build_ms = std::chrono::duration<double, std::milli>(
+            Clock::now() - build_begin).count();
+    }
+    auto index = gvs::IvfFlatIndex::load(index_path);
+    gvs::GpuStats ivf_stats;
+    const auto approximate = index.search(queries.values, queries.count, nprobe,
+                                          config, &ivf_stats);
+    const double ivf_recall = gvs::recall_at_k(exact, approximate);
+    const double score_error = mean_rank_score_error(exact, approximate);
+
+    std::vector<double> exact_latencies;
+    exact_latencies.reserve(latency_limit);
+    for (std::size_t q = 0; q < latency_limit; ++q) {
+        std::vector<float> one(queries.values.begin() + q * database.dim,
+                               queries.values.begin() + (q + 1) * database.dim);
+        gvs::GpuStats one_stats;
+        gvs::gpu_exact_search(database.values, one, database.count, 1, database.dim,
+                              database.metric, gvs::SearchConfig{k, 1}, &one_stats);
+        exact_latencies.push_back(one_stats.elapsed_ms);
+    }
+
+    const double exact_qps = exact_stats.elapsed_ms == 0.0 ? 0.0
+        : queries.count * 1000.0 / exact_stats.elapsed_ms;
+    const double ivf_qps = ivf_stats.elapsed_ms == 0.0 ? 0.0
+        : queries.count * 1000.0 / ivf_stats.elapsed_ms;
+    const double cpu_per_query = cpu_limit == 0 ? 0.0 : cpu_ms / cpu_limit;
+    const double gpu_per_query = queries.count == 0 ? 0.0 : exact_stats.elapsed_ms / queries.count;
+    const double speedup = gpu_per_query == 0.0 ? 0.0 : cpu_per_query / gpu_per_query;
+
+    std::ofstream log(args.require("log"));
+    if (!log) throw std::runtime_error("cannot open benchmark log");
+    log << std::fixed << std::setprecision(6)
+        << "{\n"
+        << "  \"status\": \"measured\",\n"
+        << "  \"device\": \"" << json_escape(exact_stats.device_name) << "\",\n"
+        << "  \"dataset\": {\"n\": " << database.count << ", \"nq\": "
+        << queries.count << ", \"dim\": " << database.dim << ", \"k\": " << k << "},\n"
+        << "  \"exact\": {\"query_ms\": " << exact_stats.elapsed_ms
+        << ", \"qps\": " << exact_qps << ", \"p50_ms\": "
+        << percentile(exact_latencies, 0.50) << ", \"p99_ms\": "
+        << percentile(exact_latencies, 0.99) << ", \"workspace_mib\": "
+        << exact_stats.peak_workspace_bytes / (1024.0 * 1024.0)
+        << ", \"cpu_gpu_recall_at_k\": " << exact_recall
+        << ", \"speedup_vs_cpu\": " << speedup << "},\n"
+        << "  \"cpu\": {\"queries_measured\": " << cpu_limit
+        << ", \"query_ms\": " << cpu_ms << "},\n"
+        << "  \"ivf_flat\": {\"index_build_ms\": " << build_ms
+        << ", \"query_ms\": " << ivf_stats.elapsed_ms << ", \"qps\": "
+        << ivf_qps << ", \"nlist\": " << index.nlist() << ", \"nprobe\": "
+        << nprobe << ", \"recall_at_k\": " << ivf_recall
+        << ", \"mean_rank_score_error\": " << score_error
+        << ", \"workspace_mib\": "
+        << ivf_stats.peak_workspace_bytes / (1024.0 * 1024.0) << "}\n"
+        << "}\n";
+    std::cout << "benchmark written to " << args.require("log") << '\n';
+}
+
+void usage() {
+    std::cout
+        << "GPU Vector Search\n"
+        << "  generate --db DB --queries Q --n N --nq NQ --dim D --metric l2|inner_product\n"
+        << "  build --db DB --index INDEX --nlist C [--iterations I]\n"
+        << "  search --db DB --queries Q --params CONFIG --index INDEX --output OUT\n"
+        << "  benchmark --db DB --queries Q --index INDEX --log JSON [--k K --batch B]\n";
+}
+
+}  // namespace
+
+int main(int argc, char** argv) {
+    try {
+        if (argc < 2) { usage(); return 1; }
+        const std::string command = argv[1];
+        const Args args(argc, argv);
+        if (command == "generate") command_generate(args);
+        else if (command == "build") command_build(args);
+        else if (command == "search") command_search(args);
+        else if (command == "benchmark") command_benchmark(args);
+        else { usage(); return 1; }
+        return 0;
+    } catch (const std::exception& error) {
+        std::cerr << "error: " << error.what() << '\n';
+        return 2;
+    }
+}
+
diff --git a/09_vector_retrieval/ccyyy1023/tests/test_search.cpp b/09_vector_retrieval/ccyyy1023/tests/test_search.cpp
new file mode 100644
index 00000000..d02e24ca
--- /dev/null
+++ b/09_vector_retrieval/ccyyy1023/tests/test_search.cpp
@@ -0,0 +1,72 @@
+#include "gvs/search.hpp"
+
+#include <algorithm>
+#include <cmath>
+#include <iostream>
+#include <stdexcept>
+#include <vector>
+
+namespace {
+
+void require(bool condition, const char* message) {
+    if (!condition) throw std::runtime_error(message);
+}
+
+void test_metric(gvs::Metric metric) {
+    constexpr std::size_t n = 257;
+    constexpr std::size_t nq = 7;
+    constexpr std::size_t dim = 33;
+    constexpr std::size_t k = 10;
+    std::vector<float> database(n * dim);
+    std::vector<float> queries(nq * dim);
+    for (std::size_t i = 0; i < database.size(); ++i)
+        database[i] = std::sin(static_cast<float>(i) * 0.013f);
+    for (std::size_t i = 0; i < queries.size(); ++i)
+        queries[i] = std::cos(static_cast<float>(i) * 0.017f);
+    const auto cpu = gvs::cpu_exact_search(database, queries, n, nq, dim, metric, k);
+    const auto gpu = gvs::gpu_exact_search(database, queries, n, nq, dim, metric,
+                                           gvs::SearchConfig{k, 4});
+    require(gvs::recall_at_k(cpu, gpu) == 1.0, "GPU exact Top-K differs from CPU");
+    for (std::size_t i = 0; i < cpu.scores.size(); ++i)
+        require(std::abs(cpu.scores[i] - gpu.scores[i]) < 2e-3f,
+                "GPU score differs from CPU tolerance");
+}
+
+void test_ivf_roundtrip() {
+    constexpr std::size_t n = 512;
+    constexpr std::size_t nq = 8;
+    constexpr std::size_t dim = 16;
+    std::vector<float> database(n * dim);
+    std::vector<float> queries(nq * dim);
+    for (std::size_t i = 0; i < n; ++i)
+        for (std::size_t d = 0; d < dim; ++d)
+            database[i * dim + d] = static_cast<float>(i % 16) * 4.0f +
+                                    static_cast<float>((i + d) % 7) * 0.01f;
+    for (std::size_t q = 0; q < nq; ++q)
+        std::copy_n(database.data() + q * 16 * dim, dim, queries.data() + q * dim);
+    gvs::IvfFlatIndex index(dim, 16, gvs::Metric::L2);
+    index.train(database, n, 8, 7);
+    index.add(database, n);
+    index.save("gvs_test_index.bin");
+    auto loaded = gvs::IvfFlatIndex::load("gvs_test_index.bin");
+    const auto exact = gvs::cpu_exact_search(database, queries, n, nq, dim,
+                                             gvs::Metric::L2, 10);
+    const auto approximate = loaded.search(queries, nq, 4, gvs::SearchConfig{10, 4});
+    require(gvs::recall_at_k(exact, approximate) >= 0.90,
+            "IVF recall is unexpectedly low on clustered test data");
+}
+
+}  // namespace
+
+int main() {
+    try {
+        test_metric(gvs::Metric::L2);
+        test_metric(gvs::Metric::InnerProduct);
+        test_ivf_roundtrip();
+        std::cout << "all correctness tests passed\n";
+        return 0;
+    } catch (const std::exception& error) {
+        std::cerr << "test failure: " << error.what() << '\n';
+        return 1;
+    }
+}