diff --git a/cpp/test/cluster/cluster_solvers.cu b/cpp/test/cluster/cluster_solvers.cu deleted file mode 100644 index cc0a381bbf..0000000000 --- a/cpp/test/cluster/cluster_solvers.cu +++ /dev/null @@ -1,104 +0,0 @@ -/* - * Copyright (c) 2020-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include -#include -#include -#include -#include - -#include - -#include -#include - -namespace raft { -namespace spectral { - -TEST(Raft, ClusterSolvers) -{ - using namespace matrix; - using index_type = int; - using value_type = double; - - raft::resources h; - - index_type maxiter{100}; - value_type tol{1.0e-10}; - unsigned long long seed{100110021003}; - - auto stream = resource::get_cuda_stream(h); - - index_type n{100}; - index_type d{10}; - index_type k{5}; - - // nullptr expected to trigger exceptions: - // - value_type* eigvecs{nullptr}; - index_type* codes{nullptr}; - - cluster_solver_config_t cfg{k, maxiter, tol, seed}; - - kmeans_solver_t cluster_solver{cfg}; - - EXPECT_ANY_THROW(cluster_solver.solve(h, n, d, eigvecs, codes)); -} - -TEST(Raft, ModularitySolvers) -{ - using namespace matrix; - using index_type = int; - using value_type = double; - - raft::resources h; - ASSERT_EQ(0, resource::get_device_id(h)); - - index_type neigvs{10}; - index_type maxiter{100}; - index_type restart_iter{10}; - value_type tol{1.0e-10}; - bool reorthog{true}; - - // nullptr expected to trigger exceptions: - // - index_type* clusters{nullptr}; - value_type* eigvals{nullptr}; - value_type* eigvecs{nullptr}; - - unsigned long long seed{100110021003}; - - eigen_solver_config_t eig_cfg{ - neigvs, maxiter, restart_iter, tol, reorthog, seed}; - lanczos_solver_t eig_solver{eig_cfg}; - - index_type k{5}; - - cluster_solver_config_t clust_cfg{k, maxiter, tol, seed}; - kmeans_solver_t cluster_solver{clust_cfg}; - - auto stream = resource::get_cuda_stream(h); - sparse_matrix_t sm{h, nullptr, nullptr, nullptr, 0, 0}; - - EXPECT_ANY_THROW(spectral::modularity_maximization( - h, sm, eig_solver, cluster_solver, clusters, eigvals, eigvecs)); - - value_type modularity{0}; - EXPECT_ANY_THROW(spectral::analyzeModularity(h, sm, k, clusters, modularity)); -} - -} // namespace spectral -} // namespace raft diff --git a/cpp/test/cluster/cluster_solvers_deprecated.cu b/cpp/test/cluster/cluster_solvers_deprecated.cu deleted file mode 100644 index 954e3e5bb6..0000000000 --- a/cpp/test/cluster/cluster_solvers_deprecated.cu +++ /dev/null @@ -1,59 +0,0 @@ -/* - * Copyright (c) 2020-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include -#include -#include - -#include - -#include -#include - -namespace raft { -namespace spectral { - -TEST(Raft, ClusterSolvers) -{ - using namespace matrix; - using index_type = int; - using value_type = double; - - raft::resources h; - - index_type maxiter{100}; - value_type tol{1.0e-10}; - unsigned long long seed{100110021003}; - - auto stream = resource::get_cuda_stream(h); - - index_type n{100}; - index_type d{10}; - index_type k{5}; - - // nullptr expected to trigger exceptions: - // - value_type* eigvecs{nullptr}; - index_type* codes{nullptr}; - - cluster_solver_config_deprecated_t cfg{k, maxiter, tol, seed}; - kmeans_solver_deprecated_t cluster_solver{cfg}; - - EXPECT_ANY_THROW(cluster_solver.solve(h, n, d, eigvecs, codes)); -} - -} // namespace spectral -} // namespace raft diff --git a/cpp/test/cluster/kmeans.cu b/cpp/test/cluster/kmeans.cu deleted file mode 100644 index 33c17c527d..0000000000 --- a/cpp/test/cluster/kmeans.cu +++ /dev/null @@ -1,363 +0,0 @@ -/* - * Copyright (c) 2022-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" - -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include - -#include - -#include -#include - -namespace raft { - -template -struct KmeansInputs { - int n_row; - int n_col; - int n_clusters; - T tol; - bool weighted; -}; - -template -void run_cluster_cost(const raft::resources& handle, - raft::device_vector_view minClusterDistance, - rmm::device_uvector& workspace, - raft::device_scalar_view clusterCost) -{ - raft::cluster::kmeans::cluster_cost( - handle, minClusterDistance, workspace, clusterCost, raft::add_op{}); -} - -template -class KmeansTest : public ::testing::TestWithParam> { - protected: - KmeansTest() - : d_labels(0, resource::get_cuda_stream(handle)), - d_labels_ref(0, resource::get_cuda_stream(handle)), - d_centroids(0, resource::get_cuda_stream(handle)), - d_sample_weight(0, resource::get_cuda_stream(handle)) - { - } - - void apiTest() - { - testparams = ::testing::TestWithParam>::GetParam(); - - auto stream = resource::get_cuda_stream(handle); - int n_samples = testparams.n_row; - int n_features = testparams.n_col; - params.n_clusters = testparams.n_clusters; - params.tol = testparams.tol; - params.n_init = 1; - params.rng_state.seed = 1; - params.oversampling_factor = 0; - - raft::random::RngState rng(params.rng_state.seed, params.rng_state.type); - - auto X = raft::make_device_matrix(handle, n_samples, n_features); - auto labels = raft::make_device_vector(handle, n_samples); - - raft::random::make_blobs(X.data_handle(), - labels.data_handle(), - n_samples, - n_features, - params.n_clusters, - stream, - true, - nullptr, - nullptr, - T(1.0), - false, - (T)-10.0f, - (T)10.0f, - (uint64_t)1234); - d_labels.resize(n_samples, stream); - d_labels_ref.resize(n_samples, stream); - d_centroids.resize(params.n_clusters * n_features, stream); - raft::copy(d_labels_ref.data(), labels.data_handle(), n_samples, stream); - rmm::device_uvector d_sample_weight(n_samples, stream); - thrust::fill( - thrust::cuda::par.on(stream), d_sample_weight.data(), d_sample_weight.data() + n_samples, 1); - auto weight_view = - raft::make_device_vector_view(d_sample_weight.data(), n_samples); - - T inertia = 0; - int n_iter = 0; - rmm::device_uvector workspace(0, stream); - rmm::device_uvector L2NormBuf_OR_DistBuf(0, stream); - rmm::device_uvector inRankCp(0, stream); - auto X_view = raft::make_const_mdspan(X.view()); - auto centroids_view = - raft::make_device_matrix_view(d_centroids.data(), params.n_clusters, n_features); - auto miniX = raft::make_device_matrix(handle, n_samples / 4, n_features); - - // Initialize kmeans on a portion of X - raft::cluster::kmeans::shuffle_and_gather( - handle, - X_view, - raft::make_device_matrix_view(miniX.data_handle(), miniX.extent(0), miniX.extent(1)), - miniX.extent(0), - params.rng_state.seed); - - raft::cluster::kmeans::init_plus_plus( - handle, params, raft::make_const_mdspan(miniX.view()), centroids_view, workspace); - - auto minClusterDistance = raft::make_device_vector(handle, n_samples); - auto minClusterAndDistance = - raft::make_device_vector, int>(handle, n_samples); - auto L2NormX = raft::make_device_vector(handle, n_samples); - auto clusterCostBefore = raft::make_device_scalar(handle, 0); - auto clusterCostAfter = raft::make_device_scalar(handle, 0); - - raft::linalg::rowNorm(L2NormX.data_handle(), - X.data_handle(), - X.extent(1), - X.extent(0), - raft::linalg::L2Norm, - true, - stream); - - raft::cluster::kmeans::min_cluster_distance(handle, - X_view, - centroids_view, - minClusterDistance.view(), - L2NormX.view(), - L2NormBuf_OR_DistBuf, - params.metric, - params.batch_samples, - params.batch_centroids, - workspace); - - run_cluster_cost(handle, minClusterDistance.view(), workspace, clusterCostBefore.view()); - - // Run a fit of kmeans - raft::cluster::kmeans::fit_main(handle, - params, - X_view, - weight_view, - centroids_view, - raft::make_host_scalar_view(&inertia), - raft::make_host_scalar_view(&n_iter), - workspace); - - // Check that the cluster cost decreased - raft::cluster::kmeans::min_cluster_distance(handle, - X_view, - centroids_view, - minClusterDistance.view(), - L2NormX.view(), - L2NormBuf_OR_DistBuf, - params.metric, - params.batch_samples, - params.batch_centroids, - workspace); - - run_cluster_cost(handle, minClusterDistance.view(), workspace, clusterCostAfter.view()); - T h_clusterCostBefore = T(0); - T h_clusterCostAfter = T(0); - raft::update_host(&h_clusterCostBefore, clusterCostBefore.data_handle(), 1, stream); - raft::update_host(&h_clusterCostAfter, clusterCostAfter.data_handle(), 1, stream); - ASSERT_TRUE(h_clusterCostAfter < h_clusterCostBefore); - - // Count samples in clusters using 2 methods and compare them - // Fill minClusterAndDistance - raft::cluster::kmeans::min_cluster_and_distance( - handle, - X_view, - raft::make_device_matrix_view( - d_centroids.data(), params.n_clusters, n_features), - minClusterAndDistance.view(), - L2NormX.view(), - L2NormBuf_OR_DistBuf, - params.metric, - params.batch_samples, - params.batch_centroids, - workspace); - raft::cluster::kmeans::KeyValueIndexOp conversion_op; - cub::TransformInputIterator, - raft::KeyValuePair*> - itr(minClusterAndDistance.data_handle(), conversion_op); - - auto sampleCountInCluster = raft::make_device_vector(handle, params.n_clusters); - auto weigthInCluster = raft::make_device_vector(handle, params.n_clusters); - auto newCentroids = raft::make_device_matrix(handle, params.n_clusters, n_features); - raft::cluster::kmeans::update_centroids(handle, - X_view, - weight_view, - raft::make_device_matrix_view( - d_centroids.data(), params.n_clusters, n_features), - itr, - weigthInCluster.view(), - newCentroids.view()); - raft::cluster::kmeans::count_samples_in_cluster(handle, - params, - X_view, - L2NormX.view(), - newCentroids.view(), - workspace, - sampleCountInCluster.view()); - - ASSERT_TRUE(devArrMatch(sampleCountInCluster.data_handle(), - weigthInCluster.data_handle(), - params.n_clusters, - CompareApprox(params.tol))); - } - - void basicTest() - { - testparams = ::testing::TestWithParam>::GetParam(); - - int n_samples = testparams.n_row; - int n_features = testparams.n_col; - params.n_clusters = testparams.n_clusters; - params.tol = testparams.tol; - params.n_init = 5; - params.rng_state.seed = 1; - params.oversampling_factor = 0; - - auto X = raft::make_device_matrix(handle, n_samples, n_features); - auto labels = raft::make_device_vector(handle, n_samples); - auto stream = resource::get_cuda_stream(handle); - - raft::random::make_blobs(X.data_handle(), - labels.data_handle(), - n_samples, - n_features, - params.n_clusters, - stream, - true, - nullptr, - nullptr, - T(1.0), - false, - (T)-10.0f, - (T)10.0f, - (uint64_t)1234); - - d_labels.resize(n_samples, stream); - d_labels_ref.resize(n_samples, stream); - d_centroids.resize(params.n_clusters * n_features, stream); - - std::optional> d_sw = std::nullopt; - auto d_centroids_view = - raft::make_device_matrix_view(d_centroids.data(), params.n_clusters, n_features); - if (testparams.weighted) { - d_sample_weight.resize(n_samples, stream); - d_sw = std::make_optional( - raft::make_device_vector_view(d_sample_weight.data(), n_samples)); - thrust::fill(thrust::cuda::par.on(stream), - d_sample_weight.data(), - d_sample_weight.data() + n_samples, - 1); - } - - raft::copy(d_labels_ref.data(), labels.data_handle(), n_samples, stream); - - T inertia = 0; - int n_iter = 0; - auto X_view = raft::make_const_mdspan(X.view()); - - raft::cluster::kmeans_fit_predict( - handle, - params, - X_view, - d_sw, - d_centroids_view, - raft::make_device_vector_view(d_labels.data(), n_samples), - raft::make_host_scalar_view(&inertia), - raft::make_host_scalar_view(&n_iter)); - - resource::sync_stream(handle, stream); - - score = raft::stats::adjusted_rand_index( - d_labels_ref.data(), d_labels.data(), n_samples, resource::get_cuda_stream(handle)); - - if (score < 1.0) { - std::stringstream ss; - ss << "Expected: " << raft::arr2Str(d_labels_ref.data(), 25, "d_labels_ref", stream); - std::cout << (ss.str().c_str()) << '\n'; - ss.str(std::string()); - ss << "Actual: " << raft::arr2Str(d_labels.data(), 25, "d_labels", stream); - std::cout << (ss.str().c_str()) << '\n'; - std::cout << "Score = " << score << '\n'; - } - } - - void SetUp() override - { - basicTest(); - apiTest(); - } - - protected: - raft::resources handle; - KmeansInputs testparams; - rmm::device_uvector d_labels; - rmm::device_uvector d_labels_ref; - rmm::device_uvector d_centroids; - rmm::device_uvector d_sample_weight; - double score; - raft::cluster::KMeansParams params; -}; - -const std::vector> inputsf2 = {{1000, 32, 5, 0.0001f, true}, - {1000, 32, 5, 0.0001f, false}, - {1000, 100, 20, 0.0001f, true}, - {1000, 100, 20, 0.0001f, false}, - {10000, 32, 10, 0.0001f, true}, - {10000, 32, 10, 0.0001f, false}, - {10000, 100, 50, 0.0001f, true}, - {10000, 100, 50, 0.0001f, false}, - {10000, 500, 100, 0.0001f, true}, - {10000, 500, 100, 0.0001f, false}}; - -const std::vector> inputsd2 = {{1000, 32, 5, 0.0001, true}, - {1000, 32, 5, 0.0001, false}, - {1000, 100, 20, 0.0001, true}, - {1000, 100, 20, 0.0001, false}, - {10000, 32, 10, 0.0001, true}, - {10000, 32, 10, 0.0001, false}, - {10000, 100, 50, 0.0001, true}, - {10000, 100, 50, 0.0001, false}, - {10000, 500, 100, 0.0001, true}, - {10000, 500, 100, 0.0001, false}}; - -typedef KmeansTest KmeansTestF; -TEST_P(KmeansTestF, Result) { ASSERT_TRUE(score == 1.0); } - -typedef KmeansTest KmeansTestD; -TEST_P(KmeansTestD, Result) { ASSERT_TRUE(score == 1.0); } - -INSTANTIATE_TEST_CASE_P(KmeansTests, KmeansTestF, ::testing::ValuesIn(inputsf2)); - -INSTANTIATE_TEST_CASE_P(KmeansTests, KmeansTestD, ::testing::ValuesIn(inputsd2)); - -} // namespace raft diff --git a/cpp/test/cluster/kmeans_balanced.cu b/cpp/test/cluster/kmeans_balanced.cu deleted file mode 100644 index 5009eaf122..0000000000 --- a/cpp/test/cluster/kmeans_balanced.cu +++ /dev/null @@ -1,240 +0,0 @@ -/* - * Copyright (c) 2022-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.h" - -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include - -#include - -#include -#include - -/* This test takes advantage of the fact that make_blobs generates balanced clusters. - * It doesn't currently test whether the algorithm can make balanced clusters with an imbalanced - * dataset. - */ - -namespace raft { - -template -struct KmeansBalancedInputs { - IdxT n_rows; - IdxT n_cols; - IdxT n_clusters; - raft::cluster::kmeans_balanced_params kb_params; - MathT tol; -}; - -template -::std::ostream& operator<<(::std::ostream& os, const KmeansBalancedInputs& p) -{ - os << "{ " << p.n_rows << ", " << p.n_cols << ", " << p.n_clusters << ", " << p.kb_params.n_iters - << static_cast(p.kb_params.metric) << '}' << std::endl; - return os; -} - -template -class KmeansBalancedTest : public ::testing::TestWithParam> { - protected: - KmeansBalancedTest() - : stream(resource::get_cuda_stream(handle)), - d_labels(0, stream), - d_labels_ref(0, stream), - d_centroids(0, stream) - { - } - - void basicTest() - { - MappingOpT op{}; - - auto p = ::testing::TestWithParam>::GetParam(); - - auto X = raft::make_device_matrix(handle, p.n_rows, p.n_cols); - auto blob_labels = raft::make_device_vector(handle, p.n_rows); - - MathT* blobs_ptr; - rmm::device_uvector blobs(0, stream); - if constexpr (!std::is_same_v) { - blobs.resize(p.n_rows * p.n_cols, stream); - blobs_ptr = blobs.data(); - } else { - blobs_ptr = X.data_handle(); - } - - raft::random::make_blobs(blobs_ptr, - blob_labels.data_handle(), - p.n_rows, - p.n_cols, - p.n_clusters, - stream, - true, - nullptr, - nullptr, - MathT{0.1}, - true, - MathT{-1}, - MathT{1}, - (uint64_t)1234); - - // Convert blobs dataset to DataT if necessary - if constexpr (!std::is_same_v) { - raft::linalg::unaryOp( - X.data_handle(), blobs.data(), p.n_rows * p.n_cols, op.reverse_op, stream); - } - - d_labels.resize(p.n_rows, stream); - d_labels_ref.resize(p.n_rows, stream); - d_centroids.resize(p.n_clusters * p.n_cols, stream); - - raft::linalg::unaryOp( - d_labels_ref.data(), blob_labels.data_handle(), p.n_rows, raft::cast_op(), stream); - - auto X_view = - raft::make_device_matrix_view(X.data_handle(), X.extent(0), X.extent(1)); - auto d_centroids_view = - raft::make_device_matrix_view(d_centroids.data(), p.n_clusters, p.n_cols); - auto d_labels_view = raft::make_device_vector_view(d_labels.data(), p.n_rows); - - raft::cluster::kmeans_balanced::fit_predict( - handle, p.kb_params, X_view, d_centroids_view, d_labels_view, op); - - resource::sync_stream(handle, stream); - - score = raft::stats::adjusted_rand_index( - d_labels_ref.data(), d_labels.data(), p.n_rows, resource::get_cuda_stream(handle)); - - if (score < 1.0) { - std::stringstream ss; - ss << "Expected: " << raft::arr2Str(d_labels_ref.data(), 25, "d_labels_ref", stream); - std::cout << (ss.str().c_str()) << '\n'; - ss.str(std::string()); - ss << "Actual: " << raft::arr2Str(d_labels.data(), 25, "d_labels", stream); - std::cout << (ss.str().c_str()) << '\n'; - std::cout << "Score = " << score << '\n'; - } - } - - void SetUp() override { basicTest(); } - - protected: - raft::handle_t handle; - cudaStream_t stream; - rmm::device_uvector d_labels; - rmm::device_uvector d_labels_ref; - rmm::device_uvector d_centroids; - double score; -}; - -template -std::vector> get_kmeans_balanced_inputs() -{ - std::vector> out; - KmeansBalancedInputs p; - p.kb_params.n_iters = 20; - p.kb_params.metric = raft::distance::DistanceType::L2Expanded; - p.tol = MathT{0.0001}; - std::vector> row_cols_k = {{1000, 32, 5}, - {1000, 100, 20}, - {10000, 32, 10}, - {10000, 100, 50}, - {10000, 500, 100}, - {1000000, 128, 10}}; - for (auto& rck : row_cols_k) { - p.n_rows = static_cast(std::get<0>(rck)); - p.n_cols = static_cast(std::get<1>(rck)); - p.n_clusters = static_cast(std::get<2>(rck)); - out.push_back(p); - } - return out; -} - -const auto inputsf_i32 = get_kmeans_balanced_inputs(); -const auto inputsd_i32 = get_kmeans_balanced_inputs(); -const auto inputsf_i64 = get_kmeans_balanced_inputs(); -const auto inputsd_i64 = get_kmeans_balanced_inputs(); - -#define KB_TEST(test_type, test_name, test_inputs) \ - typedef RAFT_DEPAREN(test_type) test_name; \ - TEST_P(test_name, Result) { ASSERT_TRUE(score == 1.0); } \ - INSTANTIATE_TEST_CASE_P(KmeansBalancedTests, test_name, ::testing::ValuesIn(test_inputs)) - -/* - * First set of tests: no conversion - */ - -KB_TEST((KmeansBalancedTest), - KmeansBalancedTestFFU32I32, - inputsf_i32); -KB_TEST((KmeansBalancedTest), - KmeansBalancedTestDDU32I32, - inputsd_i32); -KB_TEST((KmeansBalancedTest), - KmeansBalancedTestFFU32I64, - inputsf_i64); -KB_TEST((KmeansBalancedTest), - KmeansBalancedTestDDU32I64, - inputsd_i64); -KB_TEST((KmeansBalancedTest), - KmeansBalancedTestFFI32I32, - inputsf_i32); -KB_TEST((KmeansBalancedTest), - KmeansBalancedTestFFI32I64, - inputsf_i64); -KB_TEST((KmeansBalancedTest), - KmeansBalancedTestFFI64I32, - inputsf_i32); -KB_TEST((KmeansBalancedTest), - KmeansBalancedTestFFI64I64, - inputsf_i64); - -/* - * Second set of tests: integer dataset with conversion - */ - -template -struct i2f_scaler { - // Note: with a scaling factor of 42, and generating blobs with centers between -1 and 1 with a - // standard deviation of 0.1, it's statistically very unlikely that we'd overflow - const raft::compose_op, raft::cast_op> op{ - raft::div_const_op{42}, raft::cast_op{}}; - const raft::compose_op, raft::mul_const_op> reverse_op{ - raft::cast_op{}, raft::mul_const_op{42}}; - - RAFT_INLINE_FUNCTION auto operator()(const DataT& x) const { return op(x); }; -}; - -KB_TEST((KmeansBalancedTest>), - KmeansBalancedTestFI8U32I32, - inputsf_i32); -KB_TEST((KmeansBalancedTest>), - KmeansBalancedTestDI8U32I32, - inputsd_i32); - -} // namespace raft diff --git a/cpp/test/cluster/kmeans_find_k.cu b/cpp/test/cluster/kmeans_find_k.cu deleted file mode 100644 index 8e05ad3695..0000000000 --- a/cpp/test/cluster/kmeans_find_k.cu +++ /dev/null @@ -1,142 +0,0 @@ -/* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.h" - -#include -#include -#include -#include -#include -#include - -#include - -#include -#include - -namespace raft { - -template -struct KmeansFindKInputs { - int n_row; - int n_col; - int n_clusters; - T tol; - bool weighted; -}; - -template -class KmeansFindKTest : public ::testing::TestWithParam> { - protected: - KmeansFindKTest() - : stream(resource::get_cuda_stream(handle)), best_k(raft::make_host_scalar(0)) - { - } - - void basicTest() - { - testparams = ::testing::TestWithParam>::GetParam(); - - int n_samples = testparams.n_row; - int n_features = testparams.n_col; - int n_clusters = testparams.n_clusters; - - auto X = raft::make_device_matrix(handle, n_samples, n_features); - auto labels = raft::make_device_vector(handle, n_samples); - - raft::random::make_blobs(X.data_handle(), - labels.data_handle(), - n_samples, - n_features, - n_clusters, - stream, - true, - nullptr, - nullptr, - T(.001), - false, - (T)-10.0f, - (T)10.0f, - (uint64_t)1234); - - auto inertia = raft::make_host_scalar(0); - auto n_iter = raft::make_host_scalar(0); - - auto X_view = - raft::make_device_matrix_view(X.data_handle(), X.extent(0), X.extent(1)); - - raft::cluster::kmeans::find_k( - handle, X_view, best_k.view(), inertia.view(), n_iter.view(), n_clusters); - - resource::sync_stream(handle, stream); - } - - void SetUp() override { basicTest(); } - - protected: - raft::resources handle; - cudaStream_t stream; - KmeansFindKInputs testparams; - raft::host_scalar best_k; -}; - -const std::vector> inputsf2 = {{1000, 32, 8, 0.001f, true}, - {1000, 32, 8, 0.001f, false}, - {1000, 100, 20, 0.001f, true}, - {1000, 100, 20, 0.001f, false}, - {10000, 32, 10, 0.001f, true}, - {10000, 32, 10, 0.001f, false}, - {10000, 100, 50, 0.001f, true}, - {10000, 100, 50, 0.001f, false}, - {10000, 500, 100, 0.001f, true}, - {10000, 500, 100, 0.001f, false}}; - -const std::vector> inputsd2 = {{1000, 32, 5, 0.0001, true}, - {1000, 32, 5, 0.0001, false}, - {1000, 100, 20, 0.0001, true}, - {1000, 100, 20, 0.0001, false}, - {10000, 32, 10, 0.0001, true}, - {10000, 32, 10, 0.0001, false}, - {10000, 100, 50, 0.0001, true}, - {10000, 100, 50, 0.0001, false}, - {10000, 500, 100, 0.0001, true}, - {10000, 500, 100, 0.0001, false}}; - -typedef KmeansFindKTest KmeansFindKTestF; -TEST_P(KmeansFindKTestF, Result) -{ - if (best_k.view()[0] != testparams.n_clusters) { - std::cout << best_k.view()[0] << " " << testparams.n_clusters << std::endl; - } - ASSERT_TRUE(best_k.view()[0] == testparams.n_clusters); -} - -typedef KmeansFindKTest KmeansFindKTestD; -TEST_P(KmeansFindKTestD, Result) -{ - if (best_k.view()[0] != testparams.n_clusters) { - std::cout << best_k.view()[0] << " " << testparams.n_clusters << std::endl; - } - - ASSERT_TRUE(best_k.view()[0] == testparams.n_clusters); -} - -INSTANTIATE_TEST_CASE_P(KmeansFindKTests, KmeansFindKTestF, ::testing::ValuesIn(inputsf2)); - -INSTANTIATE_TEST_CASE_P(KmeansFindKTests, KmeansFindKTestD, ::testing::ValuesIn(inputsd2)); - -} // namespace raft diff --git a/cpp/test/cluster/linkage.cu b/cpp/test/cluster/linkage.cu deleted file mode 100644 index ba7ed4254e..0000000000 --- a/cpp/test/cluster/linkage.cu +++ /dev/null @@ -1,674 +0,0 @@ -/* - * Copyright (c) 2021-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -// XXX: We allow the instantiation of masked_l2_nn here: -// raft::linkage::FixConnectivitiesRedOp red_op(params.n_row); -// raft::linkage::cross_component_nn( -// handle, out_edges, data.data(), colors.data(), params.n_row, params.n_col, red_op); -// -// TODO: consider adding this to libraft.so or creating an instance in a -// separate translation unit for this test. -#undef RAFT_EXPLICIT_INSTANTIATE_ONLY - -#include "../test_utils.cuh" - -#include -#include -#include -#include -#include -#include -#include - -#include - -#include - -#include - -namespace raft { - -using namespace std; - -template -struct LinkageInputs { - IdxT n_row; - IdxT n_col; - - std::vector data; - - std::vector expected_labels; - - int n_clusters; - - bool use_knn; - - int c; -}; - -/** - * @brief kernel to calculate the values of a and b - * @param firstClusterArray: the array of classes of type T - * @param secondClusterArray: the array of classes of type T - * @param size: the size of the data points - * @param a: number of pairs of points that both the clusters have classified the same - * @param b: number of pairs of points that both the clusters have classified differently - */ -template -RAFT_KERNEL computeTheNumerator( - const T* firstClusterArray, const T* secondClusterArray, uint64_t size, uint64_t* a, uint64_t* b) -{ - // calculating the indices of pairs of datapoints compared by the current thread - uint64_t j = threadIdx.x + blockIdx.x * blockDim.x; - uint64_t i = threadIdx.y + blockIdx.y * blockDim.y; - - // thread-local variables to count a and b - uint64_t myA = 0, myB = 0; - - if (i < size && j < size && j < i) { - // checking if the pair have been classified the same by both the clusters - if (firstClusterArray[i] == firstClusterArray[j] && - secondClusterArray[i] == secondClusterArray[j]) { - ++myA; - } - - // checking if the pair have been classified differently by both the clusters - else if (firstClusterArray[i] != firstClusterArray[j] && - secondClusterArray[i] != secondClusterArray[j]) { - ++myB; - } - } - - // specialize blockReduce for a 2D block of 1024 threads of type uint64_t - typedef cub::BlockReduce - BlockReduce; - - // Allocate shared memory for blockReduce - __shared__ typename BlockReduce::TempStorage temp_storage; - - // summing up thread-local counts specific to a block - myA = BlockReduce(temp_storage).Sum(myA); - __syncthreads(); - myB = BlockReduce(temp_storage).Sum(myB); - __syncthreads(); - - // executed once per block - if (threadIdx.x == 0 && threadIdx.y == 0) { - raft::myAtomicAdd((unsigned long long int*)a, myA); - raft::myAtomicAdd((unsigned long long int*)b, myB); - } -} - -/** - * @brief Function to calculate RandIndex - * more info on rand index - * @param firstClusterArray: the array of classes of type T - * @param secondClusterArray: the array of classes of type T - * @param size: the size of the data points of type uint64_t - * @param stream: the cudaStream object - */ -template -double compute_rand_index(T* firstClusterArray, - T* secondClusterArray, - uint64_t size, - cudaStream_t stream) -{ - // rand index for size less than 2 is not defined - ASSERT(size >= 2, "Rand Index for size less than 2 not defined!"); - - // allocating and initializing memory for a and b in the GPU - rmm::device_uvector arr_buf(2, stream); - RAFT_CUDA_TRY(cudaMemsetAsync(arr_buf.data(), 0, 2 * sizeof(uint64_t), stream)); - - // kernel configuration - static const int BLOCK_DIM_Y = 16, BLOCK_DIM_X = 16; - dim3 numThreadsPerBlock(BLOCK_DIM_X, BLOCK_DIM_Y); - dim3 numBlocks(raft::ceildiv(size, numThreadsPerBlock.x), - raft::ceildiv(size, numThreadsPerBlock.y)); - - // calling the kernel - computeTheNumerator<<>>( - firstClusterArray, secondClusterArray, size, arr_buf.data(), arr_buf.data() + 1); - - // synchronizing and updating the calculated values of a and b from device to host - uint64_t ab_host[2] = {0}; - raft::update_host(ab_host, arr_buf.data(), 2, stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - - // error handling - RAFT_CUDA_TRY(cudaGetLastError()); - - // denominator - uint64_t nChooseTwo = size * (size - 1) / 2; - - // calculating the rand_index - return (double)(((double)(ab_host[0] + ab_host[1])) / (double)nChooseTwo); -} - -template -::std::ostream& operator<<(::std::ostream& os, const LinkageInputs& dims) -{ - return os; -} - -template -class LinkageTest : public ::testing::TestWithParam> { - public: - LinkageTest() - : params(::testing::TestWithParam>::GetParam()), - labels(0, resource::get_cuda_stream(handle)), - labels_ref(0, resource::get_cuda_stream(handle)) - { - } - - protected: - void basicTest() - { - auto stream = resource::get_cuda_stream(handle); - - labels.resize(params.n_row, stream); - labels_ref.resize(params.n_row, stream); - rmm::device_uvector data(params.n_row * params.n_col, stream); - - raft::copy(data.data(), params.data.data(), data.size(), stream); - raft::copy(labels_ref.data(), params.expected_labels.data(), params.n_row, stream); - - rmm::device_uvector out_children(params.n_row * 2, stream); - - auto data_view = raft::make_device_matrix_view( - data.data(), params.n_row, params.n_col); - auto dendrogram_view = - raft::make_device_matrix_view(out_children.data(), params.n_row, 2); - auto labels_view = raft::make_device_vector_view(labels.data(), params.n_row); - - if (params.use_knn) { - raft::cluster::hierarchy:: - single_linkage( - handle, - data_view, - dendrogram_view, - labels_view, - raft::distance::DistanceType::L2SqrtExpanded, - params.n_clusters, - std::make_optional(params.c)); - - } else { - raft::cluster::hierarchy:: - single_linkage( - handle, - data_view, - dendrogram_view, - labels_view, - raft::distance::DistanceType::L2SqrtExpanded, - params.n_clusters, - std::make_optional(params.c)); - } - - resource::sync_stream(handle, stream); - - score = compute_rand_index(labels.data(), labels_ref.data(), params.n_row, stream); - } - - void SetUp() override { basicTest(); } - - protected: - raft::resources handle; - - LinkageInputs params; - rmm::device_uvector labels, labels_ref; - double score; -}; - -const std::vector> linkage_inputsf2 = { - // Test n_clusters == n_points - {10, - 5, - {0.21390334, 0.50261639, 0.91036676, 0.59166485, 0.71162682, 0.10248392, 0.77782677, 0.43772379, - 0.4035871, 0.3282796, 0.47544681, 0.59862974, 0.12319357, 0.06239463, 0.28200272, 0.1345717, - 0.50498218, 0.5113505, 0.16233086, 0.62165332, 0.42281548, 0.933117, 0.41386077, 0.23264562, - 0.73325968, 0.37537541, 0.70719873, 0.14522645, 0.73279625, 0.9126674, 0.84854131, 0.28890216, - 0.85267903, 0.74703138, 0.83842071, 0.34942792, 0.27864171, 0.70911132, 0.21338564, 0.32035554, - 0.73788331, 0.46926692, 0.57570162, 0.42559178, 0.87120209, 0.22734951, 0.01847905, 0.75549396, - 0.76166195, 0.66613745}, - {9, 8, 7, 6, 5, 4, 3, 2, 1, 0}, - 10, - true, - -1}, - // // Test outlier points - {9, - 2, - {-1, -50, 3, 4, 5000, 10000, 1, 3, 4, 5, 0.000005, 0.00002, 2000000, 500000, 10, 50, 30, 5}, - {6, 0, 5, 0, 0, 4, 3, 2, 1}, - 7, - true, - -1}, - - // Test n_clusters == (n_points / 2) - {10, - 5, - {0.21390334, 0.50261639, 0.91036676, 0.59166485, 0.71162682, 0.10248392, 0.77782677, 0.43772379, - 0.4035871, 0.3282796, 0.47544681, 0.59862974, 0.12319357, 0.06239463, 0.28200272, 0.1345717, - 0.50498218, 0.5113505, 0.16233086, 0.62165332, 0.42281548, 0.933117, 0.41386077, 0.23264562, - 0.73325968, 0.37537541, 0.70719873, 0.14522645, 0.73279625, 0.9126674, 0.84854131, 0.28890216, - 0.85267903, 0.74703138, 0.83842071, 0.34942792, 0.27864171, 0.70911132, 0.21338564, 0.32035554, - 0.73788331, 0.46926692, 0.57570162, 0.42559178, 0.87120209, 0.22734951, 0.01847905, 0.75549396, - 0.76166195, 0.66613745}, - {1, 0, 4, 0, 0, 3, 2, 0, 2, 1}, - 5, - true, - -1}, - - // Test n_points == 100 - {100, - 10, - {6.26168372e-01, 9.30437651e-01, 6.02450208e-01, 2.73025296e-01, 9.53050619e-01, 3.32164396e-01, - 6.88942598e-01, 5.79163537e-01, 6.70341547e-01, 2.70140602e-02, 9.30429671e-01, 7.17721157e-01, - 9.89948537e-01, 7.75253347e-01, 1.34491522e-02, 2.48522428e-02, 3.51413378e-01, 7.64405834e-01, - 7.86373507e-01, 7.18748577e-01, 8.66998621e-01, 6.80316582e-01, 2.51288712e-01, 4.91078420e-01, - 3.76246281e-01, 4.86828710e-01, 5.67464772e-01, 5.30734742e-01, 8.99478296e-01, 7.66699088e-01, - 9.49339111e-01, 3.55248484e-01, 9.06046929e-01, 4.48407772e-01, 6.96395305e-01, 2.44277335e-01, - 7.74840000e-01, 5.21046603e-01, 4.66423971e-02, 5.12019638e-02, 8.95019614e-01, 5.28956953e-01, - 4.31536306e-01, 5.83857744e-01, 4.41787364e-01, 4.68656523e-01, 5.73971433e-01, 6.79989654e-01, - 3.19650588e-01, 6.12579596e-01, 6.49126442e-02, 8.39131142e-01, 2.85252117e-01, 5.84848929e-01, - 9.46507115e-01, 8.58440748e-01, 3.61528940e-01, 2.44215959e-01, 3.80101125e-01, 4.57128957e-02, - 8.82216988e-01, 8.31498633e-01, 7.23474381e-01, 7.75788607e-01, 1.40864146e-01, 6.62092382e-01, - 5.13985168e-01, 3.00686418e-01, 8.70109949e-01, 2.43187753e-01, 2.89391938e-01, 2.84214238e-01, - 8.70985521e-01, 8.77491176e-01, 6.72537226e-01, 3.30929686e-01, 1.85934324e-01, 9.16222614e-01, - 6.18239142e-01, 2.64768597e-01, 5.76145451e-01, 8.62961369e-01, 6.84757925e-01, 7.60549082e-01, - 1.27645356e-01, 4.51004673e-01, 3.92292980e-01, 4.63170803e-01, 4.35449330e-02, 2.17583404e-01, - 5.71832605e-02, 2.06763039e-01, 3.70116249e-01, 2.09750028e-01, 6.17283019e-01, 8.62549231e-01, - 9.84156240e-02, 2.66249156e-01, 3.87635103e-01, 2.85591012e-02, 4.24826068e-01, 4.45795088e-01, - 6.86227676e-01, 1.08848960e-01, 5.96731841e-02, 3.71770228e-01, 1.91548833e-01, 6.95136078e-01, - 9.00700636e-01, 8.76363105e-01, 2.67334632e-01, 1.80619709e-01, 7.94060419e-01, 1.42854171e-02, - 1.09372387e-01, 8.74028108e-01, 6.46403232e-01, 4.86588834e-01, 5.93446175e-02, 6.11886291e-01, - 8.83865057e-01, 3.15879821e-01, 2.27043992e-01, 9.76764951e-01, 6.15620336e-01, 9.76199360e-01, - 2.40548962e-01, 3.21795663e-01, 8.75087904e-02, 8.11234663e-01, 6.96070480e-01, 8.12062321e-01, - 1.21958818e-01, 3.44348628e-02, 8.72630414e-01, 3.06162776e-01, 1.76043529e-02, 9.45894971e-01, - 5.33896401e-01, 6.21642973e-01, 4.93062535e-01, 4.48984262e-01, 2.24560379e-01, 4.24052195e-02, - 4.43447610e-01, 8.95646149e-01, 6.05220676e-01, 1.81840491e-01, 9.70831206e-01, 2.12563586e-02, - 6.92582693e-01, 7.55946922e-01, 7.95086143e-01, 6.05328941e-01, 3.99350764e-01, 4.32846636e-01, - 9.81114529e-01, 4.98266428e-01, 6.37127930e-03, 1.59085889e-01, 6.34682067e-05, 5.59429440e-01, - 7.38827633e-01, 8.93214770e-01, 2.16494306e-01, 9.35430573e-02, 4.75665868e-02, 7.80503518e-01, - 7.86240041e-01, 7.06854594e-01, 2.13725879e-02, 7.68246091e-01, 4.50234808e-01, 5.21231104e-01, - 5.01989826e-03, 4.22081572e-02, 1.65337732e-01, 8.54134740e-01, 4.99430262e-01, 8.94525601e-01, - 1.14028379e-01, 3.69739861e-01, 1.32955599e-01, 2.65563824e-01, 2.52811151e-01, 1.44792843e-01, - 6.88449594e-01, 4.44921417e-01, 8.23296587e-01, 1.93266317e-01, 1.19033309e-01, 1.36368966e-01, - 3.42600285e-01, 5.64505195e-01, 5.57594559e-01, 7.44257892e-01, 8.38231569e-02, 4.11548847e-01, - 3.21010077e-01, 8.55081359e-01, 4.30105779e-01, 1.16229135e-01, 9.87731964e-02, 3.14712335e-01, - 4.50880592e-01, 2.72289598e-01, 6.31615256e-01, 8.97432958e-01, 4.44764250e-01, 8.03776440e-01, - 2.68767748e-02, 2.43374608e-01, 4.02141103e-01, 4.98881209e-01, 5.33173003e-01, 8.82890436e-01, - 7.16149148e-01, 4.19664401e-01, 2.29335357e-01, 2.88637806e-01, 3.44696803e-01, 6.78171906e-01, - 5.69849716e-01, 5.86454477e-01, 3.54474989e-01, 9.03876540e-01, 6.45980000e-01, 6.34887593e-01, - 7.88039746e-02, 2.04814126e-01, 7.82251754e-01, 2.43147074e-01, 7.50951808e-01, 1.72799092e-02, - 2.95349590e-01, 6.57991826e-01, 8.81214312e-01, 5.73970708e-01, 2.77610881e-01, 1.82155097e-01, - 7.69797417e-02, 6.44792402e-01, 9.46950998e-01, 7.73064845e-01, 6.04733624e-01, 5.80094567e-01, - 1.67498426e-01, 2.66514296e-01, 6.50140368e-01, 1.91170299e-01, 2.08752199e-01, 3.01664091e-01, - 9.85033484e-01, 2.92909152e-01, 8.65816607e-01, 1.85222119e-01, 2.28814559e-01, 1.34286382e-02, - 2.89234322e-01, 8.18668708e-01, 4.71706924e-01, 9.23199803e-01, 2.80879188e-01, 1.47319284e-01, - 4.13915748e-01, 9.31274932e-02, 6.66322195e-01, 9.66953974e-01, 3.19405786e-01, 6.69486551e-01, - 5.03096313e-02, 6.95225201e-01, 5.78469859e-01, 6.29481655e-01, 1.39252534e-01, 1.22564968e-01, - 6.80663678e-01, 6.34607157e-01, 6.42765834e-01, 1.57127410e-02, 2.92132086e-01, 5.24423878e-01, - 4.68676824e-01, 2.86003928e-01, 7.18608322e-01, 8.95617933e-01, 5.48844309e-01, 1.74517278e-01, - 5.24379196e-01, 2.13526524e-01, 5.88375435e-01, 9.88560185e-01, 4.17435771e-01, 6.14438688e-01, - 9.53760881e-01, 5.27151288e-01, 7.03017278e-01, 3.44448559e-01, 4.47059676e-01, 2.83414901e-01, - 1.98979011e-01, 4.24917361e-01, 5.73172761e-01, 2.32398853e-02, 1.65887230e-01, 4.05552785e-01, - 9.29665524e-01, 2.26135696e-01, 9.20563384e-01, 7.65259963e-01, 4.54820075e-01, 8.97710267e-01, - 3.78559302e-03, 9.15219382e-01, 3.55705698e-01, 6.94905124e-01, 8.58540202e-01, 3.89790666e-01, - 2.49478206e-01, 7.93679304e-01, 4.75830027e-01, 4.40425353e-01, 3.70579459e-01, 1.40578049e-01, - 1.70386675e-01, 7.04056121e-01, 4.85963102e-01, 9.68450060e-01, 6.77178001e-01, 2.65934654e-01, - 2.58915007e-01, 6.70052890e-01, 2.61945109e-01, 8.46207759e-01, 1.01928951e-01, 2.85611334e-01, - 2.45776933e-01, 2.66658783e-01, 3.71724077e-01, 4.34319025e-01, 4.24407347e-01, 7.15417683e-01, - 8.07997684e-01, 1.64296275e-01, 6.01638065e-01, 8.60606804e-02, 2.68719187e-01, 5.11764101e-01, - 9.75844338e-01, 7.81226782e-01, 2.20925515e-01, 7.18135040e-01, 9.82395577e-01, 8.39160243e-01, - 9.08058083e-01, 6.88010677e-01, 8.14271847e-01, 5.12460821e-01, 1.17311345e-01, 5.96075228e-01, - 9.17455497e-01, 2.12052706e-01, 7.04074603e-01, 8.72872565e-02, 8.76047818e-01, 6.96235046e-01, - 8.54801557e-01, 2.49729159e-01, 9.76594604e-01, 2.87386363e-01, 2.36461559e-02, 9.94075254e-01, - 4.25193986e-01, 7.61869994e-01, 5.13334255e-01, 6.44711165e-02, 8.92156689e-01, 3.55235167e-01, - 1.08154647e-01, 8.78446825e-01, 2.43833016e-01, 9.23071293e-01, 2.72724115e-01, 9.46631338e-01, - 3.74510294e-01, 4.08451278e-02, 9.78392777e-01, 3.65079221e-01, 6.37199516e-01, 5.51144906e-01, - 5.25978080e-01, 1.42803678e-01, 4.05451674e-01, 7.79788219e-01, 6.26009784e-01, 3.35249497e-01, - 1.43159543e-02, 1.80363779e-01, 5.05096904e-01, 2.82619947e-01, 5.83561392e-01, 3.10951324e-01, - 8.73223968e-01, 4.38545619e-01, 4.81348800e-01, 6.68497085e-01, 3.79345401e-01, 9.58832501e-01, - 1.89869550e-01, 2.34083070e-01, 2.94066207e-01, 5.74892667e-02, 6.92106828e-02, 9.61127686e-02, - 6.72650672e-02, 8.47345378e-01, 2.80916761e-01, 7.32177357e-03, 9.80785961e-01, 5.73192225e-02, - 8.48781331e-01, 8.83225408e-01, 7.34398275e-01, 7.70381941e-01, 6.20778343e-01, 8.96822048e-01, - 5.40732486e-01, 3.69704071e-01, 5.77305837e-01, 2.08221827e-01, 7.34275341e-01, 1.06110900e-01, - 3.49496706e-01, 8.34948910e-01, 1.56403291e-02, 6.78576376e-01, 8.96141268e-01, 5.94835119e-01, - 1.43943153e-01, 3.49618530e-01, 2.10440392e-01, 3.46585620e-01, 1.05153093e-01, 3.45446174e-01, - 2.72177079e-01, 7.07946300e-01, 4.33717726e-02, 3.31232203e-01, 3.91874320e-01, 4.76338141e-01, - 6.22777789e-01, 2.95989228e-02, 4.32855769e-01, 7.61049310e-01, 3.63279149e-01, 9.47210350e-01, - 6.43721247e-01, 6.58025802e-01, 1.05247633e-02, 5.29974442e-01, 7.30675767e-01, 4.30041079e-01, - 6.62634841e-01, 8.25936616e-01, 9.91253704e-01, 6.79399281e-01, 5.44177006e-01, 7.52876048e-01, - 3.32139049e-01, 7.98732398e-01, 7.38865223e-01, 9.16055132e-01, 6.11736493e-01, 9.63672879e-01, - 1.83778839e-01, 7.27558919e-02, 5.91602822e-01, 3.25235484e-01, 2.34741217e-01, 9.52346277e-01, - 9.18556407e-01, 9.35373324e-01, 6.89209070e-01, 2.56049054e-01, 6.17975395e-01, 7.82285691e-01, - 9.84983432e-01, 6.62322741e-01, 2.04144457e-01, 3.98446577e-01, 1.38918297e-01, 3.05919921e-01, - 3.14043787e-01, 5.91072666e-01, 7.44703771e-01, 8.92272567e-01, 9.78017873e-01, 9.01203161e-01, - 1.41526372e-01, 4.14878484e-01, 6.80683651e-01, 5.01733152e-02, 8.14635389e-01, 2.27926375e-01, - 9.03269815e-01, 8.68443745e-01, 9.86939190e-01, 7.40779486e-01, 2.61005311e-01, 3.19276232e-01, - 9.69509248e-01, 1.11908818e-01, 4.49198556e-01, 1.27056715e-01, 3.84064823e-01, 5.14591811e-01, - 2.10747488e-01, 9.53884090e-01, 8.43167950e-01, 4.51187972e-01, 3.75331782e-01, 6.23566461e-01, - 3.55290379e-01, 2.95705968e-01, 1.69622690e-01, 1.42981830e-01, 2.72180991e-01, 9.46468040e-01, - 3.70932500e-01, 9.94292830e-01, 4.62587505e-01, 7.14817405e-01, 2.45370540e-02, 3.00906377e-01, - 5.75768304e-01, 9.71448393e-01, 6.95574827e-02, 3.93693854e-01, 5.29306116e-01, 5.04694554e-01, - 6.73797120e-02, 6.76596969e-01, 5.50948898e-01, 3.24909641e-01, 7.70337719e-01, 6.51842631e-03, - 3.03264879e-01, 7.61037886e-03, 2.72289601e-01, 1.50502041e-01, 6.71103888e-02, 7.41503703e-01, - 1.92088941e-01, 2.19043977e-01, 9.09320161e-01, 2.37993569e-01, 6.18107973e-02, 8.31447852e-01, - 2.23355609e-01, 1.84789435e-01, 4.16104518e-01, 4.21573859e-01, 8.72446305e-02, 2.97294197e-01, - 4.50328256e-01, 8.72199917e-01, 2.51279916e-01, 4.86219272e-01, 7.57071329e-01, 4.85655942e-01, - 1.06187277e-01, 4.92341327e-01, 1.46017513e-01, 5.25421017e-01, 4.22637906e-01, 2.24685018e-01, - 8.72648431e-01, 5.54051490e-01, 1.80745062e-01, 2.12756336e-01, 5.20883169e-01, 7.60363654e-01, - 8.30254678e-01, 5.00003328e-01, 4.69017439e-01, 6.38105527e-01, 3.50638261e-02, 5.22217353e-02, - 9.06516882e-02, 8.52975842e-01, 1.19985883e-01, 3.74926753e-01, 6.50302066e-01, 1.98875727e-01, - 6.28362507e-02, 4.32693501e-01, 3.10500685e-01, 6.20732833e-01, 4.58503272e-01, 3.20790034e-01, - 7.91284868e-01, 7.93054570e-01, 2.93406765e-01, 8.95399023e-01, 1.06441034e-01, 7.53085241e-02, - 8.67523104e-01, 1.47963482e-01, 1.25584706e-01, 3.81545040e-02, 6.34338619e-01, 1.76368938e-02, - 5.75553531e-02, 5.31607516e-01, 2.63869588e-01, 9.41945823e-01, 9.24028838e-02, 5.21496463e-01, - 7.74866558e-01, 5.65210610e-01, 7.28015327e-02, 6.51963790e-01, 8.94727453e-01, 4.49571590e-01, - 1.29932405e-01, 8.64026259e-01, 9.92599934e-01, 7.43721560e-01, 8.87300215e-01, 1.06369925e-01, - 8.11335531e-01, 7.87734900e-01, 9.87344678e-01, 5.32502820e-01, 4.42612382e-01, 9.64041183e-01, - 1.66085871e-01, 1.12937664e-01, 5.24423470e-01, 6.54689333e-01, 4.59119726e-01, 5.22774091e-01, - 3.08722276e-02, 6.26979315e-01, 4.49754105e-01, 8.07495757e-01, 2.34199499e-01, 1.67765675e-01, - 9.22168418e-01, 3.73210378e-01, 8.04432575e-01, 5.61890354e-01, 4.47025593e-01, 6.43155678e-01, - 2.40407640e-01, 5.91631279e-01, 1.59369206e-01, 7.75799090e-01, 8.32067212e-01, 5.59791576e-02, - 6.39105224e-01, 4.85274738e-01, 2.12630838e-01, 2.81431312e-02, 7.16205363e-01, 6.83885011e-01, - 5.23869697e-01, 9.99418314e-01, 8.35331599e-01, 4.69877463e-02, 6.74712562e-01, 7.99273684e-01, - 2.77001890e-02, 5.75809742e-01, 2.78513031e-01, 8.36209905e-01, 7.25472379e-01, 4.87173943e-01, - 7.88311357e-01, 9.64676177e-01, 1.75752651e-01, 4.98112580e-01, 8.08850418e-02, 6.40981131e-01, - 4.06647450e-01, 8.46539387e-01, 2.12620694e-01, 9.11012851e-01, 8.25041445e-01, 8.90065575e-01, - 9.63626055e-01, 5.96689242e-01, 1.63372670e-01, 4.51640148e-01, 3.43026542e-01, 5.80658851e-01, - 2.82327625e-01, 4.75535418e-01, 6.27760926e-01, 8.46314115e-01, 9.61961932e-01, 3.19806094e-01, - 5.05508062e-01, 5.28102944e-01, 6.13045057e-01, 7.44714938e-01, 1.50586073e-01, 7.91878033e-01, - 4.89839179e-01, 3.10496849e-01, 8.82309038e-01, 2.86922314e-01, 4.84687559e-01, 5.20838630e-01, - 4.62955493e-01, 2.38185305e-01, 5.47259907e-02, 7.10916137e-01, 7.31887202e-01, 6.25602317e-01, - 8.77741168e-01, 4.19881322e-01, 4.81222328e-01, 1.28224501e-01, 2.46034010e-01, 3.34971854e-01, - 7.37216484e-01, 5.62134821e-02, 7.14089724e-01, 9.85549393e-01, 4.66295827e-01, 3.08722434e-03, - 4.70237690e-01, 2.66524167e-01, 7.93875484e-01, 4.54795911e-02, 8.09702944e-01, 1.47709735e-02, - 1.70082405e-01, 6.35905179e-01, 3.75379109e-01, 4.30315011e-01, 3.15788760e-01, 5.58065230e-01, - 2.24643800e-01, 2.42142981e-01, 6.57283636e-01, 3.34921891e-01, 1.26588975e-01, 7.68064155e-01, - 9.43856291e-01, 4.47518596e-01, 5.44453573e-01, 9.95764932e-01, 7.16444391e-01, 8.51019765e-01, - 1.01179183e-01, 4.45473958e-01, 4.60327322e-01, 4.96895844e-02, 4.72907738e-01, 5.58987444e-01, - 3.41027487e-01, 1.56175026e-01, 7.58283148e-01, 6.83600909e-01, 2.14623396e-01, 3.27348880e-01, - 3.92517893e-01, 6.70418431e-01, 5.16440832e-01, 8.63140348e-01, 5.73277464e-01, 3.46608058e-01, - 7.39396341e-01, 7.20852434e-01, 2.35653246e-02, 3.89935659e-01, 7.53783745e-01, 6.34563528e-01, - 8.79339335e-01, 7.41599159e-02, 5.62433904e-01, 6.15553852e-01, 4.56956324e-01, 5.20047447e-01, - 5.26845015e-02, 5.58471266e-01, 1.63632233e-01, 5.38936665e-02, 6.49593683e-01, 2.56838748e-01, - 8.99035326e-01, 7.20847756e-01, 5.68954684e-01, 7.43684755e-01, 5.70924238e-01, 3.82318724e-01, - 4.89328290e-01, 5.62208561e-01, 4.97540804e-02, 4.18011085e-01, 6.88041565e-01, 2.16234653e-01, - 7.89548214e-01, 8.46136387e-01, 8.46816189e-01, 1.73842353e-01, 6.11627842e-02, 8.44440559e-01, - 4.50646654e-01, 3.74785037e-01, 4.87196697e-01, 4.56276448e-01, 9.13284391e-01, 4.15715464e-01, - 7.13597697e-01, 1.23641270e-02, 5.10031271e-01, 4.74601930e-02, 2.55731159e-01, 3.22090006e-01, - 1.91165703e-01, 4.51170940e-01, 7.50843157e-01, 4.42420576e-01, 4.25380660e-01, 4.50667257e-01, - 6.55689206e-01, 9.68257670e-02, 1.96528793e-01, 8.97343028e-01, 4.99940904e-01, 6.65504083e-01, - 9.41828079e-01, 4.54397338e-01, 5.61893331e-01, 5.09839880e-01, 4.53117514e-01, 8.96804127e-02, - 1.74888861e-01, 6.65641378e-01, 2.81668336e-01, 1.89532742e-01, 5.61668382e-01, 8.68330157e-02, - 8.25092797e-01, 5.18106324e-01, 1.71904024e-01, 3.68385523e-01, 1.62005436e-01, 7.48507399e-01, - 9.30274827e-01, 2.38198517e-01, 9.52222901e-01, 5.23587800e-01, 6.94384557e-01, 1.09338652e-01, - 4.83356794e-01, 2.73050402e-01, 3.68027050e-01, 5.92366466e-01, 1.83192289e-01, 8.60376029e-01, - 7.13926203e-01, 8.16750052e-01, 1.57890291e-01, 6.25691951e-01, 5.24831646e-01, 1.73873797e-01, - 1.02429784e-01, 9.17488471e-01, 4.03584434e-01, 9.31170884e-01, 2.79386137e-01, 8.77745206e-01, - 2.45200576e-01, 1.28896951e-01, 3.15713052e-01, 5.27874291e-01, 2.16444335e-01, 7.03883817e-01, - 7.74738919e-02, 8.42422142e-01, 3.75598924e-01, 3.51002411e-01, 6.22752776e-01, 4.82407943e-01, - 7.43107867e-01, 9.46182666e-01, 9.44344819e-01, 3.28124763e-01, 1.06147431e-01, 1.65102684e-01, - 3.84060507e-01, 2.91057722e-01, 7.68173662e-02, 1.03543651e-01, 6.76698940e-01, 1.43141994e-01, - 7.21342202e-01, 6.69471294e-03, 9.07298311e-01, 5.57080171e-01, 8.10954489e-01, 4.11120526e-01, - 2.06407453e-01, 2.59590556e-01, 7.58512718e-01, 5.79873897e-01, 2.92875650e-01, 2.83686529e-01, - 2.42829343e-01, 9.19323719e-01, 3.46832864e-01, 3.58238858e-01, 7.42827585e-01, 2.05760059e-01, - 9.58438860e-01, 5.66326411e-01, 6.60292846e-01, 5.61095078e-02, 6.79465531e-01, 7.05118513e-01, - 4.44713264e-01, 2.09732933e-01, 5.22732436e-01, 1.74396512e-01, 5.29356748e-01, 4.38475687e-01, - 4.94036404e-01, 4.09785794e-01, 6.40025507e-01, 5.79371821e-01, 1.57726118e-01, 6.04572263e-01, - 5.41072639e-01, 5.18847173e-01, 1.97093284e-01, 8.91767002e-01, 4.29050835e-01, 8.25490570e-01, - 3.87699807e-01, 4.50705808e-01, 2.49371643e-01, 3.36074898e-01, 9.29925118e-01, 6.65393649e-01, - 9.07275994e-01, 3.73075859e-01, 4.14044139e-03, 2.37463702e-01, 2.25893784e-01, 2.46900245e-01, - 4.50350196e-01, 3.48618117e-01, 5.07193932e-01, 5.23435142e-01, 8.13611417e-01, 8.92715622e-01, - 1.02623450e-01, 3.06088345e-01, 7.80461650e-01, 2.21453645e-01, 2.01419652e-01, 2.84254457e-01, - 3.68286735e-01, 7.39358243e-01, 8.97879394e-01, 9.81599566e-01, 7.56526442e-01, 7.37645545e-01, - 4.23976657e-02, 8.25922012e-01, 2.60956996e-01, 2.90702065e-01, 8.98388344e-01, 3.03733299e-01, - 8.49071471e-01, 3.45835425e-01, 7.65458276e-01, 5.68094872e-01, 8.93770930e-01, 9.93161641e-01, - 5.63368667e-02, 4.26548945e-01, 5.46745780e-01, 5.75674571e-01, 7.94599487e-01, 7.18935553e-02, - 4.46492976e-01, 6.40240123e-01, 2.73246969e-01, 2.00465968e-01, 1.30718835e-01, 1.92492005e-01, - 1.96617189e-01, 6.61271644e-01, 8.12687657e-01, 8.66342445e-01 - - }, - {0, 9, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 7, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 6, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 5, 0, 0, 0, 0, 4, 0, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 2, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, - 10, - true, - -4}, - {10, - 5, - {0.21390334, 0.50261639, 0.91036676, 0.59166485, 0.71162682, 0.10248392, 0.77782677, 0.43772379, - 0.4035871, 0.3282796, 0.47544681, 0.59862974, 0.12319357, 0.06239463, 0.28200272, 0.1345717, - 0.50498218, 0.5113505, 0.16233086, 0.62165332, 0.42281548, 0.933117, 0.41386077, 0.23264562, - 0.73325968, 0.37537541, 0.70719873, 0.14522645, 0.73279625, 0.9126674, 0.84854131, 0.28890216, - 0.85267903, 0.74703138, 0.83842071, 0.34942792, 0.27864171, 0.70911132, 0.21338564, 0.32035554, - 0.73788331, 0.46926692, 0.57570162, 0.42559178, 0.87120209, 0.22734951, 0.01847905, 0.75549396, - 0.76166195, 0.66613745}, - {9, 8, 7, 6, 5, 4, 3, 2, 1, 0}, - 10, - false, - 5}, - // Test outlier points - {9, - 2, - {-1, -50, 3, 4, 5000, 10000, 1, 3, 4, 5, 0.000005, 0.00002, 2000000, 500000, 10, 50, 30, 5}, - {6, 0, 5, 0, 0, 4, 3, 2, 1}, - 7, - false, - 5}, - - // Test n_clusters == (n_points / 2) - {10, - 5, - {0.21390334, 0.50261639, 0.91036676, 0.59166485, 0.71162682, 0.10248392, 0.77782677, 0.43772379, - 0.4035871, 0.3282796, 0.47544681, 0.59862974, 0.12319357, 0.06239463, 0.28200272, 0.1345717, - 0.50498218, 0.5113505, 0.16233086, 0.62165332, 0.42281548, 0.933117, 0.41386077, 0.23264562, - 0.73325968, 0.37537541, 0.70719873, 0.14522645, 0.73279625, 0.9126674, 0.84854131, 0.28890216, - 0.85267903, 0.74703138, 0.83842071, 0.34942792, 0.27864171, 0.70911132, 0.21338564, 0.32035554, - 0.73788331, 0.46926692, 0.57570162, 0.42559178, 0.87120209, 0.22734951, 0.01847905, 0.75549396, - 0.76166195, 0.66613745}, - {1, 0, 4, 0, 0, 3, 2, 0, 2, 1}, - 5, - false, - 10}, - - // Test n_points == 100 - {100, - 10, - {6.26168372e-01, 9.30437651e-01, 6.02450208e-01, 2.73025296e-01, 9.53050619e-01, 3.32164396e-01, - 6.88942598e-01, 5.79163537e-01, 6.70341547e-01, 2.70140602e-02, 9.30429671e-01, 7.17721157e-01, - 9.89948537e-01, 7.75253347e-01, 1.34491522e-02, 2.48522428e-02, 3.51413378e-01, 7.64405834e-01, - 7.86373507e-01, 7.18748577e-01, 8.66998621e-01, 6.80316582e-01, 2.51288712e-01, 4.91078420e-01, - 3.76246281e-01, 4.86828710e-01, 5.67464772e-01, 5.30734742e-01, 8.99478296e-01, 7.66699088e-01, - 9.49339111e-01, 3.55248484e-01, 9.06046929e-01, 4.48407772e-01, 6.96395305e-01, 2.44277335e-01, - 7.74840000e-01, 5.21046603e-01, 4.66423971e-02, 5.12019638e-02, 8.95019614e-01, 5.28956953e-01, - 4.31536306e-01, 5.83857744e-01, 4.41787364e-01, 4.68656523e-01, 5.73971433e-01, 6.79989654e-01, - 3.19650588e-01, 6.12579596e-01, 6.49126442e-02, 8.39131142e-01, 2.85252117e-01, 5.84848929e-01, - 9.46507115e-01, 8.58440748e-01, 3.61528940e-01, 2.44215959e-01, 3.80101125e-01, 4.57128957e-02, - 8.82216988e-01, 8.31498633e-01, 7.23474381e-01, 7.75788607e-01, 1.40864146e-01, 6.62092382e-01, - 5.13985168e-01, 3.00686418e-01, 8.70109949e-01, 2.43187753e-01, 2.89391938e-01, 2.84214238e-01, - 8.70985521e-01, 8.77491176e-01, 6.72537226e-01, 3.30929686e-01, 1.85934324e-01, 9.16222614e-01, - 6.18239142e-01, 2.64768597e-01, 5.76145451e-01, 8.62961369e-01, 6.84757925e-01, 7.60549082e-01, - 1.27645356e-01, 4.51004673e-01, 3.92292980e-01, 4.63170803e-01, 4.35449330e-02, 2.17583404e-01, - 5.71832605e-02, 2.06763039e-01, 3.70116249e-01, 2.09750028e-01, 6.17283019e-01, 8.62549231e-01, - 9.84156240e-02, 2.66249156e-01, 3.87635103e-01, 2.85591012e-02, 4.24826068e-01, 4.45795088e-01, - 6.86227676e-01, 1.08848960e-01, 5.96731841e-02, 3.71770228e-01, 1.91548833e-01, 6.95136078e-01, - 9.00700636e-01, 8.76363105e-01, 2.67334632e-01, 1.80619709e-01, 7.94060419e-01, 1.42854171e-02, - 1.09372387e-01, 8.74028108e-01, 6.46403232e-01, 4.86588834e-01, 5.93446175e-02, 6.11886291e-01, - 8.83865057e-01, 3.15879821e-01, 2.27043992e-01, 9.76764951e-01, 6.15620336e-01, 9.76199360e-01, - 2.40548962e-01, 3.21795663e-01, 8.75087904e-02, 8.11234663e-01, 6.96070480e-01, 8.12062321e-01, - 1.21958818e-01, 3.44348628e-02, 8.72630414e-01, 3.06162776e-01, 1.76043529e-02, 9.45894971e-01, - 5.33896401e-01, 6.21642973e-01, 4.93062535e-01, 4.48984262e-01, 2.24560379e-01, 4.24052195e-02, - 4.43447610e-01, 8.95646149e-01, 6.05220676e-01, 1.81840491e-01, 9.70831206e-01, 2.12563586e-02, - 6.92582693e-01, 7.55946922e-01, 7.95086143e-01, 6.05328941e-01, 3.99350764e-01, 4.32846636e-01, - 9.81114529e-01, 4.98266428e-01, 6.37127930e-03, 1.59085889e-01, 6.34682067e-05, 5.59429440e-01, - 7.38827633e-01, 8.93214770e-01, 2.16494306e-01, 9.35430573e-02, 4.75665868e-02, 7.80503518e-01, - 7.86240041e-01, 7.06854594e-01, 2.13725879e-02, 7.68246091e-01, 4.50234808e-01, 5.21231104e-01, - 5.01989826e-03, 4.22081572e-02, 1.65337732e-01, 8.54134740e-01, 4.99430262e-01, 8.94525601e-01, - 1.14028379e-01, 3.69739861e-01, 1.32955599e-01, 2.65563824e-01, 2.52811151e-01, 1.44792843e-01, - 6.88449594e-01, 4.44921417e-01, 8.23296587e-01, 1.93266317e-01, 1.19033309e-01, 1.36368966e-01, - 3.42600285e-01, 5.64505195e-01, 5.57594559e-01, 7.44257892e-01, 8.38231569e-02, 4.11548847e-01, - 3.21010077e-01, 8.55081359e-01, 4.30105779e-01, 1.16229135e-01, 9.87731964e-02, 3.14712335e-01, - 4.50880592e-01, 2.72289598e-01, 6.31615256e-01, 8.97432958e-01, 4.44764250e-01, 8.03776440e-01, - 2.68767748e-02, 2.43374608e-01, 4.02141103e-01, 4.98881209e-01, 5.33173003e-01, 8.82890436e-01, - 7.16149148e-01, 4.19664401e-01, 2.29335357e-01, 2.88637806e-01, 3.44696803e-01, 6.78171906e-01, - 5.69849716e-01, 5.86454477e-01, 3.54474989e-01, 9.03876540e-01, 6.45980000e-01, 6.34887593e-01, - 7.88039746e-02, 2.04814126e-01, 7.82251754e-01, 2.43147074e-01, 7.50951808e-01, 1.72799092e-02, - 2.95349590e-01, 6.57991826e-01, 8.81214312e-01, 5.73970708e-01, 2.77610881e-01, 1.82155097e-01, - 7.69797417e-02, 6.44792402e-01, 9.46950998e-01, 7.73064845e-01, 6.04733624e-01, 5.80094567e-01, - 1.67498426e-01, 2.66514296e-01, 6.50140368e-01, 1.91170299e-01, 2.08752199e-01, 3.01664091e-01, - 9.85033484e-01, 2.92909152e-01, 8.65816607e-01, 1.85222119e-01, 2.28814559e-01, 1.34286382e-02, - 2.89234322e-01, 8.18668708e-01, 4.71706924e-01, 9.23199803e-01, 2.80879188e-01, 1.47319284e-01, - 4.13915748e-01, 9.31274932e-02, 6.66322195e-01, 9.66953974e-01, 3.19405786e-01, 6.69486551e-01, - 5.03096313e-02, 6.95225201e-01, 5.78469859e-01, 6.29481655e-01, 1.39252534e-01, 1.22564968e-01, - 6.80663678e-01, 6.34607157e-01, 6.42765834e-01, 1.57127410e-02, 2.92132086e-01, 5.24423878e-01, - 4.68676824e-01, 2.86003928e-01, 7.18608322e-01, 8.95617933e-01, 5.48844309e-01, 1.74517278e-01, - 5.24379196e-01, 2.13526524e-01, 5.88375435e-01, 9.88560185e-01, 4.17435771e-01, 6.14438688e-01, - 9.53760881e-01, 5.27151288e-01, 7.03017278e-01, 3.44448559e-01, 4.47059676e-01, 2.83414901e-01, - 1.98979011e-01, 4.24917361e-01, 5.73172761e-01, 2.32398853e-02, 1.65887230e-01, 4.05552785e-01, - 9.29665524e-01, 2.26135696e-01, 9.20563384e-01, 7.65259963e-01, 4.54820075e-01, 8.97710267e-01, - 3.78559302e-03, 9.15219382e-01, 3.55705698e-01, 6.94905124e-01, 8.58540202e-01, 3.89790666e-01, - 2.49478206e-01, 7.93679304e-01, 4.75830027e-01, 4.40425353e-01, 3.70579459e-01, 1.40578049e-01, - 1.70386675e-01, 7.04056121e-01, 4.85963102e-01, 9.68450060e-01, 6.77178001e-01, 2.65934654e-01, - 2.58915007e-01, 6.70052890e-01, 2.61945109e-01, 8.46207759e-01, 1.01928951e-01, 2.85611334e-01, - 2.45776933e-01, 2.66658783e-01, 3.71724077e-01, 4.34319025e-01, 4.24407347e-01, 7.15417683e-01, - 8.07997684e-01, 1.64296275e-01, 6.01638065e-01, 8.60606804e-02, 2.68719187e-01, 5.11764101e-01, - 9.75844338e-01, 7.81226782e-01, 2.20925515e-01, 7.18135040e-01, 9.82395577e-01, 8.39160243e-01, - 9.08058083e-01, 6.88010677e-01, 8.14271847e-01, 5.12460821e-01, 1.17311345e-01, 5.96075228e-01, - 9.17455497e-01, 2.12052706e-01, 7.04074603e-01, 8.72872565e-02, 8.76047818e-01, 6.96235046e-01, - 8.54801557e-01, 2.49729159e-01, 9.76594604e-01, 2.87386363e-01, 2.36461559e-02, 9.94075254e-01, - 4.25193986e-01, 7.61869994e-01, 5.13334255e-01, 6.44711165e-02, 8.92156689e-01, 3.55235167e-01, - 1.08154647e-01, 8.78446825e-01, 2.43833016e-01, 9.23071293e-01, 2.72724115e-01, 9.46631338e-01, - 3.74510294e-01, 4.08451278e-02, 9.78392777e-01, 3.65079221e-01, 6.37199516e-01, 5.51144906e-01, - 5.25978080e-01, 1.42803678e-01, 4.05451674e-01, 7.79788219e-01, 6.26009784e-01, 3.35249497e-01, - 1.43159543e-02, 1.80363779e-01, 5.05096904e-01, 2.82619947e-01, 5.83561392e-01, 3.10951324e-01, - 8.73223968e-01, 4.38545619e-01, 4.81348800e-01, 6.68497085e-01, 3.79345401e-01, 9.58832501e-01, - 1.89869550e-01, 2.34083070e-01, 2.94066207e-01, 5.74892667e-02, 6.92106828e-02, 9.61127686e-02, - 6.72650672e-02, 8.47345378e-01, 2.80916761e-01, 7.32177357e-03, 9.80785961e-01, 5.73192225e-02, - 8.48781331e-01, 8.83225408e-01, 7.34398275e-01, 7.70381941e-01, 6.20778343e-01, 8.96822048e-01, - 5.40732486e-01, 3.69704071e-01, 5.77305837e-01, 2.08221827e-01, 7.34275341e-01, 1.06110900e-01, - 3.49496706e-01, 8.34948910e-01, 1.56403291e-02, 6.78576376e-01, 8.96141268e-01, 5.94835119e-01, - 1.43943153e-01, 3.49618530e-01, 2.10440392e-01, 3.46585620e-01, 1.05153093e-01, 3.45446174e-01, - 2.72177079e-01, 7.07946300e-01, 4.33717726e-02, 3.31232203e-01, 3.91874320e-01, 4.76338141e-01, - 6.22777789e-01, 2.95989228e-02, 4.32855769e-01, 7.61049310e-01, 3.63279149e-01, 9.47210350e-01, - 6.43721247e-01, 6.58025802e-01, 1.05247633e-02, 5.29974442e-01, 7.30675767e-01, 4.30041079e-01, - 6.62634841e-01, 8.25936616e-01, 9.91253704e-01, 6.79399281e-01, 5.44177006e-01, 7.52876048e-01, - 3.32139049e-01, 7.98732398e-01, 7.38865223e-01, 9.16055132e-01, 6.11736493e-01, 9.63672879e-01, - 1.83778839e-01, 7.27558919e-02, 5.91602822e-01, 3.25235484e-01, 2.34741217e-01, 9.52346277e-01, - 9.18556407e-01, 9.35373324e-01, 6.89209070e-01, 2.56049054e-01, 6.17975395e-01, 7.82285691e-01, - 9.84983432e-01, 6.62322741e-01, 2.04144457e-01, 3.98446577e-01, 1.38918297e-01, 3.05919921e-01, - 3.14043787e-01, 5.91072666e-01, 7.44703771e-01, 8.92272567e-01, 9.78017873e-01, 9.01203161e-01, - 1.41526372e-01, 4.14878484e-01, 6.80683651e-01, 5.01733152e-02, 8.14635389e-01, 2.27926375e-01, - 9.03269815e-01, 8.68443745e-01, 9.86939190e-01, 7.40779486e-01, 2.61005311e-01, 3.19276232e-01, - 9.69509248e-01, 1.11908818e-01, 4.49198556e-01, 1.27056715e-01, 3.84064823e-01, 5.14591811e-01, - 2.10747488e-01, 9.53884090e-01, 8.43167950e-01, 4.51187972e-01, 3.75331782e-01, 6.23566461e-01, - 3.55290379e-01, 2.95705968e-01, 1.69622690e-01, 1.42981830e-01, 2.72180991e-01, 9.46468040e-01, - 3.70932500e-01, 9.94292830e-01, 4.62587505e-01, 7.14817405e-01, 2.45370540e-02, 3.00906377e-01, - 5.75768304e-01, 9.71448393e-01, 6.95574827e-02, 3.93693854e-01, 5.29306116e-01, 5.04694554e-01, - 6.73797120e-02, 6.76596969e-01, 5.50948898e-01, 3.24909641e-01, 7.70337719e-01, 6.51842631e-03, - 3.03264879e-01, 7.61037886e-03, 2.72289601e-01, 1.50502041e-01, 6.71103888e-02, 7.41503703e-01, - 1.92088941e-01, 2.19043977e-01, 9.09320161e-01, 2.37993569e-01, 6.18107973e-02, 8.31447852e-01, - 2.23355609e-01, 1.84789435e-01, 4.16104518e-01, 4.21573859e-01, 8.72446305e-02, 2.97294197e-01, - 4.50328256e-01, 8.72199917e-01, 2.51279916e-01, 4.86219272e-01, 7.57071329e-01, 4.85655942e-01, - 1.06187277e-01, 4.92341327e-01, 1.46017513e-01, 5.25421017e-01, 4.22637906e-01, 2.24685018e-01, - 8.72648431e-01, 5.54051490e-01, 1.80745062e-01, 2.12756336e-01, 5.20883169e-01, 7.60363654e-01, - 8.30254678e-01, 5.00003328e-01, 4.69017439e-01, 6.38105527e-01, 3.50638261e-02, 5.22217353e-02, - 9.06516882e-02, 8.52975842e-01, 1.19985883e-01, 3.74926753e-01, 6.50302066e-01, 1.98875727e-01, - 6.28362507e-02, 4.32693501e-01, 3.10500685e-01, 6.20732833e-01, 4.58503272e-01, 3.20790034e-01, - 7.91284868e-01, 7.93054570e-01, 2.93406765e-01, 8.95399023e-01, 1.06441034e-01, 7.53085241e-02, - 8.67523104e-01, 1.47963482e-01, 1.25584706e-01, 3.81545040e-02, 6.34338619e-01, 1.76368938e-02, - 5.75553531e-02, 5.31607516e-01, 2.63869588e-01, 9.41945823e-01, 9.24028838e-02, 5.21496463e-01, - 7.74866558e-01, 5.65210610e-01, 7.28015327e-02, 6.51963790e-01, 8.94727453e-01, 4.49571590e-01, - 1.29932405e-01, 8.64026259e-01, 9.92599934e-01, 7.43721560e-01, 8.87300215e-01, 1.06369925e-01, - 8.11335531e-01, 7.87734900e-01, 9.87344678e-01, 5.32502820e-01, 4.42612382e-01, 9.64041183e-01, - 1.66085871e-01, 1.12937664e-01, 5.24423470e-01, 6.54689333e-01, 4.59119726e-01, 5.22774091e-01, - 3.08722276e-02, 6.26979315e-01, 4.49754105e-01, 8.07495757e-01, 2.34199499e-01, 1.67765675e-01, - 9.22168418e-01, 3.73210378e-01, 8.04432575e-01, 5.61890354e-01, 4.47025593e-01, 6.43155678e-01, - 2.40407640e-01, 5.91631279e-01, 1.59369206e-01, 7.75799090e-01, 8.32067212e-01, 5.59791576e-02, - 6.39105224e-01, 4.85274738e-01, 2.12630838e-01, 2.81431312e-02, 7.16205363e-01, 6.83885011e-01, - 5.23869697e-01, 9.99418314e-01, 8.35331599e-01, 4.69877463e-02, 6.74712562e-01, 7.99273684e-01, - 2.77001890e-02, 5.75809742e-01, 2.78513031e-01, 8.36209905e-01, 7.25472379e-01, 4.87173943e-01, - 7.88311357e-01, 9.64676177e-01, 1.75752651e-01, 4.98112580e-01, 8.08850418e-02, 6.40981131e-01, - 4.06647450e-01, 8.46539387e-01, 2.12620694e-01, 9.11012851e-01, 8.25041445e-01, 8.90065575e-01, - 9.63626055e-01, 5.96689242e-01, 1.63372670e-01, 4.51640148e-01, 3.43026542e-01, 5.80658851e-01, - 2.82327625e-01, 4.75535418e-01, 6.27760926e-01, 8.46314115e-01, 9.61961932e-01, 3.19806094e-01, - 5.05508062e-01, 5.28102944e-01, 6.13045057e-01, 7.44714938e-01, 1.50586073e-01, 7.91878033e-01, - 4.89839179e-01, 3.10496849e-01, 8.82309038e-01, 2.86922314e-01, 4.84687559e-01, 5.20838630e-01, - 4.62955493e-01, 2.38185305e-01, 5.47259907e-02, 7.10916137e-01, 7.31887202e-01, 6.25602317e-01, - 8.77741168e-01, 4.19881322e-01, 4.81222328e-01, 1.28224501e-01, 2.46034010e-01, 3.34971854e-01, - 7.37216484e-01, 5.62134821e-02, 7.14089724e-01, 9.85549393e-01, 4.66295827e-01, 3.08722434e-03, - 4.70237690e-01, 2.66524167e-01, 7.93875484e-01, 4.54795911e-02, 8.09702944e-01, 1.47709735e-02, - 1.70082405e-01, 6.35905179e-01, 3.75379109e-01, 4.30315011e-01, 3.15788760e-01, 5.58065230e-01, - 2.24643800e-01, 2.42142981e-01, 6.57283636e-01, 3.34921891e-01, 1.26588975e-01, 7.68064155e-01, - 9.43856291e-01, 4.47518596e-01, 5.44453573e-01, 9.95764932e-01, 7.16444391e-01, 8.51019765e-01, - 1.01179183e-01, 4.45473958e-01, 4.60327322e-01, 4.96895844e-02, 4.72907738e-01, 5.58987444e-01, - 3.41027487e-01, 1.56175026e-01, 7.58283148e-01, 6.83600909e-01, 2.14623396e-01, 3.27348880e-01, - 3.92517893e-01, 6.70418431e-01, 5.16440832e-01, 8.63140348e-01, 5.73277464e-01, 3.46608058e-01, - 7.39396341e-01, 7.20852434e-01, 2.35653246e-02, 3.89935659e-01, 7.53783745e-01, 6.34563528e-01, - 8.79339335e-01, 7.41599159e-02, 5.62433904e-01, 6.15553852e-01, 4.56956324e-01, 5.20047447e-01, - 5.26845015e-02, 5.58471266e-01, 1.63632233e-01, 5.38936665e-02, 6.49593683e-01, 2.56838748e-01, - 8.99035326e-01, 7.20847756e-01, 5.68954684e-01, 7.43684755e-01, 5.70924238e-01, 3.82318724e-01, - 4.89328290e-01, 5.62208561e-01, 4.97540804e-02, 4.18011085e-01, 6.88041565e-01, 2.16234653e-01, - 7.89548214e-01, 8.46136387e-01, 8.46816189e-01, 1.73842353e-01, 6.11627842e-02, 8.44440559e-01, - 4.50646654e-01, 3.74785037e-01, 4.87196697e-01, 4.56276448e-01, 9.13284391e-01, 4.15715464e-01, - 7.13597697e-01, 1.23641270e-02, 5.10031271e-01, 4.74601930e-02, 2.55731159e-01, 3.22090006e-01, - 1.91165703e-01, 4.51170940e-01, 7.50843157e-01, 4.42420576e-01, 4.25380660e-01, 4.50667257e-01, - 6.55689206e-01, 9.68257670e-02, 1.96528793e-01, 8.97343028e-01, 4.99940904e-01, 6.65504083e-01, - 9.41828079e-01, 4.54397338e-01, 5.61893331e-01, 5.09839880e-01, 4.53117514e-01, 8.96804127e-02, - 1.74888861e-01, 6.65641378e-01, 2.81668336e-01, 1.89532742e-01, 5.61668382e-01, 8.68330157e-02, - 8.25092797e-01, 5.18106324e-01, 1.71904024e-01, 3.68385523e-01, 1.62005436e-01, 7.48507399e-01, - 9.30274827e-01, 2.38198517e-01, 9.52222901e-01, 5.23587800e-01, 6.94384557e-01, 1.09338652e-01, - 4.83356794e-01, 2.73050402e-01, 3.68027050e-01, 5.92366466e-01, 1.83192289e-01, 8.60376029e-01, - 7.13926203e-01, 8.16750052e-01, 1.57890291e-01, 6.25691951e-01, 5.24831646e-01, 1.73873797e-01, - 1.02429784e-01, 9.17488471e-01, 4.03584434e-01, 9.31170884e-01, 2.79386137e-01, 8.77745206e-01, - 2.45200576e-01, 1.28896951e-01, 3.15713052e-01, 5.27874291e-01, 2.16444335e-01, 7.03883817e-01, - 7.74738919e-02, 8.42422142e-01, 3.75598924e-01, 3.51002411e-01, 6.22752776e-01, 4.82407943e-01, - 7.43107867e-01, 9.46182666e-01, 9.44344819e-01, 3.28124763e-01, 1.06147431e-01, 1.65102684e-01, - 3.84060507e-01, 2.91057722e-01, 7.68173662e-02, 1.03543651e-01, 6.76698940e-01, 1.43141994e-01, - 7.21342202e-01, 6.69471294e-03, 9.07298311e-01, 5.57080171e-01, 8.10954489e-01, 4.11120526e-01, - 2.06407453e-01, 2.59590556e-01, 7.58512718e-01, 5.79873897e-01, 2.92875650e-01, 2.83686529e-01, - 2.42829343e-01, 9.19323719e-01, 3.46832864e-01, 3.58238858e-01, 7.42827585e-01, 2.05760059e-01, - 9.58438860e-01, 5.66326411e-01, 6.60292846e-01, 5.61095078e-02, 6.79465531e-01, 7.05118513e-01, - 4.44713264e-01, 2.09732933e-01, 5.22732436e-01, 1.74396512e-01, 5.29356748e-01, 4.38475687e-01, - 4.94036404e-01, 4.09785794e-01, 6.40025507e-01, 5.79371821e-01, 1.57726118e-01, 6.04572263e-01, - 5.41072639e-01, 5.18847173e-01, 1.97093284e-01, 8.91767002e-01, 4.29050835e-01, 8.25490570e-01, - 3.87699807e-01, 4.50705808e-01, 2.49371643e-01, 3.36074898e-01, 9.29925118e-01, 6.65393649e-01, - 9.07275994e-01, 3.73075859e-01, 4.14044139e-03, 2.37463702e-01, 2.25893784e-01, 2.46900245e-01, - 4.50350196e-01, 3.48618117e-01, 5.07193932e-01, 5.23435142e-01, 8.13611417e-01, 8.92715622e-01, - 1.02623450e-01, 3.06088345e-01, 7.80461650e-01, 2.21453645e-01, 2.01419652e-01, 2.84254457e-01, - 3.68286735e-01, 7.39358243e-01, 8.97879394e-01, 9.81599566e-01, 7.56526442e-01, 7.37645545e-01, - 4.23976657e-02, 8.25922012e-01, 2.60956996e-01, 2.90702065e-01, 8.98388344e-01, 3.03733299e-01, - 8.49071471e-01, 3.45835425e-01, 7.65458276e-01, 5.68094872e-01, 8.93770930e-01, 9.93161641e-01, - 5.63368667e-02, 4.26548945e-01, 5.46745780e-01, 5.75674571e-01, 7.94599487e-01, 7.18935553e-02, - 4.46492976e-01, 6.40240123e-01, 2.73246969e-01, 2.00465968e-01, 1.30718835e-01, 1.92492005e-01, - 1.96617189e-01, 6.61271644e-01, 8.12687657e-01, 8.66342445e-01 - - }, - {0, 9, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 7, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 6, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, - 0, 0, 0, 0, 0, 5, 0, 0, 0, 0, 4, 0, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 2, 0, - 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}, - 10, - false, - 5}}; - -typedef LinkageTest LinkageTestF_Int; -TEST_P(LinkageTestF_Int, Result) { EXPECT_TRUE(score == 1.0); } - -INSTANTIATE_TEST_CASE_P(LinkageTest, LinkageTestF_Int, ::testing::ValuesIn(linkage_inputsf2)); -} // end namespace raft diff --git a/cpp/test/cluster/spectral.cu b/cpp/test/cluster/spectral.cu deleted file mode 100644 index b8ee611f29..0000000000 --- a/cpp/test/cluster/spectral.cu +++ /dev/null @@ -1,109 +0,0 @@ -/* - * Copyright (c) 2020-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" - -#include -#include -#include - -#include - -#include -#include - -namespace raft { -namespace cluster { - -/** - * Warning: There appears to be a CUDA 12.2 bug in cusparse that causes an - * alignment issue. We've fixed the bug in our code through a workaround - * (see raft/sparse/linalg/spmm.hpp for fix). This test is meant to fail - * in the case where the fix is accidentally reverted, so that it doesn't - * break any downstream libraries that depend on RAFT - */ -TEST(Raft, Spectral) -{ - raft::handle_t handle; - - std::vector h_offsets({0, 2, 4, 7, 10, 12, 14}); - std::vector h_indices({1, 2, 0, 2, 0, 1, 3, 2, 4, 5, 3, 5, 3, 4}); - std::vector h_values( - {1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0}); - std::vector expected_clustering({1, 1, 1, 0, 0, 0}); - - int32_t n_clusters{2}; - int32_t n_eigenvectors{2}; - int32_t evs_max_it{100}; - int32_t kmean_max_it{100}; - int32_t restartIter_lanczos = 15 + n_eigenvectors; - float evs_tol{0.001}; - float kmean_tol{0.001}; - unsigned long long seed1{1234567}; - unsigned long long seed2{12345678}; - bool reorthog{false}; - - rmm::device_uvector offsets(h_offsets.size(), handle.get_stream()); - rmm::device_uvector indices(h_indices.size(), handle.get_stream()); - rmm::device_uvector values(h_indices.size(), handle.get_stream()); - rmm::device_uvector clustering(expected_clustering.size(), handle.get_stream()); - rmm::device_uvector eigenvalues(n_eigenvectors, handle.get_stream()); - rmm::device_uvector eigenvectors(n_eigenvectors * expected_clustering.size(), - handle.get_stream()); - - rmm::device_uvector exp_dev(expected_clustering.size(), handle.get_stream()); - - raft::update_device( - exp_dev.data(), expected_clustering.data(), expected_clustering.size(), handle.get_stream()); - - raft::update_device(offsets.data(), h_offsets.data(), h_offsets.size(), handle.get_stream()); - raft::update_device(indices.data(), h_indices.data(), h_indices.size(), handle.get_stream()); - raft::update_device(values.data(), h_values.data(), h_values.size(), handle.get_stream()); - - raft::spectral::matrix::sparse_matrix_t const matrix{ - handle, - offsets.data(), - indices.data(), - values.data(), - static_cast(offsets.size() - 1), - static_cast(indices.size())}; - - raft::spectral::eigen_solver_config_t eig_cfg{ - n_eigenvectors, evs_max_it, restartIter_lanczos, evs_tol, reorthog, seed1}; - raft::spectral::lanczos_solver_t eig_solver{eig_cfg}; - - raft::spectral::cluster_solver_config_t clust_cfg{ - n_clusters, kmean_max_it, kmean_tol, seed2}; - raft::spectral::kmeans_solver_t cluster_solver{clust_cfg}; - - raft::spectral::partition(handle, - matrix, - eig_solver, - cluster_solver, - clustering.data(), - eigenvalues.data(), - eigenvectors.data()); - - ASSERT_TRUE(devArrMatch(expected_clustering.data(), - exp_dev.data(), - exp_dev.size(), - 1, - raft::Compare(), - handle.get_stream())); -} - -} // namespace cluster -} // namespace raft \ No newline at end of file diff --git a/cpp/test/distance/dist_adj.cu b/cpp/test/distance/dist_adj.cu deleted file mode 100644 index a22fb7b1f9..0000000000 --- a/cpp/test/distance/dist_adj.cu +++ /dev/null @@ -1,196 +0,0 @@ -/* - * Copyright (c) 2018-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "dist_adj.cuh" - -#include -#include -#include -#include -#include -#include - -#include - -#include - -namespace raft { -namespace distance { - -template -RAFT_KERNEL naiveDistanceAdjKernel(uint8_t* dist, - const DataType* x, - const DataType* y, - int m, - int n, - int k, - DataType eps, - bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) return; - DataType acc = DataType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto diff = x[xidx] - y[yidx]; - acc += diff * diff; - } - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc <= eps; -} - -template -void naiveDistanceAdj(uint8_t* dist, - const DataType* x, - const DataType* y, - int m, - int n, - int k, - DataType eps, - bool isRowMajor, - cudaStream_t stream) -{ - static const dim3 TPB(16, 32, 1); - dim3 nblks(raft::ceildiv(m, (int)TPB.x), raft::ceildiv(n, (int)TPB.y), 1); - naiveDistanceAdjKernel<<>>(dist, x, y, m, n, k, eps, isRowMajor); - RAFT_CUDA_TRY(cudaPeekAtLastError()); -} - -template -struct DistanceAdjInputs { - DataType eps; - int m, n, k; - bool isRowMajor; - unsigned long long int seed; -}; - -template -::std::ostream& operator<<(::std::ostream& os, const DistanceAdjInputs& dims) -{ - return os; -} - -template -class DistanceAdjTest : public ::testing::TestWithParam> { - public: - DistanceAdjTest() - : params(::testing::TestWithParam>::GetParam()), - stream(resource::get_cuda_stream(handle)), - dist(params.m * params.n, stream), - dist_ref(params.m * params.n, stream) - { - } - - void SetUp() override - { - raft::random::RngState r(params.seed); - int m = params.m; - int n = params.n; - int k = params.k; - bool isRowMajor = params.isRowMajor; - - rmm::device_uvector x(m * k, stream); - rmm::device_uvector y(n * k, stream); - - uniform(handle, r, x.data(), m * k, DataType(-1.0), DataType(1.0)); - uniform(handle, r, y.data(), n * k, DataType(-1.0), DataType(1.0)); - - DataType threshold = params.eps; - - naiveDistanceAdj(dist_ref.data(), x.data(), y.data(), m, n, k, threshold, isRowMajor, stream); - size_t worksize = raft::distance:: - getWorkspaceSize( - x.data(), y.data(), m, n, k); - rmm::device_uvector workspace(worksize, stream); - - using threshold_final_op_ = threshold_final_op; - threshold_final_op_ threshold_op(threshold); - - raft::distance::distance(handle, - x.data(), - y.data(), - dist.data(), - m, - n, - k, - workspace.data(), - worksize, - threshold_op, - isRowMajor); - resource::sync_stream(handle, stream); - } - - void TearDown() override {} - - protected: - DistanceAdjInputs params; - // We use uint8_t even if the output in this test is a bool because - // cutlass doesn't support bool as output buffer yet. In cuda - // sizeof(bool) is 1 byte hence it doesn't increase - // memory consumption if we use uint8_t instead of bool. - rmm::device_uvector dist_ref; - rmm::device_uvector dist; - raft::resources handle; - cudaStream_t stream; -}; - -const std::vector> inputsf = { - {0.01f, 1024, 1024, 32, true, 1234ULL}, - {0.1f, 1024, 1024, 32, true, 1234ULL}, - {1.0f, 1024, 1024, 32, true, 1234ULL}, - {10.0f, 1024, 1024, 32, true, 1234ULL}, - {0.01f, 1024, 1024, 32, false, 1234ULL}, - {0.1f, 1024, 1024, 32, false, 1234ULL}, - {1.0f, 1024, 1024, 32, false, 1234ULL}, - {10.0f, 1024, 1024, 32, false, 1234ULL}, -}; -typedef DistanceAdjTest DistanceAdjTestF; -TEST_P(DistanceAdjTestF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch(dist_ref.data(), dist.data(), m, n, raft::Compare(), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceAdjTests, DistanceAdjTestF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.01, 1024, 1024, 32, true, 1234ULL}, - {0.1, 1024, 1024, 32, true, 1234ULL}, - {1.0, 1024, 1024, 32, true, 1234ULL}, - {10.0, 1024, 1024, 32, true, 1234ULL}, - {0.01, 1024, 1024, 32, false, 1234ULL}, - {0.1, 1024, 1024, 32, false, 1234ULL}, - {1.0, 1024, 1024, 32, false, 1234ULL}, - {10.0, 1024, 1024, 32, false, 1234ULL}, -}; -typedef DistanceAdjTest DistanceAdjTestD; -TEST_P(DistanceAdjTestD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch(dist_ref.data(), dist.data(), m, n, raft::Compare(), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceAdjTests, DistanceAdjTestD, ::testing::ValuesIn(inputsd)); - -} // namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_adj.cuh b/cpp/test/distance/dist_adj.cuh deleted file mode 100644 index 2861cb33de..0000000000 --- a/cpp/test/distance/dist_adj.cuh +++ /dev/null @@ -1,72 +0,0 @@ -/* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "dist_adj_threshold.cuh" - -#include - -#define instantiate_raft_distance_distance(DT, DataT, AccT, OutT, FinalLambda, IdxT) \ - extern template void raft::distance::distance( \ - raft::resources const& handle, \ - const DataT* x, \ - const DataT* y, \ - OutT* dist, \ - IdxT m, \ - IdxT n, \ - IdxT k, \ - void* workspace, \ - size_t worksize, \ - FinalLambda fin_op, \ - bool isRowMajor, \ - DataT metric_arg) - -instantiate_raft_distance_distance(raft::distance::DistanceType::L2Expanded, - float, - float, - uint8_t, - raft::distance::threshold_float, - int); - -instantiate_raft_distance_distance(raft::distance::DistanceType::L2Expanded, - double, - double, - uint8_t, - raft::distance::threshold_double, - int); - -#undef instantiate_raft_distance_distance - -#define instantiate_raft_distance_getWorkspaceSize(DistT, DataT, AccT, OutT, IdxT) \ - extern template size_t raft::distance::getWorkspaceSize( \ - const DataT* x, const DataT* y, IdxT m, IdxT n, IdxT k) - -instantiate_raft_distance_getWorkspaceSize( - raft::distance::DistanceType::L2Expanded, float, float, uint8_t, int); -instantiate_raft_distance_getWorkspaceSize( - raft::distance::DistanceType::L2Expanded, double, double, uint8_t, int); - -#undef instantiate_raft_distance_getWorkspaceSize - -#define instantiate_raft_distance_getWorkspaceSize(DistT, DataT, AccT, OutT, IdxT) \ - extern template size_t raft::distance::getWorkspaceSize( \ - const DataT* x, const DataT* y, IdxT m, IdxT n, IdxT k) - -instantiate_raft_distance_getWorkspaceSize( - raft::distance::DistanceType::L2Expanded, float, float, uint8_t, int); -instantiate_raft_distance_getWorkspaceSize( - raft::distance::DistanceType::L2Expanded, double, double, uint8_t, int); - -#undef instantiate_raft_distance_getWorkspaceSize diff --git a/cpp/test/distance/dist_adj_distance_instance.cu b/cpp/test/distance/dist_adj_distance_instance.cu deleted file mode 100644 index 158a5986c2..0000000000 --- a/cpp/test/distance/dist_adj_distance_instance.cu +++ /dev/null @@ -1,65 +0,0 @@ -/* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#undef RAFT_EXPLICIT_INSTANTIATE_ONLY - -#include "dist_adj_threshold.cuh" - -#include - -#include - -#define instantiate_raft_distance_distance(DT, DataT, AccT, OutT, FinalLambda, IdxT) \ - template void raft::distance::distance( \ - raft::resources const& handle, \ - const DataT* x, \ - const DataT* y, \ - OutT* dist, \ - IdxT m, \ - IdxT n, \ - IdxT k, \ - void* workspace, \ - size_t worksize, \ - FinalLambda fin_op, \ - bool isRowMajor, \ - DataT metric_arg) - -instantiate_raft_distance_distance(raft::distance::DistanceType::L2Expanded, - float, - float, - uint8_t, - raft::distance::threshold_float, - int); - -instantiate_raft_distance_distance(raft::distance::DistanceType::L2Expanded, - double, - double, - uint8_t, - raft::distance::threshold_double, - int); - -#undef instantiate_raft_distance_distance - -#define instantiate_raft_distance_getWorkspaceSize(DistT, DataT, AccT, OutT, IdxT) \ - template size_t raft::distance::getWorkspaceSize( \ - const DataT* x, const DataT* y, IdxT m, IdxT n, IdxT k) - -instantiate_raft_distance_getWorkspaceSize( - raft::distance::DistanceType::L2Expanded, float, float, uint8_t, int); -instantiate_raft_distance_getWorkspaceSize( - raft::distance::DistanceType::L2Expanded, double, double, uint8_t, int); - -#undef instantiate_raft_distance_getWorkspaceSize diff --git a/cpp/test/distance/dist_adj_threshold.cuh b/cpp/test/distance/dist_adj_threshold.cuh deleted file mode 100644 index 78663b3cd1..0000000000 --- a/cpp/test/distance/dist_adj_threshold.cuh +++ /dev/null @@ -1,36 +0,0 @@ -/* - * Copyright (c) 2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include // uint8_t - -namespace raft::distance { - -template -struct threshold_final_op { - DataT threshold_val; - - __device__ __host__ threshold_final_op() noexcept : threshold_val(0.0) {} - __device__ __host__ threshold_final_op(DataT val) noexcept : threshold_val(val) {} - __device__ __host__ OutT operator()(AccT d_val, Index g_idx) const noexcept - { - return d_val <= threshold_val; - } -}; - -using threshold_float = threshold_final_op; -using threshold_double = threshold_final_op; - -} // namespace raft::distance diff --git a/cpp/test/distance/dist_canberra.cu b/cpp/test/distance/dist_canberra.cu deleted file mode 100644 index 9b8b6c016b..0000000000 --- a/cpp/test/distance/dist_canberra.cu +++ /dev/null @@ -1,70 +0,0 @@ -/* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceCanberra : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceCanberra DistanceCanberraF; -TEST_P(DistanceCanberraF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceCanberraF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceCanberra DistanceCanberraD; -TEST_P(DistanceCanberraD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceCanberraD, ::testing::ValuesIn(inputsd)); - -class BigMatrixCanberra : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixCanberra, Result) {} - -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_correlation.cu b/cpp/test/distance/dist_correlation.cu deleted file mode 100644 index aa2866483a..0000000000 --- a/cpp/test/distance/dist_correlation.cu +++ /dev/null @@ -1,94 +0,0 @@ -/* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceCorrelation - : public DistanceTest {}; - -template -class DistanceCorrelationXequalY - : public DistanceTestSameBuffer {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceCorrelation DistanceCorrelationF; -TEST_P(DistanceCorrelationF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceCorrelationF, ::testing::ValuesIn(inputsf)); - -typedef DistanceCorrelationXequalY DistanceCorrelationXequalYF; -TEST_P(DistanceCorrelationXequalYF, Result) -{ - int m = params.m; - ASSERT_TRUE(raft::devArrMatch(dist_ref[0].data(), - dist[0].data(), - m, - m, - raft::CompareApprox(params.tolerance), - stream)); - ASSERT_TRUE(raft::devArrMatch(dist_ref[1].data(), - dist[1].data(), - m / 2, - m, - raft::CompareApprox(params.tolerance), - stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceCorrelationXequalYF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceCorrelation DistanceCorrelationD; -TEST_P(DistanceCorrelationD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceCorrelationD, ::testing::ValuesIn(inputsd)); - -class BigMatrixCorrelation - : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixCorrelation, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_cos.cu b/cpp/test/distance/dist_cos.cu deleted file mode 100644 index b792ec4039..0000000000 --- a/cpp/test/distance/dist_cos.cu +++ /dev/null @@ -1,112 +0,0 @@ -/* - * Copyright (c) 2018-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceExpCos : public DistanceTest { -}; - -template -class DistanceExpCosXequalY - : public DistanceTestSameBuffer {}; - -const std::vector> inputsf = { - {0.001f, 128, (65536 + 128) * 128, 8, true, 1234ULL}, - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, (65536 + 128) * 128, 128, 8, false, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; - -const std::vector> inputsXeqYf = { - {0.01f, 1024, 1024, 32, true, 1234ULL}, - {0.01f, 1024, 32, 1024, true, 1234ULL}, - {0.01f, 32, 1024, 1024, true, 1234ULL}, - {0.03f, 1024, 1024, 1024, true, 1234ULL}, - {0.01f, 1024, 1024, 32, false, 1234ULL}, - {0.01f, 1024, 32, 1024, false, 1234ULL}, - {0.01f, 32, 1024, 1024, false, 1234ULL}, - {0.03f, 1024, 1024, 1024, false, 1234ULL}, -}; - -typedef DistanceExpCos DistanceExpCosF; -TEST_P(DistanceExpCosF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceExpCosF, ::testing::ValuesIn(inputsf)); - -typedef DistanceExpCosXequalY DistanceExpCosXequalYF; -TEST_P(DistanceExpCosXequalYF, Result) -{ - int m = params.m; - int n = params.m; - ASSERT_TRUE(raft::devArrMatch(dist_ref[0].data(), - dist[0].data(), - m, - n, - raft::CompareApprox(params.tolerance), - stream)); - n = params.isRowMajor ? m : m / 2; - m = params.isRowMajor ? m / 2 : m; - - ASSERT_TRUE(raft::devArrMatch(dist_ref[1].data(), - dist[1].data(), - m, - n, - raft::CompareApprox(params.tolerance), - stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceExpCosXequalYF, ::testing::ValuesIn(inputsXeqYf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceExpCos DistanceExpCosD; -TEST_P(DistanceExpCosD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceExpCosD, ::testing::ValuesIn(inputsd)); - -class BigMatrixCos : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixCos, Result) {} - -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_dice.cu b/cpp/test/distance/dist_dice.cu deleted file mode 100644 index e127659dc6..0000000000 --- a/cpp/test/distance/dist_dice.cu +++ /dev/null @@ -1,112 +0,0 @@ -/* - * Copyright (c) 2018-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceExpDice : public DistanceTest { -}; - -template -class DistanceExpDiceXequalY - : public DistanceTestSameBuffer {}; - -const std::vector> inputsf = { - {0.001f, 128, (65536 + 128) * 128, 8, true, 1234ULL}, - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, (65536 + 128) * 128, 128, 8, false, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; - -const std::vector> inputsXeqYf = { - {0.01f, 1024, 1024, 32, true, 1234ULL}, - {0.01f, 1024, 32, 1024, true, 1234ULL}, - {0.01f, 32, 1024, 1024, true, 1234ULL}, - {0.03f, 1024, 1024, 1024, true, 1234ULL}, - {0.01f, 1024, 1024, 32, false, 1234ULL}, - {0.01f, 1024, 32, 1024, false, 1234ULL}, - {0.01f, 32, 1024, 1024, false, 1234ULL}, - {0.03f, 1024, 1024, 1024, false, 1234ULL}, -}; - -typedef DistanceExpDice DistanceExpDiceF; -TEST_P(DistanceExpDiceF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApproxNaN(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceExpDiceF, ::testing::ValuesIn(inputsf)); - -typedef DistanceExpDiceXequalY DistanceExpDiceXequalYF; -TEST_P(DistanceExpDiceXequalYF, Result) -{ - int m = params.m; - int n = params.m; - ASSERT_TRUE(raft::devArrMatch(dist_ref[0].data(), - dist[0].data(), - m, - n, - raft::CompareApproxNaN(params.tolerance), - stream)); - n = params.isRowMajor ? m : m / 2; - m = params.isRowMajor ? m / 2 : m; - - ASSERT_TRUE(raft::devArrMatch(dist_ref[1].data(), - dist[1].data(), - m, - n, - raft::CompareApproxNaN(params.tolerance), - stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceExpDiceXequalYF, ::testing::ValuesIn(inputsXeqYf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceExpDice DistanceExpDiceD; -TEST_P(DistanceExpDiceD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApproxNaN(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceExpDiceD, ::testing::ValuesIn(inputsd)); - -class BigMatrixDice : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixDice, Result) {} - -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_hamming.cu b/cpp/test/distance/dist_hamming.cu deleted file mode 100644 index 9529ec2eaa..0000000000 --- a/cpp/test/distance/dist_hamming.cu +++ /dev/null @@ -1,71 +0,0 @@ -/* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceHamming - : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceHamming DistanceHammingF; -TEST_P(DistanceHammingF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceHammingF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceHamming DistanceHammingD; -TEST_P(DistanceHammingD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceHammingD, ::testing::ValuesIn(inputsd)); - -class BigMatrixHamming - : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixHamming, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_hellinger.cu b/cpp/test/distance/dist_hellinger.cu deleted file mode 100644 index 93d6101a18..0000000000 --- a/cpp/test/distance/dist_hellinger.cu +++ /dev/null @@ -1,71 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceHellingerExp - : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceHellingerExp DistanceHellingerExpF; -TEST_P(DistanceHellingerExpF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceHellingerExpF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceHellingerExp DistanceHellingerExpD; -TEST_P(DistanceHellingerExpD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceHellingerExpD, ::testing::ValuesIn(inputsd)); - -class BigMatrixHellingerExp - : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixHellingerExp, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_inner_product.cu b/cpp/test/distance/dist_inner_product.cu deleted file mode 100644 index 8dd7ef0874..0000000000 --- a/cpp/test/distance/dist_inner_product.cu +++ /dev/null @@ -1,74 +0,0 @@ -/* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceInnerProduct - : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 10, 5, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceInnerProduct DistanceInnerProductF; -TEST_P(DistanceInnerProductF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceInnerProductF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceInnerProduct DistanceInnerProductD; -TEST_P(DistanceInnerProductD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceInnerProductD, ::testing::ValuesIn(inputsd)); - -class BigMatrixInnerProduct - : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixInnerProduct, Result) {} - -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_jensen_shannon.cu b/cpp/test/distance/dist_jensen_shannon.cu deleted file mode 100644 index e0e256c925..0000000000 --- a/cpp/test/distance/dist_jensen_shannon.cu +++ /dev/null @@ -1,71 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceJensenShannon - : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceJensenShannon DistanceJensenShannonF; -TEST_P(DistanceJensenShannonF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceJensenShannonF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceJensenShannon DistanceJensenShannonD; -TEST_P(DistanceJensenShannonD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceJensenShannonD, ::testing::ValuesIn(inputsd)); - -class BigMatrixJensenShannon - : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixJensenShannon, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_kl_divergence.cu b/cpp/test/distance/dist_kl_divergence.cu deleted file mode 100644 index 1f79ebcad4..0000000000 --- a/cpp/test/distance/dist_kl_divergence.cu +++ /dev/null @@ -1,71 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceKLDivergence - : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceKLDivergence DistanceKLDivergenceF; -TEST_P(DistanceKLDivergenceF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceKLDivergenceF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceKLDivergence DistanceKLDivergenceD; -TEST_P(DistanceKLDivergenceD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceKLDivergenceD, ::testing::ValuesIn(inputsd)); - -class BigMatrixKLDivergence - : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixKLDivergence, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_l1.cu b/cpp/test/distance/dist_l1.cu deleted file mode 100644 index ce62a4aeec..0000000000 --- a/cpp/test/distance/dist_l1.cu +++ /dev/null @@ -1,70 +0,0 @@ -/* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceUnexpL1 : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceUnexpL1 DistanceUnexpL1F; -TEST_P(DistanceUnexpL1F, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceUnexpL1F, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceUnexpL1 DistanceUnexpL1D; -TEST_P(DistanceUnexpL1D, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceUnexpL1D, ::testing::ValuesIn(inputsd)); - -class BigMatrixUnexpL1 : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixUnexpL1, Result) {} - -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_l2_exp.cu b/cpp/test/distance/dist_l2_exp.cu deleted file mode 100644 index 0203d9ed9d..0000000000 --- a/cpp/test/distance/dist_l2_exp.cu +++ /dev/null @@ -1,115 +0,0 @@ -/* - * Copyright (c) 2018-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceEucExpTest : public DistanceTest { -}; - -template -class DistanceEucExpTestXequalY - : public DistanceTestSameBuffer {}; - -const std::vector> inputsf = { - {0.001f, 128, (65536 + 128) * 128, 8, true, 1234ULL}, - {0.001f, 2048, 4096, 128, true, 1234ULL}, - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.003f, 1021, 1021, 1021, true, 1234ULL}, - {0.001f, (65536 + 128) * 128, 128, 8, false, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, - {0.003f, 1021, 1021, 1021, false, 1234ULL}, -}; - -const std::vector> inputsXeqYf = { - {0.01f, 2048, 4096, 128, true, 1234ULL}, - {0.01f, 1024, 1024, 32, true, 1234ULL}, - {0.01f, 1024, 32, 1024, true, 1234ULL}, - {0.01f, 32, 1024, 1024, true, 1234ULL}, - {0.03f, 1024, 1024, 1024, true, 1234ULL}, - {0.03f, 1021, 1021, 1021, true, 1234ULL}, - {0.01f, 1024, 1024, 32, false, 1234ULL}, - {0.01f, 1024, 32, 1024, false, 1234ULL}, - {0.01f, 32, 1024, 1024, false, 1234ULL}, - {0.03f, 1024, 1024, 1024, false, 1234ULL}, - {0.03f, 1021, 1021, 1021, false, 1234ULL}, -}; - -typedef DistanceEucExpTest DistanceEucExpTestF; -TEST_P(DistanceEucExpTestF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceEucExpTestF, ::testing::ValuesIn(inputsf)); - -typedef DistanceEucExpTestXequalY DistanceEucExpTestXequalYF; -TEST_P(DistanceEucExpTestXequalYF, Result) -{ - int m = params.m; - ASSERT_TRUE(raft::devArrMatch(dist_ref[0].data(), - dist[0].data(), - m, - m, - raft::CompareApprox(params.tolerance), - stream)); - ASSERT_TRUE(raft::devArrMatch(dist_ref[1].data(), - dist[1].data(), - m / 2, - m, - raft::CompareApprox(params.tolerance), - stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, - DistanceEucExpTestXequalYF, - ::testing::ValuesIn(inputsXeqYf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceEucExpTest DistanceEucExpTestD; -TEST_P(DistanceEucExpTestD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceEucExpTestD, ::testing::ValuesIn(inputsd)); - -class BigMatrixEucExp : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixEucExp, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_l2_sqrt_exp.cu b/cpp/test/distance/dist_l2_sqrt_exp.cu deleted file mode 100644 index 5bccabcc3f..0000000000 --- a/cpp/test/distance/dist_l2_sqrt_exp.cu +++ /dev/null @@ -1,74 +0,0 @@ -/* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceEucSqrtExpTest - : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 2048, 4096, 128, true, 1234ULL}, - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.003f, 1021, 1021, 1021, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, - {0.003f, 1021, 1021, 1021, false, 1234ULL}, -}; -typedef DistanceEucSqrtExpTest DistanceEucSqrtExpTestF; -TEST_P(DistanceEucSqrtExpTestF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceEucSqrtExpTestF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceEucSqrtExpTest DistanceEucSqrtExpTestD; -TEST_P(DistanceEucSqrtExpTestD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceEucSqrtExpTestD, ::testing::ValuesIn(inputsd)); - -class BigMatrixEucSqrtExp - : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixEucSqrtExp, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_l2_unexp.cu b/cpp/test/distance/dist_l2_unexp.cu deleted file mode 100644 index 19b0ff6dbf..0000000000 --- a/cpp/test/distance/dist_l2_unexp.cu +++ /dev/null @@ -1,71 +0,0 @@ -/* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceEucUnexpTest - : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceEucUnexpTest DistanceEucUnexpTestF; -TEST_P(DistanceEucUnexpTestF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceEucUnexpTestF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceEucUnexpTest DistanceEucUnexpTestD; -TEST_P(DistanceEucUnexpTestD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceEucUnexpTestD, ::testing::ValuesIn(inputsd)); - -class BigMatrixEucUnexp : public BigMatrixDistanceTest { -}; -TEST_F(BigMatrixEucUnexp, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_l_inf.cu b/cpp/test/distance/dist_l_inf.cu deleted file mode 100644 index 223d186a8d..0000000000 --- a/cpp/test/distance/dist_l_inf.cu +++ /dev/null @@ -1,70 +0,0 @@ -/* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceLinf : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceLinf DistanceLinfF; -TEST_P(DistanceLinfF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceLinfF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceLinf DistanceLinfD; -TEST_P(DistanceLinfD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceLinfD, ::testing::ValuesIn(inputsd)); - -class BigMatrixLinf : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixLinf, Result) {} - -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_lp_unexp.cu b/cpp/test/distance/dist_lp_unexp.cu deleted file mode 100644 index 9d6f5921a7..0000000000 --- a/cpp/test/distance/dist_lp_unexp.cu +++ /dev/null @@ -1,71 +0,0 @@ -/* - * Copyright (c) 2018-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceLpUnexp : public DistanceTest { -}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL, 4.0f}, - {0.001f, 1024, 32, 1024, true, 1234ULL, 3.0f}, - {0.001f, 32, 1024, 1024, true, 1234ULL, 4.0f}, - {0.003f, 1024, 1024, 1024, true, 1234ULL, 3.0f}, - {0.001f, 1024, 1024, 32, false, 1234ULL, 4.0f}, - {0.001f, 1024, 32, 1024, false, 1234ULL, 3.0f}, - {0.001f, 32, 1024, 1024, false, 1234ULL, 4.0f}, - {0.003f, 1024, 1024, 1024, false, 1234ULL, 3.0f}, -}; -typedef DistanceLpUnexp DistanceLpUnexpF; -TEST_P(DistanceLpUnexpF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceLpUnexpF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL, 4.0}, - {0.001, 1024, 32, 1024, true, 1234ULL, 3.0}, - {0.001, 32, 1024, 1024, true, 1234ULL, 4.0}, - {0.003, 1024, 1024, 1024, true, 1234ULL, 3.0}, - {0.001, 1024, 1024, 32, false, 1234ULL, 4.0}, - {0.001, 1024, 32, 1024, false, 1234ULL, 3.0}, - {0.001, 32, 1024, 1024, false, 1234ULL, 4.0}, - {0.003, 1024, 1024, 1024, false, 1234ULL, 3.0}, -}; -typedef DistanceLpUnexp DistanceLpUnexpD; -TEST_P(DistanceLpUnexpD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceLpUnexpD, ::testing::ValuesIn(inputsd)); - -class BigMatrixLpUnexp : public BigMatrixDistanceTest { -}; -TEST_F(BigMatrixLpUnexp, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/dist_russell_rao.cu b/cpp/test/distance/dist_russell_rao.cu deleted file mode 100644 index 73cf4b33a4..0000000000 --- a/cpp/test/distance/dist_russell_rao.cu +++ /dev/null @@ -1,71 +0,0 @@ -/* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "distance_base.cuh" - -namespace raft { -namespace distance { - -template -class DistanceRussellRao - : public DistanceTest {}; - -const std::vector> inputsf = { - {0.001f, 1024, 1024, 32, true, 1234ULL}, - {0.001f, 1024, 32, 1024, true, 1234ULL}, - {0.001f, 32, 1024, 1024, true, 1234ULL}, - {0.003f, 1024, 1024, 1024, true, 1234ULL}, - {0.001f, 1024, 1024, 32, false, 1234ULL}, - {0.001f, 1024, 32, 1024, false, 1234ULL}, - {0.001f, 32, 1024, 1024, false, 1234ULL}, - {0.003f, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceRussellRao DistanceRussellRaoF; -TEST_P(DistanceRussellRaoF, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceRussellRaoF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.001, 1024, 1024, 32, true, 1234ULL}, - {0.001, 1024, 32, 1024, true, 1234ULL}, - {0.001, 32, 1024, 1024, true, 1234ULL}, - {0.003, 1024, 1024, 1024, true, 1234ULL}, - {0.001, 1024, 1024, 32, false, 1234ULL}, - {0.001, 1024, 32, 1024, false, 1234ULL}, - {0.001, 32, 1024, 1024, false, 1234ULL}, - {0.003, 1024, 1024, 1024, false, 1234ULL}, -}; -typedef DistanceRussellRao DistanceRussellRaoD; -TEST_P(DistanceRussellRaoD, Result) -{ - int m = params.isRowMajor ? params.m : params.n; - int n = params.isRowMajor ? params.n : params.m; - ASSERT_TRUE(raft::devArrMatch( - dist_ref.data(), dist.data(), m, n, raft::CompareApprox(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(DistanceTests, DistanceRussellRaoD, ::testing::ValuesIn(inputsd)); - -class BigMatrixRussellRao - : public BigMatrixDistanceTest {}; -TEST_F(BigMatrixRussellRao, Result) {} -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/distance_base.cuh b/cpp/test/distance/distance_base.cuh deleted file mode 100644 index f44fb18519..0000000000 --- a/cpp/test/distance/distance_base.cuh +++ /dev/null @@ -1,708 +0,0 @@ -/* - * Copyright (c) 2018-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" - -#include // common::nvtx::range -#include // make_device_matrix_view -#include // raft::sqrt -#include -#include // raft::resources -#include -#include // raft::distance::DistanceType -#include - -#include // rmm::device_uvector - -#include - -namespace raft { -namespace distance { - -template -RAFT_KERNEL naiveDistanceKernel(DataType* dist, - const DataType* x, - const DataType* y, - int m, - int n, - int k, - raft::distance::DistanceType type, - bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) return; - DataType acc = DataType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto diff = x[xidx] - y[yidx]; - acc += diff * diff; - } - if (type == raft::distance::DistanceType::L2SqrtExpanded || - type == raft::distance::DistanceType::L2SqrtUnexpanded) - acc = raft::sqrt(acc); - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc; -} - -template -RAFT_KERNEL naiveL1_Linf_CanberraDistanceKernel(DataType* dist, - const DataType* x, - const DataType* y, - int m, - int n, - int k, - raft::distance::DistanceType type, - bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) { return; } - - DataType acc = DataType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - auto diff = (a > b) ? (a - b) : (b - a); - if (type == raft::distance::DistanceType::Linf) { - acc = raft::max(acc, diff); - } else if (type == raft::distance::DistanceType::Canberra) { - const auto add = raft::abs(a) + raft::abs(b); - // deal with potential for 0 in denominator by - // forcing 1/0 instead - acc += ((add != 0) * diff / (add + (add == 0))); - } else { - acc += diff; - } - } - - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc; -} - -template -RAFT_KERNEL naiveDiceDistanceKernel( - DataType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) { return; } - - DataType acc_a = DataType(0); - DataType acc_b = DataType(0); - DataType acc_ab = DataType(0); - - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - acc_a += a; - acc_b += b; - acc_ab += a * b; - } - - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - - // Use 1.0 - (dice dissimilarity) to calc the distance - dist[outidx] = (DataType)1.0 - (2 * acc_ab / ((acc_a) + (acc_b))); -} - -template -RAFT_KERNEL naiveCosineDistanceKernel( - DataType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) { return; } - - DataType acc_a = DataType(0); - DataType acc_b = DataType(0); - DataType acc_ab = DataType(0); - - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - acc_a += a * a; - acc_b += b * b; - acc_ab += a * b; - } - - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - - // Use 1.0 - (cosine similarity) to calc the distance - dist[outidx] = (DataType)1.0 - acc_ab / (raft::sqrt(acc_a) * raft::sqrt(acc_b)); -} - -template -RAFT_KERNEL naiveInnerProductKernel( - DataType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) { return; } - - DataType acc_ab = DataType(0); - - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - acc_ab += a * b; - } - - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc_ab; -} - -template -RAFT_KERNEL naiveHellingerDistanceKernel( - DataType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) { return; } - - DataType acc_ab = DataType(0); - - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - acc_ab += raft::sqrt(a) * raft::sqrt(b); - } - - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - - // Adjust to replace NaN in sqrt with 0 if input to sqrt is negative - acc_ab = 1 - acc_ab; - auto rectifier = (!signbit(acc_ab)); - dist[outidx] = raft::sqrt(rectifier * acc_ab); -} - -template -RAFT_KERNEL naiveLpUnexpDistanceKernel(DataType* dist, - const DataType* x, - const DataType* y, - int m, - int n, - int k, - bool isRowMajor, - DataType p) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) return; - DataType acc = DataType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - auto diff = raft::abs(a - b); - acc += raft::pow(diff, p); - } - auto one_over_p = 1 / p; - acc = raft::pow(acc, one_over_p); - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc; -} - -template -RAFT_KERNEL naiveHammingDistanceKernel( - DataType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) return; - DataType acc = DataType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - acc += (a != b); - } - acc = acc / k; - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc; -} - -template -RAFT_KERNEL naiveJensenShannonDistanceKernel( - DataType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) return; - DataType acc = DataType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - - DataType m = 0.5f * (a + b); - bool a_zero = a == 0; - bool b_zero = b == 0; - - DataType p = (!a_zero * m) / (a_zero + a); - DataType q = (!b_zero * m) / (b_zero + b); - - bool p_zero = p == 0; - bool q_zero = q == 0; - - acc += (-a * (!p_zero * log(p + p_zero))) + (-b * (!q_zero * log(q + q_zero))); - } - acc = raft::sqrt(0.5f * acc); - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc; -} - -template -RAFT_KERNEL naiveRussellRaoDistanceKernel( - OutType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) return; - OutType acc = OutType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - acc += (a * b); - } - acc = (k - acc) / k; - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc; -} - -template -RAFT_KERNEL naiveKLDivergenceDistanceKernel( - OutType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) return; - OutType acc = OutType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - bool b_zero = (b == 0); - bool a_zero = (a == 0); - acc += a * (log(a + a_zero) - log(b + b_zero)); - } - acc = 0.5f * acc; - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc; -} - -template -RAFT_KERNEL naiveCorrelationDistanceKernel( - OutType* dist, const DataType* x, const DataType* y, int m, int n, int k, bool isRowMajor) -{ - int midx = threadIdx.x + blockIdx.x * blockDim.x; - int nidx = threadIdx.y + blockIdx.y * blockDim.y; - if (midx >= m || nidx >= n) return; - OutType acc = OutType(0); - auto a_norm = DataType(0); - auto b_norm = DataType(0); - auto a_sq_norm = DataType(0); - auto b_sq_norm = DataType(0); - for (int i = 0; i < k; ++i) { - int xidx = isRowMajor ? i + midx * k : i * m + midx; - int yidx = isRowMajor ? i + nidx * k : i * n + nidx; - auto a = x[xidx]; - auto b = y[yidx]; - a_norm += a; - b_norm += b; - a_sq_norm += (a * a); - b_sq_norm += (b * b); - acc += (a * b); - } - - auto numer = k * acc - (a_norm * b_norm); - auto Q_denom = k * a_sq_norm - (a_norm * a_norm); - auto R_denom = k * b_sq_norm - (b_norm * b_norm); - - acc = 1 - (numer / raft::sqrt(Q_denom * R_denom)); - - int outidx = isRowMajor ? midx * n + nidx : midx + m * nidx; - dist[outidx] = acc; -} - -template -void naiveDistance(DataType* dist, - const DataType* x, - const DataType* y, - int m, - int n, - int k, - raft::distance::DistanceType type, - bool isRowMajor, - DataType metric_arg = 2.0f, - cudaStream_t stream = 0) -{ - static const dim3 TPB(4, 256, 1); - dim3 nblks(raft::ceildiv(m, (int)TPB.x), raft::ceildiv(n, (int)TPB.y), 1); - - switch (type) { - case raft::distance::DistanceType::Canberra: - case raft::distance::DistanceType::Linf: - case raft::distance::DistanceType::L1: - naiveL1_Linf_CanberraDistanceKernel - <<>>(dist, x, y, m, n, k, type, isRowMajor); - break; - case raft::distance::DistanceType::L2SqrtUnexpanded: - case raft::distance::DistanceType::L2Unexpanded: - case raft::distance::DistanceType::L2SqrtExpanded: - case raft::distance::DistanceType::L2Expanded: - naiveDistanceKernel - <<>>(dist, x, y, m, n, k, type, isRowMajor); - break; - case raft::distance::DistanceType::CosineExpanded: - naiveCosineDistanceKernel - <<>>(dist, x, y, m, n, k, isRowMajor); - break; - case raft::distance::DistanceType::HellingerExpanded: - naiveHellingerDistanceKernel - <<>>(dist, x, y, m, n, k, isRowMajor); - break; - case raft::distance::DistanceType::LpUnexpanded: - naiveLpUnexpDistanceKernel - <<>>(dist, x, y, m, n, k, isRowMajor, metric_arg); - break; - case raft::distance::DistanceType::HammingUnexpanded: - naiveHammingDistanceKernel - <<>>(dist, x, y, m, n, k, isRowMajor); - break; - case raft::distance::DistanceType::InnerProduct: - naiveInnerProductKernel<<>>(dist, x, y, m, n, k, isRowMajor); - break; - case raft::distance::DistanceType::JensenShannon: - naiveJensenShannonDistanceKernel - <<>>(dist, x, y, m, n, k, isRowMajor); - break; - case raft::distance::DistanceType::RusselRaoExpanded: - naiveRussellRaoDistanceKernel - <<>>(dist, x, y, m, n, k, isRowMajor); - break; - case raft::distance::DistanceType::KLDivergence: - naiveKLDivergenceDistanceKernel - <<>>(dist, x, y, m, n, k, isRowMajor); - break; - case raft::distance::DistanceType::CorrelationExpanded: - naiveCorrelationDistanceKernel - <<>>(dist, x, y, m, n, k, isRowMajor); - break; - case raft::distance::DistanceType::DiceExpanded: - naiveDiceDistanceKernel<<>>(dist, x, y, m, n, k, isRowMajor); - break; - default: FAIL() << "should be here\n"; - } - RAFT_CUDA_TRY(cudaPeekAtLastError()); -} - -template -struct DistanceInputs { - DataType tolerance; - int m, n, k; - bool isRowMajor; - unsigned long long int seed; - DataType metric_arg = 2.0f; -}; - -template -::std::ostream& operator<<(::std::ostream& os, const DistanceInputs& dims) -{ - return os; -} - -// TODO: Remove when mdspan-based raft::runtime::distance::pairwise_distance is -// implemented. -// -// Context: -// https://github.com/rapidsai/raft/issues/1338 -template -constexpr bool layout_to_row_major(); - -template <> -constexpr bool layout_to_row_major() -{ - return true; -} -template <> -constexpr bool layout_to_row_major() -{ - return false; -} - -template -void distanceLauncher(raft::resources const& handle, - DataType* x, - DataType* y, - DataType* dist, - DataType* dist2, - int m, - int n, - int k, - DistanceInputs& params, - DataType threshold, - DataType metric_arg = 2.0f) -{ - auto x_v = make_device_matrix_view(x, m, k); - auto y_v = make_device_matrix_view(y, n, k); - auto dist_v = make_device_matrix_view(dist, m, n); - - raft::distance::distance( - handle, x_v, y_v, dist_v, metric_arg); -} - -template -class DistanceTest : public ::testing::TestWithParam> { - public: - DistanceTest() - : params(::testing::TestWithParam>::GetParam()), - stream(resource::get_cuda_stream(handle)), - x(params.m * params.k, stream), - y(params.n * params.k, stream), - dist_ref(params.m * params.n, stream), - dist(params.m * params.n, stream), - dist2(params.m * params.n, stream) - { - } - - void SetUp() override - { - auto testInfo = testing::UnitTest::GetInstance()->current_test_info(); - common::nvtx::range fun_scope("test::%s/%s", testInfo->test_suite_name(), testInfo->name()); - - raft::random::RngState r(params.seed); - int m = params.m; - int n = params.n; - int k = params.k; - DataType metric_arg = params.metric_arg; - bool isRowMajor = params.isRowMajor; - if (distanceType == raft::distance::DistanceType::HellingerExpanded || - distanceType == raft::distance::DistanceType::JensenShannon || - distanceType == raft::distance::DistanceType::KLDivergence) { - // Hellinger works only on positive numbers - uniform(handle, r, x.data(), m * k, DataType(0.0), DataType(1.0)); - uniform(handle, r, y.data(), n * k, DataType(0.0), DataType(1.0)); - } else if (distanceType == raft::distance::DistanceType::RusselRaoExpanded || - distanceType == raft::distance::DistanceType::DiceExpanded) { - uniform(handle, r, x.data(), m * k, DataType(0.0), DataType(1.0)); - uniform(handle, r, y.data(), n * k, DataType(0.0), DataType(1.0)); - // Russel rao works on boolean values. - bernoulli(handle, r, x.data(), m * k, 0.5f); - bernoulli(handle, r, y.data(), n * k, 0.5f); - } else { - uniform(handle, r, x.data(), m * k, DataType(-1.0), DataType(1.0)); - uniform(handle, r, y.data(), n * k, DataType(-1.0), DataType(1.0)); - } - naiveDistance( - dist_ref.data(), x.data(), y.data(), m, n, k, distanceType, isRowMajor, metric_arg, stream); - - DataType threshold = -10000.f; - - if (isRowMajor) { - distanceLauncher(handle, - x.data(), - y.data(), - dist.data(), - dist2.data(), - m, - n, - k, - params, - threshold, - metric_arg); - - } else { - distanceLauncher(handle, - x.data(), - y.data(), - dist.data(), - dist2.data(), - m, - n, - k, - params, - threshold, - metric_arg); - } - resource::sync_stream(handle, stream); - } - - protected: - raft::resources handle; - cudaStream_t stream; - - DistanceInputs params; - rmm::device_uvector x, y, dist_ref, dist, dist2; -}; - -/* - * This test suite verifies the path when X and Y are same buffer, - * distance metrics which requires norms like L2 expanded/cosine/correlation - * takes a more optimal path in such case to skip norm calculation for Y buffer. - * It may happen that though both X and Y are same buffer but user passes - * different dimensions for them like in case of tiled_brute_force_knn. - */ -template -class DistanceTestSameBuffer : public ::testing::TestWithParam> { - public: - using dev_vector = rmm::device_uvector; - DistanceTestSameBuffer() - : params(::testing::TestWithParam>::GetParam()), - stream(resource::get_cuda_stream(handle)), - x(params.m * params.k, stream), - dist_ref({dev_vector(params.m * params.m, stream), dev_vector(params.m * params.m, stream)}), - dist({dev_vector(params.m * params.m, stream), dev_vector(params.m * params.m, stream)}), - dist2({dev_vector(params.m * params.m, stream), dev_vector(params.m * params.m, stream)}) - { - } - - void SetUp() override - { - auto testInfo = testing::UnitTest::GetInstance()->current_test_info(); - common::nvtx::range fun_scope("test::%s/%s", testInfo->test_suite_name(), testInfo->name()); - - raft::random::RngState r(params.seed); - int m = params.m; - int n = params.m; - int k = params.k; - DataType metric_arg = params.metric_arg; - bool isRowMajor = params.isRowMajor; - if (distanceType == raft::distance::DistanceType::HellingerExpanded || - distanceType == raft::distance::DistanceType::JensenShannon || - distanceType == raft::distance::DistanceType::KLDivergence) { - // Hellinger works only on positive numbers - uniform(handle, r, x.data(), m * k, DataType(0.0), DataType(1.0)); - } else if (distanceType == raft::distance::DistanceType::RusselRaoExpanded || - distanceType == raft::distance::DistanceType::DiceExpanded) { - uniform(handle, r, x.data(), m * k, DataType(0.0), DataType(1.0)); - // Russel rao works on boolean values. - bernoulli(handle, r, x.data(), m * k, 0.5f); - } else { - uniform(handle, r, x.data(), m * k, DataType(-1.0), DataType(1.0)); - } - - for (int i = 0; i < 2; i++) { - // both X and Y are same buffer but when i = 1 - // different dimensions for x & y is passed. - m = m / (i + 1); - naiveDistance(dist_ref[i].data(), - x.data(), - x.data(), - m, - n, - k, - distanceType, - isRowMajor, - metric_arg, - stream); - - DataType threshold = -10000.f; - - if (isRowMajor) { - distanceLauncher(handle, - x.data(), - x.data(), - dist[i].data(), - dist2[i].data(), - m, - n, - k, - params, - threshold, - metric_arg); - - } else { - distanceLauncher(handle, - x.data(), - x.data(), - dist[i].data(), - dist2[i].data(), - m, - n, - k, - params, - threshold, - metric_arg); - } - } - resource::sync_stream(handle, stream); - } - - protected: - raft::resources handle; - cudaStream_t stream; - - DistanceInputs params; - dev_vector x; - static const int N = 2; - std::array dist_ref, dist, dist2; -}; - -template -class BigMatrixDistanceTest : public ::testing::Test { - public: - BigMatrixDistanceTest() - : x(m * k, resource::get_cuda_stream(handle)), - dist(std::size_t(m) * m, resource::get_cuda_stream(handle)){}; - void SetUp() override - { - auto testInfo = testing::UnitTest::GetInstance()->current_test_info(); - common::nvtx::range fun_scope("test::%s/%s", testInfo->test_suite_name(), testInfo->name()); - - void pairwise_distance(raft::resources const& handle, - float* x, - float* y, - float* dists, - int m, - int n, - int k, - raft::distance::DistanceType metric, - bool isRowMajor, - float metric_arg); - constexpr bool row_major = true; - constexpr float metric_arg = 0.0f; - raft::distance::distance( - handle, x.data(), x.data(), dist.data(), m, n, k, row_major, metric_arg); - RAFT_CUDA_TRY(cudaStreamSynchronize(resource::get_cuda_stream(handle))); - } - - protected: - raft::resources handle; - int m = 48000; - int n = 48000; - int k = 1; - rmm::device_uvector x, dist; -}; -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/fused_cosine_nn.cu b/cpp/test/distance/fused_cosine_nn.cu deleted file mode 100644 index d4d632e1dc..0000000000 --- a/cpp/test/distance/fused_cosine_nn.cu +++ /dev/null @@ -1,420 +0,0 @@ -/* - * Copyright (c) 2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#undef RAFT_EXPLICIT_INSTANTIATE_ONLY // Search with filter instantiation - -#include "../test_utils.cuh" - -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -namespace raft { -namespace distance { - -template -struct RaftKVPMinReduce { - typedef raft::KeyValuePair KVP; - - DI KVP operator()(LabelT rit, const KVP& a, const KVP& b) { return b.value < a.value ? b : a; } - - DI KVP operator()(const KVP& a, const KVP& b) { return b.value < a.value ? b : a; } - -}; // KVPMinReduce - -template -__global__ void naiveCosKernel(raft::KeyValuePair* min, - DataT* x, - DataT* y, - int m, - int n, - int k, - int* workspace, - DataT maxVal) -{ - int midx = threadIdx.y + blockIdx.y * blockDim.y; - int nidx = threadIdx.x + blockIdx.x * blockDim.x; - DataT acc_a = DataT(0); - DataT acc_b = DataT(0); - DataT acc_ab = DataT(0); - // if (midx >= m || nidx >= n) { return; } - - for (int i = 0; i < k; ++i) { - int xidx = i + midx * k; - int yidx = i + nidx * k; - auto a = x[xidx]; - auto b = y[yidx]; - acc_a += a * a; - acc_b += b * b; - acc_ab += a * b; - } - - // Use 1.0 - (cosine similarity) to calc the distance - DataT acc = maxVal; - if (midx < m || nidx < n) { acc = (DataT)1.0 - acc_ab / (raft::sqrt(acc_a) * raft::sqrt(acc_b)); } - - ReduceOpT redOp; - typedef cub::WarpReduce> WarpReduce; - __shared__ typename WarpReduce::TempStorage temp[NWARPS]; - int warpId = threadIdx.x / raft::WarpSize; - raft::KeyValuePair tmp; - tmp.key = nidx; - tmp.value = midx >= m || nidx >= n ? maxVal : acc; - tmp = WarpReduce(temp[warpId]).Reduce(tmp, RaftKVPMinReduce()); - if (threadIdx.x % raft::WarpSize == 0 && midx < m) { - while (atomicCAS(workspace + midx, 0, 1) == 1) - ; - __threadfence(); - redOp(midx, min + midx, tmp); - __threadfence(); - atomicCAS(workspace + midx, 1, 0); - } -} - -template -void naive(raft::KeyValuePair* min, - DataT* x, - DataT* y, - int m, - int n, - int k, - int* workspace, - cudaStream_t stream) -{ - static const dim3 TPB(32, 16, 1); - dim3 nblks(raft::ceildiv(n, (int)TPB.x), raft::ceildiv(m, (int)TPB.y), 1); - RAFT_CUDA_TRY(cudaMemsetAsync(workspace, 0, sizeof(int) * m, stream)); - auto blks = raft::ceildiv(m, 256); - MinAndDistanceReduceOp op; - detail::initKernel, int> - <<>>(min, m, std::numeric_limits::max(), op); - RAFT_CUDA_TRY(cudaGetLastError()); - naiveCosKernel, 16> - <<>>(min, x, y, m, n, k, workspace, std::numeric_limits::max()); - RAFT_CUDA_TRY(cudaGetLastError()); -} - -template -struct Inputs { - DataT tolerance; - int m, n, k; - unsigned long long int seed; - - friend std::ostream& operator<<(std::ostream& os, const Inputs& p) - { - return os << "m: " << p.m - << ", " - "n: " - << p.n - << ", " - "k: " - << p.k - << ", " - "seed: " - << p.seed - << ", " - "tol: " - << p.tolerance; - } -}; - -template -class FusedCosineNNTest : public ::testing::TestWithParam> { - public: - FusedCosineNNTest() - : params(::testing::TestWithParam>::GetParam()), - stream(resource::get_cuda_stream(handle)), - x(params.m * params.k, stream), - y(params.n * params.k, stream), - xn(params.m, stream), - yn(params.n, stream), - min(params.m, stream), - min_ref(params.m, stream), - workspace(params.m * sizeof(int), stream) - { - } - - protected: - void SetUp() override - { - raft::random::RngState r(params.seed); - int m = params.m; - int n = params.n; - int k = params.k; - uniform(handle, r, x.data(), m * k, DataT(-1.0), DataT(1.0)); - uniform(handle, r, y.data(), n * k, DataT(-1.0), DataT(1.0)); - generateGoldenResult(); - raft::linalg::rowNorm( - xn.data(), x.data(), k, m, raft::linalg::L2Norm, true, stream, raft::sqrt_op{}); - raft::linalg::rowNorm( - yn.data(), y.data(), k, n, raft::linalg::L2Norm, true, stream, raft::sqrt_op{}); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - } - - protected: - raft::resources handle; - cudaStream_t stream; - Inputs params; - rmm::device_uvector x; - rmm::device_uvector y; - rmm::device_uvector xn; - rmm::device_uvector yn; - rmm::device_uvector> min; - rmm::device_uvector> min_ref; - rmm::device_uvector workspace; - - virtual void generateGoldenResult() - { - int m = params.m; - int n = params.n; - int k = params.k; - naive(min_ref.data(), x.data(), y.data(), m, n, k, (int*)workspace.data(), stream); - } - - void runTest(raft::KeyValuePair* out) - { - int m = params.m; - int n = params.n; - int k = params.k; - raft::distance::DistanceType metric = raft::distance::DistanceType::CosineExpanded; - constexpr bool init_out_buffer = true; - fusedDistanceNNMinReduce, int>(out, - x.data(), - y.data(), - xn.data(), - yn.data(), - m, - n, - k, - (void*)workspace.data(), - false, - init_out_buffer, - true, - metric, - 0.0f, - stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - } -}; - -template -struct CompareApproxAbsKVP { - typedef typename raft::KeyValuePair KVP; - CompareApproxAbsKVP(T eps_) : eps(eps_) {} - bool operator()(const KVP& a, const KVP& b) const - { - T diff = std::abs(std::abs(a.value) - std::abs(b.value)); - T m = std::max(std::abs(a.value), std::abs(b.value)); - T ratio = m >= eps ? diff / m : diff; - return (ratio <= eps); - } - - private: - T eps; -}; - -template -struct CompareExactKVP { - typedef typename raft::KeyValuePair KVP; - bool operator()(const KVP& a, const KVP& b) const - { - if (a.value != b.value) return false; - return true; - } -}; - -template -::testing::AssertionResult devArrMatch(const raft::KeyValuePair* expected, - const raft::KeyValuePair* actual, - size_t size, - L eq_compare, - cudaStream_t stream = 0) -{ - typedef typename raft::KeyValuePair KVP; - std::shared_ptr exp_h(new KVP[size]); - std::shared_ptr act_h(new KVP[size]); - raft::update_host(exp_h.get(), expected, size, stream); - raft::update_host(act_h.get(), actual, size, stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - for (size_t i(0); i < size; ++i) { - auto exp = exp_h.get()[i]; - auto act = act_h.get()[i]; - if (!eq_compare(exp, act)) { - return ::testing::AssertionFailure() - << "actual=" << act.key << "," << act.value << " != expected=" << exp.key << "," - << exp.value << " @" << i; - } - } - return ::testing::AssertionSuccess(); -} - -const std::vector> inputsf = { - {0.001f, 32, 32, 32, 1234ULL}, - {0.001f, 32, 64, 32, 1234ULL}, - {0.001f, 64, 32, 32, 1234ULL}, - {0.001f, 64, 64, 32, 1234ULL}, - {0.001f, 128, 32, 32, 1234ULL}, - {0.001f, 128, 64, 32, 1234ULL}, - {0.001f, 128, 128, 64, 1234ULL}, - {0.001f, 64, 128, 128, 1234ULL}, - - {0.001f, 32, 32, 34, 1234ULL}, - {0.001f, 32, 64, 34, 1234ULL}, - {0.001f, 64, 32, 34, 1234ULL}, - {0.001f, 64, 64, 34, 1234ULL}, - {0.001f, 128, 32, 34, 1234ULL}, - {0.001f, 128, 64, 34, 1234ULL}, - {0.001f, 128, 128, 66, 1234ULL}, - {0.001f, 64, 128, 130, 1234ULL}, - - {0.001f, 32, 32, 33, 1234ULL}, - {0.001f, 32, 64, 33, 1234ULL}, - {0.001f, 64, 32, 33, 1234ULL}, - {0.001f, 64, 64, 33, 1234ULL}, - {0.001f, 128, 32, 33, 1234ULL}, - {0.001f, 128, 64, 33, 1234ULL}, - {0.001f, 128, 128, 65, 1234ULL}, - {0.001f, 64, 128, 129, 1234ULL}, - {0.006f, 1805, 134, 2, 1234ULL}, - {0.006f, 8192, 1024, 64, 1234ULL}, - {0.006f, 8192, 1025, 64, 1234ULL}, - - // Repeat with smaller values of k - {0.006f, 32, 32, 1, 1234ULL}, - {0.001f, 32, 64, 2, 1234ULL}, - {0.001f, 64, 32, 3, 1234ULL}, - {0.001f, 64, 64, 4, 1234ULL}, - {0.001f, 128, 32, 5, 1234ULL}, - {0.001f, 128, 64, 6, 1234ULL}, - {0.001f, 128, 128, 7, 1234ULL}, - {0.001f, 64, 128, 8, 1234ULL}, - - {0.001f, 32, 32, 9, 1234ULL}, - {0.001f, 32, 64, 10, 1234ULL}, - {0.001f, 64, 32, 11, 1234ULL}, - {0.001f, 64, 64, 12, 1234ULL}, - {0.001f, 128, 32, 13, 1234ULL}, - {0.001f, 128, 64, 14, 1234ULL}, - {0.001f, 128, 128, 15, 1234ULL}, - {0.001f, 64, 128, 16, 1234ULL}, - - {0.001f, 32, 32, 17, 1234ULL}, - {0.001f, 32, 64, 18, 1234ULL}, - {0.001f, 64, 32, 19, 1234ULL}, - {0.001f, 64, 64, 20, 1234ULL}, - {0.001f, 128, 32, 21, 1234ULL}, - {0.001f, 128, 64, 22, 1234ULL}, - {0.001f, 128, 128, 23, 1234ULL}, - {0.00001, 64, 128, 24, 1234ULL}, - {0.001f, 1805, 134, 25, 1234ULL}, - {0.006f, 8192, 1024, 25, 1234ULL}, - {0.006f, 8192, 1024, 66, 1234ULL}, -}; -typedef FusedCosineNNTest FusedCosineNNTestF; -TEST_P(FusedCosineNNTestF, Result) -{ - runTest(min.data()); - ASSERT_TRUE(devArrMatch( - min_ref.data(), min.data(), params.m, CompareApproxAbsKVP(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(FusedCosineNNTests, FusedCosineNNTestF, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.00001, 32, 32, 32, 1234ULL}, {0.00001, 32, 64, 32, 1234ULL}, - {0.00001, 64, 32, 32, 1234ULL}, {0.00001, 64, 64, 32, 1234ULL}, - {0.00001, 128, 32, 32, 1234ULL}, {0.00001, 128, 64, 32, 1234ULL}, - {0.00001, 128, 128, 64, 1234ULL}, {0.00001, 64, 128, 128, 1234ULL}, - - {0.00001, 32, 32, 34, 1234ULL}, {0.00001, 32, 64, 34, 1234ULL}, - {0.00001, 64, 32, 34, 1234ULL}, {0.00001, 64, 64, 34, 1234ULL}, - {0.00001, 128, 32, 34, 1234ULL}, {0.00001, 128, 64, 34, 1234ULL}, - {0.00001, 128, 128, 66, 1234ULL}, {0.00001, 64, 128, 130, 1234ULL}, - - {0.00001, 32, 32, 33, 1234ULL}, {0.00001, 32, 64, 33, 1234ULL}, - {0.00001, 64, 32, 33, 1234ULL}, {0.00001, 64, 64, 33, 1234ULL}, - {0.00001, 128, 32, 33, 1234ULL}, {0.00001, 128, 64, 33, 1234ULL}, - {0.00001, 128, 128, 65, 1234ULL}, {0.00001, 64, 128, 129, 1234ULL}, - - {0.00001, 1805, 134, 2, 1234ULL}, {0.00001, 8192, 1024, 25, 1234ULL}, -}; -typedef FusedCosineNNTest FusedCosineNNTestD; -TEST_P(FusedCosineNNTestD, Result) -{ - runTest(min.data()); - ASSERT_TRUE(devArrMatch( - min_ref.data(), min.data(), params.m, CompareApproxAbsKVP(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(FusedCosineNNTests, FusedCosineNNTestD, ::testing::ValuesIn(inputsd)); - -/// This is to test output determinism of the prim -template -class FusedCosineNNDetTest : public FusedCosineNNTest { - public: - FusedCosineNNDetTest() : stream(resource::get_cuda_stream(handle)), min1(0, stream) {} - - void SetUp() override - { - FusedCosineNNTest::SetUp(); - int m = this->params.m; - min1.resize(m, stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - } - - void TearDown() override { FusedCosineNNTest::TearDown(); } - - protected: - raft::resources handle; - cudaStream_t stream; - - rmm::device_uvector> min1; - - static const int NumRepeats = 3; - - void generateGoldenResult() override {} -}; - -typedef FusedCosineNNDetTest FusedCosineNNDetTestF; -TEST_P(FusedCosineNNDetTestF, Result) -{ - runTest(min.data()); // assumed to be golden - for (int i = 0; i < NumRepeats; ++i) { - runTest(min1.data()); - ASSERT_TRUE(devArrMatch(min.data(), min1.data(), params.m, CompareExactKVP(), stream)); - cudaMemsetAsync(min1.data(), 0, sizeof(*min.data()) * params.m, stream); - } -} -INSTANTIATE_TEST_CASE_P(FusedCosineNNDetTests, FusedCosineNNDetTestF, ::testing::ValuesIn(inputsf)); - -typedef FusedCosineNNDetTest FusedCosineNNDetTestD; -TEST_P(FusedCosineNNDetTestD, Result) -{ - runTest(min.data()); // assumed to be golden - for (int i = 0; i < NumRepeats; ++i) { - runTest(min1.data()); - ASSERT_TRUE(devArrMatch(min.data(), min1.data(), params.m, CompareExactKVP(), stream)); - } -} -INSTANTIATE_TEST_CASE_P(FusedCosineNNDetTests, FusedCosineNNDetTestD, ::testing::ValuesIn(inputsd)); - -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/fused_l2_nn.cu b/cpp/test/distance/fused_l2_nn.cu deleted file mode 100644 index 6fd8f15808..0000000000 --- a/cpp/test/distance/fused_l2_nn.cu +++ /dev/null @@ -1,437 +0,0 @@ -/* - * Copyright (c) 2021-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" - -#include -#include -#include -#include -#include -#include -#include - -#include - -namespace raft { -namespace distance { - -template -struct RaftKVPMinReduce { - typedef raft::KeyValuePair KVP; - - DI KVP operator()(LabelT rit, const KVP& a, const KVP& b) { return b.value < a.value ? b : a; } - - DI KVP operator()(const KVP& a, const KVP& b) { return b.value < a.value ? b : a; } - -}; // KVPMinReduce - -template -RAFT_KERNEL naiveKernel(raft::KeyValuePair* min, - DataT* x, - DataT* y, - int m, - int n, - int k, - int* workspace, - DataT maxVal) -{ - int midx = threadIdx.y + blockIdx.y * blockDim.y; - int nidx = threadIdx.x + blockIdx.x * blockDim.x; - DataT acc = DataT(0); - for (int i = 0; i < k; ++i) { - int xidx = i + midx * k; - int yidx = i + nidx * k; - auto diff = midx >= m || nidx >= n ? DataT(0) : x[xidx] - y[yidx]; - acc += diff * diff; - } - - if (Sqrt) { acc = raft::sqrt(acc); } - ReduceOpT redOp; - typedef cub::WarpReduce> WarpReduce; - __shared__ typename WarpReduce::TempStorage temp[NWARPS]; - int warpId = threadIdx.x / raft::WarpSize; - raft::KeyValuePair tmp; - tmp.key = nidx; - tmp.value = midx >= m || nidx >= n ? maxVal : acc; - tmp = WarpReduce(temp[warpId]).Reduce(tmp, RaftKVPMinReduce()); - if (threadIdx.x % raft::WarpSize == 0 && midx < m) { - while (atomicCAS(workspace + midx, 0, 1) == 1) - ; - __threadfence(); - redOp(midx, min + midx, tmp); - __threadfence(); - atomicCAS(workspace + midx, 1, 0); - } -} - -template -void naive(raft::KeyValuePair* min, - DataT* x, - DataT* y, - int m, - int n, - int k, - int* workspace, - cudaStream_t stream) -{ - static const dim3 TPB(32, 16, 1); - dim3 nblks(raft::ceildiv(n, (int)TPB.x), raft::ceildiv(m, (int)TPB.y), 1); - RAFT_CUDA_TRY(cudaMemsetAsync(workspace, 0, sizeof(int) * m, stream)); - auto blks = raft::ceildiv(m, 256); - MinAndDistanceReduceOp op; - detail::initKernel, int> - <<>>(min, m, std::numeric_limits::max(), op); - RAFT_CUDA_TRY(cudaGetLastError()); - naiveKernel, 16> - <<>>(min, x, y, m, n, k, workspace, std::numeric_limits::max()); - RAFT_CUDA_TRY(cudaGetLastError()); -} - -template -struct Inputs { - DataT tolerance; - int m, n, k; - unsigned long long int seed; - - friend std::ostream& operator<<(std::ostream& os, const Inputs& p) - { - return os << "m: " << p.m - << ", " - "n: " - << p.n - << ", " - "k: " - << p.k - << ", " - "seed: " - << p.seed - << ", " - "tol: " - << p.tolerance; - } -}; - -template -class FusedL2NNTest : public ::testing::TestWithParam> { - public: - FusedL2NNTest() - : params(::testing::TestWithParam>::GetParam()), - stream(resource::get_cuda_stream(handle)), - x(params.m * params.k, stream), - y(params.n * params.k, stream), - xn(params.m, stream), - yn(params.n, stream), - min(params.m, stream), - min_ref(params.m, stream), - workspace(params.m * sizeof(int), stream) - { - } - - protected: - void SetUp() override - { - raft::random::RngState r(params.seed); - int m = params.m; - int n = params.n; - int k = params.k; - uniform(handle, r, x.data(), m * k, DataT(-1.0), DataT(1.0)); - uniform(handle, r, y.data(), n * k, DataT(-1.0), DataT(1.0)); - generateGoldenResult(); - raft::linalg::rowNorm(xn.data(), x.data(), k, m, raft::linalg::L2Norm, true, stream); - raft::linalg::rowNorm(yn.data(), y.data(), k, n, raft::linalg::L2Norm, true, stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - } - - protected: - raft::resources handle; - cudaStream_t stream; - Inputs params; - rmm::device_uvector x; - rmm::device_uvector y; - rmm::device_uvector xn; - rmm::device_uvector yn; - rmm::device_uvector> min; - rmm::device_uvector> min_ref; - rmm::device_uvector workspace; - - virtual void generateGoldenResult() - { - int m = params.m; - int n = params.n; - int k = params.k; - naive(min_ref.data(), x.data(), y.data(), m, n, k, (int*)workspace.data(), stream); - } - - void runTest(raft::KeyValuePair* out) - { - int m = params.m; - int n = params.n; - int k = params.k; - - const bool init_out_buffer = true; - fusedL2NNMinReduce, int>(out, - x.data(), - y.data(), - xn.data(), - yn.data(), - m, - n, - k, - (void*)workspace.data(), - Sqrt, - init_out_buffer, - stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - } -}; - -template -struct CompareApproxAbsKVP { - typedef typename raft::KeyValuePair KVP; - CompareApproxAbsKVP(T eps_) : eps(eps_) {} - bool operator()(const KVP& a, const KVP& b) const - { - T diff = std::abs(std::abs(a.value) - std::abs(b.value)); - T m = std::max(std::abs(a.value), std::abs(b.value)); - T ratio = m >= eps ? diff / m : diff; - return (ratio <= eps); - } - - private: - T eps; -}; - -template -struct CompareExactKVP { - typedef typename raft::KeyValuePair KVP; - bool operator()(const KVP& a, const KVP& b) const - { - if (a.value != b.value) return false; - return true; - } -}; - -template -::testing::AssertionResult devArrMatch(const raft::KeyValuePair* expected, - const raft::KeyValuePair* actual, - size_t size, - L eq_compare, - cudaStream_t stream = 0) -{ - typedef typename raft::KeyValuePair KVP; - std::shared_ptr exp_h(new KVP[size]); - std::shared_ptr act_h(new KVP[size]); - raft::update_host(exp_h.get(), expected, size, stream); - raft::update_host(act_h.get(), actual, size, stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - for (size_t i(0); i < size; ++i) { - auto exp = exp_h.get()[i]; - auto act = act_h.get()[i]; - if (!eq_compare(exp, act)) { - return ::testing::AssertionFailure() - << "actual=" << act.key << "," << act.value << " != expected=" << exp.key << "," - << exp.value << " @" << i; - } - } - return ::testing::AssertionSuccess(); -} - -const std::vector> inputsf = { - {0.001f, 32, 32, 32, 1234ULL}, - {0.001f, 32, 64, 32, 1234ULL}, - {0.001f, 64, 32, 32, 1234ULL}, - {0.001f, 64, 64, 32, 1234ULL}, - {0.001f, 128, 32, 32, 1234ULL}, - {0.001f, 128, 64, 32, 1234ULL}, - {0.001f, 128, 128, 64, 1234ULL}, - {0.001f, 64, 128, 128, 1234ULL}, - - {0.001f, 32, 32, 34, 1234ULL}, - {0.001f, 32, 64, 34, 1234ULL}, - {0.001f, 64, 32, 34, 1234ULL}, - {0.001f, 64, 64, 34, 1234ULL}, - {0.001f, 128, 32, 34, 1234ULL}, - {0.001f, 128, 64, 34, 1234ULL}, - {0.001f, 128, 128, 66, 1234ULL}, - {0.001f, 64, 128, 130, 1234ULL}, - - {0.001f, 32, 32, 33, 1234ULL}, - {0.001f, 32, 64, 33, 1234ULL}, - {0.001f, 64, 32, 33, 1234ULL}, - {0.001f, 64, 64, 33, 1234ULL}, - {0.001f, 128, 32, 33, 1234ULL}, - {0.001f, 128, 64, 33, 1234ULL}, - {0.001f, 128, 128, 65, 1234ULL}, - {0.001f, 64, 128, 129, 1234ULL}, - {0.006f, 1805, 134, 2, 1234ULL}, - {0.006f, 8192, 1024, 64, 1234ULL}, - {0.006f, 8192, 1025, 64, 1234ULL}, - - // Repeat with smaller values of k - {0.006f, 32, 32, 1, 1234ULL}, - {0.001f, 32, 64, 2, 1234ULL}, - {0.001f, 64, 32, 3, 1234ULL}, - {0.001f, 64, 64, 4, 1234ULL}, - {0.001f, 128, 32, 5, 1234ULL}, - {0.001f, 128, 64, 6, 1234ULL}, - {0.001f, 128, 128, 7, 1234ULL}, - {0.001f, 64, 128, 8, 1234ULL}, - - {0.001f, 32, 32, 9, 1234ULL}, - {0.001f, 32, 64, 10, 1234ULL}, - {0.001f, 64, 32, 11, 1234ULL}, - {0.001f, 64, 64, 12, 1234ULL}, - {0.001f, 128, 32, 13, 1234ULL}, - {0.001f, 128, 64, 14, 1234ULL}, - {0.001f, 128, 128, 15, 1234ULL}, - {0.001f, 64, 128, 16, 1234ULL}, - - {0.001f, 32, 32, 17, 1234ULL}, - {0.001f, 32, 64, 18, 1234ULL}, - {0.001f, 64, 32, 19, 1234ULL}, - {0.001f, 64, 64, 20, 1234ULL}, - {0.001f, 128, 32, 21, 1234ULL}, - {0.001f, 128, 64, 22, 1234ULL}, - {0.001f, 128, 128, 23, 1234ULL}, - {0.00001, 64, 128, 24, 1234ULL}, - {0.001f, 1805, 134, 25, 1234ULL}, - {0.006f, 8192, 1024, 25, 1234ULL}, - {0.006f, 8192, 1024, 66, 1234ULL}, -}; -typedef FusedL2NNTest FusedL2NNTestF_Sq; -TEST_P(FusedL2NNTestF_Sq, Result) -{ - runTest(min.data()); - ASSERT_TRUE(devArrMatch( - min_ref.data(), min.data(), params.m, CompareApproxAbsKVP(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(FusedL2NNTests, FusedL2NNTestF_Sq, ::testing::ValuesIn(inputsf)); -typedef FusedL2NNTest FusedL2NNTestF_Sqrt; -TEST_P(FusedL2NNTestF_Sqrt, Result) -{ - runTest(min.data()); - ASSERT_TRUE(devArrMatch( - min_ref.data(), min.data(), params.m, CompareApproxAbsKVP(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(FusedL2NNTests, FusedL2NNTestF_Sqrt, ::testing::ValuesIn(inputsf)); - -const std::vector> inputsd = { - {0.00001, 32, 32, 32, 1234ULL}, {0.00001, 32, 64, 32, 1234ULL}, - {0.00001, 64, 32, 32, 1234ULL}, {0.00001, 64, 64, 32, 1234ULL}, - {0.00001, 128, 32, 32, 1234ULL}, {0.00001, 128, 64, 32, 1234ULL}, - {0.00001, 128, 128, 64, 1234ULL}, {0.00001, 64, 128, 128, 1234ULL}, - - {0.00001, 32, 32, 34, 1234ULL}, {0.00001, 32, 64, 34, 1234ULL}, - {0.00001, 64, 32, 34, 1234ULL}, {0.00001, 64, 64, 34, 1234ULL}, - {0.00001, 128, 32, 34, 1234ULL}, {0.00001, 128, 64, 34, 1234ULL}, - {0.00001, 128, 128, 66, 1234ULL}, {0.00001, 64, 128, 130, 1234ULL}, - - {0.00001, 32, 32, 33, 1234ULL}, {0.00001, 32, 64, 33, 1234ULL}, - {0.00001, 64, 32, 33, 1234ULL}, {0.00001, 64, 64, 33, 1234ULL}, - {0.00001, 128, 32, 33, 1234ULL}, {0.00001, 128, 64, 33, 1234ULL}, - {0.00001, 128, 128, 65, 1234ULL}, {0.00001, 64, 128, 129, 1234ULL}, - - {0.00001, 1805, 134, 2, 1234ULL}, //{0.00001, 8192, 1024, 25, 1234ULL}, -}; -typedef FusedL2NNTest FusedL2NNTestD_Sq; -TEST_P(FusedL2NNTestD_Sq, Result) -{ - runTest(min.data()); - ASSERT_TRUE(devArrMatch( - min_ref.data(), min.data(), params.m, CompareApproxAbsKVP(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(FusedL2NNTests, FusedL2NNTestD_Sq, ::testing::ValuesIn(inputsd)); -typedef FusedL2NNTest FusedL2NNTestD_Sqrt; -TEST_P(FusedL2NNTestD_Sqrt, Result) -{ - runTest(min.data()); - ASSERT_TRUE(devArrMatch( - min_ref.data(), min.data(), params.m, CompareApproxAbsKVP(params.tolerance), stream)); -} -INSTANTIATE_TEST_CASE_P(FusedL2NNTests, FusedL2NNTestD_Sqrt, ::testing::ValuesIn(inputsd)); - -/// This is to test output determinism of the prim -template -class FusedL2NNDetTest : public FusedL2NNTest { - public: - FusedL2NNDetTest() : stream(resource::get_cuda_stream(handle)), min1(0, stream) {} - - void SetUp() override - { - FusedL2NNTest::SetUp(); - int m = this->params.m; - min1.resize(m, stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - } - - void TearDown() override { FusedL2NNTest::TearDown(); } - - protected: - raft::resources handle; - cudaStream_t stream; - - rmm::device_uvector> min1; - - static const int NumRepeats = 3; - - void generateGoldenResult() override {} -}; - -typedef FusedL2NNDetTest FusedL2NNDetTestF_Sq; -TEST_P(FusedL2NNDetTestF_Sq, Result) -{ - runTest(min.data()); // assumed to be golden - for (int i = 0; i < NumRepeats; ++i) { - runTest(min1.data()); - ASSERT_TRUE(devArrMatch(min.data(), min1.data(), params.m, CompareExactKVP(), stream)); - } -} -INSTANTIATE_TEST_CASE_P(FusedL2NNDetTests, FusedL2NNDetTestF_Sq, ::testing::ValuesIn(inputsf)); -typedef FusedL2NNDetTest FusedL2NNDetTestF_Sqrt; -TEST_P(FusedL2NNDetTestF_Sqrt, Result) -{ - runTest(min.data()); // assumed to be golden - for (int i = 0; i < NumRepeats; ++i) { - runTest(min1.data()); - ASSERT_TRUE(devArrMatch(min.data(), min1.data(), params.m, CompareExactKVP(), stream)); - } -} -INSTANTIATE_TEST_CASE_P(FusedL2NNDetTests, FusedL2NNDetTestF_Sqrt, ::testing::ValuesIn(inputsf)); - -typedef FusedL2NNDetTest FusedL2NNDetTestD_Sq; -TEST_P(FusedL2NNDetTestD_Sq, Result) -{ - runTest(min.data()); // assumed to be golden - for (int i = 0; i < NumRepeats; ++i) { - runTest(min1.data()); - ASSERT_TRUE(devArrMatch(min.data(), min1.data(), params.m, CompareExactKVP(), stream)); - } -} -INSTANTIATE_TEST_CASE_P(FusedL2NNDetTests, FusedL2NNDetTestD_Sq, ::testing::ValuesIn(inputsd)); -typedef FusedL2NNDetTest FusedL2NNDetTestD_Sqrt; -TEST_P(FusedL2NNDetTestD_Sqrt, Result) -{ - runTest(min.data()); // assumed to be golden - for (int i = 0; i < NumRepeats; ++i) { - runTest(min1.data()); - ASSERT_TRUE(devArrMatch(min.data(), min1.data(), params.m, CompareExactKVP(), stream)); - } -} -INSTANTIATE_TEST_CASE_P(FusedL2NNDetTests, FusedL2NNDetTestD_Sqrt, ::testing::ValuesIn(inputsd)); - -} // end namespace distance -} // end namespace raft diff --git a/cpp/test/distance/gram.cu b/cpp/test/distance/gram.cu deleted file mode 100644 index e911a25ff1..0000000000 --- a/cpp/test/distance/gram.cu +++ /dev/null @@ -1,174 +0,0 @@ -/* - * Copyright (c) 2019-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "gram_base.cuh" - -#include -#include -#include -#include -#include - -#include - -#include - -#include -#include - -namespace raft::distance::kernels { - -struct GramMatrixInputs { - int n1; // feature vectors in matrix 1 - int n2; // featuer vectors in matrix 2 - int n_cols; // number of elements in a feature vector - bool is_row_major; - KernelParams kernel; - int ld1; - int ld2; - int ld_out; - // We will generate random input using the dimensions given here. - // The reference output is calculated by a custom kernel. -}; - -std::ostream& operator<<(std::ostream& os, const GramMatrixInputs& p) -{ - std::vector kernel_names{"linear", "poly", "rbf", "tanh"}; - os << "/" << p.n1 << "x" << p.n2 << "x" << p.n_cols << "/" - << (p.is_row_major ? "RowMajor/" : "ColMajor/") << kernel_names[p.kernel.kernel] << "/ld_" - << p.ld1 << "x" << p.ld2 << "x" << p.ld_out; - return os; -} - -const std::vector inputs = { - {42, 137, 2, false, {KernelType::LINEAR}}, - {42, 137, 2, true, {KernelType::LINEAR}}, - {42, 137, 2, false, {KernelType::LINEAR}, 64, 179, 181}, - {42, 137, 2, true, {KernelType::LINEAR}, 64, 179, 181}, - {137, 42, 2, false, {KernelType::POLYNOMIAL, 2, 0.5, 2.4}}, - {137, 42, 2, true, {KernelType::POLYNOMIAL, 2, 0.5, 2.4}}, - {137, 42, 2, false, {KernelType::POLYNOMIAL, 2, 0.5, 2.4}, 159, 73, 144}, - {137, 42, 2, true, {KernelType::POLYNOMIAL, 2, 0.5, 2.4}, 159, 73, 144}, - {42, 137, 2, false, {KernelType::TANH, 0, 0.5, 2.4}}, - {42, 137, 2, true, {KernelType::TANH, 0, 0.5, 2.4}}, - {42, 137, 2, false, {KernelType::TANH, 0, 0.5, 2.4}, 64, 155, 49}, - {42, 137, 2, true, {KernelType::TANH, 0, 0.5, 2.4}, 64, 155, 143}, - {3, 4, 2, false, {KernelType::RBF, 0, 0.5}}, - {42, 137, 2, false, {KernelType::RBF, 0, 0.5}}, - {42, 137, 2, true, {KernelType::RBF, 0, 0.5}}, - // Distance kernel does not support LD parameter yet. - //{42, 137, 2, false, {KernelType::RBF, 0, 0.5}, 64, 155, 49}, - // {42, 137, 2, true, {KernelType::RBF, 0, 0.5}, 64, 155, 143}, -}; - -template -class GramMatrixTest : public ::testing::TestWithParam { - protected: - GramMatrixTest() - : params(GetParam()), - handle(), - x1(0, resource::get_cuda_stream(handle)), - x2(0, resource::get_cuda_stream(handle)), - gram(0, resource::get_cuda_stream(handle)), - gram_host(0) - { - auto stream = resource::get_cuda_stream(handle); - - if (params.ld1 == 0) { params.ld1 = params.is_row_major ? params.n_cols : params.n1; } - if (params.ld2 == 0) { params.ld2 = params.is_row_major ? params.n_cols : params.n2; } - if (params.ld_out == 0) { params.ld_out = params.is_row_major ? params.n2 : params.n1; } - // Derive the size of the output from the offset of the last element. - size_t size = get_offset(params.n1 - 1, params.n_cols - 1, params.ld1, params.is_row_major) + 1; - x1.resize(size, stream); - size = get_offset(params.n2 - 1, params.n_cols - 1, params.ld2, params.is_row_major) + 1; - x2.resize(size, stream); - size = get_offset(params.n1 - 1, params.n2 - 1, params.ld_out, params.is_row_major) + 1; - - gram.resize(size, stream); - RAFT_CUDA_TRY(cudaMemsetAsync(gram.data(), 0, gram.size() * sizeof(math_t), stream)); - gram_host.resize(gram.size()); - std::fill(gram_host.begin(), gram_host.end(), 0); - - raft::random::RngState rng(42137ULL); - raft::random::uniform(handle, rng, x1.data(), x1.size(), math_t(0), math_t(1)); - raft::random::uniform(handle, rng, x2.data(), x2.size(), math_t(0), math_t(1)); - } - - ~GramMatrixTest() override {} - - void runTest() - { - std::unique_ptr> kernel = - std::unique_ptr>(KernelFactory::create(params.kernel)); - - auto x1_span = - params.is_row_major - ? raft::make_device_strided_matrix_view( - x1.data(), params.n1, params.n_cols, params.ld1) - : raft::make_device_strided_matrix_view( - x1.data(), params.n1, params.n_cols, params.ld1); - auto x2_span = - params.is_row_major - ? raft::make_device_strided_matrix_view( - x2.data(), params.n2, params.n_cols, params.ld2) - : raft::make_device_strided_matrix_view( - x2.data(), params.n2, params.n_cols, params.ld2); - auto out_span = - params.is_row_major - ? raft::make_device_strided_matrix_view( - gram.data(), params.n1, params.n2, params.ld_out) - : raft::make_device_strided_matrix_view( - gram.data(), params.n1, params.n2, params.ld_out); - - (*kernel)(handle, x1_span, x2_span, out_span); - - auto stream = resource::get_cuda_stream(handle); - naiveGramMatrixKernel(params.n1, - params.n2, - params.n_cols, - x1, - x2, - gram_host.data(), - params.ld1, - params.ld2, - params.ld_out, - params.is_row_major, - params.kernel, - stream, - handle); - - ASSERT_TRUE(raft::devArrMatchHost( - gram_host.data(), gram.data(), gram.size(), raft::CompareApprox(1e-6f), stream)); - } - - GramMatrixInputs params; - raft::resources handle; - - rmm::device_uvector x1; - rmm::device_uvector x2; - rmm::device_uvector gram; - - std::vector gram_host; -}; - -typedef GramMatrixTest GramMatrixTestFloat; -typedef GramMatrixTest GramMatrixTestDouble; - -TEST_P(GramMatrixTestFloat, Gram) { runTest(); } - -INSTANTIATE_TEST_SUITE_P(GramMatrixTests, GramMatrixTestFloat, ::testing::ValuesIn(inputs)); -}; // end namespace raft::distance::kernels diff --git a/cpp/test/distance/gram_base.cuh b/cpp/test/distance/gram_base.cuh deleted file mode 100644 index 170bcb76c1..0000000000 --- a/cpp/test/distance/gram_base.cuh +++ /dev/null @@ -1,90 +0,0 @@ -/* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include -#include -#include -#include -#include - -#include - -#include -#include - -namespace raft { -namespace distance { -namespace kernels { - -// Get the offset of element [i,k]. -HDI int get_offset(int i, int k, int ld, bool is_row_major) -{ - return is_row_major ? i * ld + k : i + k * ld; -} - -// Calculate the Gram matrix on the host. -template -void naiveGramMatrixKernel(int n1, - int n2, - int n_cols, - const rmm::device_uvector& x1, - const rmm::device_uvector& x2, - math_t* gram_host, - int ld1, - int ld2, - int ld_out, - bool is_row_major, - KernelParams kernel, - cudaStream_t stream, - const raft::resources& handle) -{ - std::vector x1_host(x1.size()); - raft::update_host(x1_host.data(), x1.data(), x1.size(), stream); - std::vector x2_host(x2.size()); - raft::update_host(x2_host.data(), x2.data(), x2.size(), stream); - resource::sync_stream(handle, stream); - - for (int i = 0; i < n1; i++) { - for (int j = 0; j < n2; j++) { - float d = 0; - for (int k = 0; k < n_cols; k++) { - if (kernel.kernel == KernelType::RBF) { - math_t diff = x1_host[get_offset(i, k, ld1, is_row_major)] - - x2_host[get_offset(j, k, ld2, is_row_major)]; - d += diff * diff; - } else { - d += x1_host[get_offset(i, k, ld1, is_row_major)] * - x2_host[get_offset(j, k, ld2, is_row_major)]; - } - } - int idx = get_offset(i, j, ld_out, is_row_major); - math_t v = 0; - switch (kernel.kernel) { - case (KernelType::LINEAR): gram_host[idx] = d; break; - case (KernelType::POLYNOMIAL): - v = kernel.gamma * d + kernel.coef0; - gram_host[idx] = std::pow(v, kernel.degree); - break; - case (KernelType::TANH): gram_host[idx] = std::tanh(kernel.gamma * d + kernel.coef0); break; - case (KernelType::RBF): gram_host[idx] = exp(-kernel.gamma * d); break; - } - } - } -} - -} // namespace kernels -} // namespace distance -} // namespace raft diff --git a/cpp/test/distance/masked_nn.cu b/cpp/test/distance/masked_nn.cu deleted file mode 100644 index 34fa07c45b..0000000000 --- a/cpp/test/distance/masked_nn.cu +++ /dev/null @@ -1,438 +0,0 @@ -/* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.h" - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include - -namespace raft::distance::masked_nn { - -// The adjacency pattern determines what distances get computed. -enum AdjacencyPattern { - checkerboard = 0, // adjacency matrix looks like a checkerboard (half the distances are computed) - checkerboard_4 = 1, // checkerboard with tiles of size 4x4 - checkerboard_64 = 2, // checkerboard with tiles of size 64x64 - all_true = 3, // no distance computations can be skipped - all_false = 4 // all distance computations can be skipped -}; - -// Kernels: -// - init_adj: to initialize the adjacency kernel with a specific adjacency pattern -// - referenceKernel: to produce the ground-truth output - -RAFT_KERNEL init_adj(AdjacencyPattern pattern, - int n, - raft::device_matrix_view adj, - raft::device_vector_view group_idxs) -{ - int m = adj.extent(0); - int num_groups = adj.extent(1); - - for (int idx_m = blockIdx.y * blockDim.y + threadIdx.y; idx_m < m; - idx_m += blockDim.y * gridDim.y) { - for (int idx_g = blockIdx.x * blockDim.x + threadIdx.x; idx_g < num_groups; - idx_g += blockDim.x * gridDim.x) { - switch (pattern) { - case checkerboard: adj(idx_m, idx_g) = (idx_m + idx_g) % 2; break; - case checkerboard_4: adj(idx_m, idx_g) = (idx_m / 4 + idx_g) % 2; break; - case checkerboard_64: adj(idx_m, idx_g) = (idx_m / 64 + idx_g) % 2; break; - case all_true: adj(idx_m, idx_g) = true; break; - case all_false: adj(idx_m, idx_g) = false; break; - default: assert(false && "unknown pattern"); - } - } - } - // Each group is of size n / num_groups. - // - // - group_idxs[j] indicates the start of group j + 1 (i.e. is the inclusive - // scan of the group lengths) - // - // - The first group always starts at index zero, so we do not store it. - // - // - The group_idxs[num_groups - 1] should always equal n. - - if (blockIdx.y == 0 && threadIdx.y == 0) { - const int g_stride = blockDim.x * gridDim.x; - for (int idx_g = blockIdx.x * blockDim.x + threadIdx.x; idx_g < num_groups; idx_g += g_stride) { - group_idxs(idx_g) = (idx_g + 1) * (n / num_groups); - } - group_idxs(num_groups - 1) = n; - } -} - -template -__launch_bounds__(32 * NWARPS, 2) RAFT_KERNEL referenceKernel(raft::KeyValuePair* min, - DataT* x, - DataT* y, - bool* adj, - int* group_idxs, - int m, - int n, - int k, - int num_groups, - bool sqrt, - int* workspace, - DataT maxVal) -{ - const int m_stride = blockDim.y * gridDim.y; - const int m_offset = threadIdx.y + blockIdx.y * blockDim.y; - const int n_stride = blockDim.x * gridDim.x; - const int n_offset = threadIdx.x + blockIdx.x * blockDim.x; - - for (int m_grid = 0; m_grid < m; m_grid += m_stride) { - for (int n_grid = 0; n_grid < n; n_grid += n_stride) { - int midx = m_grid + m_offset; - int nidx = n_grid + n_offset; - - // Do a reverse linear search to determine the group index. - int group_idx = 0; - for (int i = num_groups; 0 <= i; --i) { - if (nidx < group_idxs[i]) { group_idx = i; } - } - const bool include_dist = adj[midx * num_groups + group_idx] && midx < m && nidx < n; - - // Compute L2 metric. - DataT acc = DataT(0); - for (int i = 0; i < k; ++i) { - int xidx = i + midx * k; - int yidx = i + nidx * k; - auto diff = x[xidx] - y[yidx]; - acc += diff * diff; - } - if (sqrt) { acc = raft::sqrt(acc); } - ReduceOpT redOp; - typedef cub::WarpReduce> WarpReduce; - __shared__ typename WarpReduce::TempStorage temp[NWARPS]; - int warpId = threadIdx.x / raft::WarpSize; - raft::KeyValuePair tmp; - tmp.key = include_dist ? nidx : -1; - tmp.value = include_dist ? acc : maxVal; - tmp = WarpReduce(temp[warpId]).Reduce(tmp, raft::distance::KVPMinReduce{}); - if (threadIdx.x % raft::WarpSize == 0 && midx < m) { - while (atomicCAS(workspace + midx, 0, 1) == 1) - ; - __threadfence(); - redOp(midx, min + midx, tmp); - __threadfence(); - atomicCAS(workspace + midx, 1, 0); - } - __syncthreads(); - } - } -} - -// Structs -// - Params: holds parameters for test case -// - Inputs: holds the inputs to the functions under test (x, y, adj, group_idxs). Is generated from -// the inputs. -struct Params { - double tolerance; - int m, n, k, num_groups; - bool sqrt; - unsigned long long int seed; - AdjacencyPattern pattern; -}; - -inline auto operator<<(std::ostream& os, const Params& p) -> std::ostream& -{ - os << "m: " << p.m << ", n: " << p.n << ", k: " << p.k << ", num_groups: " << p.num_groups - << ", sqrt: " << p.sqrt << ", seed: " << p.seed << ", tol: " << p.tolerance; - return os; -} - -template -struct Inputs { - using IdxT = int; - - raft::device_matrix x, y; - raft::device_matrix adj; - raft::device_vector group_idxs; - - Inputs(const raft::handle_t& handle, const Params& p) - : x{raft::make_device_matrix(handle, p.m, p.k)}, - y{raft::make_device_matrix(handle, p.n, p.k)}, - adj{raft::make_device_matrix(handle, p.m, p.num_groups)}, - group_idxs{raft::make_device_vector(handle, p.num_groups)} - { - // Initialize x, y - raft::random::RngState r(p.seed); - uniform(handle, r, x.data_handle(), p.m * p.k, DataT(-1.0), DataT(1.0)); - uniform(handle, r, y.data_handle(), p.n * p.k, DataT(-1.0), DataT(1.0)); - - // Initialize adj, group_idxs. - dim3 block(32, 32); - dim3 grid(10, 10); - init_adj<<>>( - p.pattern, p.n, adj.view(), group_idxs.view()); - RAFT_CUDA_TRY(cudaGetLastError()); - } -}; - -template > -auto reference(const raft::handle_t& handle, Inputs inp, const Params& p) - -> raft::device_vector -{ - int m = inp.x.extent(0); - int n = inp.y.extent(0); - int k = inp.x.extent(1); - int num_groups = inp.group_idxs.extent(0); - - if (m == 0 || n == 0 || k == 0 || num_groups == 0) { - return raft::make_device_vector(handle, 0); - } - - // Initialize workspace - auto stream = resource::get_cuda_stream(handle); - rmm::device_uvector workspace(p.m * sizeof(int), stream); - RAFT_CUDA_TRY(cudaMemsetAsync(workspace.data(), 0, sizeof(int) * m, stream)); - - // Initialize output - auto out = raft::make_device_vector(handle, m); - auto blks = raft::ceildiv(m, 256); - MinAndDistanceReduceOp op; - raft::distance::detail::initKernel, int> - <<>>(out.data_handle(), m, std::numeric_limits::max(), op); - RAFT_CUDA_TRY(cudaGetLastError()); - - // Launch reference kernel - const int nwarps = 16; - static const dim3 TPB(32, nwarps, 1); - dim3 nblks(1, 200, 1); - referenceKernel - <<>>(out.data_handle(), - inp.x.data_handle(), - inp.y.data_handle(), - inp.adj.data_handle(), - inp.group_idxs.data_handle(), - m, - n, - k, - num_groups, - p.sqrt, - (int*)workspace.data(), - std::numeric_limits::max()); - RAFT_CUDA_TRY(cudaGetLastError()); - - return out; -} - -template > -auto run_masked_nn(const raft::handle_t& handle, Inputs inp, const Params& p) - -> raft::device_vector -{ - // Compute norms: - auto x_norm = raft::make_device_vector(handle, p.m); - auto y_norm = raft::make_device_vector(handle, p.n); - - raft::linalg::norm(handle, - std::as_const(inp.x).view(), - x_norm.view(), - raft::linalg::L2Norm, - raft::linalg::Apply::ALONG_ROWS); - raft::linalg::norm(handle, - std::as_const(inp.y).view(), - y_norm.view(), - raft::linalg::L2Norm, - raft::linalg::Apply::ALONG_ROWS); - - // Create parameters for masked_l2_nn - using IdxT = int; - using RedOpT = MinAndDistanceReduceOp; - using PairRedOpT = raft::distance::KVPMinReduce; - using ParamT = raft::distance::masked_l2_nn_params; - - bool init_out = true; - ParamT masked_l2_params{RedOpT{}, PairRedOpT{}, p.sqrt, init_out}; - - // Create output - auto out = raft::make_device_vector(handle, p.m); - - // Launch kernel - raft::distance::masked_l2_nn(handle, - masked_l2_params, - inp.x.view(), - inp.y.view(), - x_norm.view(), - y_norm.view(), - inp.adj.view(), - inp.group_idxs.view(), - out.view()); - - resource::sync_stream(handle); - - return out; -} - -template -struct CompareApproxAbsKVP { - typedef typename raft::KeyValuePair KVP; - CompareApproxAbsKVP(T eps_) : eps(eps_) {} - bool operator()(const KVP& a, const KVP& b) const - { - T diff = raft::abs(raft::abs(a.value) - raft::abs(b.value)); - T m = std::max(raft::abs(a.value), raft::abs(b.value)); - T ratio = m >= eps ? diff / m : diff; - return (ratio <= eps); - } - - private: - T eps; -}; - -template -::testing::AssertionResult devArrMatch(const raft::KeyValuePair* expected, - const raft::KeyValuePair* actual, - size_t size, - L eq_compare, - cudaStream_t stream = 0) -{ - typedef typename raft::KeyValuePair KVP; - std::shared_ptr exp_h(new KVP[size]); - std::shared_ptr act_h(new KVP[size]); - raft::update_host(exp_h.get(), expected, size, stream); - raft::update_host(act_h.get(), actual, size, stream); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - for (size_t i(0); i < size; ++i) { - auto exp = exp_h.get()[i]; - auto act = act_h.get()[i]; - if (!eq_compare(exp, act)) { - return ::testing::AssertionFailure() - << "actual=" << act.key << "," << act.value << " != expected=" << exp.key << "," - << exp.value << " @" << i; - } - } - return ::testing::AssertionSuccess(); -} - -inline auto gen_params() -> std::vector -{ - // Regular powers of two - auto regular = raft::util::itertools::product({0.001f}, // tolerance - {32, 64, 512}, // m - {32, 64, 512}, // n - {8, 32}, // k - {2, 32}, // num_groups - {true, false}, // sqrt - {1234ULL}, // seed - {AdjacencyPattern::all_true, - AdjacencyPattern::checkerboard, - AdjacencyPattern::checkerboard_64, - AdjacencyPattern::all_false}); - - // Irregular sizes to check tiling and bounds checking - auto irregular = raft::util::itertools::product({0.001f}, // tolerance - {511, 512, 513}, // m - {127, 128, 129}, // n - {5}, // k - {3, 9}, // num_groups - {true, false}, // sqrt - {1234ULL}, // seed - {AdjacencyPattern::all_true, - AdjacencyPattern::checkerboard, - AdjacencyPattern::checkerboard_64}); - - regular.insert(regular.end(), irregular.begin(), irregular.end()); - - return regular; -} - -class MaskedL2NNTest : public ::testing::TestWithParam { - // Empty. -}; - -// -TEST_P(MaskedL2NNTest, ReferenceCheckFloat) -{ - using DataT = float; - - // Get parameters; create handle and input data. - Params p = GetParam(); - raft::handle_t handle{}; - Inputs inputs{handle, p}; - - // Calculate reference and test output - auto out_reference = reference(handle, inputs, p); - auto out_fast = run_masked_nn(handle, inputs, p); - - // Check for differences. - ASSERT_TRUE(devArrMatch(out_reference.data_handle(), - out_fast.data_handle(), - p.m, - CompareApproxAbsKVP(p.tolerance), - resource::get_cuda_stream(handle))); -} - -// This test checks whether running the masked_l2_nn twice returns the same -// output. -TEST_P(MaskedL2NNTest, DeterminismCheck) -{ - using DataT = float; - - // Get parameters; create handle and input data. - Params p = GetParam(); - raft::handle_t handle{}; - Inputs inputs{handle, p}; - - // Calculate reference and test output - auto out1 = run_masked_nn(handle, inputs, p); - auto out2 = run_masked_nn(handle, inputs, p); - - // Check for differences. - ASSERT_TRUE(devArrMatch(out1.data_handle(), - out2.data_handle(), - p.m, - CompareApproxAbsKVP(p.tolerance), - resource::get_cuda_stream(handle))); -} - -TEST_P(MaskedL2NNTest, ReferenceCheckDouble) -{ - using DataT = double; - - // Get parameters; create handle and input data. - Params p = GetParam(); - raft::handle_t handle{}; - Inputs inputs{handle, p}; - - // Calculate reference and test output - auto out_reference = reference(handle, inputs, p); - auto out_fast = run_masked_nn(handle, inputs, p); - - // Check for differences. - ASSERT_TRUE(devArrMatch(out_reference.data_handle(), - out_fast.data_handle(), - p.m, - CompareApproxAbsKVP(p.tolerance), - resource::get_cuda_stream(handle))); -} - -INSTANTIATE_TEST_CASE_P(MaskedL2NNTests, MaskedL2NNTest, ::testing::ValuesIn(gen_params())); - -} // end namespace raft::distance::masked_nn diff --git a/cpp/test/distance/masked_nn_compress_to_bits.cu b/cpp/test/distance/masked_nn_compress_to_bits.cu deleted file mode 100644 index 2512af5e4f..0000000000 --- a/cpp/test/distance/masked_nn_compress_to_bits.cu +++ /dev/null @@ -1,220 +0,0 @@ -/* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include "../test_utils.cuh" -#include "../test_utils.h" - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include - -namespace raft::distance::masked_nn::compress_to_bits { - -/** - * @brief Transpose and decompress 2D bitfield to boolean matrix - * - * Inverse operation of compress_to_bits - * - * @tparam T - * - * @parameter[in] in An `m x n` bitfield matrix. Row major. - * @parameter in_rows The number of rows of `in`, i.e. `m`. - * @parameter in_cols The number of cols of `in`, i.e. `n`. - * - * @parameter[out] out An `(m * bits_per_elem) x n` boolean matrix. - */ -template ::value>> -RAFT_KERNEL decompress_bits_kernel(const T* in, int in_rows, int in_cols, bool* out) -{ - constexpr int bits_per_element = 8 * sizeof(T); - - const size_t i = threadIdx.y + blockIdx.y * blockDim.y; - const size_t j = threadIdx.x + blockIdx.x * blockDim.x; - - if (in_rows <= i || in_cols <= j) { return; } - - const size_t out_rows = in_rows * bits_per_element; - const size_t out_cols = in_cols; - const size_t out_i = i * bits_per_element; - const size_t out_j = j; - - if (out_rows <= out_i && out_cols <= out_j) { return; } - - T bitfield = in[i * in_cols + j]; - for (int bitpos = 0; bitpos < bits_per_element; ++bitpos) { - bool bit = ((T(1) << bitpos) & bitfield) != 0; - out[(out_i + bitpos) * out_cols + out_j] = bit; - } -} - -/** - * @brief Transpose and decompress 2D bitfield to boolean matrix - * - * Inverse operation of compress_to_bits - * - * @tparam T - * - * @parameter[in] in An `m x n` bitfield matrix. Row major. - * @parameter in_rows The number of rows of `in`, i.e. `m`. - * @parameter in_cols The number of cols of `in`, i.e. `n`. - * - * @parameter[out] out An `n x (m * bits_per_elem)` boolean matrix. - */ -template ::value>> -void decompress_bits(const raft::handle_t& handle, const T* in, int in_rows, int in_cols, bool* out) -{ - auto stream = resource::get_cuda_stream(handle); - dim3 grid(raft::ceildiv(in_cols, 32), raft::ceildiv(in_rows, 32)); - dim3 block(32, 32); - decompress_bits_kernel<<>>(in, in_rows, in_cols, out); - RAFT_CUDA_TRY(cudaGetLastError()); -} - -// Params holds parameters for test case -struct Params { - int m, n; -}; - -inline auto operator<<(std::ostream& os, const Params& p) -> std::ostream& -{ - return os << "m: " << p.m << ", n: " << p.n; -} - -// Check that the following holds -// -// decompress(compress(x)) == x -// -// for 2D boolean matrices x. -template -void check_invertible(const Params& p) -{ - using raft::distance::detail::compress_to_bits; - constexpr int bits_per_elem = sizeof(T) * 8; - - // Make m and n that are safe to ceildiv. - int m = raft::round_up_safe(p.m, bits_per_elem); - int n = p.n; - - // Generate random input - raft::handle_t handle{}; - raft::random::RngState r(1ULL); - auto in = raft::make_device_matrix(handle, m, n); - raft::random::bernoulli(handle, r, in.data_handle(), m * n, 0.5f); - - int tmp_m = raft::ceildiv(m, bits_per_elem); - int out_m = tmp_m * bits_per_elem; - - auto tmp = raft::make_device_matrix(handle, tmp_m, n); - auto out = raft::make_device_matrix(handle, out_m, n); - - resource::sync_stream(handle); - RAFT_CUDA_TRY(cudaGetLastError()); - - ASSERT_EQ(in.extent(0), out.extent(0)) << "M does not match"; - ASSERT_EQ(in.extent(1), out.extent(1)) << "N does not match"; - - compress_to_bits(handle, in.view(), tmp.view()); - resource::sync_stream(handle); - RAFT_CUDA_TRY(cudaGetLastError()); - - decompress_bits(handle, tmp.data_handle(), tmp.extent(0), tmp.extent(1), out.data_handle()); - resource::sync_stream(handle); - RAFT_CUDA_TRY(cudaGetLastError()); - - // Check for differences. - ASSERT_TRUE(raft::devArrMatch(in.data_handle(), - out.data_handle(), - in.extent(0) * in.extent(1), - raft::Compare(), - resource::get_cuda_stream(handle))); - resource::sync_stream(handle); - RAFT_CUDA_TRY(cudaGetLastError()); -} - -void check_all_true(const Params& p) -{ - using raft::distance::detail::compress_to_bits; - using T = uint64_t; - constexpr int bits_per_elem = sizeof(T) * 8; - - // Make m and n that are safe to ceildiv. - int m = raft::round_up_safe(p.m, bits_per_elem); - int n = p.n; - - raft::handle_t handle{}; - raft::random::RngState r(1ULL); - auto in = raft::make_device_matrix(handle, m, n); - raft::matrix::fill(handle, in.view(), true); - - int tmp_m = raft::ceildiv(m, bits_per_elem); - auto tmp = raft::make_device_matrix(handle, tmp_m, n); - resource::sync_stream(handle); - RAFT_CUDA_TRY(cudaGetLastError()); - - compress_to_bits(handle, in.view(), tmp.view()); - resource::sync_stream(handle); - RAFT_CUDA_TRY(cudaGetLastError()); - - auto expected = raft::make_device_matrix(handle, tmp_m, n); - raft::matrix::fill(handle, expected.view(), ~T(0)); - - // Check for differences. - ASSERT_TRUE(raft::devArrMatch(expected.data_handle(), - tmp.data_handle(), - tmp.extent(0) * tmp.extent(1), - raft::Compare(), - resource::get_cuda_stream(handle))); - resource::sync_stream(handle); - RAFT_CUDA_TRY(cudaGetLastError()); -} - -class CompressToBitsTest : public ::testing::TestWithParam { - // Empty. -}; - -TEST_P(CompressToBitsTest, CheckTrue64) { check_all_true(GetParam()); } - -TEST_P(CompressToBitsTest, CheckInvertible64) -{ - using T = uint64_t; - check_invertible(GetParam()); -} - -TEST_P(CompressToBitsTest, CheckInvertible32) -{ - using T = uint32_t; - check_invertible(GetParam()); -} - -std::vector params = raft::util::itertools::product( - {1, 3, 32, 33, 63, 64, 65, 128, 10013}, {1, 3, 32, 33, 63, 64, 65, 13001}); - -INSTANTIATE_TEST_CASE_P(CompressToBits, CompressToBitsTest, ::testing::ValuesIn(params)); - -} // namespace raft::distance::masked_nn::compress_to_bits \ No newline at end of file