diff --git a/c/include/cuvs/preprocessing/quantize/scalar.h b/c/include/cuvs/preprocessing/quantize/scalar.h index 9724ddbfa8..68175d98ca 100644 --- a/c/include/cuvs/preprocessing/quantize/scalar.h +++ b/c/include/cuvs/preprocessing/quantize/scalar.h @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -19,15 +19,24 @@ extern "C" { * @defgroup preprocessing_c_scalar C API for Scalar Quantizer * @{ */ +/** Select global or per-dimension scalar quantization bounds. */ +enum cuvsScalarQuantizationMode { + CUVS_SCALAR_GLOBAL = 0, + CUVS_SCALAR_PER_DIMENSION = 1 +}; + /** - * @brief Scalar quantizer parameters. + * @brief Scalar quantizer parameters. Allocate with cuvsScalarQuantizerParamsCreate. */ struct cuvsScalarQuantizerParams { - /* + /** * specifies how many outliers at top & bottom will be ignored * needs to be within range of (0, 1] */ float quantile; + + /** Quantization mode. Defaults to CUVS_SCALAR_GLOBAL. */ + enum cuvsScalarQuantizationMode mode; }; typedef struct cuvsScalarQuantizerParams* cuvsScalarQuantizerParams_t; @@ -53,6 +62,11 @@ CUVS_EXPORT cuvsError_t cuvsScalarQuantizerParamsDestroy(cuvsScalarQuantizerPara * * The quantization is performed by a linear mapping of an interval in the * float data type to the full range of the quantized int type. + * Allocate with cuvsScalarQuantizerCreate and release with cuvsScalarQuantizerDestroy; + * do not copy the handle's contents. The implementation also owns any per-dimension bounds. + * In global mode, min_ and max_ are the trained bounds. In per-dimension mode, they are + * numeric_limits::min() and numeric_limits::max() for the training data type T, + * and are not used for transformation. */ typedef struct { double min_; @@ -70,7 +84,7 @@ typedef cuvsScalarQuantizer* cuvsScalarQuantizer_t; CUVS_EXPORT cuvsError_t cuvsScalarQuantizerCreate(cuvsScalarQuantizer_t* quantizer); /** - * @brief De-allocate Scalar Quantizer + * @brief De-allocate Scalar Quantizer after outstanding stream work completes. * * @param[in] quantizer * @return cuvsError_t @@ -80,10 +94,15 @@ CUVS_EXPORT cuvsError_t cuvsScalarQuantizerDestroy(cuvsScalarQuantizer_t quantiz /** * @brief Trains a scalar quantizer to be used later for quantizing the dataset. * + * Supports float16, float32, and float64 in global or per-dimension mode. + * Per-dimension mode applies the quantile independently to each dimension using every row. + * Successful retraining replaces the previous bounds; failed training preserves them. + * Complete outstanding stream work using the quantizer before retraining. + * * @param[in] res raft resource * @param[in] params configure scalar quantizer, e.g. quantile * @param[in] dataset a row-major host or device matrix - * @param[out] quantizer trained scalar quantizer + * @param[in,out] quantizer scalar quantizer allocated by cuvsScalarQuantizerCreate */ CUVS_EXPORT cuvsError_t cuvsScalarQuantizerTrain(cuvsResources_t res, cuvsScalarQuantizerParams_t params, @@ -93,6 +112,9 @@ CUVS_EXPORT cuvsError_t cuvsScalarQuantizerTrain(cuvsResources_t res, /** * @brief Applies quantization transform to given dataset * + * Per-dimension mode requires device matrices with the same input type and width as training. + * Its output codes are int8 in [-127, 127]. The quantizer must outlive stream work using its bounds. + * * @param[in] res raft resource * @param[in] quantizer a scalar quantizer * @param[in] dataset a row-major host or device matrix to transform @@ -107,7 +129,8 @@ CUVS_EXPORT cuvsError_t cuvsScalarQuantizerTransform(cuvsResources_t res, * @brief Perform inverse quantization step on previously quantized dataset * * Note that depending on the chosen data types train dataset the conversion is - * not lossless. + * not lossless. Per-dimension mode requires device matrices with the same output type and + * width as training. The quantizer must outlive stream work using its bounds. * * @param[in] res raft resource * @param[in] quantizer a scalar quantizer diff --git a/c/src/preprocessing/quantize/scalar.cpp b/c/src/preprocessing/quantize/scalar.cpp index 99ab090fa8..7b7eca81bd 100644 --- a/c/src/preprocessing/quantize/scalar.cpp +++ b/c/src/preprocessing/quantize/scalar.cpp @@ -1,9 +1,11 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include +#include +#include #include #include @@ -15,49 +17,40 @@ namespace { +namespace scalar = cuvs::preprocessing::quantize::scalar; + +// Keep the public min/max layout while owning per-dimension bounds behind the factory handle. +struct scalar_quantizer : cuvsScalarQuantizer { + uintptr_t addr = 0; + DLDataType dtype{}; +}; + template -void _train(cuvsResources_t res, - cuvsScalarQuantizerParams params, - DLManagedTensor* dataset_tensor, - cuvsScalarQuantizer_t quantizer) +scalar::quantizer _train(cuvsResources_t res, + scalar::params params, + DLManagedTensor* dataset_tensor) { auto dataset = dataset_tensor->dl_tensor; - auto res_ptr = reinterpret_cast(res); - - auto quantizer_params = cuvs::preprocessing::quantize::scalar::params(); - quantizer_params.quantile = params.quantile; - - cuvs::preprocessing::quantize::scalar::quantizer ret; - if (cuvs::core::is_dlpack_device_compatible(dataset)) { using mdspan_type = raft::device_matrix_view; - auto mds = cuvs::core::from_dlpack(dataset_tensor); - ret = cuvs::preprocessing::quantize::scalar::train(*res_ptr, quantizer_params, mds); + return scalar::train(*res_ptr, params, cuvs::core::from_dlpack(dataset_tensor)); } else if (cuvs::core::is_dlpack_host_compatible(dataset)) { using mdspan_type = raft::host_matrix_view; - auto mds = cuvs::core::from_dlpack(dataset_tensor); - ret = cuvs::preprocessing::quantize::scalar::train(*res_ptr, quantizer_params, mds); + return scalar::train(*res_ptr, params, cuvs::core::from_dlpack(dataset_tensor)); } else { RAFT_FAIL("dataset must be accessible on host or device memory"); } - - quantizer->min_ = ret.min_; - quantizer->max_ = ret.max_; } template void _transform(cuvsResources_t res, - cuvsScalarQuantizer_t quantizer_, + const scalar::quantizer& quantizer, DLManagedTensor* dataset_tensor, DLManagedTensor* out_tensor) { auto res_ptr = reinterpret_cast(res); - cuvs::preprocessing::quantize::scalar::quantizer quantizer; - quantizer.min_ = quantizer_->min_; - quantizer.max_ = quantizer_->max_; - auto dataset = dataset_tensor->dl_tensor; if (cuvs::core::is_dlpack_device_compatible(dataset)) { using mdspan_type = raft::device_matrix_view; @@ -85,16 +78,12 @@ void _transform(cuvsResources_t res, template void _inverse_transform(cuvsResources_t res, - cuvsScalarQuantizer_t quantizer_, + const scalar::quantizer& quantizer, DLManagedTensor* dataset_tensor, DLManagedTensor* out_tensor) { auto res_ptr = reinterpret_cast(res); - cuvs::preprocessing::quantize::scalar::quantizer quantizer; - quantizer.min_ = quantizer_->min_; - quantizer.max_ = quantizer_->max_; - auto dataset = dataset_tensor->dl_tensor; if (cuvs::core::is_dlpack_device_compatible(dataset)) { using mdspan_type = raft::device_matrix_view; @@ -119,12 +108,59 @@ void _inverse_transform(cuvsResources_t res, RAFT_FAIL("dataset must be accessible on host or device memory"); } } + +template +void _transform(cuvsResources_t res, + cuvsScalarQuantizer_t quantizer, + DLManagedTensor* dataset, + DLManagedTensor* out) +{ + _transform(res, + scalar::quantizer{static_cast(quantizer->min_), static_cast(quantizer->max_)}, + dataset, + out); +} + +template +void _inverse_transform(cuvsResources_t res, + cuvsScalarQuantizer_t quantizer, + DLManagedTensor* dataset, + DLManagedTensor* out) +{ + _inverse_transform(res, + scalar::quantizer{static_cast(quantizer->min_), static_cast(quantizer->max_)}, + dataset, + out); +} + +template +void dispatch_float(DLDataType dtype, Fn&& fn) +{ + RAFT_EXPECTS(dtype.code == kDLFloat && dtype.lanes == 1, "Expected a scalar floating-point dtype"); + switch (dtype.bits) { + case 16: fn.template operator()(); break; + case 32: fn.template operator()(); break; + case 64: fn.template operator()(); break; + default: RAFT_FAIL("Unsupported floating-point dtype bits: %d", dtype.bits); + } +} + +void destroy_per_dim(scalar_quantizer* quantizer) +{ + if (quantizer->addr) { + dispatch_float(quantizer->dtype, [&]() { + delete reinterpret_cast*>(quantizer->addr); + }); + quantizer->addr = 0; + } +} + } // namespace extern "C" cuvsError_t cuvsScalarQuantizerParamsCreate(cuvsScalarQuantizerParams_t* params) { return cuvs::core::translate_exceptions( - [=] { *params = new cuvsScalarQuantizerParams{.quantile = 0.99}; }); + [=] { *params = new cuvsScalarQuantizerParams{.quantile = 0.99, .mode = CUVS_SCALAR_GLOBAL}; }); } extern "C" cuvsError_t cuvsScalarQuantizerParamsDestroy(cuvsScalarQuantizerParams_t params) @@ -134,12 +170,21 @@ extern "C" cuvsError_t cuvsScalarQuantizerParamsDestroy(cuvsScalarQuantizerParam extern "C" cuvsError_t cuvsScalarQuantizerCreate(cuvsScalarQuantizer_t* quantizer) { - return cuvs::core::translate_exceptions([=] { *quantizer = new cuvsScalarQuantizer{}; }); + return cuvs::core::translate_exceptions([=] { + RAFT_EXPECTS(quantizer != nullptr, "quantizer must not be null"); + *quantizer = new scalar_quantizer{}; + }); } extern "C" cuvsError_t cuvsScalarQuantizerDestroy(cuvsScalarQuantizer_t quantizer) { - return cuvs::core::translate_exceptions([=] { delete quantizer; }); + return cuvs::core::translate_exceptions([=] { + if (quantizer) { + auto q = static_cast(quantizer); + destroy_per_dim(q); + delete q; + } + }); } extern "C" cuvsError_t cuvsScalarQuantizerTrain(cuvsResources_t res, @@ -148,18 +193,34 @@ extern "C" cuvsError_t cuvsScalarQuantizerTrain(cuvsResources_t res, cuvsScalarQuantizer_t quantizer) { return cuvs::core::translate_exceptions([=] { - auto dataset = dataset_tensor->dl_tensor; - if (dataset.dtype.code == kDLFloat && dataset.dtype.bits == 32) { - _train(res, *params, dataset_tensor, quantizer); - } else if (dataset.dtype.code == kDLFloat && dataset.dtype.bits == 16) { - _train(res, *params, dataset_tensor, quantizer); - } else if (dataset.dtype.code == kDLFloat && dataset.dtype.bits == 64) { - _train(res, *params, dataset_tensor, quantizer); - } else { - RAFT_FAIL("Unsupported dataset DLtensor dtype: %d and bits: %d", - dataset.dtype.code, - dataset.dtype.bits); + RAFT_EXPECTS(params && dataset_tensor && quantizer, + "params, dataset, and quantizer must not be null"); + scalar::params cpp_params; + cpp_params.quantile = params->quantile; + switch (params->mode) { + case CUVS_SCALAR_GLOBAL: cpp_params.range = scalar::quantization_mode::GLOBAL; break; + case CUVS_SCALAR_PER_DIMENSION: + cpp_params.range = scalar::quantization_mode::PER_DIMENSION; + break; + default: RAFT_FAIL("Unsupported scalar quantization mode"); } + auto q = static_cast(quantizer); + dispatch_float(dataset_tensor->dl_tensor.dtype, [&]() { + auto trained = _train(res, cpp_params, dataset_tensor); + if (auto const* bounds = + std::get_if::global_type>(&trained.value)) { + destroy_per_dim(q); + q->min_ = bounds->min; + q->max_ = bounds->max; + } else { + auto owned = std::make_unique>(std::move(trained)); + destroy_per_dim(q); + q->addr = reinterpret_cast(owned.release()); + q->dtype = dataset_tensor->dl_tensor.dtype; + q->min_ = static_cast(cuda::std::numeric_limits::min()); + q->max_ = static_cast(cuda::std::numeric_limits::max()); + } + }); }); } @@ -169,18 +230,16 @@ extern "C" cuvsError_t cuvsScalarQuantizerTransform(cuvsResources_t res, DLManagedTensor* out_tensor) { return cuvs::core::translate_exceptions([=] { - auto dataset = dataset_tensor->dl_tensor; - if (dataset.dtype.code == kDLFloat && dataset.dtype.bits == 32) { - _transform(res, quantizer, dataset_tensor, out_tensor); - } else if (dataset.dtype.code == kDLFloat && dataset.dtype.bits == 16) { - _transform(res, quantizer, dataset_tensor, out_tensor); - } else if (dataset.dtype.code == kDLFloat && dataset.dtype.bits == 64) { - _transform(res, quantizer, dataset_tensor, out_tensor); - } else { - RAFT_FAIL("Unsupported dataset DLtensor dtype: %d and bits: %d", - dataset.dtype.code, - dataset.dtype.bits); - } + RAFT_EXPECTS(quantizer && dataset_tensor && out_tensor, + "quantizer, dataset, and out must not be null"); + auto q = static_cast(quantizer); + dispatch_float(q->addr ? q->dtype : dataset_tensor->dl_tensor.dtype, [&]() { + if (q->addr) { + _transform(res, *reinterpret_cast*>(q->addr), dataset_tensor, out_tensor); + } else { + _transform(res, quantizer, dataset_tensor, out_tensor); + } + }); }); } @@ -190,16 +249,14 @@ cuvsError_t cuvsScalarQuantizerInverseTransform(cuvsResources_t res, DLManagedTensor* out) { return cuvs::core::translate_exceptions([=] { - auto dtype = out->dl_tensor.dtype; - if (dtype.code == kDLFloat && dtype.bits == 32) { - _inverse_transform(res, quantizer, dataset, out); - } else if (dtype.code == kDLFloat && dtype.bits == 16) { - _inverse_transform(res, quantizer, dataset, out); - } else if (dtype.code == kDLFloat && dtype.bits == 64) { - _inverse_transform(res, quantizer, dataset, out); - } else { - RAFT_FAIL( - "Unsupported output dataset DLtensor dtype: %d and bits: %d", dtype.code, dtype.bits); - } + RAFT_EXPECTS(quantizer && dataset && out, "quantizer, dataset, and out must not be null"); + auto q = static_cast(quantizer); + dispatch_float(q->addr ? q->dtype : out->dl_tensor.dtype, [&]() { + if (q->addr) { + _inverse_transform(res, *reinterpret_cast*>(q->addr), dataset, out); + } else { + _inverse_transform(res, quantizer, dataset, out); + } + }); }); } diff --git a/c/tests/CMakeLists.txt b/c/tests/CMakeLists.txt index cfceb58f33..e52d1f2208 100644 --- a/c/tests/CMakeLists.txt +++ b/c/tests/CMakeLists.txt @@ -95,6 +95,7 @@ ConfigureTest( ) ConfigureTest(NAME PCA_C_TEST PATH preprocessing/pca_c.cu) +ConfigureTest(NAME SCALAR_QUANTIZATION_C_TEST PATH preprocessing/scalar_c.cu) if(BUILD_CAGRA_HNSWLIB) ConfigureTest(NAME HNSW_C_TEST PATH neighbors/ann_hnsw_c.cu) diff --git a/c/tests/preprocessing/scalar_c.cu b/c/tests/preprocessing/scalar_c.cu new file mode 100644 index 0000000000..3ed1d52d51 --- /dev/null +++ b/c/tests/preprocessing/scalar_c.cu @@ -0,0 +1,163 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "../../src/core/interop.hpp" +#include +#include + +#include +#include +#include +#include + +#include +#include + +namespace { + +template +class ScalarQuantizerC : public ::testing::Test { + protected: + cuvsResources_t res{}; + cudaStream_t stream{}; + cuvsScalarQuantizerParams_t params{}; + cuvsScalarQuantizer_t quantizer{}; + std::vector values{0.f, 10.f, 9.f, 254.f, 137.f, 9.f}; + std::vector host_codes = std::vector(values.size()); + rmm::device_uvector input{0, stream}; + rmm::device_uvector codes{0, stream}; + DLManagedTensor host{}, device{}, encoded{}, host_encoded{}; + + void SetUp() override + { + ASSERT_EQ(cuvsResourcesCreate(&res), CUVS_SUCCESS); + ASSERT_EQ(cuvsStreamGet(res, &stream), CUVS_SUCCESS); + ASSERT_EQ(cuvsScalarQuantizerParamsCreate(¶ms), CUVS_SUCCESS); + EXPECT_EQ(params->mode, CUVS_SCALAR_GLOBAL); + params->quantile = 1.f; + ASSERT_EQ(cuvsScalarQuantizerCreate(&quantizer), CUVS_SUCCESS); + input.resize(values.size(), stream); + codes.resize(values.size(), stream); + raft::update_device(input.data(), values.data(), values.size(), stream); + cuvs::core::to_dlpack(raft::make_host_matrix_view(values.data(), 2, 3), &host); + cuvs::core::to_dlpack(raft::make_device_matrix_view(input.data(), 2, 3), &device); + cuvs::core::to_dlpack(raft::make_device_matrix_view(codes.data(), 2, 3), &encoded); + cuvs::core::to_dlpack(raft::make_host_matrix_view(host_codes.data(), 2, 3), &host_encoded); + ASSERT_EQ(cuvsStreamSync(res), CUVS_SUCCESS); + } + + void TearDown() override + { + EXPECT_EQ(cuvsStreamSync(res), CUVS_SUCCESS); + EXPECT_EQ(cuvsScalarQuantizerDestroy(quantizer), CUVS_SUCCESS); + EXPECT_EQ(cuvsScalarQuantizerParamsDestroy(params), CUVS_SUCCESS); + for (auto tensor : {&host, &device, &encoded, &host_encoded}) { + if (tensor->deleter) { tensor->deleter(tensor); } + } + EXPECT_EQ(cuvsResourcesDestroy(res), CUVS_SUCCESS); + } + + cuvsError_t train(cuvsScalarQuantizationMode mode, DLManagedTensor& data) + { + params->mode = mode; + return cuvsScalarQuantizerTrain(res, params, &data, quantizer); + } + + void check_per_dimension_round_trip() + { + ASSERT_EQ(cuvsScalarQuantizerTransform(res, quantizer, &device, &encoded), CUVS_SUCCESS); + ASSERT_EQ(cuvsScalarQuantizerInverseTransform(res, quantizer, &encoded, &device), CUVS_SUCCESS); + std::vector decoded(values.size()); + raft::update_host(host_codes.data(), codes.data(), host_codes.size(), stream); + raft::update_host(decoded.data(), input.data(), decoded.size(), stream); + ASSERT_EQ(cuvsStreamSync(res), CUVS_SUCCESS); + EXPECT_EQ(host_codes, (std::vector{-127, -127, 0, 127, 127, 0})); + EXPECT_EQ(decoded, values); + } +}; + +} // namespace + +using ScalarTypes = ::testing::Types; +TYPED_TEST_SUITE(ScalarQuantizerC, ScalarTypes); + +TYPED_TEST(ScalarQuantizerC, PerDimensionRoundTrip) +{ + ASSERT_EQ(this->train(CUVS_SCALAR_PER_DIMENSION, this->device), CUVS_SUCCESS); + EXPECT_EQ(this->quantizer->min_, double(cuda::std::numeric_limits::min())); + EXPECT_EQ(this->quantizer->max_, double(cuda::std::numeric_limits::max())); + this->check_per_dimension_round_trip(); +} + +TYPED_TEST(ScalarQuantizerC, GlobalBoundsAndRoundTrip) +{ + std::fill_n(this->values.begin(), 3, TypeParam{-3}); + std::fill_n(this->values.begin() + 3, 3, TypeParam{7}); + auto original = this->values; + ASSERT_EQ(this->train(CUVS_SCALAR_GLOBAL, this->host), CUVS_SUCCESS); + EXPECT_EQ(this->quantizer->min_, -3); + EXPECT_EQ(this->quantizer->max_, 7); + ASSERT_EQ(cuvsScalarQuantizerTransform(this->res, this->quantizer, &this->host, &this->host_encoded), CUVS_SUCCESS); + EXPECT_EQ(this->host_codes, (std::vector{-128, -128, -128, 127, 127, 127})); + ASSERT_EQ(cuvsScalarQuantizerInverseTransform(this->res, this->quantizer, &this->host_encoded, &this->host), CUVS_SUCCESS); + EXPECT_EQ(this->values, original); + this->quantizer->min_ = 1; + this->quantizer->max_ = 5; + ASSERT_EQ(cuvsScalarQuantizerInverseTransform(this->res, this->quantizer, &this->host_encoded, &this->host), CUVS_SUCCESS); + EXPECT_EQ(this->values, (std::vector{1.f, 1.f, 1.f, 5.f, 5.f, 5.f})); +} + +using ScalarQuantizerCFloat = ScalarQuantizerC; +TEST_F(ScalarQuantizerCFloat, HostRetraining) +{ + ASSERT_EQ(train(CUVS_SCALAR_PER_DIMENSION, device), CUVS_SUCCESS); + ASSERT_EQ(train(CUVS_SCALAR_PER_DIMENSION, host), CUVS_SUCCESS); + check_per_dimension_round_trip(); +} + +TEST_F(ScalarQuantizerCFloat, RetrainingChangesMode) +{ + ASSERT_EQ(train(CUVS_SCALAR_GLOBAL, host), CUVS_SUCCESS); + ASSERT_EQ(train(CUVS_SCALAR_PER_DIMENSION, host), CUVS_SUCCESS); + check_per_dimension_round_trip(); + ASSERT_EQ(train(CUVS_SCALAR_GLOBAL, host), CUVS_SUCCESS); + EXPECT_EQ(quantizer->min_, 0); + EXPECT_EQ(quantizer->max_, 254); + EXPECT_EQ(cuvsScalarQuantizerTransform(res, quantizer, &host, &host_encoded), CUVS_SUCCESS); +} + +TEST_F(ScalarQuantizerCFloat, InvalidTrainingPreservesQuantizer) +{ + ASSERT_EQ(train(CUVS_SCALAR_PER_DIMENSION, host), CUVS_SUCCESS); + params->quantile = 0; + EXPECT_EQ(train(CUVS_SCALAR_PER_DIMENSION, host), CUVS_ERROR); + params->quantile = 1; + EXPECT_EQ(train(static_cast(2), host), CUVS_ERROR); + check_per_dimension_round_trip(); +} + +TEST_F(ScalarQuantizerCFloat, RejectsMismatchedTypes) +{ + ASSERT_EQ(train(CUVS_SCALAR_PER_DIMENSION, host), CUVS_SUCCESS); + device.dl_tensor.dtype.bits = 64; + EXPECT_EQ(cuvsScalarQuantizerTransform(res, quantizer, &device, &encoded), CUVS_ERROR); + device.dl_tensor.dtype.bits = 32; + encoded.dl_tensor.dtype.code = kDLUInt; + EXPECT_EQ(cuvsScalarQuantizerTransform(res, quantizer, &device, &encoded), CUVS_ERROR); +} + +TEST_F(ScalarQuantizerCFloat, RejectsMismatchedWidth) +{ + ASSERT_EQ(train(CUVS_SCALAR_PER_DIMENSION, host), CUVS_SUCCESS); + device.dl_tensor.shape[1] = 2; + EXPECT_EQ(cuvsScalarQuantizerTransform(res, quantizer, &device, &encoded), CUVS_ERROR); +} + +TEST_F(ScalarQuantizerCFloat, RejectsHostPerDimensionTransforms) +{ + ASSERT_EQ(train(CUVS_SCALAR_PER_DIMENSION, host), CUVS_SUCCESS); + EXPECT_EQ(cuvsScalarQuantizerTransform(res, quantizer, &host, &host_encoded), CUVS_ERROR); + EXPECT_EQ(cuvsScalarQuantizerInverseTransform(res, quantizer, &host_encoded, &host), CUVS_ERROR); +} diff --git a/cpp/include/cuvs/core/dataset/sq.hpp b/cpp/include/cuvs/core/dataset/sq.hpp index bc875784f2..7f636c33d8 100644 --- a/cpp/include/cuvs/core/dataset/sq.hpp +++ b/cpp/include/cuvs/core/dataset/sq.hpp @@ -16,12 +16,13 @@ #include #include +#include namespace CUVS_EXPORT cuvs { /** * Scalar-quantized dataset specification: every component of a vector becomes one code, so the - * manifest is only the range the codes are mapped back onto and the dimension is still a property + * manifest holds the quantizer and the dimension is still a property * of the data. */ template @@ -33,10 +34,10 @@ struct sq_spec { using math_type = MathT; - // Nothing to allocate: the whole manifest is the interval the codes are dequantized into, - // which is exactly what the scalar quantizer holds. - using manifest_type = cuvs::preprocessing::quantize::scalar::quantizer; - using manifest_view_type = manifest_type; + // The manifest owns either global bounds or per-dimension bounds. + using manifest_type = cuvs::preprocessing::quantize::scalar::quantizer; + using manifest_view_type = typename manifest_type::view_type; + using const_manifest_view_type = typename manifest_type::const_view_type; template struct apply : storage_spec { @@ -45,21 +46,28 @@ struct sq_spec { using code_type = int8_t; // NOTE: re-exported rather than declared here, so that they do not belong to one instantiation // of `apply`; qualified lookup into `apply` does not reach the enclosing spec on its own. - using math_type = sq_spec::math_type; - using manifest_type = sq_spec::manifest_type; - using manifest_view_type = sq_spec::manifest_view_type; + using math_type = sq_spec::math_type; + using manifest_type = sq_spec::manifest_type; + using manifest_view_type = sq_spec::manifest_view_type; + using const_manifest_view_type = sq_spec::const_manifest_view_type; // NOTE: every accessor is inherited; the dimension is still one code per component, so even // `get_dim` is the dense one. - [[nodiscard]] static auto get_manifest_view(const manifest_type& manifest) noexcept + /** Borrow mutable per-dimension bounds from the owning manifest. */ + [[nodiscard]] static auto get_manifest_view(manifest_type& manifest) noexcept -> manifest_view_type { - return manifest; + return manifest.view(); } - // NOTE: the manifest is its own view here, read-only or not, so one overload serves both the - // owning dataset and the view. + /** Borrow read-only quantizer parameters; the manifest must outlive use of the view. */ [[nodiscard]] static auto get_manifest_const_view(const manifest_type& manifest) noexcept - -> manifest_view_type + -> const_manifest_view_type + { + return manifest.view(); + } + /** Convert a mutable view to read-only; its owner must outlive stream work using it. */ + [[nodiscard]] static auto get_manifest_const_view(manifest_view_type manifest) noexcept + -> const_manifest_view_type { return manifest; } @@ -75,4 +83,75 @@ template using host_sq_dataset = sq_dataset>>; +namespace preprocessing::quantize::scalar { + +/** + * @defgroup scalar_dataset Scalar-quantized dataset construction + * @{ + */ +/** + * @brief Train and encode a device corpus with the selected quantization mode. + * For separate training and encoding, use train and transform. + * + * @param[in] res Resources used for training and allocation. + * @param[in] params Scalar quantizer parameters. + * @param[in] dataset Device corpus to train on and encode. + * @return Device dataset owning the quantized codes and quantizer. + */ +template +inline device_sq_dataset make_sq_dataset( + raft::resources const& res, params params, raft::device_matrix_view dataset) +{ + auto quantizer = train(res, params, dataset); + auto codes = raft::make_device_matrix(res, dataset.extent(0), dataset.extent(1)); + transform(res, quantizer, dataset, codes.view()); + return device_sq_dataset{std::move(codes), std::move(quantizer)}; +} + +/** + * @brief Train and encode a host corpus in GPU batches, returning device-resident int8 codes. + * Uses one quantizer across all batches and measures the maximum error over every row. + * Configure params.device_buffer_samples to bound input staging and an auxiliary stream pool + * on res to enable double buffering. All int8 codes must fit in device memory. + * + * @param[in] res Resources used for training, allocation, and transfer streams. + * @param[in] params Scalar quantizer and host batching parameters. + * @param[in] dataset Row-major host corpus; pageable and pinned input are supported. + * @return Device dataset owning the quantized codes and quantizer. + */ +device_sq_dataset make_sq_dataset( + raft::resources const& res, params params, raft::host_matrix_view dataset); + +/** + * @brief Train and encode a host corpus in GPU batches, returning device-resident int8 codes. + * Uses one quantizer across all batches and measures the maximum error over every row. + * Configure params.device_buffer_samples to bound input staging and an auxiliary stream pool + * on res to enable double buffering. All int8 codes must fit in device memory. + * + * @param[in] res Resources used for training, allocation, and transfer streams. + * @param[in] params Scalar quantizer and host batching parameters. + * @param[in] dataset Row-major host corpus; pageable and pinned input are supported. + * @return Device dataset owning the quantized codes and quantizer. + */ +device_sq_dataset make_sq_dataset( + raft::resources const& res, params params, raft::host_matrix_view dataset); + +/** + * @brief Train and encode a host corpus in GPU batches, returning device-resident int8 codes. + * Uses one quantizer across all batches and measures the maximum error over every row. + * Configure params.device_buffer_samples to bound input staging and an auxiliary stream pool + * on res to enable double buffering. All int8 codes must fit in device memory. + * + * @param[in] res Resources used for training, allocation, and transfer streams. + * @param[in] params Scalar quantizer and host batching parameters. + * @param[in] dataset Row-major host corpus; pageable and pinned input are supported. + * @return Device dataset owning the quantized codes and quantizer. + */ +device_sq_dataset make_sq_dataset( + raft::resources const& res, params params, raft::host_matrix_view dataset); + +/** @} */ + +} // namespace preprocessing::quantize::scalar + } // namespace CUVS_EXPORT cuvs diff --git a/cpp/include/cuvs/preprocessing/quantize/scalar.hpp b/cpp/include/cuvs/preprocessing/quantize/scalar.hpp index 9cb0afb537..1fc452132b 100644 --- a/cpp/include/cuvs/preprocessing/quantize/scalar.hpp +++ b/cpp/include/cuvs/preprocessing/quantize/scalar.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -14,6 +14,11 @@ #include #include +#include +#include +#include +#include + namespace CUVS_EXPORT cuvs { namespace preprocessing { namespace quantize { @@ -24,6 +29,9 @@ namespace scalar { * @{ */ +/** Select global or per-dimension quantization bounds. */ +enum class quantization_mode { GLOBAL, PER_DIMENSION }; + /** * @brief quantizer parameters. */ @@ -33,23 +41,135 @@ struct params { * Needs to be within range of (0, 1]. */ float quantile = 0.99; + /** + * Per-dimension training uses every row of a nonempty, finite dataset. + * It applies the quantile independently to each dimension to select lower and upper bounds. + * Transforms derive each dimension's shift from the midpoint of its bounds, rounded to the input + * type. Each dimension's scale fits its bounds around that shift into the symmetric range of + * quantized values. + * + * For a zero-width interval, its scale is one. Otherwise, each scale is at least + * the smallest positive normal value of the scale type. + */ + quantization_mode range = quantization_mode::GLOBAL; + /** + * Maximum rows staged on device for host training and host dataset build. + * Default: 65536. Zero processes the entire host input in one batch. Ignored for device input. + * Multiple batches use two input buffers when res has an auxiliary stream pool, and one + * otherwise. Pageable input uses batch-sized pinned staging buffers, without registering + * the full input. Each input buffer uses at most device_buffer_samples * dim * sizeof(T) + * bytes, in addition to reduction and encoding workspace. Substitute n_rows when it is zero. + * + * Per-dimension host training with quantile < 1 selects exact quantiles using one column of + * host scratch (n_rows * sizeof(T) bytes). Global training retains its sampled quantiles. + */ + int64_t device_buffer_samples = 1 << 16; }; /** - * @brief Defines and stores scalar for quantisation upon training - * - * The quantization is performed by a linear mapping of an interval in the - * float data type to the full range of the quantized int type. - * - * @tparam T data element type - * + * Non-owning scalar quantizer bounds, with the constness semantics of an mdspan: + * quantizer_view can modify per-dimension bounds; quantizer_view cannot. + * Global bounds and max_error are copied. The owner must outlive stream work using the vectors. + */ +template +struct quantizer_view { + using value_type = std::remove_const_t; + using scale_type = std::conditional_t, float, value_type>; + using vector_view_type = raft::device_vector_view; + /** Global bounds. Transforms derive an affine mapping in double precision. + * Encode (x - shift) / scale, rounding to nearest with ties away from zero and clamping + * to the full quantized range. Decode shift + scale * code, saturating to finite input limits. + * Equal bounds represent a constant range; every code decodes to that value. */ + struct global_type { + /** Lower bound shared by all dimensions. */ + value_type min = value_type{0}; + /** Upper bound shared by all dimensions. */ + value_type max = value_type{0}; + }; + /** Per-dimension bounds. Transforms derive midpoint shifts and scales from these bounds. + * Encode (x - shift) / scale, rounding to nearest with ties to even and + * clamping to the symmetric range of quantized values. + * Decode shift + scale * code, saturating to finite limits of the input type. */ + struct per_dim_type { + /** One finite lower bound per dimension, with mins[j] <= maxs[j]. */ + vector_view_type mins; + /** One finite upper bound per dimension; must have the same extent as mins. */ + vector_view_type maxs; + }; + /** Global bounds or borrowed per-dimension bounds. */ + std::variant value; + /** Maximum training-row L2 reconstruction error; see quantizer::max_error. */ + scale_type max_error = INFINITY; + + /** A mutable view decays to the read-only one, the way an mdspan does. */ + template + requires(!std::is_const_v && std::is_same_v) + operator quantizer_view() const noexcept // NOLINT(google-explicit-constructor) + { + using const_view_type = quantizer_view; + if (auto const* range = std::get_if(&value)) { + return {typename const_view_type::global_type{range->min, range->max}, max_error}; + } + auto const& bounds = std::get(value); + return {typename const_view_type::per_dim_type{raft::make_const_mdspan(bounds.mins), + raft::make_const_mdspan(bounds.maxs)}, + max_error}; + } +}; + +/** + * Owns global bounds or device bounds for each dimension. + * Move-only; copy view() to borrow the parameters. */ template struct quantizer { - /** Minimum value of the quantization range. */ - T min_; - /** Maximum value of the quantization range. */ - T max_; + static_assert(!std::is_const_v, "An owning quantizer cannot hold const elements."); + + using view_type = quantizer_view; + using const_view_type = quantizer_view; + using scale_type = typename view_type::scale_type; + using global_type = typename view_type::global_type; + /** Owning per-dimension bounds. */ + struct per_dim_type { + raft::device_vector mins; + raft::device_vector maxs; + }; + /** Global bounds or owning per-dimension bounds. */ + std::variant value; + /** + * Upper bound on the L2 error between a training row and its int8 encode/decode result. + * train measures every input row, including clipped outliers, and rounds the maximum up + * with a relative 1e-6 margin. This bound applies only to the training data, not later queries. + * It is infinity when constructed from parameters alone or when the error is not finite. + */ + scale_type max_error = INFINITY; + + /** Construct default global bounds. */ + quantizer() : value{global_type{}} {} + /** Construct from finite global bounds with min <= max. */ + quantizer(T min, T max); + /** Copy nonempty device bounds of equal length, with finite min <= max. */ + quantizer(raft::resources const& res, + raft::device_vector_view mins, + raft::device_vector_view maxs); + + /** Borrow mutable per-dimension bounds; the owner must outlive stream work using them. */ + [[nodiscard]] auto view() noexcept -> view_type + { + if (auto const* range = std::get_if(&value)) { return {*range, max_error}; } + auto& bounds = std::get(value); + return {typename view_type::per_dim_type{bounds.mins.view(), bounds.maxs.view()}, max_error}; + } + /** Borrow read-only bounds; the owner must outlive stream work using per-dimension bounds. */ + [[nodiscard]] auto view() const noexcept -> const_view_type + { + if (auto const* range = std::get_if(&value)) { + return {typename const_view_type::global_type{range->min, range->max}, max_error}; + } + auto const& bounds = std::get(value); + return {typename const_view_type::per_dim_type{bounds.mins.view(), bounds.maxs.view()}, + max_error}; + } }; /** @@ -117,20 +237,23 @@ void transform(raft::resources const& res, raft::device_matrix_view out); /** - * @brief Applies quantization transform to given dataset + * @brief Applies global quantization to a host dataset + * + * Throws raft::logic_error for per-dimension parameters. Use device transforms or host build + * for per-dimension quantization. * * Usage example: * @code{.cpp} * raft::handle_t handle; * cuvs::preprocessing::quantize::scalar::params params; - * auto quantizer = cuvs::preprocessing::quantize::scalar::train(handle, params, + * auto quantizer = cuvs::preprocessing::quantize::scalar::train(handle, params, * dataset); auto quantized_dataset = raft::make_host_matrix(samples, features); * cuvs::preprocessing::quantize::scalar::transform(handle, quantizer, dataset, * quantized_dataset.view()); * @endcode * * @param[in] res raft resource - * @param[in] quantizer a scalar quantizer + * @param[in] quantizer a scalar quantizer with global parameters * @param[in] dataset a row-major matrix view on host * @param[out] out a row-major matrix view on host * @@ -168,7 +291,10 @@ void inverse_transform(raft::resources const& res, raft::device_matrix_view out); /** - * @brief Perform inverse quantization step on previously quantized dataset + * @brief Applies inverse global quantization to a host dataset + * + * Throws raft::logic_error for per-dimension parameters. Use device inverse_transform + * for per-dimension quantization. * * Note that depending on the chosen data types train dataset the conversion is * not lossless. @@ -179,11 +305,11 @@ void inverse_transform(raft::resources const& res, * cuvs::preprocessing::quantize::scalar::transform(handle, quantizer, dataset, * quantized_dataset.view()); auto dataset_revert = raft::make_host_matrix(samples, * features); cuvs::preprocessing::quantize::scalar::inverse_transform(handle, quantizer, - * dataset_revert.view()); + * raft::make_const_mdspan(quantized_dataset.view()), dataset_revert.view()); * @endcode * * @param[in] res raft resource - * @param[in] quantizer a scalar quantizer + * @param[in] quantizer a scalar quantizer with global parameters * @param[in] dataset a row-major matrix view on host * @param[out] out a row-major matrix view on host * @@ -258,20 +384,23 @@ void transform(raft::resources const& res, raft::device_matrix_view out); /** - * @brief Applies quantization transform to given dataset + * @brief Applies global quantization to a host dataset + * + * Throws raft::logic_error for per-dimension parameters. Use device transforms or host build + * for per-dimension quantization. * * Usage example: * @code{.cpp} * raft::handle_t handle; * cuvs::preprocessing::quantize::scalar::params params; - * auto quantizer = cuvs::preprocessing::quantize::scalar::train(handle, params, + * auto quantizer = cuvs::preprocessing::quantize::scalar::train(handle, params, * dataset); auto quantized_dataset = raft::make_host_matrix(samples, features); * cuvs::preprocessing::quantize::scalar::transform(handle, quantizer, dataset, * quantized_dataset.view()); * @endcode * * @param[in] res raft resource - * @param[in] quantizer a scalar quantizer + * @param[in] quantizer a scalar quantizer with global parameters * @param[in] dataset a row-major matrix view on host * @param[out] out a row-major matrix view on host * @@ -308,7 +437,10 @@ void inverse_transform(raft::resources const& res, raft::device_matrix_view out); /** - * @brief Perform inverse quantization step on previously quantized dataset + * @brief Applies inverse global quantization to a host dataset + * + * Throws raft::logic_error for per-dimension parameters. Use device inverse_transform + * for per-dimension quantization. * * Note that depending on the chosen data types train dataset the conversion is * not lossless. @@ -319,11 +451,11 @@ void inverse_transform(raft::resources const& res, * cuvs::preprocessing::quantize::scalar::transform(handle, quantizer, dataset, * quantized_dataset.view()); auto dataset_revert = raft::make_host_matrix(samples, * features); cuvs::preprocessing::quantize::scalar::inverse_transform(handle, quantizer, - * dataset_revert.view()); + * raft::make_const_mdspan(quantized_dataset.view()), dataset_revert.view()); * @endcode * * @param[in] res raft resource - * @param[in] quantizer a scalar quantizer + * @param[in] quantizer a scalar quantizer with global parameters * @param[in] dataset a row-major matrix view on host * @param[out] out a row-major matrix view on host * @@ -398,20 +530,23 @@ void transform(raft::resources const& res, raft::device_matrix_view out); /** - * @brief Applies quantization transform to given dataset + * @brief Applies global quantization to a host dataset + * + * Throws raft::logic_error for per-dimension parameters. Use device transforms or host build + * for per-dimension quantization. * * Usage example: * @code{.cpp} * raft::handle_t handle; * cuvs::preprocessing::quantize::scalar::params params; - * auto quantizer = cuvs::preprocessing::quantize::scalar::train(handle, params, + * auto quantizer = cuvs::preprocessing::quantize::scalar::train(handle, params, * dataset); auto quantized_dataset = raft::make_host_matrix(samples, features); * cuvs::preprocessing::quantize::scalar::transform(handle, quantizer, dataset, * quantized_dataset.view()); * @endcode * * @param[in] res raft resource - * @param[in] quantizer a scalar quantizer + * @param[in] quantizer a scalar quantizer with global parameters * @param[in] dataset a row-major matrix view on host * @param[out] out a row-major matrix view on host * @@ -448,7 +583,10 @@ void inverse_transform(raft::resources const& res, raft::device_matrix_view out); /** - * @brief Perform inverse quantization step on previously quantized dataset + * @brief Applies inverse global quantization to a host dataset + * + * Throws raft::logic_error for per-dimension parameters. Use device inverse_transform + * for per-dimension quantization. * * Note that depending on the chosen data types train dataset the conversion is * not lossless. @@ -459,11 +597,11 @@ void inverse_transform(raft::resources const& res, * cuvs::preprocessing::quantize::scalar::transform(handle, quantizer, dataset, * quantized_dataset.view()); auto dataset_revert = raft::make_host_matrix(samples, * features); cuvs::preprocessing::quantize::scalar::inverse_transform(handle, quantizer, - * dataset_revert.view()); + * raft::make_const_mdspan(quantized_dataset.view()), dataset_revert.view()); * @endcode * * @param[in] res raft resource - * @param[in] quantizer a scalar quantizer + * @param[in] quantizer a scalar quantizer with global parameters * @param[in] dataset a row-major matrix view on host * @param[out] out a row-major matrix view on host * diff --git a/cpp/src/preprocessing/quantize/detail/scalar.cuh b/cpp/src/preprocessing/quantize/detail/scalar.cuh index 555aa7707b..d4549eecd9 100644 --- a/cpp/src/preprocessing/quantize/detail/scalar.cuh +++ b/cpp/src/preprocessing/quantize/detail/scalar.cuh @@ -5,6 +5,9 @@ #pragma once +#include "../../../cluster/detail/kmeans_batch_loader.cuh" +#include "../../../neighbors/detail/ann_utils.cuh" + #include #include #include @@ -13,11 +16,20 @@ #include #include #include +#include + +#include +#include +#include +#include #include #include #include +#include +#include #include +#include namespace cuvs::preprocessing::quantize::scalar::detail { @@ -38,39 +50,145 @@ struct fp_lt_op { _RAFT_HOST_DEVICE bool operator()(const T& a, const T& b) const { return fp_lt(a, b); } }; -template +template +typename quantizer::global_type train(T min, T max) +{ + double const lo = static_cast(min); + double const hi = static_cast(max); + RAFT_EXPECTS(std::isfinite(lo) && std::isfinite(hi) && lo <= hi, + "Bounds must be finite with min <= max"); + return {min, max}; +} + +template struct quantize_op { - const T min_; - const T max_; - const QuantI q_type_min_ = std::numeric_limits::min(); - const QuantI q_type_max_ = std::numeric_limits::max(); - const TempT scalar_; - const TempT offset_; - - constexpr explicit quantize_op(T min, T max) - : min_(min), - max_(max), - scalar_(static_cast(max_) > static_cast(min_) - ? ((static_cast(q_type_max_) - static_cast(q_type_min_)) / - (static_cast(max_) - static_cast(min_))) - : static_cast(1)), - offset_(static_cast(q_type_min_) - static_cast(min_) * scalar_) + double shift; + double scale; + + explicit quantize_op(typename quantizer_view::global_type bounds) { + auto const range = detail::train(bounds.min, bounds.max); + double const lo = static_cast(range.min); + double const hi = static_cast(range.max); + if (lo == hi) { + shift = lo; + scale = 0; + return; + } + constexpr double q_min = std::numeric_limits::min(); + constexpr double q_max = std::numeric_limits::max(); + // Divide the endpoints separately if subtracting them would overflow. + double const width = hi - lo; + scale = std::max( + std::isfinite(width) ? width / (q_max - q_min) : hi / (q_max - q_min) - lo / (q_max - q_min), + std::numeric_limits::min()); + shift = lo * 0.5 + hi * 0.5 - (q_min + q_max) * 0.5 * scale; } constexpr RAFT_INLINE_FUNCTION QuantI operator()(const T& x) const { - if (!fp_lt(min_, x)) return q_type_min_; - if (!fp_lt(x, max_)) return q_type_max_; - return static_cast(lroundf(scalar_ * static_cast(x) + offset_)); + constexpr double q_min = std::numeric_limits::min(); + constexpr double q_max = std::numeric_limits::max(); + if (scale == 0) { return static_cast(static_cast(x) <= shift ? q_min : q_max); } + double const code = (static_cast(x) - shift) / scale; + return static_cast(lroundf(fmin(fmax(code, q_min), q_max))); } constexpr RAFT_INLINE_FUNCTION T operator()(const QuantI& x) const { - return static_cast((static_cast(x) - offset_) / scalar_); + double const value = fma(scale, static_cast(x), shift); + double const limit = static_cast(cuda::std::numeric_limits::max()); + return static_cast(fmin(fmax(value, -limit), limit)); + } +}; + +template +struct per_dim_quantize_op { + using compute_t = typename quantizer_view::scale_type; + compute_t shift; + compute_t scale; + + RAFT_INLINE_FUNCTION per_dim_quantize_op(T min, T max) + { + compute_t const lo = static_cast(min); + compute_t const hi = static_cast(max); + // Halve before adding to avoid overflow, and round the midpoint to the input type. + shift = static_cast( + static_cast(lo == hi ? lo : lo * compute_t{0.5} + hi * compute_t{0.5})); + compute_t const half_range = fmax(hi - shift, shift - lo); + constexpr compute_t q_max = std::numeric_limits::max(); + scale = half_range > 0 ? fmax(half_range / q_max, std::numeric_limits::min()) + : compute_t{1}; + } + + RAFT_INLINE_FUNCTION QuantI operator()(const T& x) const + { + compute_t const inv = compute_t{1} / scale; + compute_t const code = rint((static_cast(x) - shift) * inv); + constexpr compute_t q_max = std::numeric_limits::max(); + constexpr compute_t q_min = std::max(std::numeric_limits::min(), -q_max); + return static_cast(fmin(fmax(code, q_min), q_max)); + } + + RAFT_INLINE_FUNCTION T operator()(const QuantI& x) const + { + compute_t const value = fma(scale, static_cast(x), shift); + compute_t const limit = static_cast(cuda::std::numeric_limits::max()); + return static_cast(fmin(fmax(value, -limit), limit)); } }; +template +RAFT_INLINE_FUNCTION double reconstruction_error(T x, int64_t j, Op op) +{ + double const value = static_cast(x); + double decoded; + if constexpr (std::is_same_v::per_dim_type>) { + auto const quantize = + per_dim_quantize_op{op.mins.data_handle()[j], op.maxs.data_handle()[j]}; + decoded = static_cast(quantize(quantize(x))); + } else { + decoded = static_cast(op(op(x))); + } + return isfinite(value) && isfinite(decoded) ? fabs(value - decoded) + : std::numeric_limits::infinity(); +} + +template +auto round_error_up(double error) -> typename quantizer::scale_type +{ + using bound_t = typename quantizer::scale_type; + // Apply a conservative margin, then round upwards in the public result type. + auto const bound = static_cast(error * (1.0 + 1e-6)); + return error > 0 ? std::nextafter(bound, std::numeric_limits::infinity()) : bound; +} + +template +void max_reconstruction_error(raft::resources const& res, + raft::device_matrix_view dataset, + Op op, + raft::device_vector_view errors, + raft::device_vector_view maximum, + bool inplace) +{ + raft::linalg::reduce( + res, + dataset, + errors, + 0.0, + false, + [op] __device__(T x, int64_t j) { return reconstruction_error(x, j, op); }, + [] __device__(double a, double b) { return hypot(a, b); }); + raft::linalg::reduce( + res, + raft::make_device_matrix_view(errors.data_handle(), 1, errors.size()), + maximum, + 0.0, + inplace, + raft::identity_op{}, + raft::max_op{}); +} + template std::tuple quantile_min_max( raft::resources const& res, @@ -129,58 +247,72 @@ std::tuple quantile_min_max( } template -cuvs::preprocessing::quantize::scalar::quantizer train( - raft::resources const& res, - const cuvs::preprocessing::quantize::scalar::params params, - raft::device_matrix_view dataset) +typename quantizer::per_dim_type train(raft::resources const& res, + raft::device_vector&& minima, + raft::device_vector&& maxima) { - RAFT_EXPECTS(params.quantile > 0.0 && params.quantile <= 1.0, - "quantile for scalar quantization needs to be within (0, 1] but is %f", - params.quantile); - - auto [min, max] = detail::quantile_min_max(res, dataset, params.quantile); - - RAFT_LOG_DEBUG("quantizer train min=%lf max=%lf.", double(min), double(max)); - - return cuvs::preprocessing::quantize::scalar::quantizer{min, max}; -} - -template -cuvs::preprocessing::quantize::scalar::quantizer train( - raft::resources const& res, - const cuvs::preprocessing::quantize::scalar::params params, - raft::host_matrix_view dataset) -{ - RAFT_EXPECTS(params.quantile > 0.0 && params.quantile <= 1.0, - "quantile for scalar quantization needs to be within (0, 1] but is %f", - params.quantile); - - auto [min, max] = detail::quantile_min_max(res, dataset, params.quantile); - - RAFT_LOG_DEBUG("quantizer train min=%lf max=%lf.", double(min), double(max)); - - return cuvs::preprocessing::quantize::scalar::quantizer{min, max}; + using compute_t = typename quantizer_view::scale_type; + auto const dim = minima.extent(0); + auto invalid = raft::make_device_vector(res, 1); + auto const* max_ptr = maxima.data_handle(); + raft::linalg::reduce( + res, + raft::make_device_matrix_view(minima.data_handle(), 1, dim), + invalid.view(), + 0, + false, + [max_ptr] __device__(T min, int64_t j) { + compute_t const lo = static_cast(min); + compute_t const hi = static_cast(max_ptr[j]); + return static_cast(!isfinite(lo) || !isfinite(hi) || lo > hi); + }, + raft::max_op{}); + int invalid_bounds = 0; + raft::copy(res, + raft::make_host_scalar_view(&invalid_bounds), + raft::make_device_scalar_view(invalid.data_handle())); + raft::resource::sync_stream(res); + RAFT_EXPECTS(invalid_bounds == 0, "Bounds must be finite with min <= max"); + return {std::move(minima), std::move(maxima)}; } -template -void transform(raft::resources const& res, - const cuvs::preprocessing::quantize::scalar::quantizer& quantizer, - raft::device_matrix_view dataset, - raft::device_matrix_view out) +// Encoding and decoding share the traversal; the input type selects the operator overload. +template +void apply_quantizer(raft::resources const& res, + quantizer_view quantizer, + raft::device_matrix_view dataset, + raft::device_matrix_view out) { - raft::linalg::map(res, - out, - quantize_op(quantizer.min_, quantizer.max_), - raft::make_const_mdspan(dataset)); + if (auto const* bounds = + std::get_if::per_dim_type>(&quantizer.value)) { + auto const dim = dataset.extent(1); + RAFT_EXPECTS(dim > 0 && dim == bounds->mins.extent(0) && dim == bounds->maxs.extent(0), + "Input width must match the nonempty min and max vectors"); + RAFT_EXPECTS(dataset.extent(0) == out.extent(0) && dim == out.extent(1), + "Input and output shapes must match"); + auto const* input = dataset.data_handle(); + auto const* mins = bounds->mins.data_handle(); + auto const* maxs = bounds->maxs.data_handle(); + raft::linalg::map_offset(res, out, [input, mins, maxs, dim] __device__(int64_t i) { + return per_dim_quantize_op{mins[i % dim], maxs[i % dim]}(input[i]); + }); + } else { + auto const& range = std::get::global_type>(quantizer.value); + raft::linalg::map(res, out, quantize_op{range}, dataset); + } } -template -void transform(raft::resources const& res, - const cuvs::preprocessing::quantize::scalar::quantizer& quantizer, - raft::host_matrix_view dataset, - raft::host_matrix_view out) +template +void apply_quantizer(raft::resources const& res, + quantizer_view quantizer, + raft::host_matrix_view dataset, + raft::host_matrix_view out) { - auto main_op = quantize_op(quantizer.min_, quantizer.max_); + RAFT_EXPECTS( + std::holds_alternative::global_type>(quantizer.value), + "Host quantization requires global parameters"); + auto const& range = std::get::global_type>(quantizer.value); + auto main_op = quantize_op{range}; size_t n_elements = dataset.extent(0) * dataset.extent(1); // In dry-run mode the views may be unbacked or alias the shared probe buffer. @@ -192,34 +324,174 @@ void transform(raft::resources const& res, } } -template -void inverse_transform(raft::resources const& res, - const cuvs::preprocessing::quantize::scalar::quantizer& quantizer, - raft::device_matrix_view dataset, - raft::device_matrix_view out) -{ - raft::linalg::map(res, - out, - quantize_op(quantizer.min_, quantizer.max_), - raft::make_const_mdspan(dataset)); -} - -template -void inverse_transform(raft::resources const& res, - const cuvs::preprocessing::quantize::scalar::quantizer& quantizer, - raft::host_matrix_view dataset, - raft::host_matrix_view out) +// The optional output lets dataset construction encode during the error-measurement pass. +// Host and device inputs share the same batch processing and corpus-wide error reduction. +template +quantizer train( + raft::resources const& res, + params params, + raft::mdspan, raft::row_major, Accessor> dataset, + std::optional> out = std::nullopt) { - auto main_op = quantize_op(quantizer.min_, quantizer.max_); - size_t n_elements = dataset.extent(0) * dataset.extent(1); - - // In dry-run mode the views may be unbacked or alias the shared probe buffer. - if (raft::resource::get_dry_run_flag(res)) { return; } - -#pragma omp parallel for - for (size_t i = 0; i < n_elements; ++i) { - out.data_handle()[i] = main_op(dataset.data_handle()[i]); + using compute_t = typename quantizer::scale_type; + auto const rows = dataset.extent(0); + auto const dim = dataset.extent(1); + RAFT_EXPECTS(rows > 0 && dim > 0 && dataset.data_handle() != nullptr, + "Training requires nonempty rows and dimensions"); + RAFT_EXPECTS(params.quantile > 0.f && params.quantile <= 1.f, "Quantile must be within (0, 1]"); + constexpr bool data_on_device = raft::is_device_mdspan_v; + if constexpr (!data_on_device) { + RAFT_EXPECTS(params.device_buffer_samples >= 0, "device_buffer_samples must be nonnegative"); } + auto const batch_rows = data_on_device || params.device_buffer_samples == 0 + ? rows + : std::min(rows, params.device_buffer_samples); + auto [copy_stream, prefetch] = cuvs::spatial::knn::detail::utils::get_prefetch_stream(res); + cuvs::cluster::kmeans::detail::kmeans_batch_loader batches( + res, + dataset, + batch_rows, + copy_stream, + raft::resource::get_large_workspace_resource_ref(res), + prefetch); + bool const dry_run = raft::resource::get_dry_run_flag(res); + auto for_each_batch = [&](auto&& consume, bool repeat) { + if (dry_run) { + // Account for the largest batch's workspace without submitting loader copies. + consume(raft::make_device_matrix_view(nullptr, batch_rows, dim), + int64_t{0}); + return; + } + batches.start(); + for (std::size_t pos = 0; pos < batches.num_batches(); ++pos) { + auto batch = batches.acquire(pos); + if (pos + 1 < batches.num_batches() || repeat) { + batches.prefetch((pos + 1) % batches.num_batches()); + } + consume(raft::make_device_matrix_view(batch.data(), batch.size(), dim), + static_cast(batch.offset())); + auto const next = pos + (prefetch ? 2 : 1); + if (next < batches.num_batches() || repeat) { + batches.recycle(batch, next % batches.num_batches()); + } else { + batches.release(batch); + } + } + }; + + quantizer result; + if (params.range == quantization_mode::GLOBAL) { + auto [min, max] = quantile_min_max(res, dataset, params.quantile); + result.value = detail::train(min, max); + } else { + auto minima = raft::make_device_vector(res, dim); + auto maxima = raft::make_device_vector(res, dim); + if (params.quantile == 1.f) { + for_each_batch( + [&](auto input, int64_t offset) { + raft::linalg::reduce( + res, + input, + minima.view(), + static_cast(std::numeric_limits::infinity()), + offset != 0, + raft::identity_op{}, + raft::min_op{}); + raft::linalg::reduce( + res, + input, + maxima.view(), + static_cast(-std::numeric_limits::infinity()), + offset != 0, + raft::identity_op{}, + raft::max_op{}); + }, + true); + } else if constexpr (data_on_device) { + auto column = raft::make_device_vector(res, rows); + auto stream = raft::resource::get_cuda_stream(res).get(); + size_t workspace_bytes = 0; + RAFT_CUDA_TRY(cub::DeviceMergeSort::SortKeys( + nullptr, workspace_bytes, column.data_handle(), rows, fp_lt_op{}, stream)); + auto workspace = raft::make_device_vector(res, workspace_bytes); + auto const upper = static_cast(std::ceil((0.5 + 0.5 * params.quantile) * rows)) - 1; + auto const lower = rows - upper - 1; + auto const* input = dataset.data_handle(); + for (int64_t j = 0; j < dim; ++j) { + raft::linalg::map_offset( + res, column.view(), [input, dim, j] __device__(int64_t i) { return input[i * dim + j]; }); + if (!raft::resource::get_dry_run_flag(res)) { + RAFT_CUDA_TRY(cub::DeviceMergeSort::SortKeys(workspace.data_handle(), + workspace_bytes, + column.data_handle(), + rows, + fp_lt_op{}, + stream)); + } + raft::copy(res, + raft::make_device_scalar_view(minima.data_handle() + j), + raft::make_device_scalar_view(column.data_handle() + lower)); + raft::copy(res, + raft::make_device_scalar_view(maxima.data_handle() + j), + raft::make_device_scalar_view(column.data_handle() + upper)); + } + } else { + // Select over the complete column, not independently fitted batch quantiles. + std::vector mins(dim, T{0}), maxs(dim, T{0}); + if (!dry_run) { + std::vector column(rows); + auto const upper = + static_cast(std::ceil((0.5 + 0.5 * params.quantile) * rows)) - 1; + auto const lower = rows - upper - 1; + for (int64_t j = 0; j < dim; ++j) { + for (int64_t i = 0; i < rows; ++i) { + column[i] = dataset(i, j); + } + std::nth_element(column.begin(), column.begin() + upper, column.end(), fp_lt_op{}); + maxs[j] = column[upper]; + std::nth_element( + column.begin(), column.begin() + lower, column.begin() + upper + 1, fp_lt_op{}); + mins[j] = column[lower]; + } + } + raft::copy( + res, minima.view(), raft::make_host_vector_view(mins.data(), dim)); + raft::copy( + res, maxima.view(), raft::make_host_vector_view(maxs.data(), dim)); + raft::resource::sync_stream(res); + } + result.value = train(res, std::move(minima), std::move(maxima)); + } + + auto errors = raft::make_device_vector(res, batch_rows); + auto maximum = raft::make_device_vector(res, 1); + quantizer_view view = result.view(); + for_each_batch( + [&](auto input, int64_t offset) { + auto batch_errors = + raft::make_device_vector_view(errors.data_handle(), input.extent(0)); + if (auto const* bounds = + std::get_if::per_dim_type>(&view.value)) { + max_reconstruction_error( + res, input, *bounds, batch_errors, maximum.view(), offset != 0); + } else { + auto const& range = std::get::global_type>(view.value); + max_reconstruction_error( + res, input, quantize_op{range}, batch_errors, maximum.view(), offset != 0); + } + if (out) { + auto batch_codes = raft::make_device_matrix_view( + out->data_handle() + offset * dim, input.extent(0), dim); + apply_quantizer(res, view, input, batch_codes); + } + }, + false); + double error = std::numeric_limits::infinity(); + raft::copy( + res, raft::make_host_scalar_view(&error), raft::make_device_scalar_view(maximum.data_handle())); + raft::resource::sync_stream(res); + result.max_error = round_error_up(error); + return result; } } // namespace cuvs::preprocessing::quantize::scalar::detail diff --git a/cpp/src/preprocessing/quantize/scalar.cu b/cpp/src/preprocessing/quantize/scalar.cu index 012d36b932..e4efcf16ea 100644 --- a/cpp/src/preprocessing/quantize/scalar.cu +++ b/cpp/src/preprocessing/quantize/scalar.cu @@ -1,55 +1,86 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ +#include +#include + #include "./detail/scalar.cuh" #include namespace cuvs::preprocessing::quantize::scalar { -#define CUVS_INST_QUANTIZATION(T, QuantI) \ - auto train(raft::resources const& res, \ - const params params, \ - raft::device_matrix_view dataset) -> quantizer \ - { \ - return detail::train(res, params, dataset); \ - } \ - auto train(raft::resources const& res, \ - const params params, \ - raft::host_matrix_view dataset) -> quantizer \ - { \ - return detail::train(res, params, dataset); \ - } \ - void transform(raft::resources const& res, \ - const quantizer& quantizer, \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view out) \ - { \ - detail::transform(res, quantizer, dataset, out); \ - } \ - void transform(raft::resources const& res, \ - const quantizer& quantizer, \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view out) \ - { \ - detail::transform(res, quantizer, dataset, out); \ - } \ - void inverse_transform(raft::resources const& res, \ - const quantizer& quantizer, \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view out) \ - { \ - detail::inverse_transform(res, quantizer, dataset, out); \ - } \ - void inverse_transform(raft::resources const& res, \ - const quantizer& quantizer, \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view out) \ - { \ - detail::inverse_transform(res, quantizer, dataset, out); \ - } \ +template +quantizer::quantizer(T min, T max) : value{detail::train(min, max)} +{ +} + +template +quantizer::quantizer(raft::resources const& res, + raft::device_vector_view mins, + raft::device_vector_view maxs) +{ + RAFT_EXPECTS(mins.extent(0) > 0 && mins.extent(0) == maxs.extent(0), + "Bounds must be nonempty and have equal lengths"); + auto minima = raft::make_device_vector(res, mins.extent(0)); + auto maxima = raft::make_device_vector(res, maxs.extent(0)); + raft::copy(res, minima.view(), mins); + raft::copy(res, maxima.view(), maxs); + value = detail::train(res, std::move(minima), std::move(maxima)); +} + +#define CUVS_INST_QUANTIZATION(T, QuantI) \ + auto train(raft::resources const& res, \ + const params params, \ + raft::device_matrix_view dataset) -> quantizer \ + { \ + return detail::train(res, params, dataset); \ + } \ + auto train(raft::resources const& res, \ + const params params, \ + raft::host_matrix_view dataset) -> quantizer \ + { \ + return detail::train(res, params, dataset); \ + } \ + void transform(raft::resources const& res, \ + const quantizer& quantizer, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view out) \ + { \ + detail::apply_quantizer(res, quantizer.view(), dataset, out); \ + } \ + void transform(raft::resources const& res, \ + const quantizer& quantizer, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view out) \ + { \ + detail::apply_quantizer(res, quantizer.view(), dataset, out); \ + } \ + void inverse_transform(raft::resources const& res, \ + const quantizer& quantizer, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view out) \ + { \ + detail::apply_quantizer(res, quantizer.view(), dataset, out); \ + } \ + void inverse_transform(raft::resources const& res, \ + const quantizer& quantizer, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view out) \ + { \ + detail::apply_quantizer(res, quantizer.view(), dataset, out); \ + } \ + auto make_sq_dataset( \ + raft::resources const& res, params params, raft::host_matrix_view dataset) \ + -> device_sq_dataset \ + { \ + auto codes = \ + raft::make_device_matrix(res, dataset.extent(0), dataset.extent(1)); \ + auto quantizer = detail::train(res, params, dataset, codes.view()); \ + return {std::move(codes), std::move(quantizer)}; \ + } \ template struct quantizer; CUVS_INST_QUANTIZATION(double, int8_t); diff --git a/cpp/tests/preprocessing/scalar_quantization.cu b/cpp/tests/preprocessing/scalar_quantization.cu index e053c99d56..03022edbab 100644 --- a/cpp/tests/preprocessing/scalar_quantization.cu +++ b/cpp/tests/preprocessing/scalar_quantization.cu @@ -3,9 +3,13 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "../../src/preprocessing/quantize/detail/scalar.cuh" #include "../test_utils.cuh" +#include #include +#include #include +#include #include #include #include @@ -14,6 +18,65 @@ namespace cuvs::preprocessing::quantize::scalar { +// Compute row norms in extended precision without squaring extreme residuals. +template +long double expectReconstructionErrorBound( + const T* original, const T* reconstructed, int64_t rows, int64_t dim, long double bound) +{ + long double maximum = 0; + for (int64_t row = 0; row < rows; ++row) { + SCOPED_TRACE(row); + long double error = 0; + for (int64_t col = 0; col < dim; ++col) { + auto i = row * dim + col; + auto delta = static_cast(static_cast(original[i])) - + static_cast(static_cast(reconstructed[i])); + error = std::hypot(error, delta); + } + EXPECT_GE(bound, error); + maximum = std::max(maximum, error); + } + return maximum; +} + +template +auto makeDeviceInput(raft::resources const& res, std::vector const& values, int64_t dim) +{ + auto input = raft::make_device_matrix(res, values.size() / dim, dim); + raft::update_device( + input.data_handle(), values.data(), values.size(), raft::resource::get_cuda_stream(res)); + // The caller may pass a temporary host vector. + raft::resource::sync_stream(res); + return input; +} + +template +std::vector copyToHost(raft::resources const& res, const T* data, size_t size) +{ + std::vector values(size); + raft::update_host(values.data(), data, size, raft::resource::get_cuda_stream(res)); + raft::resource::sync_stream(res); + return values; +} + +template +void expectDeviceValues(raft::resources const& res, + const T* device_values, + std::vector const& expected) +{ + EXPECT_EQ(copyToHost(res, device_values, expected.size()), expected); +} + +template +std::vector makeShiftedDataset(int64_t rows, int64_t dim) +{ + std::vector values(rows * dim); + for (int64_t i = 0; i < rows * dim; ++i) { + values[i] = static_cast((i / dim - rows / 2) * .17 + (i % dim % 5) * 20.0); + } + return values; +} + template struct QuantizationInputs { cuvs::preprocessing::quantize::scalar::params quantization_params; @@ -115,7 +178,8 @@ class QuantizationTest : public ::testing::TestWithParam> }, cuvs::alloc_behavior::DATA_DRIVEN, min_train_alloc); - std::cerr << "Q1: min = " << (double)quantizer_1.min_ << ", max = " << (double)quantizer_1.max_ + auto bounds_1 = std::get::global_type>(quantizer_1.value); + std::cerr << "Q1: min = " << (double)bounds_1.min << ", max = " << (double)bounds_1.max << std::endl; { @@ -167,6 +231,12 @@ class QuantizationTest : public ::testing::TestWithParam> h, quantizer_1, quantized_input_h_const_view, re_transformed_input_h.view()); }, cuvs::alloc_behavior::NO_ALLOCATIONS); + // Also covers the rows outside the range-fitting sample on the large input. + expectReconstructionErrorBound(host_input_.data(), + re_transformed_input_h.data_handle(), + rows_, + cols_, + quantizer_1.max_error); auto quantized_input_d_const_view = raft::make_device_matrix_view( quantized_input_d.data_handle(), rows_, cols_); @@ -203,13 +273,14 @@ class QuantizationTest : public ::testing::TestWithParam> }, cuvs::alloc_behavior::DATA_DRIVEN, min_train_alloc); - std::cerr << "Q2: min = " << (double)quantizer_2.min_ << ", max = " << (double)quantizer_2.max_ + auto bounds_2 = std::get::global_type>(quantizer_2.value); + std::cerr << "Q2: min = " << (double)bounds_2.min << ", max = " << (double)bounds_2.max << std::endl; // check both quantizers are the same (valid if sampling is identical) if (input_.size() <= 1000000) { - ASSERT_TRUE((double)quantizer_1.min_ == (double)quantizer_2.min_); - ASSERT_TRUE((double)quantizer_1.max_ == (double)quantizer_2.max_); + ASSERT_TRUE((double)bounds_1.min == (double)bounds_2.min); + ASSERT_TRUE((double)bounds_1.max == (double)bounds_2.max); } { @@ -290,6 +361,561 @@ class QuantizationTest : public ::testing::TestWithParam> std::vector host_input_; }; +using bounds_view = quantizer_view::per_dim_type; + +template +void expectPerDimensionBounds(raft::resources const& res, + quantizer const& owner, + std::vector const& mins, + std::vector const& maxs) +{ + auto bounds = std::get::per_dim_type>(owner.view().value); + expectDeviceValues(res, bounds.mins.data_handle(), mins); + expectDeviceValues(res, bounds.maxs.data_handle(), maxs); +} + +void testPerDimensionRange(std::vector const& values, + int64_t dim, + std::vector const& expected, + float tolerance = 0.f) +{ + raft::resources res; + auto input = makeDeviceInput(res, values, dim); + auto owner = train( + res, params{1.f, quantization_mode::PER_DIMENSION}, raft::make_const_mdspan(input.view())); + auto codes = raft::make_device_matrix(res, input.extent(0), dim); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + expectDeviceValues(res, codes.data_handle(), expected); + inverse_transform(res, owner, raft::make_const_mdspan(codes.view()), input.view()); + auto decoded = copyToHost(res, input.data_handle(), values.size()); + for (size_t i = 0; i < values.size(); ++i) { + EXPECT_TRUE(std::isfinite(decoded[i])); + EXPECT_NEAR(decoded[i], values[i], tolerance) << i; + } + expectReconstructionErrorBound( + values.data(), decoded.data(), input.extent(0), dim, owner.max_error); + if (tolerance == 0.f) { + auto bounds = std::get(std::as_const(owner).view().value); + EXPECT_EQ(copyToHost(res, bounds.mins.data_handle(), dim), + copyToHost(res, bounds.maxs.data_handle(), dim)); + EXPECT_EQ(owner.max_error, 0.f); + } +} + +TEST(PerDimensionRangeTest, IndependentRanges) +{ + testPerDimensionRange( + {-127.f, -127.f / 128.f, 9.f, 1.f, 1.f / 128.f, 9.f, 127.f, 127.f / 128.f, 9.f}, + 3, + {-127, -127, 0, 1, 1, 0, 127, 127, 0}, + 1e-6f); +} + +TEST(PerDimensionRangeTest, ConstantDimensions) +{ + testPerDimensionRange({10.f, -7.f, 10.f, -7.f}, 2, {0, 0, 0, 0}); +} + +TEST(PerDimensionRangeTest, ConstantSubnormalDimensions) +{ + auto tiny = std::numeric_limits::denorm_min(); + testPerDimensionRange({tiny, -tiny}, 2, {0, 0}); +} + +TEST(PerDimensionRangeTest, SymmetricExtremeBounds) +{ + auto limit = std::numeric_limits::max(); + testPerDimensionRange({-limit, limit}, 1, {-127, 127}, limit / 254.f); +} + +TEST(PerDimensionRangeTest, AsymmetricExtremeBounds) +{ + auto limit = std::numeric_limits::max(); + testPerDimensionRange({-std::ldexp(124.f, 104), limit}, + 1, + {-127, 127}, + (std::ldexp(124.f, 103) + limit * 0.5f) / 254.f); +} + +template +class PerDimensionQuantizationTest : public ::testing::Test {}; +using PerDimensionTypes = ::testing::Types; +TYPED_TEST_SUITE(PerDimensionQuantizationTest, PerDimensionTypes); + +template +void testTrainingReconstructionError(params p) +{ + raft::resources res; + // A partial warp plus a second column per lane, with clipped outliers in different rows. + constexpr int64_t rows = 9, dim = 35; + auto values = makeShiftedDataset(rows, dim); + values.front() = T{-300.f}; + values.back() = T{500.f}; + auto input = makeDeviceInput(res, values, dim); + auto codes = raft::make_device_matrix(res, rows, dim); + auto decoded = raft::make_device_matrix(res, rows, dim); + auto owner = train(res, p, raft::make_const_mdspan(input.view())); + auto view = std::as_const(owner).view(); + EXPECT_EQ(owner.max_error, view.max_error); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + inverse_transform(res, owner, raft::make_const_mdspan(codes.view()), decoded.view()); + auto reconstructed = copyToHost(res, decoded.data_handle(), values.size()); + auto maximum = + expectReconstructionErrorBound(values.data(), reconstructed.data(), rows, dim, owner.max_error); + EXPECT_GT(maximum, 0); + EXPECT_NEAR(static_cast(owner.max_error), + static_cast(maximum), + static_cast(maximum) * 2e-6); + if (p.quantile < 1.f) { EXPECT_GT(owner.max_error, 100); } + if (p.range == quantization_mode::GLOBAL) { + auto host = train(res, + params{p.quantile}, + raft::make_host_matrix_view(values.data(), rows, dim)); + auto const host_bounds = std::get::global_type>(host.value); + auto const bounds = std::get::global_type>(owner.value); + EXPECT_EQ(host_bounds.min, bounds.min); + EXPECT_EQ(host_bounds.max, bounds.max); + if (p.quantile == 1.f) { + EXPECT_EQ(bounds.min, values.front()); + EXPECT_EQ(bounds.max, values.back()); + } + EXPECT_GE(static_cast(host.max_error), maximum); + EXPECT_NEAR(static_cast(host.max_error), + static_cast(owner.max_error), + static_cast(maximum) * 1e-6); + } +} +TYPED_TEST(PerDimensionQuantizationTest, GlobalTrainingReconstructionError) +{ + testTrainingReconstructionError(params{1.f, quantization_mode::GLOBAL}); +} + +TYPED_TEST(PerDimensionQuantizationTest, GlobalClippedReconstructionError) +{ + testTrainingReconstructionError(params{.5f, quantization_mode::GLOBAL}); +} + +TYPED_TEST(PerDimensionQuantizationTest, PerDimensionTrainingReconstructionError) +{ + testTrainingReconstructionError(params{1.f, quantization_mode::PER_DIMENSION}); +} + +TYPED_TEST(PerDimensionQuantizationTest, PerDimensionClippedReconstructionError) +{ + testTrainingReconstructionError(params{.5f, quantization_mode::PER_DIMENSION}); +} + +void testDoubleReconstructionError(double scale) +{ + raft::resources res; + std::vector values{-scale, -scale, scale * 0.3, scale * 0.3, scale, scale}; + auto input = makeDeviceInput(res, values, 2); + auto owner = train( + res, params{1.f, quantization_mode::PER_DIMENSION}, raft::make_const_mdspan(input.view())); + auto codes = raft::make_device_matrix(res, 3, 2); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + inverse_transform(res, owner, raft::make_const_mdspan(codes.view()), input.view()); + auto decoded = copyToHost(res, input.data_handle(), values.size()); + auto maximum = + expectReconstructionErrorBound(values.data(), decoded.data(), 3, 2, owner.max_error); + EXPECT_GT(maximum, 0); + EXPECT_TRUE(std::isfinite(owner.max_error)); + EXPECT_NEAR(owner.max_error, maximum, maximum * 2e-6); +} + +// Squaring these residuals would underflow or overflow, although their L2 norm is finite. +TEST(PerDimensionRangeTest, ReconstructionErrorUnderflow) { testDoubleReconstructionError(1e-300); } +TEST(PerDimensionRangeTest, ReconstructionErrorOverflow) { testDoubleReconstructionError(1e200); } + +TYPED_TEST(PerDimensionQuantizationTest, CorpusAndQueriesShareQuantizer) +{ + using T = TypeParam; + raft::resources res; + auto stream = raft::resource::get_cuda_stream(res); + std::vector corpus{0.f, 10.f, 9.f, 254.f, 137.f, 9.f}; + auto input = makeDeviceInput(res, corpus, 3); + auto owner = train( + res, params{.range = quantization_mode::PER_DIMENSION}, raft::make_const_mdspan(input.view())); + auto quantizer = std::as_const(owner).view(); + EXPECT_EQ(owner.max_error, 0); + EXPECT_EQ(quantizer.max_error, 0); + expectPerDimensionBounds(res, owner, {0.f, 10.f, 9.f}, {254.f, 137.f, 9.f}); + // The corpus manifest, rather than query extrema, controls ties and saturation. + std::vector queries{127.5f, 73.75f, 1000.f, 128.5f, 72.75f, -1000.f}; + raft::update_device(input.data_handle(), queries.data(), queries.size(), stream); + auto codes = raft::make_device_matrix(res, 2, 3); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + expectDeviceValues(res, codes.data_handle(), {0, 0, 127, 2, -2, -127}); + inverse_transform(res, owner, raft::make_const_mdspan(codes.view()), input.view()); + expectDeviceValues(res, input.data_handle(), {127.f, 73.5f, 136.f, 129.f, 72.5f, -118.f}); +} + +TEST(PerDimensionRangeTest, AdjacentHalfBounds) +{ + raft::resources res; + std::vector values{half{1.f}, half{1.f + 1.f / 1024.f}}; + auto input = makeDeviceInput(res, values, 1); + auto quantizer = train( + res, params{1.f, quantization_mode::PER_DIMENSION}, raft::make_const_mdspan(input.view())); + auto codes = raft::make_device_matrix(res, 2, 1); + transform(res, quantizer, raft::make_const_mdspan(input.view()), codes.view()); + inverse_transform(res, quantizer, raft::make_const_mdspan(codes.view()), input.view()); + expectDeviceValues(res, input.data_handle(), values); +} + +TEST(PerDimensionRangeTest, PerDimensionQuantiles) +{ + raft::resources res; + constexpr int64_t rows = 1000, dim = 3; + std::vector corpus(rows * dim); + for (int64_t i = 0; i < rows; ++i) { + corpus[i * dim] = static_cast(i); + corpus[i * dim + 1] = 10000.f + 2.f * static_cast(rows - i - 1); + corpus[i * dim + 2] = 9.f; + } + auto input = makeDeviceInput(res, corpus, dim); + auto check_bounds = [&](params p, std::vector mins, std::vector maxs) { + auto owner = train(res, p, raft::make_const_mdspan(input.view())); + expectPerDimensionBounds(res, owner, mins, maxs); + }; + check_bounds( + params{.range = quantization_mode::PER_DIMENSION}, {4.f, 10008.f, 9.f}, {995.f, 11990.f, 9.f}); + check_bounds( + params{1.f, quantization_mode::PER_DIMENSION}, {0.f, 10000.f, 9.f}, {999.f, 11998.f, 9.f}); + check_bounds( + params{.5f, quantization_mode::PER_DIMENSION}, {250.f, 10500.f, 9.f}, {749.f, 11498.f, 9.f}); + for (float q : {0.f, -1.f, 1.1f, std::numeric_limits::quiet_NaN()}) { + EXPECT_THROW( + train( + res, params{q, quantization_mode::PER_DIMENSION}, raft::make_const_mdspan(input.view())), + raft::logic_error); + } +} + +TYPED_TEST(PerDimensionQuantizationTest, CustomParametersAndValidation) +{ + using T = TypeParam; + raft::resources res; + auto mins = raft::make_device_vector(res, 2); + auto maxs = raft::make_device_vector(res, 2); + std::vector lower{-154.f, -104.f}, upper{354.f, -96.f}; + auto stream = raft::resource::get_cuda_stream(res); + raft::update_device(mins.data_handle(), lower.data(), lower.size(), stream); + raft::update_device(maxs.data_handle(), upper.data(), upper.size(), stream); + scalar::quantizer owner{ + res, raft::make_const_mdspan(mins.view()), raft::make_const_mdspan(maxs.view())}; + auto quantizer = std::as_const(owner).view(); + EXPECT_TRUE(std::isinf(owner.max_error)); + EXPECT_TRUE(std::isinf(quantizer.max_error)); + EXPECT_TRUE(std::isinf((scalar::quantizer{T{-1}, T{1}}.view().max_error))); + expectPerDimensionBounds(res, owner, lower, upper); + auto input = makeDeviceInput(res, {104.f, -106.f}, 2); + auto codes = raft::make_device_matrix(res, 1, 2); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + expectDeviceValues(res, codes.data_handle(), {2, -127}); + scalar::quantizer wrong_bounds; + wrong_bounds.value = typename scalar::quantizer::per_dim_type{ + raft::make_device_vector(res, 2), raft::make_device_vector(res, 1)}; + EXPECT_THROW(transform(res, wrong_bounds, raft::make_const_mdspan(input.view()), codes.view()), + raft::logic_error); + auto wrong_shape = raft::make_device_matrix(res, 2, 1); + EXPECT_THROW(transform(res, owner, raft::make_const_mdspan(input.view()), wrong_shape.view()), + raft::logic_error); + EXPECT_THROW( + (scalar::quantizer{res, + raft::make_const_mdspan(mins.view()), + raft::make_device_vector_view(maxs.data_handle(), 1)}), + raft::logic_error); + EXPECT_THROW((scalar::quantizer{ + res, raft::make_const_mdspan(maxs.view()), raft::make_const_mdspan(mins.view())}), + raft::logic_error); + EXPECT_THROW(train(res, + params{.range = quantization_mode::PER_DIMENSION}, + raft::make_device_matrix_view(nullptr, 0, 2)), + raft::logic_error); +} + +TYPED_TEST(PerDimensionQuantizationTest, DirectPerDimensionBounds) +{ + using T = TypeParam; + raft::resources res; + auto stream = raft::resource::get_cuda_stream(res); + auto mins = raft::make_device_vector(res, 2); + auto maxs = raft::make_device_vector(res, 2); + std::vector lower{-127.f, -254.f}, upper{127.f, 254.f}, values{1.f, 2.f}; + raft::update_device(mins.data_handle(), lower.data(), lower.size(), stream); + raft::update_device(maxs.data_handle(), upper.data(), upper.size(), stream); + quantizer owner; + owner.value = typename quantizer::per_dim_type{std::move(mins), std::move(maxs)}; + auto input = raft::make_device_matrix(res, 1, 2); + auto codes = raft::make_device_matrix(res, 1, 2); + raft::update_device(input.data_handle(), values.data(), values.size(), stream); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + expectDeviceValues(res, codes.data_handle(), {1, 1}); + inverse_transform(res, owner, raft::make_const_mdspan(codes.view()), input.view()); + expectDeviceValues(res, input.data_handle(), values); + + // Transforms use the owner's current bounds. + lower = {T{-254.f}, T{-508.f}}; + upper = {T{254.f}, T{508.f}}; + auto mutable_view = owner.view(); + quantizer_view read_only = mutable_view; + auto bounds = std::get::per_dim_type>(mutable_view.value); + auto const_bounds = std::get::per_dim_type>(read_only.value); + static_assert(std::is_same_v, T*>); + static_assert( + std::is_same_v, const T*>); + static_assert(!std::is_convertible_v, quantizer_view>); + EXPECT_EQ(bounds.mins.data_handle(), const_bounds.mins.data_handle()); + EXPECT_EQ(bounds.maxs.data_handle(), const_bounds.maxs.data_handle()); + raft::update_device(bounds.mins.data_handle(), lower.data(), lower.size(), stream); + raft::update_device(bounds.maxs.data_handle(), upper.data(), upper.size(), stream); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + expectDeviceValues(res, codes.data_handle(), {0, 0}); +} + +TYPED_TEST(PerDimensionQuantizationTest, DirectGlobalBounds) +{ + using T = TypeParam; + raft::resources res; + auto input = raft::make_host_matrix(1, 2); + input(0, 0) = T{1}; + input(0, 1) = T{2}; + auto codes = raft::make_host_matrix(1, 2); + quantizer owner; + owner.value = typename quantizer::global_type{T{-128}, T{127}}; + quantizer_view read_only = owner.view(); + auto copied_bounds = std::get::global_type>(read_only.value); + EXPECT_EQ(double(copied_bounds.min), -128.0); + EXPECT_EQ(double(copied_bounds.max), 127.0); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + EXPECT_EQ(codes(0, 0), 1); + EXPECT_EQ(codes(0, 1), 2); + auto& bounds = std::get::global_type>(owner.value); + bounds.min = T{-256}; + bounds.max = T{254}; + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + EXPECT_EQ(codes(0, 0), 1); + EXPECT_EQ(codes(0, 1), 1); + inverse_transform(res, owner, raft::make_const_mdspan(codes.view()), input.view()); + EXPECT_EQ(static_cast(input(0, 0)), 2); + EXPECT_EQ(static_cast(input(0, 1)), 2); +} + +TYPED_TEST(PerDimensionQuantizationTest, HostTransformsRejectPerDimensionQuantizer) +{ + using T = TypeParam; + raft::resources res; + auto values = raft::make_host_matrix(1, 1); + values(0, 0) = T{1}; + auto quantizer = train( + res, params{1.f, quantization_mode::PER_DIMENSION}, raft::make_const_mdspan(values.view())); + auto codes = raft::make_host_matrix(1, 1); + EXPECT_THROW(transform(res, quantizer, raft::make_const_mdspan(values.view()), codes.view()), + raft::logic_error); + EXPECT_THROW( + inverse_transform(res, quantizer, raft::make_const_mdspan(codes.view()), values.view()), + raft::logic_error); +} + +TEST(PerDimensionRangeTest, QuantizedTypeLimits) +{ + using QuantI = int16_t; + raft::resources res; + std::vector values{-100000.f, -32767.f, 0.f, 32767.f, 100000.f}; + auto input = makeDeviceInput(res, values, 1); + auto codes = raft::make_device_matrix(res, values.size(), 1); + auto owner = detail::train(res, + params{.5f, quantization_mode::PER_DIMENSION}, + raft::make_const_mdspan(input.view()), + codes.view()); + expectDeviceValues( + res, std::get(std::as_const(owner).view().value).maxs.data_handle(), {32767.f}); + std::vector expected{-32767, -32767, 0, 32767, 32767}; + expectDeviceValues(res, codes.data_handle(), expected); + detail::apply_quantizer( + res, std::as_const(owner).view(), raft::make_const_mdspan(codes.view()), input.view()); + auto decoded = copyToHost(res, input.data_handle(), values.size()); + EXPECT_EQ(decoded, (std::vector(expected.begin(), expected.end()))); + auto maximum = expectReconstructionErrorBound( + values.data(), decoded.data(), values.size(), 1, owner.max_error); + EXPECT_NEAR(owner.max_error, maximum, maximum * 2e-6); +} + +template +void testGlobalBounds(T min, T max) +{ + raft::resources res; + auto owner = quantizer{min, max}; + auto bounds = std::get::global_type>(owner.value); + EXPECT_EQ(bounds.min, min); + EXPECT_EQ(bounds.max, max); + auto input = makeDeviceInput(res, {T{-100}, min, max, T{100}}, 4); + auto codes = raft::make_device_matrix(res, 1, 4); + transform(res, owner, raft::make_const_mdspan(input.view()), codes.view()); + inverse_transform(res, owner, raft::make_const_mdspan(codes.view()), input.view()); + expectDeviceValues(res, input.data_handle(), {min, min, max, max}); +} + +TYPED_TEST(PerDimensionQuantizationTest, GlobalBounds) { testGlobalBounds(-3, 7); } +TYPED_TEST(PerDimensionQuantizationTest, ConstantGlobalBounds) +{ + testGlobalBounds(5, 5); +} + +TEST(QuantizationDatasetTest, GlobalRangeDataset) +{ + raft::resources res; + auto input = makeDeviceInput(res, {-1.f, 1.f}, 2); + auto encoded = make_sq_dataset(res, params{1.f}, raft::make_const_mdspan(input.view())); + typename decltype(encoded)::const_view_type read_only = encoded.view(); + EXPECT_EQ(read_only.const_manifest_view().max_error, 0.f); + EXPECT_EQ(encoded.n_rows(), 1); + EXPECT_EQ(encoded.dim(), 2u); + expectDeviceValues(res, encoded.data_const_view().data_handle(), {-128, 127}); + auto range = + std::get::global_type>(encoded.const_manifest_view().value); + EXPECT_EQ(range.min, -1.f); + EXPECT_EQ(range.max, 1.f); +} + +TEST(QuantizationDatasetTest, PerDimensionDataset) +{ + raft::resources res; + auto input = makeDeviceInput(res, {0.f, 10.f, 254.f, 14.f}, 2); + auto encoded = make_sq_dataset( + res, params{.range = quantization_mode::PER_DIMENSION}, raft::make_const_mdspan(input.view())); + typename decltype(encoded)::const_view_type read_only = encoded.view(); + EXPECT_EQ(read_only.const_manifest_view().max_error, 0.f); + EXPECT_EQ(encoded.n_rows(), 2); + EXPECT_EQ(encoded.dim(), 2u); + expectDeviceValues(res, encoded.data_const_view().data_handle(), {-127, -127, 127, 127}); + auto codes = raft::make_device_matrix(res, 2, 2); + auto quantizer = train( + res, params{.range = quantization_mode::PER_DIMENSION}, raft::make_const_mdspan(input.view())); + transform(res, quantizer, raft::make_const_mdspan(input.view()), codes.view()); + expectDeviceValues( + res, codes.data_handle(), copyToHost(res, encoded.data_const_view().data_handle(), 4)); +} + +template +void testHostDataset(params p, bool prefetch = false, bool use_pinned = false) +{ + constexpr int64_t rows = 37, dim = 35; + auto values = makeShiftedDataset(rows, dim); + // The largest clipping error precedes the final batch to check the running maximum. + values.front() = T{-1200.f}; + values.back() = T{500.f}; + raft::resources res; + if (prefetch) { + raft::resource::set_cuda_stream_pool(res, std::make_shared(1)); + } + auto input = makeDeviceInput(res, values, dim); + auto pinned = raft::make_pinned_matrix(res, rows, dim); + if (use_pinned) { std::copy(values.begin(), values.end(), pinned.data_handle()); } + // A negative staging limit is ignored for device input. + auto expected = + make_sq_dataset(res, params{p.quantile, p.range, -1}, raft::make_const_mdspan(input.view())); + auto host_input = raft::make_host_matrix_view( + use_pinned ? pinned.data_handle() : values.data(), rows, dim); + auto actual = make_sq_dataset(res, p, host_input); + EXPECT_EQ(actual.n_rows(), rows); + EXPECT_EQ(actual.dim(), dim); + expectDeviceValues(res, + actual.data_const_view().data_handle(), + copyToHost(res, expected.data_const_view().data_handle(), values.size())); + auto a = actual.const_manifest_view(); + auto e = expected.const_manifest_view(); + EXPECT_NEAR(double(a.max_error), double(e.max_error), double(e.max_error) * 1e-6); + if (p.range == quantization_mode::PER_DIMENSION) { + auto actual_bounds = std::get::per_dim_type>(a.value); + auto expected_bounds = std::get::per_dim_type>(e.value); + expectDeviceValues(res, + actual_bounds.maxs.data_handle(), + copyToHost(res, expected_bounds.maxs.data_handle(), dim)); + expectDeviceValues(res, + actual_bounds.mins.data_handle(), + copyToHost(res, expected_bounds.mins.data_handle(), dim)); + } + // Standalone training must agree without allocating corpus codes. + auto trained = train(res, p, host_input); + EXPECT_NEAR(double(trained.max_error), double(a.max_error), double(a.max_error) * 1e-6); +} + +TYPED_TEST(PerDimensionQuantizationTest, HostDatasetPartialBatch) +{ + testHostDataset(params{1.f, quantization_mode::PER_DIMENSION, 7}); +} + +struct HostDatasetInputs { + const char* name; + params quantization; + bool prefetch = false; + bool pinned = false; +}; +class HostDatasetTest : public ::testing::TestWithParam {}; +TEST_P(HostDatasetTest, MatchesDeviceDataset) +{ + auto const& p = GetParam(); + testHostDataset(p.quantization, p.prefetch, p.pinned); +} + +INSTANTIATE_TEST_SUITE_P( + HostDatasetCases, + HostDatasetTest, + ::testing::Values( + HostDatasetInputs{"SingleRowBatch", {1.f, quantization_mode::PER_DIMENSION, 1}}, + HostDatasetInputs{"FullBatch", {1.f, quantization_mode::PER_DIMENSION, 37}}, + HostDatasetInputs{"UnlimitedBatch", {1.f, quantization_mode::PER_DIMENSION, 0}}, + HostDatasetInputs{"Pinned", {1.f, quantization_mode::PER_DIMENSION, 7}, false, true}, + HostDatasetInputs{"Prefetch", {1.f, quantization_mode::PER_DIMENSION, 7}, true}, + HostDatasetInputs{"PinnedPrefetch", {1.f, quantization_mode::PER_DIMENSION, 7}, true, true}, + HostDatasetInputs{"ClippedQuantile", {.5f, quantization_mode::PER_DIMENSION, 7}}, + HostDatasetInputs{"DefaultQuantile", {.99f, quantization_mode::PER_DIMENSION, 7}}, + HostDatasetInputs{"GlobalRange", {1.f, quantization_mode::GLOBAL, 7}}, + HostDatasetInputs{"GlobalClippedQuantile", {.5f, quantization_mode::GLOBAL, 7}}), + [](auto const& info) { return info.param.name; }); + +void testHostBuildWorkspace(float quantile, bool prefetch = false) +{ + constexpr int64_t rows = 4097, dim = 35, batch_rows = 128; + std::vector values(rows * dim, 1.f); + auto input = raft::make_host_matrix_view(values.data(), rows, dim); + raft::resources res; + if (prefetch) { + raft::resource::set_cuda_stream_pool(res, std::make_shared(1)); + } + auto action = [&](raft::resources const& h) { + auto encoded = + make_sq_dataset(h, params{quantile, quantization_mode::PER_DIMENSION, batch_rows}, input); + if (!raft::resource::get_dry_run_flag(h)) { + EXPECT_EQ(encoded.const_manifest_view().max_error, 0.f); + } + }; + execute_with_dry_run_check(res, action, alloc_behavior::ARGUMENT_DRIVEN); + raft::memory_stats_resources stats(res); + action(stats); + auto peak = stats.get_bytes_peak(); + // Includes all int8 codes; excludes even one full FP32 corpus copy. + EXPECT_LT( + peak.device_workspace + peak.device_large_workspace + peak.device_global + peak.device_managed, + values.size() * sizeof(float)); +} + +TEST(QuantizationDatasetTest, HostBuildBatchWorkspace) { testHostBuildWorkspace(1.f); } +TEST(QuantizationDatasetTest, HostBuildClippedWorkspace) { testHostBuildWorkspace(.5f); } +TEST(QuantizationDatasetTest, HostBuildPrefetchWorkspace) { testHostBuildWorkspace(1.f, true); } + +TEST(QuantizationDatasetTest, HostBuildRejectsNegativeBatchSize) +{ + std::vector values{1.f}; + EXPECT_THROW( + make_sq_dataset(raft::resources{}, + params{1.f, quantization_mode::PER_DIMENSION, -1}, + raft::make_host_matrix_view(values.data(), 1, 1)), + raft::logic_error); +} + template const std::vector> inputs = { {{1.0}, 5, 5, T(0.0), T(1.0)}, diff --git a/fern/pages/c_api/c-api-preprocessing-quantize-scalar.md b/fern/pages/c_api/c-api-preprocessing-quantize-scalar.md index 86aacd7bce..efc3e4e5e7 100644 --- a/fern/pages/c_api/c-api-preprocessing-quantize-scalar.md +++ b/fern/pages/c_api/c-api-preprocessing-quantize-scalar.md @@ -8,15 +8,44 @@ _Source header: `cuvs/preprocessing/quantize/scalar.h`_ ## C API for Scalar Quantizer + +### cuvsScalarQuantizationMode + +Select global or per-dimension scalar quantization bounds. + +```c +enum cuvsScalarQuantizationMode { + CUVS_SCALAR_GLOBAL = 0, + CUVS_SCALAR_PER_DIMENSION = 1 +}; +``` + +**Values** + +| Name | Value | +| --- | --- | +| `CUVS_SCALAR_GLOBAL` | `0` | +| `CUVS_SCALAR_PER_DIMENSION` | `1` | + ### cuvsScalarQuantizerParams -Scalar quantizer parameters. +Scalar quantizer parameters. Allocate with cuvsScalarQuantizerParamsCreate. ```c -struct cuvsScalarQuantizerParams; +struct cuvsScalarQuantizerParams { + float quantile; + enum cuvsScalarQuantizationMode mode; +}; ``` +**Fields** + +| Name | Type | Description | +| --- | --- | --- | +| `quantile` | `float` | specifies how many outliers at top & bottom will be ignored needs to be within range of (0, 1] | +| `mode` | [`enum cuvsScalarQuantizationMode`](/api-reference/c-api-preprocessing-quantize-scalar#cuvsscalarquantizationmode) | Quantization mode.
Defaults to CUVS_SCALAR_GLOBAL. | + ### cuvsScalarQuantizerParamsCreate @@ -60,7 +89,7 @@ cuvsError_t cuvsScalarQuantizerParamsDestroy(cuvsScalarQuantizerParams_t params) Defines and stores scalar for quantisation upon training -The quantization is performed by a linear mapping of an interval in the float data type to the full range of the quantized int type. +The quantization is performed by a linear mapping of an interval in the float data type to the full range of the quantized int type. Allocate with cuvsScalarQuantizerCreate and release with cuvsScalarQuantizerDestroy; do not copy the handle's contents. The implementation also owns any per-dimension bounds. In global mode, min_ and max_ are the trained bounds. In per-dimension mode, they are numeric_limits<T>::min() and numeric_limits<T>::max() for the training data type T, and are not used for transformation. ```c typedef struct { @@ -98,7 +127,7 @@ cuvsError_t cuvsScalarQuantizerCreate(cuvsScalarQuantizer_t* quantizer); ### cuvsScalarQuantizerDestroy -De-allocate Scalar Quantizer +De-allocate Scalar Quantizer after outstanding stream work completes. ```c cuvsError_t cuvsScalarQuantizerDestroy(cuvsScalarQuantizer_t quantizer); @@ -126,6 +155,8 @@ DLManagedTensor* dataset, cuvsScalarQuantizer_t quantizer); ``` +Supports float16, float32, and float64 in global or per-dimension mode. Per-dimension mode applies the quantile independently to each dimension using every row. Successful retraining replaces the previous bounds; failed training preserves them. Complete outstanding stream work using the quantizer before retraining. + **Parameters** | Name | Direction | Type | Description | @@ -133,7 +164,7 @@ cuvsScalarQuantizer_t quantizer); | `res` | in | [`cuvsResources_t`](/api-reference/c-api-core-c-api#cuvsresources-t) | raft resource | | `params` | in | [`cuvsScalarQuantizerParams_t`](/api-reference/c-api-preprocessing-quantize-scalar#cuvsscalarquantizerparams) | configure scalar quantizer, e.g. quantile | | `dataset` | in | `DLManagedTensor*` | a row-major host or device matrix | -| `quantizer` | out | [`cuvsScalarQuantizer_t`](/api-reference/c-api-preprocessing-quantize-scalar#cuvsscalarquantizer) | trained scalar quantizer | +| `quantizer` | in,out | [`cuvsScalarQuantizer_t`](/api-reference/c-api-preprocessing-quantize-scalar#cuvsscalarquantizer) | scalar quantizer allocated by cuvsScalarQuantizerCreate | **Returns** @@ -151,6 +182,8 @@ DLManagedTensor* dataset, DLManagedTensor* out); ``` +Per-dimension mode requires device matrices with the same input type and width as training. Its output codes are int8 in [-127, 127]. The quantizer must outlive stream work using its bounds. + **Parameters** | Name | Direction | Type | Description | @@ -176,7 +209,7 @@ DLManagedTensor* dataset, DLManagedTensor* out); ``` -Note that depending on the chosen data types train dataset the conversion is not lossless. +Note that depending on the chosen data types train dataset the conversion is not lossless. Per-dimension mode requires device matrices with the same output type and width as training. The quantizer must outlive stream work using its bounds. **Parameters** diff --git a/fern/pages/cpp_api/cpp-api-core-dataset-sq.md b/fern/pages/cpp_api/cpp-api-core-dataset-sq.md index a26e57b857..9a5a308033 100644 --- a/fern/pages/cpp_api/cpp-api-core-dataset-sq.md +++ b/fern/pages/cpp_api/cpp-api-core-dataset-sq.md @@ -13,9 +13,105 @@ _Source header: `cuvs/core/dataset/sq.hpp`_ Scalar-quantized dataset specification: every component of a vector becomes one code, so the -manifest is only the range the codes are mapped back onto and the dimension is still a property of the data. +manifest holds the quantizer and the dimension is still a property of the data. ```cpp template struct sq_spec; ``` + +## Scalar-quantized dataset construction + + +### preprocessing::quantize::scalar::make_sq_dataset + +Train and encode a device corpus with the selected quantization mode. For separate training and encoding, use train and transform. + +```cpp +template +inline device_sq_dataset make_sq_dataset( +raft::resources const& res, params params, raft::device_matrix_view dataset); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | in | `raft::resources const&` | Resources used for training and allocation. | +| `params` | in | `params` | Scalar quantizer parameters. | +| `dataset` | in | `raft::device_matrix_view` | Device corpus to train on and encode. | + +**Returns** + +`inline device_sq_dataset` + +Device dataset owning the quantized codes and quantizer. + +**Additional overload:** `preprocessing::quantize::scalar::make_sq_dataset` + +Train and encode a host corpus in GPU batches, returning device-resident int8 codes. Uses one quantizer across all batches and measures the maximum error over every row. Configure params.device_buffer_samples to bound input staging and an auxiliary stream pool on res to enable double buffering. All int8 codes must fit in device memory. + +```cpp +device_sq_dataset make_sq_dataset( +raft::resources const& res, params params, raft::host_matrix_view dataset); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | in | `raft::resources const&` | Resources used for training, allocation, and transfer streams. | +| `params` | in | `params` | Scalar quantizer and host batching parameters. | +| `dataset` | in | `raft::host_matrix_view` | Row-major host corpus; pageable and pinned input are supported. | + +**Returns** + +`device_sq_dataset` + +Device dataset owning the quantized codes and quantizer. + +**Additional overload:** `preprocessing::quantize::scalar::make_sq_dataset` + +Train and encode a host corpus in GPU batches, returning device-resident int8 codes. Uses one quantizer across all batches and measures the maximum error over every row. Configure params.device_buffer_samples to bound input staging and an auxiliary stream pool on res to enable double buffering. All int8 codes must fit in device memory. + +```cpp +device_sq_dataset make_sq_dataset( +raft::resources const& res, params params, raft::host_matrix_view dataset); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | in | `raft::resources const&` | Resources used for training, allocation, and transfer streams. | +| `params` | in | `params` | Scalar quantizer and host batching parameters. | +| `dataset` | in | `raft::host_matrix_view` | Row-major host corpus; pageable and pinned input are supported. | + +**Returns** + +`device_sq_dataset` + +Device dataset owning the quantized codes and quantizer. + +**Additional overload:** `preprocessing::quantize::scalar::make_sq_dataset` + +Train and encode a host corpus in GPU batches, returning device-resident int8 codes. Uses one quantizer across all batches and measures the maximum error over every row. Configure params.device_buffer_samples to bound input staging and an auxiliary stream pool on res to enable double buffering. All int8 codes must fit in device memory. + +```cpp +device_sq_dataset make_sq_dataset( +raft::resources const& res, params params, raft::host_matrix_view dataset); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | in | `raft::resources const&` | Resources used for training, allocation, and transfer streams. | +| `params` | in | `params` | Scalar quantizer and host batching parameters. | +| `dataset` | in | `raft::host_matrix_view` | Row-major host corpus; pageable and pinned input are supported. | + +**Returns** + +`device_sq_dataset` + +Device dataset owning the quantized codes and quantizer. diff --git a/fern/pages/cpp_api/cpp-api-preprocessing-quantize-scalar.md b/fern/pages/cpp_api/cpp-api-preprocessing-quantize-scalar.md index 29638673f1..a4ad8c4ae0 100644 --- a/fern/pages/cpp_api/cpp-api-preprocessing-quantize-scalar.md +++ b/fern/pages/cpp_api/cpp-api-preprocessing-quantize-scalar.md @@ -8,6 +8,25 @@ _Source header: `cuvs/preprocessing/quantize/scalar.hpp`_ ## Scalar quantizer utilities + +### preprocessing::quantize::scalar::quantization_mode + +Select global or per-dimension quantization bounds. + +```cpp +enum class quantization_mode { + GLOBAL, + PER_DIMENSION +}; +``` + +**Values** + +| Name | Value | +| --- | --- | +| `GLOBAL` | `` | +| `PER_DIMENSION` | `` | + ### preprocessing::quantize::scalar::params @@ -16,6 +35,8 @@ quantizer parameters. ```cpp struct params { float quantile; + quantization_mode range; + int64_t device_buffer_samples; }; ``` @@ -24,19 +45,106 @@ struct params { | Name | Type | Description | | --- | --- | --- | | `quantile` | `float` | Specifies how many outliers at top & bottom will be ignored. Needs to be within range of (0, 1]. | +| `range` | [`quantization_mode`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantization-mode) | Per-dimension training uses every row of a nonempty, finite dataset. It applies the quantile independently to each dimension to select lower and upper bounds. Transforms derive each dimension's shift from the midpoint of its bounds, rounded to the input type. Each dimension's scale fits its bounds around that shift into the symmetric range of quantized values.

For a zero-width interval, its scale is one. Otherwise, each scale is at least the smallest positive normal value of the scale type. | +| `device_buffer_samples` | `int64_t` | Maximum rows staged on device for host training and host dataset build.
Default: 65536. Zero processes the entire host input in one batch. Ignored for device input. Multiple batches use two input buffers when res has an auxiliary stream pool, and one otherwise. Pageable input uses batch-sized pinned staging buffers, without registering the full input. Each input buffer uses at most device_buffer_samples * dim * sizeof(T) bytes, in addition to reduction and encoding workspace. Substitute n_rows when it is zero.

Per-dimension host training with quantile < 1 selects exact quantiles using one column of host scratch (n_rows * sizeof(T) bytes). Global training retains its sampled quantiles. | + + +### preprocessing::quantize::scalar::quantizer_view + +Non-owning scalar quantizer bounds, with the constness semantics of an mdspan: + +quantizer_view<T> can modify per-dimension bounds; quantizer_view<const T> cannot. Global bounds and max_error are copied. The owner must outlive stream work using the vectors. + +```cpp +template +struct quantizer_view { + std::variant value; + scale_type max_error; +}; +``` + +**Fields** + +| Name | Type | Description | +| --- | --- | --- | +| `value` | [`std::variant`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer-view-global-type) | Global bounds or borrowed per-dimension bounds. | +| `max_error` | `scale_type` | Maximum training-row L2 reconstruction error; see quantizer::max_error. | + + +### preprocessing::quantize::scalar::quantizer_view::global_type + +Global bounds. Transforms derive an affine mapping in double precision. + +Encode (x - shift) / scale, rounding to nearest with ties away from zero and clamping to the full quantized range. Decode shift + scale * code, saturating to finite input limits. Equal bounds represent a constant range; every code decodes to that value. + +```cpp +struct global_type { + value_type min; + value_type max; +}; +``` + +**Fields** + +| Name | Type | Description | +| --- | --- | --- | +| `min` | `value_type` | Lower bound shared by all dimensions. | +| `max` | `value_type` | Upper bound shared by all dimensions. | + + +### preprocessing::quantize::scalar::quantizer_view::per_dim_type + +Per-dimension bounds. Transforms derive midpoint shifts and scales from these bounds. + +Encode (x - shift) / scale, rounding to nearest with ties to even and clamping to the symmetric range of quantized values. Decode shift + scale * code, saturating to finite limits of the input type. + +```cpp +struct per_dim_type { + vector_view_type mins; + vector_view_type maxs; +}; +``` + +**Fields** + +| Name | Type | Description | +| --- | --- | --- | +| `mins` | `vector_view_type` | One finite lower bound per dimension, with mins[j] <= maxs[j]. | +| `maxs` | `vector_view_type` | One finite upper bound per dimension; must have the same extent as mins. | + + +### preprocessing::quantize::scalar::quantizer_view::requires + +A mutable view decays to the read-only one, the way an mdspan does. + +```cpp +template +requires(!std::is_const_v && std::is_same_v) +operator quantizer_view() const noexcept // NOLINT(google-explicit-constructor); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `arg1` | | `!std::is_const_v && std::is_same_v` | | + +**Returns** + +`void` ### preprocessing::quantize::scalar::quantizer -Defines and stores scalar for quantisation upon training +Owns global bounds or device bounds for each dimension. -The quantization is performed by a linear mapping of an interval in the float data type to the full range of the quantized int type. +Move-only; copy view() to borrow the parameters. ```cpp template struct quantizer { - T min_; - T max_; + std::variant value; + scale_type max_error; }; ``` @@ -44,8 +152,106 @@ struct quantizer { | Name | Type | Description | | --- | --- | --- | -| `min_` | `T` | Minimum value of the quantization range. | -| `max_` | `T` | Maximum value of the quantization range. | +| `value` | [`std::variant`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer-view-global-type) | Global bounds or owning per-dimension bounds. | +| `max_error` | `scale_type` | Upper bound on the L2 error between a training row and its int8 encode/decode result. train measures every input row, including clipped outliers, and rounds the maximum up with a relative 1e-6 margin. This bound applies only to the training data, not later queries. It is infinity when constructed from parameters alone or when the error is not finite. | + + +### preprocessing::quantize::scalar::quantizer::per_dim_type + +Owning per-dimension bounds. + +```cpp +struct per_dim_type { + raft::device_vector mins; + raft::device_vector maxs; +}; +``` + +**Fields** + +| Name | Type | Description | +| --- | --- | --- | +| `mins` | `raft::device_vector` | | +| `maxs` | `raft::device_vector` | | + + +### preprocessing::quantize::scalar::quantizer::quantizer + +Construct default global bounds. + +```cpp +quantizer(); +``` + +**Returns** + +`void` + +**Additional overload:** `preprocessing::quantize::scalar::quantizer::quantizer` + +Construct from finite global bounds with min <= max. + +```cpp +quantizer(T min, T max); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `min` | | `T` | | +| `max` | | `T` | | + +**Returns** + +`void` + +**Additional overload:** `preprocessing::quantize::scalar::quantizer::quantizer` + +Copy nonempty device bounds of equal length, with finite min <= max. + +```cpp +quantizer(raft::resources const& res, +raft::device_vector_view mins, +raft::device_vector_view maxs); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | | `raft::resources const&` | | +| `mins` | | `raft::device_vector_view` | | +| `maxs` | | `raft::device_vector_view` | | + +**Returns** + +`void` + + +### preprocessing::quantize::scalar::quantizer::view + +Borrow mutable per-dimension bounds; the owner must outlive stream work using them. + +```cpp +[[nodiscard]] auto view() noexcept -> view_type; +``` + +**Returns** + +`view_type` + +**Additional overload:** `preprocessing::quantize::scalar::quantizer::view` + +Borrow read-only bounds; the owner must outlive stream work using per-dimension bounds. + +```cpp +[[nodiscard]] auto view() const noexcept -> const_view_type; +``` + +**Returns** + +`const_view_type` ### preprocessing::quantize::scalar::train @@ -125,7 +331,7 @@ Usage example: **Additional overload:** `preprocessing::quantize::scalar::transform` -Applies quantization transform to given dataset +Applies global quantization to a host dataset ```cpp void transform(raft::resources const& res, @@ -134,6 +340,8 @@ raft::host_matrix_view dataset, raft::host_matrix_view out); ``` +Throws raft::logic_error for per-dimension parameters. Use device transforms or host build for per-dimension quantization. + Usage example: **Parameters** @@ -141,7 +349,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resource | -| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer | +| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer with global parameters | | `dataset` | in | `raft::host_matrix_view` | a row-major matrix view on host | | `out` | out | `raft::host_matrix_view` | a row-major matrix view on host | @@ -180,7 +388,7 @@ Usage example: **Additional overload:** `preprocessing::quantize::scalar::inverse_transform` -Perform inverse quantization step on previously quantized dataset +Applies inverse global quantization to a host dataset ```cpp void inverse_transform(raft::resources const& res, @@ -189,6 +397,8 @@ raft::host_matrix_view dataset, raft::host_matrix_view out); ``` +Throws raft::logic_error for per-dimension parameters. Use device inverse_transform for per-dimension quantization. + Note that depending on the chosen data types train dataset the conversion is not lossless. Usage example: @@ -198,7 +408,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resource | -| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer | +| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer with global parameters | | `dataset` | in | `raft::host_matrix_view` | a row-major matrix view on host | | `out` | out | `raft::host_matrix_view` | a row-major matrix view on host | @@ -282,7 +492,7 @@ Usage example: **Additional overload:** `preprocessing::quantize::scalar::transform` -Applies quantization transform to given dataset +Applies global quantization to a host dataset ```cpp void transform(raft::resources const& res, @@ -291,6 +501,8 @@ raft::host_matrix_view dataset, raft::host_matrix_view out); ``` +Throws raft::logic_error for per-dimension parameters. Use device transforms or host build for per-dimension quantization. + Usage example: **Parameters** @@ -298,7 +510,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resource | -| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer | +| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer with global parameters | | `dataset` | in | `raft::host_matrix_view` | a row-major matrix view on host | | `out` | out | `raft::host_matrix_view` | a row-major matrix view on host | @@ -336,7 +548,7 @@ Usage example: **Additional overload:** `preprocessing::quantize::scalar::inverse_transform` -Perform inverse quantization step on previously quantized dataset +Applies inverse global quantization to a host dataset ```cpp void inverse_transform(raft::resources const& res, @@ -345,6 +557,8 @@ raft::host_matrix_view dataset, raft::host_matrix_view out); ``` +Throws raft::logic_error for per-dimension parameters. Use device inverse_transform for per-dimension quantization. + Note that depending on the chosen data types train dataset the conversion is not lossless. Usage example: @@ -354,7 +568,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resource | -| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer | +| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer with global parameters | | `dataset` | in | `raft::host_matrix_view` | a row-major matrix view on host | | `out` | out | `raft::host_matrix_view` | a row-major matrix view on host | @@ -438,7 +652,7 @@ Usage example: **Additional overload:** `preprocessing::quantize::scalar::transform` -Applies quantization transform to given dataset +Applies global quantization to a host dataset ```cpp void transform(raft::resources const& res, @@ -447,6 +661,8 @@ raft::host_matrix_view dataset, raft::host_matrix_view out); ``` +Throws raft::logic_error for per-dimension parameters. Use device transforms or host build for per-dimension quantization. + Usage example: **Parameters** @@ -454,7 +670,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resource | -| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer | +| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer with global parameters | | `dataset` | in | `raft::host_matrix_view` | a row-major matrix view on host | | `out` | out | `raft::host_matrix_view` | a row-major matrix view on host | @@ -492,7 +708,7 @@ Usage example: **Additional overload:** `preprocessing::quantize::scalar::inverse_transform` -Perform inverse quantization step on previously quantized dataset +Applies inverse global quantization to a host dataset ```cpp void inverse_transform(raft::resources const& res, @@ -501,6 +717,8 @@ raft::host_matrix_view dataset, raft::host_matrix_view out); ``` +Throws raft::logic_error for per-dimension parameters. Use device inverse_transform for per-dimension quantization. + Note that depending on the chosen data types train dataset the conversion is not lossless. Usage example: @@ -510,7 +728,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resource | -| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer | +| `quantizer` | in | [`const quantizer&`](/api-reference/cpp-api-preprocessing-quantize-scalar#preprocessing-quantize-scalar-quantizer) | a scalar quantizer with global parameters | | `dataset` | in | `raft::host_matrix_view` | a row-major matrix view on host | | `out` | out | `raft::host_matrix_view` | a row-major matrix view on host |