Skip to content

Fix fp16 overflow #1084

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
hhy3 wants to merge 1 commit into
base: main
Choose a base branch
from
Open

Fix fp16 overflow #1084

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 4 additions & 4 deletions cpp/bench/ann/src/cuvs/cuvs_benchmark.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,8 +75,8 @@ auto create_algo(const std::string& algo_name,
if constexpr (std::is_same_v<T, uint8_t>) {}

#ifdef CUVS_ANN_BENCH_USE_CUVS_IVF_FLAT
if constexpr (std::is_same_v<T, float> || std::is_same_v<T, uint8_t> ||
std::is_same_v<T, int8_t>) {
if constexpr (std::is_same_v<T, float> || std::is_same_v<T, half> ||
std::is_same_v<T, uint8_t> || std::is_same_v<T, int8_t>) {
if (algo_name == "raft_ivf_flat" || algo_name == "cuvs_ivf_flat") {
typename cuvs::bench::cuvs_ivf_flat<T, int64_t>::build_param param;
parse_build_param<T, int64_t>(conf, param);
Expand Down Expand Up @@ -141,8 +141,8 @@ auto create_search_param(const std::string& algo_name, const nlohmann::json& con
}
#endif
#ifdef CUVS_ANN_BENCH_USE_CUVS_IVF_FLAT
if constexpr (std::is_same_v<T, float> || std::is_same_v<T, uint8_t> ||
std::is_same_v<T, int8_t>) {
if constexpr (std::is_same_v<T, float> || std::is_same_v<T, half> ||
std::is_same_v<T, uint8_t> || std::is_same_v<T, int8_t>) {
if (algo_name == "raft_ivf_flat" || algo_name == "cuvs_ivf_flat") {
auto param =
std::make_unique<typename cuvs::bench::cuvs_ivf_flat<T, int64_t>::search_param>();
Expand Down
1 change: 1 addition & 0 deletions cpp/bench/ann/src/cuvs/cuvs_ivf_flat.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@

namespace cuvs::bench {
template class cuvs_ivf_flat<float, int64_t>;
template class cuvs_ivf_flat<half, int64_t>;
template class cuvs_ivf_flat<uint8_t, int64_t>;
template class cuvs_ivf_flat<int8_t, int64_t>;
} // namespace cuvs::bench
2 changes: 1 addition & 1 deletion cpp/src/neighbors/detail/ann_utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -136,7 +136,7 @@ struct config<float> {
};
template <>
struct config<half> {
using value_t = half;
using value_t = float;

@achirkin achirkin Jul 4, 2025

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please check whether this is used outside the IVF-Flat. Changing the accumulation type like this can have a drastic impact on performance.

@hhy3 hhy3 Jul 4, 2025

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I did simple benchmark and it showed no significant differences. I'll use cuvs ann-bench to get a more detailed benchmark results later

@cjnolet cjnolet Jul 22, 2025

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@hhy3 any updates here? We're about to begin burndown for 25.08 release. Should we consider this for 25.08 or push to 25.10 (October)?

@hhy3 hhy3 Jul 23, 2025

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@cjnolet hi, push it to 25.10, thx

@hhy3 hhy3 Apr 11, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@cjnolet sorry for my late reply. I just benchmark the performance, and it shows some regression on performance, especially when nprobe is small, but significantly improves recall:

IVF-Flat FP16 AccT=float vs AccT=half Benchmark Results

  • Dataset: cohere-768-angular-fp16 (1M vectors, 768 dims, inner_product)
  • GPU: NVIDIA A100-PCIE-40GB
  • nlist=4096, ratio=10, niter=20, k=100

Latency (bs=64)

nprobe half latency (ms) float latency (ms) change half recall float recall
32 4.02 5.04 +25% 0.815 0.849
64 6.22 6.89 +11% 0.867 0.913
128 9.87 10.2 +3% 0.897 0.955

Throughput (bs=1000, threads:1)

nprobe half QPS float QPS change half recall float recall
32 24.7k 22.8k -8% 0.816 0.850
64 13.8k 13.2k -4% 0.868 0.913
128 7.87k 7.84k -0.4% 0.898 0.956

static constexpr double kDivisor = 1.0;
};
template <>
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/neighbors/ivf_flat/jit_lto_kernels/interleaved_scan_matrix.json
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,10 +32,10 @@
},
{
"data_type": "__half",
"acc_type": "__half",
"acc_type": "float",
"veclen": ["1", "8"],
"type_abbrev": "h",
"acc_abbrev": "h"
"acc_abbrev": "f"
},
{
"data_type": "uint8_t",
Expand Down
12 changes: 6 additions & 6 deletions cpp/src/neighbors/ivf_flat/jit_lto_kernels/ivf_flat_interleaved_scan_kernel.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,8 +82,8 @@ struct loadAndComputeDist {
for (int k = 0; k < Veclen; ++k) {
compute_dist<Veclen, T, AccT>(dist, queryRegs[k], encV[k]);
if constexpr (ComputeNorm) {
norm_query += queryRegs[k] * queryRegs[k];
norm_data += encV[k] * encV[k];
norm_query += (AccT)(queryRegs[k] * queryRegs[k]);
norm_data += (AccT)(encV[k] * encV[k]);
}
}
}
Expand Down Expand Up @@ -117,8 +117,8 @@ struct loadAndComputeDist {
T q = raft::shfl(queryReg, d + k, raft::WarpSize);
compute_dist<Veclen, T, AccT>(dist, q, encV[k]);
if constexpr (ComputeNorm) {
norm_query += q * q;
norm_data += encV[k] * encV[k];
norm_query += (AccT)(q * q);
norm_data += (AccT)(encV[k] * encV[k]);
}
}
}
Expand All @@ -143,8 +143,8 @@ struct loadAndComputeDist {
T q = raft::shfl(queryReg, d + k, raft::WarpSize);
compute_dist<Veclen, T, AccT>(dist, q, enc[k]);
if constexpr (ComputeNorm) {
norm_query += q * q;
norm_data += enc[k] * enc[k];
norm_query += (AccT)(q * q);
norm_data += (AccT)(enc[k] * enc[k]);
}
}
}
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/neighbors/ivf_flat/jit_lto_kernels/metric_matrix.json
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,10 +19,10 @@
},
{
"data_type": "__half",
"acc_type": "__half",
"acc_type": "float",
"veclen": ["1", "8"],
"type_abbrev": "h",
"acc_abbrev": "h"
"acc_abbrev": "f"
},
{
"data_type": "uint8_t",
Expand Down
Loading