From f94780c2adb579a788c30397ce81eaae42e1d64c Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 12 Jul 2021 16:43:35 -0400 Subject: [PATCH] Remaining sparse semiring distances (#261) This PR is intended to be merged after #207 (hash table strategy) has been merged. This PR introduces the following distances: - Hamming - Jensen-Shannon - Russell-Rao - KL-Divergence - Correlation Most of the changes here are from #207 and will be reviewed in that PR. The only files that need to be reviewed for this PR are `sparse/distance/l2_distance.cuh`, `sparse/distance/bin_distance.cuh`, `sparse/distance/lp_distances.cuh`, and their corresponding gtests: `test/sparse/distance.cuh` Authors: - Corey J. Nolet (https://github.com/cjnolet) - Divye Gala (https://github.com/divyegala) Approvers: - Dante Gama Dessavre (https://github.com/dantegd) URL: https://github.com/rapidsai/raft/pull/261 --- .../raft/sparse/distance/bin_distance.cuh | 25 +-- cpp/include/raft/sparse/distance/distance.cuh | 27 ++- .../raft/sparse/distance/ip_distance.cuh | 7 +- .../raft/sparse/distance/l2_distance.cuh | 180 ++++++++++++++++-- .../raft/sparse/distance/lp_distance.cuh | 97 +++++++++- cpp/test/sparse/distance.cu | 96 ++++++++++ 6 files changed, 387 insertions(+), 45 deletions(-) diff --git a/cpp/include/raft/sparse/distance/bin_distance.cuh b/cpp/include/raft/sparse/distance/bin_distance.cuh index e5ac85e909..f3109556b7 100644 --- a/cpp/include/raft/sparse/distance/bin_distance.cuh +++ b/cpp/include/raft/sparse/distance/bin_distance.cuh @@ -24,7 +24,6 @@ #include #include -#include #include #include @@ -87,8 +86,8 @@ void compute_bin_distance(value_t *out, const value_idx *Q_coo_rows, value_idx R_nnz, value_idx m, value_idx n, std::shared_ptr alloc, cudaStream_t stream, expansion_f expansion_func) { - raft::mr::device::buffer Q_norms(alloc, stream, m); - raft::mr::device::buffer R_norms(alloc, stream, n); + rmm::device_uvector Q_norms(m, stream); + rmm::device_uvector R_norms(n, stream); CUDA_CHECK( cudaMemsetAsync(Q_norms.data(), 0, Q_norms.size() * sizeof(value_t))); CUDA_CHECK( @@ -113,8 +112,7 @@ class jaccard_expanded_distances_t : public distances_t { explicit jaccard_expanded_distances_t( const distances_config_t &config) : config_(&config), - workspace(config.handle.get_device_allocator(), - config.handle.get_stream(), 0), + workspace(0, config.handle.get_stream()), ip_dists(config) {} void compute(value_t *out_dists) { @@ -123,9 +121,8 @@ class jaccard_expanded_distances_t : public distances_t { value_idx *b_indices = ip_dists.b_rows_coo(); value_t *b_data = ip_dists.b_data_coo(); - raft::mr::device::buffer search_coo_rows( - config_->handle.get_device_allocator(), config_->handle.get_stream(), - config_->a_nnz); + rmm::device_uvector search_coo_rows( + config_->a_nnz, config_->handle.get_stream()); raft::sparse::convert::csr_to_coo(config_->a_indptr, config_->a_nrows, search_coo_rows.data(), config_->a_nnz, config_->handle.get_stream()); @@ -150,7 +147,7 @@ class jaccard_expanded_distances_t : public distances_t { private: const distances_config_t *config_; - raft::mr::device::buffer workspace; + rmm::device_uvector workspace; ip_distances_t ip_dists; }; @@ -164,8 +161,7 @@ class dice_expanded_distances_t : public distances_t { explicit dice_expanded_distances_t( const distances_config_t &config) : config_(&config), - workspace(config.handle.get_device_allocator(), - config.handle.get_stream(), 0), + workspace(0, config.handle.get_stream()), ip_dists(config) {} void compute(value_t *out_dists) { @@ -174,9 +170,8 @@ class dice_expanded_distances_t : public distances_t { value_idx *b_indices = ip_dists.b_rows_coo(); value_t *b_data = ip_dists.b_data_coo(); - raft::mr::device::buffer search_coo_rows( - config_->handle.get_device_allocator(), config_->handle.get_stream(), - config_->a_nnz); + rmm::device_uvector search_coo_rows( + config_->a_nnz, config_->handle.get_stream()); raft::sparse::convert::csr_to_coo(config_->a_indptr, config_->a_nrows, search_coo_rows.data(), config_->a_nnz, config_->handle.get_stream()); @@ -197,7 +192,7 @@ class dice_expanded_distances_t : public distances_t { private: const distances_config_t *config_; - raft::mr::device::buffer workspace; + rmm::device_uvector workspace; ip_distances_t ip_dists; }; diff --git a/cpp/include/raft/sparse/distance/distance.cuh b/cpp/include/raft/sparse/distance/distance.cuh index 0cd0be11be..a1974b3666 100644 --- a/cpp/include/raft/sparse/distance/distance.cuh +++ b/cpp/include/raft/sparse/distance/distance.cuh @@ -56,7 +56,12 @@ static const std::unordered_set supportedDistance{ raft::distance::DistanceType::JaccardExpanded, raft::distance::DistanceType::CosineExpanded, raft::distance::DistanceType::HellingerExpanded, - raft::distance::DistanceType::DiceExpanded}; + raft::distance::DistanceType::DiceExpanded, + raft::distance::DistanceType::CorrelationExpanded, + raft::distance::DistanceType::RusselRaoExpanded, + raft::distance::DistanceType::HammingUnexpanded, + raft::distance::DistanceType::JensenShannon, + raft::distance::DistanceType::KLDivergence}; /** * Compute pairwise distances between A and B, using the provided @@ -120,6 +125,26 @@ void pairwiseDistance(value_t *out, case raft::distance::DistanceType::DiceExpanded: dice_expanded_distances_t(input_config).compute(out); break; + case raft::distance::DistanceType::CorrelationExpanded: + correlation_expanded_distances_t(input_config) + .compute(out); + break; + case raft::distance::DistanceType::RusselRaoExpanded: + russelrao_expanded_distances_t(input_config) + .compute(out); + break; + case raft::distance::DistanceType::HammingUnexpanded: + hamming_unexpanded_distances_t(input_config) + .compute(out); + break; + case raft::distance::DistanceType::JensenShannon: + jensen_shannon_unexpanded_distances_t(input_config) + .compute(out); + break; + case raft::distance::DistanceType::KLDivergence: + kl_divergence_unexpanded_distances_t(input_config) + .compute(out); + break; default: THROW("Unsupported distance: %d", metric); diff --git a/cpp/include/raft/sparse/distance/ip_distance.cuh b/cpp/include/raft/sparse/distance/ip_distance.cuh index bf45fe0d8e..882ccba027 100644 --- a/cpp/include/raft/sparse/distance/ip_distance.cuh +++ b/cpp/include/raft/sparse/distance/ip_distance.cuh @@ -23,7 +23,6 @@ #include #include -#include #include #include @@ -47,9 +46,7 @@ class ip_distances_t : public distances_t { * @param[in] config specifies inputs, outputs, and sizes */ ip_distances_t(const distances_config_t &config) - : config_(&config), - coo_rows_b(config.handle.get_device_allocator(), - config.handle.get_stream(), config.b_nnz) { + : config_(&config), coo_rows_b(config.b_nnz, config.handle.get_stream()) { raft::sparse::convert::csr_to_coo(config_->b_indptr, config_->b_nrows, coo_rows_b.data(), config_->b_nnz, config_->handle.get_stream()); @@ -74,7 +71,7 @@ class ip_distances_t : public distances_t { private: const distances_config_t *config_; - raft::mr::device::buffer coo_rows_b; + rmm::device_uvector coo_rows_b; }; }; // END namespace distance }; // END namespace sparse diff --git a/cpp/include/raft/sparse/distance/l2_distance.cuh b/cpp/include/raft/sparse/distance/l2_distance.cuh index f73e23d94b..8886d4c9df 100644 --- a/cpp/include/raft/sparse/distance/l2_distance.cuh +++ b/cpp/include/raft/sparse/distance/l2_distance.cuh @@ -24,7 +24,8 @@ #include #include #include -#include + +#include #include #include @@ -81,6 +82,35 @@ __global__ void compute_euclidean_warp_kernel( C[(size_t)i * n_cols + j] = val * (fabs(val) >= 0.0001); } +template +__global__ void compute_correlation_warp_kernel( + value_t *__restrict__ C, const value_t *__restrict__ Q_sq_norms, + const value_t *__restrict__ R_sq_norms, const value_t *__restrict__ Q_norms, + const value_t *__restrict__ R_norms, value_idx n_rows, value_idx n_cols, + value_idx n) { + value_idx tid = blockDim.x * blockIdx.x + threadIdx.x; + value_idx i = tid / n_cols; + value_idx j = tid % n_cols; + + if (i >= n_rows || j >= n_cols) return; + + value_t dot = C[(size_t)i * n_cols + j]; + value_t Q_l1 = Q_norms[i]; + value_t R_l1 = R_norms[j]; + + value_t Q_l2 = Q_sq_norms[i]; + value_t R_l2 = R_sq_norms[j]; + + value_t numer = n * dot - (Q_l1 * R_l1); + value_t Q_denom = n * Q_l2 - (Q_l1 * Q_l1); + value_t R_denom = n * R_l2 - (R_l1 * R_l1); + + value_t val = 1 - (numer / sqrt(Q_denom * R_denom)); + + // correct for small instabilities + C[(size_t)i * n_cols + j] = val * (fabs(val) >= 0.0001); +} + template void compute_euclidean(value_t *C, const value_t *Q_sq_norms, @@ -100,8 +130,8 @@ void compute_l2(value_t *out, const value_idx *Q_coo_rows, value_idx R_nnz, value_idx m, value_idx n, std::shared_ptr alloc, cudaStream_t stream, expansion_f expansion_func) { - raft::mr::device::buffer Q_sq_norms(alloc, stream, m); - raft::mr::device::buffer R_sq_norms(alloc, stream, n); + rmm::device_uvector Q_sq_norms(m, stream); + rmm::device_uvector R_sq_norms(n, stream); CUDA_CHECK( cudaMemsetAsync(Q_sq_norms.data(), 0, Q_sq_norms.size() * sizeof(value_t))); CUDA_CHECK( @@ -116,6 +146,55 @@ void compute_l2(value_t *out, const value_idx *Q_coo_rows, expansion_func); } +template +void compute_correlation(value_t *C, const value_t *Q_sq_norms, + const value_t *R_sq_norms, const value_t *Q_norms, + const value_t *R_norms, value_idx n_rows, + value_idx n_cols, value_idx n, cudaStream_t stream) { + int blocks = raft::ceildiv((size_t)n_rows * n_cols, tpb); + compute_correlation_warp_kernel<<>>( + C, Q_sq_norms, R_sq_norms, Q_norms, R_norms, n_rows, n_cols, n); +} + +template +void compute_corr(value_t *out, const value_idx *Q_coo_rows, + const value_t *Q_data, value_idx Q_nnz, + const value_idx *R_coo_rows, const value_t *R_data, + value_idx R_nnz, value_idx m, value_idx n, value_idx n_cols, + std::shared_ptr alloc, + cudaStream_t stream) { + // sum_sq for std dev + rmm::device_uvector Q_sq_norms(m, stream); + rmm::device_uvector R_sq_norms(n, stream); + + // sum for mean + rmm::device_uvector Q_norms(m, stream); + rmm::device_uvector R_norms(n, stream); + + CUDA_CHECK( + cudaMemsetAsync(Q_sq_norms.data(), 0, Q_sq_norms.size() * sizeof(value_t))); + CUDA_CHECK( + cudaMemsetAsync(R_sq_norms.data(), 0, R_sq_norms.size() * sizeof(value_t))); + + CUDA_CHECK( + cudaMemsetAsync(Q_norms.data(), 0, Q_norms.size() * sizeof(value_t))); + CUDA_CHECK( + cudaMemsetAsync(R_norms.data(), 0, R_norms.size() * sizeof(value_t))); + + compute_row_norm_kernel<<>>( + Q_sq_norms.data(), Q_coo_rows, Q_data, Q_nnz); + compute_row_norm_kernel<<>>( + R_sq_norms.data(), R_coo_rows, R_data, R_nnz); + + compute_row_sum_kernel<<>>( + Q_norms.data(), Q_coo_rows, Q_data, Q_nnz); + compute_row_sum_kernel<<>>( + R_norms.data(), R_coo_rows, R_data, R_nnz); + + compute_correlation(out, Q_sq_norms.data(), R_sq_norms.data(), Q_norms.data(), + R_norms.data(), m, n, n_cols, stream); +} + /** * L2 distance using the expanded form: sum(x_k)^2 + sum(y_k)^2 - 2 * sum(x_k * y_k) * The expanded form is more efficient for sparse data. @@ -133,9 +212,8 @@ class l2_expanded_distances_t : public distances_t { value_idx *b_indices = ip_dists.b_rows_coo(); value_t *b_data = ip_dists.b_data_coo(); - raft::mr::device::buffer search_coo_rows( - config_->handle.get_device_allocator(), config_->handle.get_stream(), - config_->a_nnz); + rmm::device_uvector search_coo_rows( + config_->a_nnz, config_->handle.get_stream()); raft::sparse::convert::csr_to_coo(config_->a_indptr, config_->a_nrows, search_coo_rows.data(), config_->a_nnz, config_->handle.get_stream()); @@ -183,6 +261,39 @@ class l2_sqrt_expanded_distances_t ~l2_sqrt_expanded_distances_t() = default; }; +template +class correlation_expanded_distances_t : public distances_t { + public: + explicit correlation_expanded_distances_t( + const distances_config_t &config) + : config_(&config), ip_dists(config) {} + + void compute(value_t *out_dists) { + ip_dists.compute(out_dists); + + value_idx *b_indices = ip_dists.b_rows_coo(); + value_t *b_data = ip_dists.b_data_coo(); + + rmm::device_uvector search_coo_rows( + config_->a_nnz, config_->handle.get_stream()); + raft::sparse::convert::csr_to_coo(config_->a_indptr, config_->a_nrows, + search_coo_rows.data(), config_->a_nnz, + config_->handle.get_stream()); + + compute_corr(out_dists, search_coo_rows.data(), config_->a_data, + config_->a_nnz, b_indices, b_data, config_->b_nnz, + config_->a_nrows, config_->b_nrows, config_->b_ncols, + config_->handle.get_device_allocator(), + config_->handle.get_stream()); + } + + ~correlation_expanded_distances_t() = default; + + protected: + const distances_config_t *config_; + ip_distances_t ip_dists; +}; + /** * Cosine distance using the expanded form: 1 - ( sum(x_k * y_k) / (sqrt(sum(x_k)^2) * sqrt(sum(y_k)^2))) * The expanded form is more efficient for sparse data. @@ -193,8 +304,7 @@ class cosine_expanded_distances_t : public distances_t { explicit cosine_expanded_distances_t( const distances_config_t &config) : config_(&config), - workspace(config.handle.get_device_allocator(), - config.handle.get_stream(), 0), + workspace(0, config.handle.get_stream()), ip_dists(config) {} void compute(value_t *out_dists) { @@ -203,9 +313,8 @@ class cosine_expanded_distances_t : public distances_t { value_idx *b_indices = ip_dists.b_rows_coo(); value_t *b_data = ip_dists.b_data_coo(); - raft::mr::device::buffer search_coo_rows( - config_->handle.get_device_allocator(), config_->handle.get_stream(), - config_->a_nnz); + rmm::device_uvector search_coo_rows( + config_->a_nnz, config_->handle.get_stream()); raft::sparse::convert::csr_to_coo(config_->a_indptr, config_->a_nrows, search_coo_rows.data(), config_->a_nnz, config_->handle.get_stream()); @@ -229,7 +338,7 @@ class cosine_expanded_distances_t : public distances_t { private: const distances_config_t *config_; - raft::mr::device::buffer workspace; + rmm::device_uvector workspace; ip_distances_t ip_dists; }; @@ -247,14 +356,11 @@ class hellinger_expanded_distances_t : public distances_t { public: explicit hellinger_expanded_distances_t( const distances_config_t &config) - : config_(&config), - workspace(config.handle.get_device_allocator(), - config.handle.get_stream(), 0) {} + : config_(&config), workspace(0, config.handle.get_stream()) {} void compute(value_t *out_dists) { - raft::mr::device::buffer coo_rows( - config_->handle.get_device_allocator(), config_->handle.get_stream(), - max(config_->b_nnz, config_->a_nnz)); + rmm::device_uvector coo_rows(max(config_->b_nnz, config_->a_nnz), + config_->handle.get_stream()); raft::sparse::convert::csr_to_coo(config_->b_indptr, config_->b_nrows, coo_rows.data(), config_->b_nnz, @@ -279,7 +385,43 @@ class hellinger_expanded_distances_t : public distances_t { private: const distances_config_t *config_; - raft::mr::device::buffer workspace; + rmm::device_uvector workspace; +}; + +template +class russelrao_expanded_distances_t : public distances_t { + public: + explicit russelrao_expanded_distances_t( + const distances_config_t &config) + : config_(&config), + workspace(0, config.handle.get_stream()), + ip_dists(config) {} + + void compute(value_t *out_dists) { + ip_dists.compute(out_dists); + + value_t n_cols = config_->a_ncols; + value_t n_cols_inv = 1.0 / n_cols; + raft::linalg::unaryOp( + out_dists, out_dists, config_->a_nrows * config_->b_nrows, + [=] __device__(value_t input) { return (n_cols - input) * n_cols_inv; }, + config_->handle.get_stream()); + + auto exec_policy = rmm::exec_policy(config_->handle.get_stream()); + auto diags = thrust::counting_iterator(0); + value_idx b_nrows = config_->b_nrows; + thrust::for_each(exec_policy, diags, diags + config_->a_nrows, + [=] __device__(value_idx input) { + out_dists[input * b_nrows + input] = 0.0; + }); + } + + ~russelrao_expanded_distances_t() = default; + + private: + const distances_config_t *config_; + rmm::device_uvector workspace; + ip_distances_t ip_dists; }; }; // END namespace distance diff --git a/cpp/include/raft/sparse/distance/lp_distance.cuh b/cpp/include/raft/sparse/distance/lp_distance.cuh index 653dc55683..885d55ee50 100644 --- a/cpp/include/raft/sparse/distance/lp_distance.cuh +++ b/cpp/include/raft/sparse/distance/lp_distance.cuh @@ -23,8 +23,7 @@ #include #include -#include -#include +#include #include #include @@ -44,9 +43,8 @@ template *config_, product_f product_func, accum_f accum_func, write_f write_func) { - raft::mr::device::buffer coo_rows( - config_->handle.get_device_allocator(), config_->handle.get_stream(), - max(config_->b_nnz, config_->a_nnz)); + rmm::device_uvector coo_rows(max(config_->b_nnz, config_->a_nnz), + config_->handle.get_stream()); raft::sparse::convert::csr_to_coo(config_->b_indptr, config_->b_nrows, coo_rows.data(), config_->b_nnz, @@ -185,6 +183,95 @@ class lp_unexpanded_distances_t : public distances_t { const distances_config_t *config_; value_t p; }; + +template +class hamming_unexpanded_distances_t : public distances_t { + public: + explicit hamming_unexpanded_distances_t( + const distances_config_t &config) + : config_(&config) {} + + void compute(value_t *out_dists) { + unexpanded_lp_distances(out_dists, config_, NotEqual(), + Sum(), AtomicAdd()); + + value_t n_cols = 1.0 / config_->a_ncols; + raft::linalg::unaryOp( + out_dists, out_dists, config_->a_nrows * config_->b_nrows, + [=] __device__(value_t input) { return input * n_cols; }, + config_->handle.get_stream()); + } + + private: + const distances_config_t *config_; +}; + +template +class jensen_shannon_unexpanded_distances_t : public distances_t { + public: + explicit jensen_shannon_unexpanded_distances_t( + const distances_config_t &config) + : config_(&config) {} + + void compute(value_t *out_dists) { + unexpanded_lp_distances( + out_dists, config_, + [] __device__(value_t a, value_t b) { + value_t m = 0.5f * (a + b); + bool a_zero = a == 0; + bool b_zero = b == 0; + + value_t x = (!a_zero * m) / (a_zero + a); + value_t y = (!b_zero * m) / (b_zero + b); + + bool x_zero = x == 0; + bool y_zero = y == 0; + + return (-a * (!x_zero * log(x + x_zero))) + + (-b * (!y_zero * log(y + y_zero))); + }, + Sum(), AtomicAdd()); + + raft::linalg::unaryOp( + out_dists, out_dists, config_->a_nrows * config_->b_nrows, + [=] __device__(value_t input) { return sqrt(0.5 * input); }, + config_->handle.get_stream()); + } + + private: + const distances_config_t *config_; +}; + +template +class kl_divergence_unexpanded_distances_t : public distances_t { + public: + explicit kl_divergence_unexpanded_distances_t( + const distances_config_t &config) + : config_(&config) {} + + void compute(value_t *out_dists) { + rmm::device_uvector coo_rows(max(config_->b_nnz, config_->a_nnz), + config_->handle.get_stream()); + + raft::sparse::convert::csr_to_coo(config_->b_indptr, config_->b_nrows, + coo_rows.data(), config_->b_nnz, + config_->handle.get_stream()); + + balanced_coo_pairwise_generalized_spmv( + out_dists, *config_, coo_rows.data(), + [] __device__(value_t a, value_t b) { return a * log(a / b); }, Sum(), + AtomicAdd()); + + raft::linalg::unaryOp( + out_dists, out_dists, config_->a_nrows * config_->b_nrows, + [=] __device__(value_t input) { return 0.5 * input; }, + config_->handle.get_stream()); + } + + private: + const distances_config_t *config_; +}; + }; // END namespace distance }; // END namespace sparse }; // END namespace raft diff --git a/cpp/test/sparse/distance.cu b/cpp/test/sparse/distance.cu index 9c2f9a4e27..0589637061 100644 --- a/cpp/test/sparse/distance.cu +++ b/cpp/test/sparse/distance.cu @@ -768,6 +768,102 @@ const std::vector> inputs_i32_f = { }, raft::distance::DistanceType::L1, 0.0}, + {5, + {0, 3, 8, 12, 16, 20, 25, 30, 35, 40, 45}, + {0, 3, 4, 0, 1, 2, 3, 4, 1, 2, 3, 4, 0, 2, 3, 4, 0, 1, 3, 4, 0, 1, 2, + 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4, 0, 1, 2, 3, 4}, + {0.70862347, 0.8232774, 0.12108795, 0.84527547, 0.94937088, 0.03258545, + 0.99584118, 0.76835667, 0.34426657, 0.2357925, 0.01274851, 0.11422017, + 0.3437756, 0.31967718, 0.5956055, 0.31610373, 0.04147273, 0.03724415, + 0.21515727, 0.04751052, 0.50283183, 0.99957274, 0.01395933, 0.96032529, + 0.88438711, 0.46095378, 0.27432481, 0.54294211, 0.54280225, 0.59503329, + 0.61364678, 0.22837736, 0.56609561, 0.29809423, 0.76736686, 0.56460608, + 0.98165371, 0.02140123, 0.19881268, 0.26057815, 0.31648823, 0.89874295, + 0.27366735, 0.5119944, 0.11416134}, + {// dense output + 0., 0.48769777, 1.88014197, 0.26127048, 0.26657011, 0.7874794, + 0.76962708, 1.122858, 1.1232498, 1.08166081, 0.48769777, 0., + 1.31332116, 0.98318907, 0.42661815, 0.09279052, 1.35187836, 1.38429055, + 0.40658897, 0.56136388, 1.88014197, 1.31332116, 0., 1.82943642, + 1.54826077, 1.05918884, 1.59360067, 1.34698954, 0.60215168, 0.46993848, + 0.26127048, 0.98318907, 1.82943642, 0., 0.29945563, 1.08494093, + 0.22934281, 0.82801925, 1.74288748, 1.50610116, 0.26657011, 0.42661815, + 1.54826077, 0.29945563, 0., 0.45060069, 0.77814948, 1.45245711, + 1.18328348, 0.82486987, 0.7874794, 0.09279052, 1.05918884, 1.08494093, + 0.45060069, 0., 1.29899154, 1.40683824, 0.48505269, 0.53862363, + 0.76962708, 1.35187836, 1.59360067, 0.22934281, 0.77814948, 1.29899154, + 0., 0.33202426, 1.92108999, 1.88812175, 1.122858, 1.38429055, + 1.34698954, 0.82801925, 1.45245711, 1.40683824, 0.33202426, 0., + 1.47318624, 1.92660889, 1.1232498, 0.40658897, 0.60215168, 1.74288748, + 1.18328348, 0.48505269, 1.92108999, 1.47318624, 0., 0.24992619, + 1.08166081, 0.56136388, 0.46993848, 1.50610116, 0.82486987, 0.53862363, + 1.88812175, 1.92660889, 0.24992619, 0.}, + raft::distance::DistanceType::CorrelationExpanded, + 0.0}, + {5, + {0, 1, 2, 4, 4, 5, 6, 7, 9, 9, 10}, + {1, 4, 0, 4, 1, 3, 0, 1, 3, 0}, + {1., 1., 1., 1., 1., 1., 1., 1., 1., 1.}, + {// dense output + 0., 1., 1., 1., 0.8, 1., 1., 0.8, 1., 1., 1., 0., 0.8, 1., 1., 1., 1., + 1., 1., 1., 1., 0.8, 0., 1., 1., 1., 0.8, 1., 1., 0.8, 1., 1., 1., 0., + 1., 1., 1., 1., 1., 1., 0.8, 1., 1., 1., 0., 1., 1., 0.8, 1., 1., 1., + 1., 1., 1., 1., 0., 1., 0.8, 1., 1., 1., 1., 0.8, 1., 1., 1., 0., 1., + 1., 0.8, 0.8, 1., 1., 1., 0.8, 0.8, 1., 0., 1., 1., 1., 1., 1., 1., 1., + 1., 1., 1., 0., 1., 1., 1., 0.8, 1., 1., 1., 0.8, 1., 1., 0.}, + raft::distance::DistanceType::RusselRaoExpanded, + 0.0}, + {5, + {0, 1, 1, 3, 3, 4, 4, 6, 9, 10, 10}, + {0, 3, 4, 4, 2, 3, 0, 2, 3, 2}, + {1., 1., 1., 1., 1., 1., 1., 1., 1., 1.}, + {// dense output + 0., 0.2, 0.6, 0.2, 0.4, 0.2, 0.6, 0.4, 0.4, 0.2, 0.2, 0., 0.4, 0., 0.2, + 0., 0.4, 0.6, 0.2, 0., 0.6, 0.4, 0., 0.4, 0.2, 0.4, 0.4, 0.6, 0.6, 0.4, + 0.2, 0., 0.4, 0., 0.2, 0., 0.4, 0.6, 0.2, 0., 0.4, 0.2, 0.2, 0.2, 0., + 0.2, 0.6, 0.8, 0.4, 0.2, 0.2, 0., 0.4, 0., 0.2, 0., 0.4, 0.6, 0.2, 0., + 0.6, 0.4, 0.4, 0.4, 0.6, 0.4, 0., 0.2, 0.2, 0.4, 0.4, 0.6, 0.6, 0.6, 0.8, + 0.6, 0.2, 0., 0.4, 0.6, 0.4, 0.2, 0.6, 0.2, 0.4, 0.2, 0.2, 0.4, 0., 0.2, + 0.2, 0., 0.4, 0., 0.2, 0., 0.4, 0.6, 0.2, 0.}, + raft::distance::DistanceType::HammingUnexpanded, + 0.0}, + {3, + {0, 1, 2}, + {0, 1}, + {1.0, 1.0}, + {0.0, 0.83255, 0.83255, 0.0}, + raft::distance::DistanceType::JensenShannon, + 0.0}, + {2, + {0, 1, 3}, + {0, 0, 1}, + {1.0, 0.5, 0.5}, + {0, 0.4645014, 0.4645014, 0}, + raft::distance::DistanceType::JensenShannon, + 0.0}, + {3, + {0, 1, 2}, + {0, 0}, + {1.0, 1.0}, + {0.0, 0.0, 0.0, 0.0}, + raft::distance::DistanceType::JensenShannon, + 0.0}, + + {3, + {0, 1, 2}, + {0, 1}, + {1.0, 1.0}, + {0.0, 1.0, 1.0, 0.0}, + raft::distance::DistanceType::DiceExpanded, + 0.0}, + {3, + {0, 1, 3}, + {0, 0, 1}, + {1.0, 1.0, 1.0}, + {0, 0.333333, 0.333333, 0}, + raft::distance::DistanceType::DiceExpanded, + 0.0}, + }; typedef SparseDistanceTest SparseDistanceTestF;