diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index b972fc1226..8ede967064 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -467,9 +467,8 @@ if(NOT BUILD_CPU_ONLY) # Keep cuVS as lightweight as possible. Only CUDA libs and rmm should be used in global target. target_link_libraries( cuvs - PUBLIC rmm::rmm raft::raft - PRIVATE nvidia::cutlass::cutlass ${CUVS_CTK_MATH_DEPENDENCIES} - $ + PUBLIC rmm::rmm raft::raft ${CUVS_CTK_MATH_DEPENDENCIES} + PRIVATE nvidia::cutlass::cutlass $ ) endif() @@ -573,8 +572,8 @@ if(BUILD_C_LIBRARY) target_link_libraries( cuvs_c - PUBLIC cuvs::cuvs - PRIVATE raft::raft ${CUVS_CTK_MATH_DEPENDENCIES} + PUBLIC cuvs::cuvs ${CUVS_CTK_MATH_DEPENDENCIES} + PRIVATE raft::raft ) # ensure CUDA symbols aren't relocated to the middle of the debug build binaries diff --git a/cpp/src/cluster/detail/mst.cuh b/cpp/src/cluster/detail/mst.cuh index de88b954e2..5804b8b5d9 100644 --- a/cpp/src/cluster/detail/mst.cuh +++ b/cpp/src/cluster/detail/mst.cuh @@ -16,8 +16,8 @@ #pragma once +#include "../../sparse/neighbors/cross_component_nn.cuh" #include -#include #include #include #include @@ -86,7 +86,7 @@ void connect_knn_graph( static constexpr size_t default_row_batch_size = 4096; static constexpr size_t default_col_batch_size = 16; - raft::sparse::neighbors::cross_component_nn(handle, + cuvs::sparse::neighbors::cross_component_nn(handle, connected_edges, X, color, @@ -166,14 +166,14 @@ void build_sorted_mst( handle, indptr, indices, pw_dists, (value_idx)m, nnz, color, stream, false, true); int iters = 1; - int n_components = raft::sparse::neighbors::get_n_components(color, m, stream); + int n_components = cuvs::sparse::neighbors::get_n_components(color, m, stream); while (n_components > 1 && iters < max_iter) { connect_knn_graph(handle, X, mst_coo, m, n, color, reduction_op); iters++; - n_components = raft::sparse::neighbors::get_n_components(color, m, stream); + n_components = cuvs::sparse::neighbors::get_n_components(color, m, stream); } /** diff --git a/cpp/src/cluster/detail/single_linkage.cuh b/cpp/src/cluster/detail/single_linkage.cuh index 837acf797e..8b90336b71 100644 --- a/cpp/src/cluster/detail/single_linkage.cuh +++ b/cpp/src/cluster/detail/single_linkage.cuh @@ -81,7 +81,7 @@ void single_linkage(raft::resources const& handle, * 2. Construct MST, sorted by weights */ rmm::device_uvector color(m, stream); - raft::sparse::neighbors::FixConnectivitiesRedOp op(m); + cuvs::sparse::neighbors::FixConnectivitiesRedOp op(m); detail::build_sorted_mst(handle, X, indptr.data(), diff --git a/cpp/src/distance/detail/compress_to_bits.cuh b/cpp/src/distance/detail/compress_to_bits.cuh index 9ce47774a6..0f02631169 100644 --- a/cpp/src/distance/detail/compress_to_bits.cuh +++ b/cpp/src/distance/detail/compress_to_bits.cuh @@ -15,6 +15,7 @@ */ #pragma once +#include #include #include #include @@ -100,7 +101,7 @@ void compress_to_bits(raft::resources const& handle, raft::device_matrix_view in, raft::device_matrix_view out) { - auto stream = resource::get_cuda_stream(handle); + auto stream = raft::resource::get_cuda_stream(handle); constexpr int bits_per_element = 8 * sizeof(T); RAFT_EXPECTS(raft::ceildiv(in.extent(0), bits_per_element) == out.extent(0), diff --git a/cpp/src/distance/detail/masked_distance_base.cuh b/cpp/src/distance/detail/masked_distance_base.cuh index d92052c84c..2c41ee3be0 100644 --- a/cpp/src/distance/detail/masked_distance_base.cuh +++ b/cpp/src/distance/detail/masked_distance_base.cuh @@ -14,7 +14,7 @@ * limitations under the License. */ #pragma once -#include "../pairwise_distance_base.cuh" +#include "pairwise_distance_base.cuh" #include #include diff --git a/cpp/src/distance/detail/masked_nn.cuh b/cpp/src/distance/detail/masked_nn.cuh index 6520b1e2e9..e10d2b7c69 100644 --- a/cpp/src/distance/detail/masked_nn.cuh +++ b/cpp/src/distance/detail/masked_nn.cuh @@ -16,9 +16,9 @@ #pragma once -#include "../compress_to_bits.cuh" -#include "../fused_distance_nn/fused_l2_nn.cuh" -#include "../masked_distance_base.cuh" +#include "compress_to_bits.cuh" +#include "fused_distance_nn/fused_l2_nn.cuh" +#include "masked_distance_base.cuh" #include #include #include @@ -251,14 +251,14 @@ void masked_l2_nn_impl(raft::resources const& handle, bool sqrt, bool initOutBuffer) { - typedef typename linalg::Policy4x4::Policy P; + typedef typename raft::linalg::Policy4x4::Policy P; static_assert(P::Mblk == 64, "masked_l2_nn_impl only supports a policy with 64 rows per block."); // Get stream and workspace memory resource rmm::mr::device_memory_resource* ws_mr = - dynamic_cast(resource::get_workspace_resource(handle)); - auto stream = resource::get_cuda_stream(handle); + dynamic_cast(raft::resource::get_workspace_resource(handle)); + auto stream = raft::resource::get_cuda_stream(handle); // Acquire temporary buffers and initialize to zero: // 1) Adjacency matrix bitfield diff --git a/cpp/src/distance/detail/pairwise_matrix/dispatch-ext.cuh b/cpp/src/distance/detail/pairwise_matrix/dispatch-ext.cuh index 3e8402f1f9..bc8189c709 100644 --- a/cpp/src/distance/detail/pairwise_matrix/dispatch-ext.cuh +++ b/cpp/src/distance/detail/pairwise_matrix/dispatch-ext.cuh @@ -22,7 +22,7 @@ #include // raft::identity_op #include // RAFT_EXPLICIT -#ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY +#ifdef CUVS_EXPLICIT_INSTANTIATE_ONLY namespace cuvs::distance::detail { diff --git a/cpp/src/distance/detail/pairwise_matrix/dispatch.cuh b/cpp/src/distance/detail/pairwise_matrix/dispatch.cuh index 4a52b7ebe7..06b039c3a1 100644 --- a/cpp/src/distance/detail/pairwise_matrix/dispatch.cuh +++ b/cpp/src/distance/detail/pairwise_matrix/dispatch.cuh @@ -15,7 +15,7 @@ */ #pragma once -#ifndef RAFT_EXPLICIT_INSTANTIATE_ONLY +#ifndef CUVS_EXPLICIT_INSTANTIATE_ONLY #include "dispatch-inl.cuh" #endif diff --git a/cpp/src/distance/distance-ext.cuh b/cpp/src/distance/distance-ext.cuh index e574c11b0c..148951afad 100644 --- a/cpp/src/distance/distance-ext.cuh +++ b/cpp/src/distance/distance-ext.cuh @@ -24,7 +24,7 @@ #include // rmm::device_uvector -#ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY +#ifdef CUVS_EXPLICIT_INSTANTIATE_ONLY namespace cuvs { namespace distance { diff --git a/cpp/src/distance/distance.cuh b/cpp/src/distance/distance.cuh index de70cd4691..b5bfc07cb2 100644 --- a/cpp/src/distance/distance.cuh +++ b/cpp/src/distance/distance.cuh @@ -15,7 +15,7 @@ */ #pragma once -#ifndef RAFT_EXPLICIT_INSTANTIATE_ONLY +#ifndef CUVS_EXPLICIT_INSTANTIATE_ONLY #include "distance-inl.cuh" #endif diff --git a/cpp/src/distance/masked_nn.cuh b/cpp/src/distance/masked_nn.cuh new file mode 100644 index 0000000000..82dd189196 --- /dev/null +++ b/cpp/src/distance/masked_nn.cuh @@ -0,0 +1,201 @@ +/* + * 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. + */ + +#ifndef __MASKED_L2_NN_H +#define __MASKED_L2_NN_H + +#pragma once + +#include "detail/masked_nn.cuh" +#include "fused_distance_nn_helpers.cuh" +#include +#include + +#include + +#include + +namespace cuvs { +namespace distance { +/** + * \defgroup masked_nn Masked 1-nearest neighbors + * @{ + */ + +/** + * @brief Parameter struct for masked_l2_nn function + * + * @tparam ReduceOpT Type of reduction operator in the epilogue. + * @tparam KVPReduceOpT Type of Reduction operation on key value pairs. + * + * Usage example: + * @code{.cpp} + * #include + * + * using IdxT = int; + * using DataT = float; + * using RedOpT = cuvs::distance::MinAndDistanceReduceOp; + * using PairRedOpT = cuvs::distance::KVPMinReduce; + * using ParamT = cuvs::distance::masked_l2_nn_params; + * + * bool init_out = true; + * bool sqrt = false; + * + * ParamT masked_l2_params{RedOpT{}, PairRedOpT{}, sqrt, init_out}; + * @endcode + * + * Prescribes how to reduce a distance to an intermediate type (`redOp`), and + * how to reduce two intermediate types (`pairRedOp`). Typically, a distance is + * mapped to an (index, value) pair and (index, value) pair with the lowest + * value (distance) is selected. + * + * In addition, prescribes whether to compute the square root of the distance + * (`sqrt`) and whether to initialize the output buffer (`initOutBuffer`). + */ +template +struct masked_l2_nn_params { + /** Reduction operator in the epilogue */ + ReduceOpT redOp; + /** Reduction operation on key value pairs */ + KVPReduceOpT pairRedOp; + /** Whether the output `minDist` should contain L2-sqrt */ + bool sqrt; + /** Whether to initialize the output buffer before the main kernel launch */ + bool initOutBuffer; +}; + +/** + * @brief Masked L2 distance and 1-nearest-neighbor computation in a single call. + * + * This function enables faster computation of nearest neighbors if the + * computation of distances between certain point pairs can be skipped. + * + * We use an adjacency matrix that describes which distances to calculate. The + * points in `y` are divided into groups, and the adjacency matrix indicates + * whether to compute distances between points in `x` and groups in `y`. In other + * words, if `adj[i,k]` is true then distance between point `x_i`, and points in + * `group_k` will be calculated. + * + * **Performance considerations** + * + * The points in `x` are processed in tiles of `M` points (`M` is currently 64, + * but may change in the future). As a result, the largest compute time + * reduction occurs if all `M` points can skip a group. If only part of the `M` + * points can skip a group, then at most a minor compute time reduction and a + * modest energy use reduction can be expected. + * + * The points in `y` are also grouped into tiles of `N` points (`N` is currently + * 64, but may change in the future). As a result, group sizes should be larger + * than `N` to avoid wasting computational resources. If the group sizes are + * evenly divisible by `N`, then the computation is most efficient, although for + * larger group sizes this effect is minor. + * + * + * **Comparison to SDDM** + * + * [SDDMM](https://ieeexplore.ieee.org/document/8638042) (sampled dense-dense + * matrix multiplication) is a matrix-matrix multiplication where only part of + * the output is computed. Compared to masked_l2_nn, there are a few differences: + * + * - The output of masked_l2_nn is a single vector (of nearest neighbors) and not + * a sparse matrix. + * + * - The sampling in masked_l2_nn is expressed through intermediate "groups" + rather than a CSR format. + * + * @tparam DataT data type + * @tparam OutT output type to either store 1-NN indices and their minimum + * distances or store only the min distances. Accordingly, one + * has to pass an appropriate `ReduceOpT` + * @tparam IdxT indexing arithmetic type + * @tparam ReduceOpT A struct to perform the final needed reduction operation + * and also to initialize the output array elements with the + * appropriate initial value needed for reduction. + * + * @param handle RAFT handle for managing expensive resources + * @param params Parameter struct specifying the reduction operations. + * @param[in] x First matrix. Row major. Dim = `m x k`. + * (on device). + * @param[in] y Second matrix. Row major. Dim = `n x k`. + * (on device). + * @param[in] x_norm L2 squared norm of `x`. Length = `m`. (on device). + * @param[in] y_norm L2 squared norm of `y`. Length = `n`. (on device) + * @param[in] adj A boolean adjacency matrix indicating for each + * row of `x` and each group in `y` whether to compute the + * distance. Dim = `m x num_groups`. + * @param[in] group_idxs An array containing the *end* indices of each group + * in `y`. The value of group_idxs[j] indicates the + * start of group j + 1, i.e., it is the inclusive + * scan of the group lengths. The first group is + * always assumed to start at index 0 and the last + * group typically ends at index `n`. Length = + * `num_groups`. + * @param[out] out will contain the reduced output (Length = `m`) + * (on device) + */ +template +void masked_l2_nn(raft::resources const& handle, + cuvs::distance::masked_l2_nn_params params, + raft::device_matrix_view x, + raft::device_matrix_view y, + raft::device_vector_view x_norm, + raft::device_vector_view y_norm, + raft::device_matrix_view adj, + raft::device_vector_view group_idxs, + raft::device_vector_view out) +{ + IdxT m = x.extent(0); + IdxT n = y.extent(0); + IdxT k = x.extent(1); + IdxT num_groups = group_idxs.extent(0); + + // Match k dimension of x, y + RAFT_EXPECTS(x.extent(1) == y.extent(1), "Dimension of vectors in x and y must be equal."); + // Match x, x_norm and y, y_norm + RAFT_EXPECTS(m == x_norm.extent(0), "Length of `x_norm` must match input `x`."); + RAFT_EXPECTS(n == y_norm.extent(0), "Length of `y_norm` must match input `y` "); + // Match adj to x and group_idxs + RAFT_EXPECTS(m == adj.extent(0), "#rows in `adj` must match input `x`."); + RAFT_EXPECTS(num_groups == adj.extent(1), "#cols in `adj` must match length of `group_idxs`."); + // NOTE: We do not check if all indices in group_idxs actually points *inside* y. + + // If there is no work to be done, return immediately. + if (m == 0 || n == 0 || k == 0 || num_groups == 0) { return; } + + detail::masked_l2_nn_impl(handle, + out.data_handle(), + x.data_handle(), + y.data_handle(), + x_norm.data_handle(), + y_norm.data_handle(), + adj.data_handle(), + group_idxs.data_handle(), + num_groups, + m, + n, + k, + params.redOp, + params.pairRedOp, + params.sqrt, + params.initOutBuffer); +} + +/** @} */ + +} // namespace distance +} // namespace cuvs + +#endif diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh b/cpp/src/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh index 63f3838d09..495ec6a4d8 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh @@ -24,7 +24,7 @@ namespace cuvs::neighbors::cagra::detail { namespace multi_cta_search { -#ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY +#ifdef CUVS_EXPLICIT_INSTANTIATE_ONLY template #endif -#ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY +#ifdef CUVS_EXPLICIT_INSTANTIATE_ONLY namespace cuvs::neighbors::detail { diff --git a/cpp/src/neighbors/refine-ext.cuh b/cpp/src/neighbors/refine-ext.cuh index 5b686e1abf..2f40842871 100644 --- a/cpp/src/neighbors/refine-ext.cuh +++ b/cpp/src/neighbors/refine-ext.cuh @@ -24,7 +24,7 @@ #include // int64_t -#ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY +#ifdef CUVS_EXPLICIT_INSTANTIATE_ONLY namespace cuvs::neighbors { diff --git a/cpp/src/neighbors/refine.cuh b/cpp/src/neighbors/refine.cuh index 15f2b02928..b3a3cdb4eb 100644 --- a/cpp/src/neighbors/refine.cuh +++ b/cpp/src/neighbors/refine.cuh @@ -15,7 +15,7 @@ */ #pragma once -#ifndef RAFT_EXPLICIT_INSTANTIATE_ONLY +#ifndef CUVS_EXPLICIT_INSTANTIATE_ONLY #include "refine-inl.cuh" #endif diff --git a/cpp/src/sparse/neighbors/cross_component_nn.cuh b/cpp/src/sparse/neighbors/cross_component_nn.cuh new file mode 100644 index 0000000000..36a0c79f59 --- /dev/null +++ b/cpp/src/sparse/neighbors/cross_component_nn.cuh @@ -0,0 +1,99 @@ +/* + * 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. + */ + +#pragma once + +#include "detail/cross_component_nn.cuh" +#include +#include +#include + +namespace cuvs::sparse::neighbors { + +template +using FixConnectivitiesRedOp = detail::FixConnectivitiesRedOp; + +/** + * Gets the number of unique components from array of + * colors or labels. This does not assume the components are + * drawn from a monotonically increasing set. + * @tparam value_idx + * @param[in] colors array of components + * @param[in] n_rows size of components array + * @param[in] stream cuda stream for which to order cuda operations + * @return total number of components + */ +template +value_idx get_n_components(value_idx* colors, size_t n_rows, cudaStream_t stream) +{ + return detail::get_n_components(colors, n_rows, stream); +} + +/** + * Connects the components of an otherwise unconnected knn graph + * by computing a 1-nn to neighboring components of each data point + * (e.g. component(nn) != component(self)) and reducing the results to + * include the set of smallest destination components for each source + * component. The result will not necessarily contain + * n_components^2 - n_components number of elements because many components + * will likely not be contained in the neighborhoods of 1-nns. + * @tparam value_idx + * @tparam value_t + * @param[in] handle raft handle + * @param[out] out output edge list containing nearest cross-component + * edges. + * @param[in] X original (row-major) dense matrix for which knn graph should be constructed. + * @param[in] orig_colors array containing component number for each row of X + * @param[in] n_rows number of rows in X + * @param[in] n_cols number of cols in X + * @param[in] reduction_op reduction operation for computing nearest neighbors. The reduction + * operation must have `gather` and `scatter` functions defined + * @param[in] row_batch_size the batch size for computing nearest neighbors. This parameter controls + * the number of samples for which the nearest neighbors are computed at once. Therefore, it affects + * the memory consumption mainly by reducing the size of the adjacency matrix for masked nearest + * neighbors computation + * @param[in] col_batch_size the input data is sorted and 'unsorted' based on color. An additional + * scratch space buffer of shape (n_rows, col_batch_size) is created for this. Usually, this + * parameter affects the memory consumption more drastically than the row_batch_size with a marginal + * increase in compute time as the col_batch_size is reduced + * @param[in] metric distance metric + */ +template +void cross_component_nn( + raft::resources const& handle, + raft::sparse::COO& out, + const value_t* X, + const value_idx* orig_colors, + size_t n_rows, + size_t n_cols, + red_op reduction_op, + size_t row_batch_size = 0, + size_t col_batch_size = 0, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2SqrtExpanded) +{ + detail::cross_component_nn(handle, + out, + X, + orig_colors, + n_rows, + n_cols, + reduction_op, + row_batch_size, + col_batch_size, + metric); +} + +}; // end namespace cuvs::sparse::neighbors \ No newline at end of file diff --git a/cpp/src/sparse/neighbors/detail/cross_component_nn.cuh b/cpp/src/sparse/neighbors/detail/cross_component_nn.cuh new file mode 100644 index 0000000000..43bb6d8a81 --- /dev/null +++ b/cpp/src/sparse/neighbors/detail/cross_component_nn.cuh @@ -0,0 +1,542 @@ +/* + * 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. + */ +#pragma once + +#include "../../../distance/masked_nn.cuh" +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace cuvs::sparse::neighbors::detail { + +/** + * Base functor with reduction ops for performing masked 1-nn + * computation. + * @tparam value_idx + * @tparam value_t + */ +template +struct FixConnectivitiesRedOp { + value_idx m; + + // default constructor for cutlass + DI FixConnectivitiesRedOp() : m(0) {} + + FixConnectivitiesRedOp(value_idx m_) : m(m_){}; + + typedef typename raft::KeyValuePair KVP; + DI void operator()(value_idx rit, KVP* out, const KVP& other) const + { + if (rit < m && other.value < out->value) { + out->key = other.key; + out->value = other.value; + } + } + + DI KVP operator()(value_idx rit, const KVP& a, const KVP& b) const + { + if (rit < m && a.value < b.value) { + return a; + } else + return b; + } + + DI void init(value_t* out, value_t maxVal) const { *out = maxVal; } + DI void init(KVP* out, value_t maxVal) const + { + out->key = -1; + out->value = maxVal; + } + + DI void init_key(value_t& out, value_idx idx) const { return; } + DI void init_key(KVP& out, value_idx idx) const { out.key = idx; } + + DI value_t get_value(KVP& out) const { return out.value; } + + DI value_t get_value(value_t& out) const { return out; } + + /** The gather and scatter ensure that operator() is still consistent after rearranging the data. + * TODO (tarang-jain): refactor cross_component_nn API to separate out the gather and scatter + * functions from the reduction op. Reference: https://github.com/rapidsai/raft/issues/1614 */ + void gather(const raft::resources& handle, value_idx* map) {} + + void scatter(const raft::resources& handle, value_idx* map) {} +}; + +/** + * Assumes 3-iterator tuple containing COO rows, cols, and + * a cub keyvalue pair object. Sorts the 3 arrays in + * ascending order: row->col->keyvaluepair + */ +struct TupleComp { + template + __host__ __device__ bool operator()(const one& t1, const two& t2) + { + // sort first by each sample's color, + if (thrust::get<0>(t1) < thrust::get<0>(t2)) return true; + if (thrust::get<0>(t1) > thrust::get<0>(t2)) return false; + + // then by the color of each sample's closest neighbor, + if (thrust::get<1>(t1) < thrust::get<1>(t2)) return true; + if (thrust::get<1>(t1) > thrust::get<1>(t2)) return false; + + // then sort by value in descending order + return thrust::get<2>(t1).value < thrust::get<2>(t2).value; + } +}; + +template +struct CubKVPMinReduce { + 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 + +/** + * Gets the number of unique components from array of + * colors or labels. This does not assume the components are + * drawn from a monotonically increasing set. + * @tparam value_idx + * @param[in] colors array of components + * @param[in] n_rows size of components array + * @param[in] stream cuda stream for which to order cuda operations + * @return total number of components + */ +template +value_idx get_n_components(value_idx* colors, size_t n_rows, cudaStream_t stream) +{ + rmm::device_uvector map_ids(0, stream); + int num_clusters = raft::label::getUniquelabels(map_ids, colors, n_rows, stream); + return num_clusters; +} + +/** + * Functor to look up a component for a vertex + * @tparam value_idx + * @tparam value_t + */ +template +struct LookupColorOp { + value_idx* colors; + + LookupColorOp(value_idx* colors_) : colors(colors_) {} + + DI value_idx + + operator()(const raft::KeyValuePair& kvp) + { + return colors[kvp.key]; + } +}; + +/** + * Compute the cross-component 1-nearest neighbors for each row in X using + * the given array of components + * @tparam value_idx + * @tparam value_t + * @param[in] handle raft handle + * @param[out] kvp mapping of closest neighbor vertex and distance for each vertex in the given + * array of components + * @param[out] nn_colors components of nearest neighbors for each vertex + * @param[in] colors components of each vertex + * @param[in] X original dense data + * @param[in] n_rows number of rows in original dense data + * @param[in] n_cols number of columns in original dense data + * @param[in] row_batch_size row batch size for computing nearest neighbors + * @param[in] col_batch_size column batch size for sorting and 'unsorting' + * @param[in] reduction_op reduction operation for computing nearest neighbors + */ +template +void perform_1nn(raft::resources const& handle, + raft::KeyValuePair* kvp, + value_idx* nn_colors, + value_idx* colors, + const value_t* X, + size_t n_rows, + size_t n_cols, + size_t row_batch_size, + size_t col_batch_size, + red_op reduction_op) +{ + auto stream = raft::resource::get_cuda_stream(handle); + auto exec_policy = raft::resource::get_thrust_policy(handle); + + auto sort_plan = raft::make_device_vector(handle, (value_idx)n_rows); + raft::linalg::map_offset(handle, sort_plan.view(), [] __device__(value_idx idx) { return idx; }); + + thrust::sort_by_key( + raft::resource::get_thrust_policy(handle), colors, colors + n_rows, sort_plan.data_handle()); + + // Modify the reduction operation based on the sort plan. + reduction_op.gather(handle, sort_plan.data_handle()); + + auto X_mutable_view = + raft::make_device_matrix_view(const_cast(X), n_rows, n_cols); + auto sort_plan_const_view = + raft::make_device_vector_view(sort_plan.data_handle(), n_rows); + raft::matrix::gather(handle, X_mutable_view, sort_plan_const_view, (value_idx)col_batch_size); + + // Get the number of unique components from the array of colors + value_idx n_components = get_n_components(colors, n_rows, stream); + + // colors_group_idxs is an array containing the *end* indices of each color + // component in colors. That is, the value of colors_group_idxs[j] indicates + // the start of color j + 1, i.e., it is the inclusive scan of the sizes of + // the color components. + auto colors_group_idxs = raft::make_device_vector(handle, n_components + 1); + raft::sparse::convert::sorted_coo_to_csr( + colors, n_rows, colors_group_idxs.data_handle(), n_components + 1, stream); + + auto group_idxs_view = raft::make_device_vector_view( + colors_group_idxs.data_handle() + 1, n_components); + + auto x_norm = raft::make_device_vector(handle, (value_idx)n_rows); + raft::linalg::rowNorm( + x_norm.data_handle(), X, n_cols, n_rows, raft::linalg::L2Norm, true, stream); + + auto adj = raft::make_device_matrix(handle, row_batch_size, n_components); + using OutT = raft::KeyValuePair; + using ParamT = cuvs::distance::masked_l2_nn_params; + + bool apply_sqrt = true; + bool init_out_buffer = true; + ParamT params{reduction_op, reduction_op, apply_sqrt, init_out_buffer}; + + auto X_full_view = raft::make_device_matrix_view(X, n_rows, n_cols); + + size_t n_batches = raft::ceildiv(n_rows, row_batch_size); + + for (size_t bid = 0; bid < n_batches; bid++) { + size_t batch_offset = bid * row_batch_size; + size_t rows_per_batch = min(row_batch_size, n_rows - batch_offset); + + auto X_batch_view = raft::make_device_matrix_view( + X + batch_offset * n_cols, rows_per_batch, n_cols); + + auto x_norm_batch_view = raft::make_device_vector_view( + x_norm.data_handle() + batch_offset, rows_per_batch); + + auto mask_op = [colors, + n_components = raft::util::FastIntDiv(n_components), + batch_offset] __device__(value_idx idx) { + value_idx row = idx / n_components; + value_idx col = idx % n_components; + return colors[batch_offset + row] != col; + }; + + auto adj_vector_view = raft::make_device_vector_view( + adj.data_handle(), rows_per_batch * n_components); + + raft::linalg::map_offset(handle, adj_vector_view, mask_op); + + auto adj_view = raft::make_device_matrix_view( + adj.data_handle(), rows_per_batch, n_components); + + auto kvp_view = + raft::make_device_vector_view, value_idx>( + kvp + batch_offset, rows_per_batch); + + cuvs::distance::masked_l2_nn(handle, + params, + X_batch_view, + X_full_view, + x_norm_batch_view, + x_norm.view(), + adj_view, + group_idxs_view, + kvp_view); + } + + // Transform the keys so that they correctly point to the unpermuted indices. + thrust::transform(exec_policy, + kvp, + kvp + n_rows, + kvp, + [sort_plan = sort_plan.data_handle()] __device__(OutT KVP) { + OutT res; + res.value = KVP.value; + res.key = sort_plan[KVP.key]; + return res; + }); + + // Undo permutation of the rows of X by scattering in place. + raft::matrix::scatter(handle, X_mutable_view, sort_plan_const_view, (value_idx)col_batch_size); + + // Undo permutation of the key-value pair and color vectors. This is not done + // inplace, so using two temporary vectors. + auto tmp_colors = raft::make_device_vector(handle, n_rows); + auto tmp_kvp = raft::make_device_vector(handle, n_rows); + + thrust::scatter(exec_policy, kvp, kvp + n_rows, sort_plan.data_handle(), tmp_kvp.data_handle()); + thrust::scatter( + exec_policy, colors, colors + n_rows, sort_plan.data_handle(), tmp_colors.data_handle()); + reduction_op.scatter(handle, sort_plan.data_handle()); + + raft::copy_async(colors, tmp_colors.data_handle(), n_rows, stream); + raft::copy_async(kvp, tmp_kvp.data_handle(), n_rows, stream); + + LookupColorOp extract_colors_op(colors); + thrust::transform(exec_policy, kvp, kvp + n_rows, nn_colors, extract_colors_op); +} + +/** + * Sort nearest neighboring components wrt component of source vertices + * @tparam value_idx + * @tparam value_t + * @param[inout] colors components array of source vertices + * @param[inout] nn_colors nearest neighbors components array + * @param[inout] kvp nearest neighbor source vertex / distance array + * @param[inout] src_indices array of source vertex indices which will become arg_sort + * indices + * @param n_rows number of components in `colors` + * @param stream stream for which to order CUDA operations + */ +template +void sort_by_color(raft::resources const& handle, + value_idx* colors, + value_idx* nn_colors, + raft::KeyValuePair* kvp, + value_idx* src_indices, + size_t n_rows) +{ + auto exec_policy = raft::resource::get_thrust_policy(handle); + thrust::counting_iterator arg_sort_iter(0); + thrust::copy(exec_policy, arg_sort_iter, arg_sort_iter + n_rows, src_indices); + + auto keys = thrust::make_zip_iterator( + thrust::make_tuple(colors, nn_colors, (raft::KeyValuePair*)kvp)); + auto vals = thrust::make_zip_iterator(thrust::make_tuple(src_indices)); + // get all the colors in contiguous locations so we can map them to warps. + thrust::sort_by_key(exec_policy, keys, keys + n_rows, vals, TupleComp()); +} + +template +RAFT_KERNEL min_components_by_color_kernel(value_idx* out_rows, + value_idx* out_cols, + value_t* out_vals, + const value_idx* out_index, + const value_idx* indices, + const raft::KeyValuePair* kvp, + size_t nnz) +{ + size_t tid = blockDim.x * blockIdx.x + threadIdx.x; + + if (tid >= nnz) return; + + int idx = out_index[tid]; + + if ((tid == 0 || (out_index[tid - 1] != idx))) { + out_rows[idx] = indices[tid]; + out_cols[idx] = kvp[tid].key; + out_vals[idx] = kvp[tid].value; + } +} + +/** + * Computes the min set of unique components that neighbor the + * components of each source vertex. + * @tparam value_idx + * @tparam value_t + * @param[out] coo output edge list + * @param[in] out_index output indptr for ordering edge list + * @param[in] indices indices of source vertices for each component + * @param[in] kvp indices and distances of each destination vertex for each component + * @param[in] n_colors number of components + * @param[in] stream cuda stream for which to order cuda operations + */ +template +void min_components_by_color(raft::sparse::COO& coo, + const value_idx* out_index, + const value_idx* indices, + const raft::KeyValuePair* kvp, + size_t nnz, + cudaStream_t stream) +{ + /** + * Arrays should be ordered by: colors_indptr->colors_n->kvp.value + * so the last element of each column in the input CSR should be + * the min. + */ + min_components_by_color_kernel<<>>( + coo.rows(), coo.cols(), coo.vals(), out_index, indices, kvp, nnz); +} + +/** + * Connects the components of an otherwise unconnected knn graph + * by computing a 1-nn to neighboring components of each data point + * (e.g. component(nn) != component(self)) and reducing the results to + * include the set of smallest destination components for each source + * component. The result will not necessarily contain + * n_components^2 - n_components number of elements because many components + * will likely not be contained in the neighborhoods of 1-nns. + * @tparam value_idx + * @tparam value_t + * @param[in] handle raft handle + * @param[out] out output edge list containing nearest cross-component + * edges. + * @param[in] X original (row-major) dense matrix for which knn graph should be constructed. + * @param[in] orig_colors array containing component number for each row of X + * @param[in] n_rows number of rows in X + * @param[in] n_cols number of cols in X + * @param[in] reduction_op reduction operation for computing nearest neighbors. The reduction + * operation must have `gather` and `scatter` functions defined + * @param[in] row_batch_size the batch size for computing nearest neighbors. This parameter controls + * the number of samples for which the nearest neighbors are computed at once. Therefore, it affects + * the memory consumption mainly by reducing the size of the adjacency matrix for masked nearest + * neighbors computation. default 0 indicates that no batching is done + * @param[in] col_batch_size the input data is sorted and 'unsorted' based on color. An additional + * scratch space buffer of shape (n_rows, col_batch_size) is created for this. Usually, this + * parameter affects the memory consumption more drastically than the col_batch_size with a marginal + * increase in compute time as the col_batch_size is reduced. default 0 indicates that no batching + * is done + * @param[in] metric distance metric + */ +template +void cross_component_nn( + raft::resources const& handle, + raft::sparse::COO& out, + const value_t* X, + const value_idx* orig_colors, + size_t n_rows, + size_t n_cols, + red_op reduction_op, + size_t row_batch_size, + size_t col_batch_size, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2SqrtExpanded) +{ + auto stream = raft::resource::get_cuda_stream(handle); + + RAFT_EXPECTS(metric == cuvs::distance::DistanceType::L2SqrtExpanded, + "Fixing connectivities for an unconnected k-NN graph only " + "supports L2SqrtExpanded currently."); + + if (row_batch_size == 0 || row_batch_size > n_rows) { row_batch_size = n_rows; } + + if (col_batch_size == 0 || col_batch_size > n_cols) { col_batch_size = n_cols; } + + rmm::device_uvector colors(n_rows, stream); + + // Normalize colors so they are drawn from a monotonically increasing set + constexpr bool zero_based = true; + raft::label::make_monotonic( + colors.data(), const_cast(orig_colors), n_rows, stream, zero_based); + + /** + * First compute 1-nn for all colors where the color of each data point + * is guaranteed to be != color of its nearest neighbor. + */ + rmm::device_uvector nn_colors(n_rows, stream); + rmm::device_uvector> temp_inds_dists(n_rows, stream); + rmm::device_uvector src_indices(n_rows, stream); + + perform_1nn(handle, + temp_inds_dists.data(), + nn_colors.data(), + colors.data(), + X, + n_rows, + n_cols, + row_batch_size, + col_batch_size, + reduction_op); + + /** + * Sort data points by color (neighbors are not sorted) + */ + // max_color + 1 = number of connected components + // sort nn_colors by key w/ original colors + sort_by_color( + handle, colors.data(), nn_colors.data(), temp_inds_dists.data(), src_indices.data(), n_rows); + + /** + * Take the min for any duplicate colors + */ + // Compute mask of duplicates + rmm::device_uvector out_index(n_rows + 1, stream); + raft::sparse::op::compute_duplicates_mask( + out_index.data(), colors.data(), nn_colors.data(), n_rows, stream); + + thrust::exclusive_scan(raft::resource::get_thrust_policy(handle), + out_index.data(), + out_index.data() + out_index.size(), + out_index.data()); + + // compute final size + value_idx size = 0; + raft::update_host(&size, out_index.data() + (out_index.size() - 1), 1, stream); + raft::resource::sync_stream(handle, stream); + + size++; + + raft::sparse::COO min_edges(stream); + min_edges.allocate(size, n_rows, n_rows, true, stream); + + min_components_by_color( + min_edges, out_index.data(), src_indices.data(), temp_inds_dists.data(), n_rows, stream); + + /** + * Symmetrize resulting edge list + */ + raft::sparse::linalg::symmetrize( + handle, min_edges.rows(), min_edges.cols(), min_edges.vals(), n_rows, n_rows, size, out); +} + +}; // end namespace cuvs::sparse::neighbors::detail diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt index 903c28f8d8..fac4c62444 100644 --- a/cpp/test/CMakeLists.txt +++ b/cpp/test/CMakeLists.txt @@ -52,6 +52,7 @@ function(ConfigureTest) GTest::gtest GTest::gtest_main Threads::Threads + ${CUVS_CTK_MATH_DEPENDENCIES} $ $ $<$:cuvs::c_api> @@ -166,6 +167,8 @@ if(BUILD_TESTS) test/distance/dist_l_inf.cu test/distance/dist_lp_unexp.cu test/distance/dist_russell_rao.cu + test/distance/masked_nn.cu + test/sparse/neighbors/cross_component_nn.cu GPUS 1 PERCENT diff --git a/cpp/test/cluster/linkage.cu b/cpp/test/cluster/linkage.cu index 6ea6bd61b9..0f2461fa72 100644 --- a/cpp/test/cluster/linkage.cu +++ b/cpp/test/cluster/linkage.cu @@ -21,7 +21,7 @@ // // TODO: consider adding this to libraft.so or creating an instance in a // separate translation unit for this test. -#undef RAFT_EXPLICIT_INSTANTIATE_ONLY +#undef CUVS_EXPLICIT_INSTANTIATE_ONLY #include "../test_utils.cuh" diff --git a/cpp/test/distance/masked_nn.cu b/cpp/test/distance/masked_nn.cu new file mode 100644 index 0000000000..a8f2f5163f --- /dev/null +++ b/cpp/test/distance/masked_nn.cu @@ -0,0 +1,438 @@ +/* + * 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 +#include +#include + +#include "../../src/distance/masked_nn.cuh" +#include "../test_utils.h" + +#include + +#include + +namespace cuvs::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, cuvs::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 = raft::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); + cuvs::distance::MinAndDistanceReduceOp op; + cuvs::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 = cuvs::distance::MinAndDistanceReduceOp; + using PairRedOpT = cuvs::distance::KVPMinReduce; + using ParamT = cuvs::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 + cuvs::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()); + + raft::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), + raft::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), + raft::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), + raft::resource::get_cuda_stream(handle))); +} + +INSTANTIATE_TEST_CASE_P(MaskedL2NNTests, MaskedL2NNTest, ::testing::ValuesIn(gen_params())); + +} // end namespace cuvs::distance::masked_nn diff --git a/cpp/test/sparse/neighbors/cross_component_nn.cu b/cpp/test/sparse/neighbors/cross_component_nn.cu new file mode 100644 index 0000000000..d931fedc46 --- /dev/null +++ b/cpp/test/sparse/neighbors/cross_component_nn.cu @@ -0,0 +1,1039 @@ +/* + * 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. + */ + +// XXX: We allow the instantiation of masked_l2_nn here: +// raft::sparse::neighbors::FixConnectivitiesRedOp red_op(params.n_row); +// cuvs::sparse::neighbors::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. +// +// TODO: edge case testing. Reference: https://github.com/rapidsai/raft/issues/1669 + +#include "../../test_utils.cuh" + +#include "../../../src/cluster/single_linkage.cuh" +#include "../../../src/sparse/neighbors/cross_component_nn.cuh" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include + +#include + +namespace cuvs { +namespace sparse { + +using namespace std; + +template +struct ConnectComponentsInputs { + value_idx n_row; + value_idx n_col; + std::vector data; + + int c; +}; + +template +class ConnectComponentsTest + : public ::testing::TestWithParam> { + protected: + void basicTest() + { + raft::resources handle; + + auto stream = raft::resource::get_cuda_stream(handle); + + params = ::testing::TestWithParam>::GetParam(); + + raft::sparse::COO out_edges(stream); + raft::sparse::COO out_edges_batched(stream); + + rmm::device_uvector data(params.n_row * params.n_col, stream); + + raft::copy( + data.data(), params.data.data(), data.size(), raft::resource::get_cuda_stream(handle)); + + rmm::device_uvector indptr(params.n_row + 1, stream); + + /** + * 1. Construct knn graph + */ + raft::sparse::COO knn_graph_coo(stream); + + raft::sparse::neighbors::knn_graph(handle, + data.data(), + params.n_row, + params.n_col, + raft::distance::DistanceType::L2SqrtExpanded, + knn_graph_coo, + params.c); + + raft::sparse::convert::sorted_coo_to_csr( + knn_graph_coo.rows(), knn_graph_coo.nnz, indptr.data(), params.n_row + 1, stream); + + /** + * 2. Construct MST, sorted by weights + */ + rmm::device_uvector colors(params.n_row, stream); + + auto mst_coo = raft::mst::mst(handle, + indptr.data(), + knn_graph_coo.cols(), + knn_graph_coo.vals(), + params.n_row, + knn_graph_coo.nnz, + colors.data(), + stream, + false, + true); + + /** + * 3. cross_component_nn to fix connectivities + */ + cuvs::sparse::neighbors::FixConnectivitiesRedOp red_op(params.n_row); + cuvs::sparse::neighbors::cross_component_nn(handle, + out_edges, + data.data(), + colors.data(), + params.n_row, + params.n_col, + red_op, + params.n_row, + params.n_col); + + cuvs::sparse::neighbors::cross_component_nn(handle, + out_edges_batched, + data.data(), + colors.data(), + params.n_row, + params.n_col, + red_op, + params.n_row / 2, + params.n_col / 2); + + ASSERT_TRUE(out_edges.nnz == out_edges_batched.nnz); + + ASSERT_TRUE( + devArrMatch(out_edges.rows(), out_edges_batched.rows(), out_edges.nnz, Compare())); + + ASSERT_TRUE( + devArrMatch(out_edges.cols(), out_edges_batched.cols(), out_edges.nnz, Compare())); + + ASSERT_TRUE(devArrMatch( + out_edges.vals(), out_edges_batched.vals(), out_edges.nnz, CompareApprox(1e-4))); + + /** + * Construct final edge list + */ + rmm::device_uvector indptr2(params.n_row + 1, stream); + + raft::sparse::convert::sorted_coo_to_csr( + out_edges.rows(), out_edges.nnz, indptr2.data(), params.n_row + 1, stream); + + auto output_mst = raft::mst::mst(handle, + indptr2.data(), + out_edges.cols(), + out_edges.vals(), + params.n_row, + out_edges.nnz, + colors.data(), + stream, + false, + false); + + raft::resource::sync_stream(handle, stream); + + // The sum of edges for both MST runs should be n_rows - 1 + final_edges = output_mst.n_edges + mst_coo.n_edges; + } + + void SetUp() override { basicTest(); } + + void TearDown() override {} + + protected: + ConnectComponentsInputs params; + + value_idx final_edges; +}; + +const std::vector> fix_conn_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}, + -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 + + }, + -4}}; + +typedef ConnectComponentsTest ConnectComponentsTestF_Int; +TEST_P(ConnectComponentsTestF_Int, Result) +{ + /** + * Verify the src & dst vertices on each edge have different colors + */ + EXPECT_TRUE(final_edges == params.n_row - 1); +} + +INSTANTIATE_TEST_CASE_P(ConnectComponentsTest, + ConnectComponentsTestF_Int, + ::testing::ValuesIn(fix_conn_inputsf2)); + +template +struct MutualReachabilityFixConnectivitiesRedOp { + value_t* core_dists; + value_idx m; + + DI MutualReachabilityFixConnectivitiesRedOp() : m(0) {} + + MutualReachabilityFixConnectivitiesRedOp(value_t* core_dists_, value_idx m_) + : core_dists(core_dists_), m(m_){}; + + typedef typename raft::KeyValuePair KVP; + DI void operator()(value_idx rit, KVP* out, const KVP& other) const + { + if (rit < m && other.value < std::numeric_limits::max()) { + value_t core_dist_rit = core_dists[rit]; + value_t core_dist_other = max(core_dist_rit, max(core_dists[other.key], other.value)); + + value_t core_dist_out; + if (out->key > -1) { + core_dist_out = max(core_dist_rit, max(core_dists[out->key], out->value)); + } else { + core_dist_out = out->value; + } + + bool smaller = core_dist_other < core_dist_out; + out->key = smaller ? other.key : out->key; + out->value = smaller ? core_dist_other : core_dist_out; + } + } + + DI KVP operator()(value_idx rit, const KVP& a, const KVP& b) const + { + if (rit < m && a.key > -1) { + value_t core_dist_rit = core_dists[rit]; + value_t core_dist_a = max(core_dist_rit, max(core_dists[a.key], a.value)); + + value_t core_dist_b; + if (b.key > -1) { + core_dist_b = max(core_dist_rit, max(core_dists[b.key], b.value)); + } else { + core_dist_b = b.value; + } + + return core_dist_a < core_dist_b ? KVP(a.key, core_dist_a) : KVP(b.key, core_dist_b); + } + + return b; + } + + DI void init(value_t* out, value_t maxVal) const { *out = maxVal; } + DI void init(KVP* out, value_t maxVal) const + { + out->key = -1; + out->value = maxVal; + } + + DI void init_key(value_t& out, value_idx idx) const { return; } + DI void init_key(KVP& out, value_idx idx) const { out.key = idx; } + + DI value_t get_value(KVP& out) const { return out.value; } + DI value_t get_value(value_t& out) const { return out; } + + void gather(const raft::resources& handle, value_idx* map) + { + auto tmp_core_dists = raft::make_device_vector(handle, m); + thrust::gather(raft::resource::get_thrust_policy(handle), + map, + map + m, + core_dists, + tmp_core_dists.data_handle()); + raft::copy_async( + core_dists, tmp_core_dists.data_handle(), m, raft::resource::get_cuda_stream(handle)); + } + + void scatter(const raft::resources& handle, value_idx* map) + { + auto tmp_core_dists = raft::make_device_vector(handle, m); + thrust::scatter(raft::resource::get_thrust_policy(handle), + core_dists, + core_dists + m, + map, + tmp_core_dists.data_handle()); + raft::copy_async( + core_dists, tmp_core_dists.data_handle(), m, raft::resource::get_cuda_stream(handle)); + } +}; + +template +struct ConnectComponentsMutualReachabilityInputs { + value_idx n_row; + value_idx n_col; + std::vector data; + std::vector core_dists; + std::vector colors; + std::vector expected_rows; + std::vector expected_cols; + std::vector expected_vals; +}; + +template +class ConnectComponentsEdgesTest + : public ::testing::TestWithParam> { + protected: + void basicTest() + { + raft::resources handle; + + auto stream = raft::resource::get_cuda_stream(handle); + + params = ::testing::TestWithParam< + ConnectComponentsMutualReachabilityInputs>::GetParam(); + + raft::sparse::COO out_edges_unbatched( + raft::resource::get_cuda_stream(handle)); + raft::sparse::COO out_edges_batched( + raft::resource::get_cuda_stream(handle)); + + rmm::device_uvector data(params.n_row * params.n_col, + raft::resource::get_cuda_stream(handle)); + rmm::device_uvector core_dists(params.n_row, raft::resource::get_cuda_stream(handle)); + rmm::device_uvector colors(params.n_row, raft::resource::get_cuda_stream(handle)); + + raft::copy( + data.data(), params.data.data(), data.size(), raft::resource::get_cuda_stream(handle)); + raft::copy(core_dists.data(), + params.core_dists.data(), + core_dists.size(), + raft::resource::get_cuda_stream(handle)); + raft::copy( + colors.data(), params.colors.data(), colors.size(), raft::resource::get_cuda_stream(handle)); + + /** + * 3. cross_component_nn to fix connectivities + */ + MutualReachabilityFixConnectivitiesRedOp red_op(core_dists.data(), + params.n_row); + + cuvs::sparse::neighbors::cross_component_nn(handle, + out_edges_unbatched, + data.data(), + colors.data(), + params.n_row, + params.n_col, + red_op, + params.n_row, + params.n_col); + + cuvs::sparse::neighbors::cross_component_nn(handle, + out_edges_batched, + data.data(), + colors.data(), + params.n_row, + params.n_col, + red_op, + 11, + 1); + + ASSERT_TRUE(out_edges_unbatched.nnz == out_edges_batched.nnz && + out_edges_unbatched.nnz == params.expected_rows.size()); + + ASSERT_TRUE(devArrMatch(out_edges_unbatched.rows(), + params.expected_rows.data(), + out_edges_unbatched.nnz, + Compare())); + + ASSERT_TRUE(devArrMatch(out_edges_unbatched.cols(), + params.expected_cols.data(), + out_edges_unbatched.nnz, + Compare())); + + ASSERT_TRUE(devArrMatch(out_edges_unbatched.vals(), + params.expected_vals.data(), + out_edges_unbatched.nnz, + CompareApprox(1e-4))); + + ASSERT_TRUE(devArrMatch(out_edges_batched.rows(), + params.expected_rows.data(), + out_edges_batched.nnz, + Compare())); + + ASSERT_TRUE(devArrMatch(out_edges_batched.cols(), + params.expected_cols.data(), + out_edges_batched.nnz, + Compare())); + + ASSERT_TRUE(devArrMatch(out_edges_batched.vals(), + params.expected_vals.data(), + out_edges_batched.nnz, + CompareApprox(1e-4))); + } + + void SetUp() override { basicTest(); } + + void TearDown() override {} + + protected: + ConnectComponentsMutualReachabilityInputs params; +}; + +const std::vector> mr_fix_conn_inputsf2 = { + {100, + 2, + {-7.72642, -8.39496, 5.4534, 0.742305, -2.97867, 9.55685, 6.04267, 0.571319, -6.52184, + -6.31932, 3.64934, 1.40687, -2.17793, 9.98983, 4.42021, 2.33028, 4.73696, 2.94181, + -3.66019, 9.38998, -3.05358, 9.12521, -6.65217, -5.57297, -6.35769, -6.58313, -3.61553, + 7.81808, -1.77073, 9.18565, -7.95052, -6.39764, -6.60294, -6.05293, -2.58121, 10.0178, + -7.76348, -6.72638, -6.40639, -6.95294, -2.97262, 8.54856, -6.95673, -6.53896, -7.32614, + -6.02371, -2.1478, 10.5523, -2.54502, 10.5789, -2.96984, 10.0714, 3.22451, 1.55252, + -6.25396, -7.73727, -7.85431, -6.09303, -8.11658, -8.20057, -7.55965, -6.64786, 4.936, + 2.23423, 4.44752, 2.27472, -5.72103, -7.70079, -0.929985, 9.78172, -3.10984, 8.72259, + -2.44167, 7.58954, -2.18511, 8.6292, 5.55528, 2.30192, 4.73164, -0.0143992, -8.2573, + -7.81793, -2.98837, 8.82863, 4.60517, 0.804492, -3.83738, 9.21115, -2.62485, 8.71318, + 3.57758, 2.44676, -8.48711, -6.69548, -6.70645, -6.49479, -6.86663, -5.42658, 3.83139, + 1.47141, 2.02013, 2.79507, 4.64499, 1.73858, -1.69667, 10.3705, -6.61974, -6.09829, + -6.05757, -4.98332, -7.10309, -6.16611, -3.52203, 9.32853, -2.26724, 7.10101, 6.11777, + 1.4549, -4.23412, 8.452, -6.58655, -7.59446, 3.93783, 1.64551, -7.12502, -7.63385, + 2.72111, 1.94666, -7.14428, -4.15994, -6.66553, -8.12585, 4.70011, 4.43641, -7.76914, + -7.69592, 4.11012, 2.48644, 4.89743, 1.89872, 4.29716, 1.17089, -6.62913, -6.53366, + -8.07093, -6.22356, -2.16558, 7.25125, 4.73953, 1.46969, -5.91625, -6.46733, 5.43091, + 1.06378, -6.82142, -8.02308, 6.52606, 2.14775, 3.08922, 2.04173, -2.14756, 8.36917, + 3.85663, 1.65111, -1.68665, 7.79344, -5.01385, -6.40628, -2.52269, 7.95658, -2.30033, + 7.05462, -1.04355, 8.78851, 3.72045, 3.5231, -3.98772, 8.29444, 4.24777, 0.509655, + 4.72693, 1.67416, 5.7827, 2.7251, -3.41722, 7.60198, 5.22674, 4.16363, -3.1109, + 10.8666, -3.18612, 9.62596, -1.4782, 9.94557, 4.47859, 2.37722, -5.79658, -5.82631, + -3.34842, 8.70507}, + {0.978428, 1.01917, 0.608673, 1.45629, 0.310713, 0.689461, 0.701126, 0.63296, 0.774788, + 0.701648, 0.513282, 0.757651, 0.45638, 0.973111, 0.901396, 0.613692, 0.482497, 0.688143, + 0.72428, 0.666345, 0.58232, 0.554756, 0.710315, 0.903611, 0.694115, 0.796099, 0.639759, + 0.798998, 0.639839, 1.30727, 0.663729, 0.57476, 0.571348, 1.14662, 1.26518, 0.485068, + 0.78207, 0.791621, 1.01678, 1.28509, 1.14715, 0.381395, 0.850507, 0.788511, 0.588341, + 0.878516, 0.928669, 0.405874, 0.776421, 0.612274, 1.84963, 0.57476, 0.95226, 0.488078, + 1.24868, 0.515136, 0.589378, 0.903632, 1.01678, 1.09964, 0.666345, 0.713265, 0.877168, + 1.10053, 1.96887, 1.03574, 2.03728, 0.969553, 0.774788, 0.586338, 0.65168, 0.435472, + 0.664396, 0.790584, 0.678637, 0.715964, 0.865494, 0.978428, 1.59242, 0.861109, 0.833259, + 0.65168, 0.903632, 1.49599, 0.76347, 0.960453, 1.1848, 1.37398, 0.928957, 1.07848, + 0.661798, 1.21104, 1.04579, 1.89047, 1.24288, 0.529553, 0.903611, 0.620897, 0.882467, + 0.647189}, + {0, 1, 2, 1, 0, 1, 2, 1, 1, 2, 2, 0, 0, 2, 2, 0, 0, 2, 0, 0, 2, 0, 0, 2, 2, + 2, 1, 0, 0, 0, 0, 1, 1, 0, 2, 2, 2, 2, 1, 1, 0, 2, 1, 2, 2, 1, 0, 0, 0, 1, + 1, 1, 2, 0, 0, 0, 2, 2, 1, 2, 0, 1, 0, 1, 0, 0, 1, 0, 1, 1, 1, 0, 0, 2, 1, + 0, 1, 0, 1, 1, 2, 1, 2, 0, 2, 2, 2, 1, 2, 1, 1, 1, 2, 1, 2, 2, 2, 1, 0, 2}, + {50, 54, 57, 63, 82, 87}, + {57, 63, 50, 54, 87, 82}, + {6.0764, 11.1843, 6.0764, 11.1843, 6.89004, 6.89004}}, + {1000, + 2, + {-6.59634, -7.13901, -6.13753, -6.58082, 5.19821, 2.04918, -2.96856, 8.16444, + -2.76879, 7.51114, -6.82261, -6.61152, 5.02008, 2.58376, 5.55621, 2.31966, + 4.86379, 3.33731, 5.84639, 1.15623, -2.17159, 8.60241, -4.97844, -6.94077, + -2.31014, 8.41407, 5.5582, 0.402669, 5.25265, 0.919754, 5.85298, 2.11489, + -3.29245, 8.69222, -1.9621, 8.81209, -1.53408, 8.86723, -2.18227, 8.79519, + 4.60519, 2.20738, -6.4759, -6.9043, -7.18766, -6.10045, -9.00148, -7.48793, + 4.01674, 1.41769, -2.45347, 10.1085, -3.20892, 9.22827, -3.18612, 9.62596, + 4.81977, 3.36517, 4.90693, 2.8628, -6.44269, -5.68946, -8.30144, -5.37878, + 4.61485, 2.79094, -1.98726, 9.31127, -3.66019, 9.38998, -6.58607, -8.23669, + -7.46015, -6.29153, 4.08468, 3.85433, -6.36842, -5.50645, -6.83602, -5.18506, + -0.627173, 10.3597, 3.98846, 1.48928, -2.9968, 8.58173, -7.2144, -7.28376, + -0.660242, 10.1409, -4.23528, -8.38308, -3.15984, 8.52716, -2.40987, 9.76567, + -8.7548, -6.76508, 4.56971, 0.312209, -7.5487, -5.8402, -1.6096, 9.32159, + 5.04813, 0.270586, -7.6525, -6.47306, -1.79758, 7.88964, -9.0153, -3.74236, + -3.5715, 9.48788, -1.65154, 8.85435, -3.47412, 9.70034, 6.31245, 2.39219, + 4.03851, 2.29295, -3.17098, 9.86672, -6.90693, -7.81338, -6.22373, -6.68537, + -3.22204, 9.12072, -0.365254, 9.6482, -7.76712, -7.31757, 4.15669, 3.54716, + 4.1937, 0.083629, -3.03896, 9.52755, -6.29293, -7.35501, -2.95926, 9.63714, + 4.02709, 1.58547, 4.56828, 1.93595, 5.6242, 1.75918, -7.36237, -7.83344, + 5.32177, 3.81988, -2.43183, 8.153, -1.97939, 10.4559, -3.49492, 9.51833, + 3.39602, 1.28026, -2.42215, 8.71528, -3.57682, 8.87191, -2.77385, 11.7345, + 5.71351, 0.946654, -6.50253, -6.90937, 4.08239, 0.603367, -5.64134, -6.85884, + -2.76177, 7.7665, -2.25165, 8.93984, -3.49071, 9.47639, -1.06792, 7.57842, + 5.15754, 1.24743, 3.63574, 1.20537, -6.07969, -8.49642, 4.12227, 2.19696, + -7.17144, -8.4433, -1.92234, 11.2047, 3.23237, 1.19535, 3.85389, 0.641937, + 4.82665, 1.21779, -7.68923, -6.45605, -7.00816, -8.76196, -5.12894, 9.83619, + -5.66247, -5.35879, 3.05598, 2.73358, 6.06038, 1.40242, -1.69568, 7.78342, + 5.13391, 2.23384, -2.96984, 10.0714, -5.36618, -6.2493, 5.55896, 1.6829, + 3.55882, 2.58911, 5.36155, 0.844118, -0.0634456, 9.14351, 4.88368, 1.40909, + -7.04675, -6.59753, -7.78333, -6.55575, 5.39881, 2.25436, -2.85189, 8.64285, + -2.22821, 8.39159, 3.88591, 1.69249, -7.55481, -7.02463, 4.60032, 2.65467, + -6.90615, -7.76198, -6.76005, -7.85318, 4.15044, 3.01733, -7.18884, -7.63227, + 4.68874, 2.01376, 3.51716, 2.35558, -3.81367, 9.68396, 4.42644, 3.4639, + 4.81758, 0.637825, -6.20705, -4.98023, -1.68603, 9.0876, -4.99504, -5.33687, + -1.77073, 9.18565, 4.86433, 3.02027, 4.20538, 1.664, 4.59042, 2.64799, + -3.09856, 9.86389, -3.02306, 7.95507, -6.32402, -6.79053, -7.67205, -7.18807, + -8.10918, -6.38341, -1.67979, 6.80315, 4.00249, 3.16219, -2.54391, 7.84561, + -3.22764, 8.80084, -2.63712, 8.05875, -2.41744, 7.02672, -6.71117, -5.56251, + 5.18348, 1.60256, -7.40824, -6.29375, -4.22233, 10.3682, 4.8509, 1.87646, + -2.99456, 9.09616, 5.1332, 2.15801, -2.27358, 9.78515, -6.73874, -8.64855, + 4.96124, 2.39509, -3.70949, 8.67978, -4.13674, 9.06237, 2.80367, 2.48116, + -0.876786, 7.58414, -3.7005, 9.67084, 6.48652, 0.903085, 6.28189, 2.98299, + -6.07922, -6.12582, -5.67921, -7.537, 4.55014, 3.41329, -1.63688, 9.19763, + -4.02439, 10.3812, 5.23053, 3.08187, -2.2951, 7.76855, -6.24491, -5.77041, + 6.02415, 2.53708, -6.91286, -7.08823, 4.83193, 1.66405, -7.07454, -5.74634, + -2.09576, 10.8911, 3.29543, 1.05452, -3.49973, 8.44799, 5.2922, 0.396778, + -2.54502, 10.5789, -6.38865, -6.14523, -1.75221, 8.09212, -9.30387, -5.99606, + -2.98113, 10.1032, -6.2017, -7.36802, 4.63628, 0.814805, -1.81905, 8.61307, + 4.88926, 3.55062, 3.08325, 2.57918, -2.51717, 10.4942, -5.75358, -6.9315, + 6.36742, 2.40949, 5.74806, 0.933264, 4.74408, 1.91058, -7.41496, -6.97064, + -2.98414, 8.36096, 6.72825, 1.83358, -2.95349, 9.39159, -3.35599, 7.49944, + 6.18738, 3.76905, -3.17182, 9.58488, 5.17863, 1.0525, -3.0397, 8.43847, + -2.23874, 8.96405, 3.04689, 2.41364, 6.14064, 2.82339, -6.33334, -6.87369, + -7.92444, -8.84647, 3.65129, 0.86958, 5.29842, 3.98337, -2.06538, 9.78892, + -6.89494, -6.30082, -2.52144, 8.11703, -8.11398, -7.47257, 5.3381, 2.36666, + -6.93452, -6.59456, -7.50634, -6.01772, 6.23438, 1.12621, -2.15218, 8.32138, + -7.04777, -7.3522, -2.52771, 8.72563, -2.77907, 8.03552, 4.29123, 1.62391, + -8.07551, -6.43551, -3.28202, 8.77747, -2.21308, 9.27534, -8.25153, -8.49367, + -3.54644, 8.82395, -8.05867, -5.69243, 4.46681, 1.98875, 3.8362, 3.61229, + -6.96231, -7.00186, 5.18993, 1.00483, -5.35116, -6.37227, 5.23298, 1.66362, + -5.68306, -7.03864, -9.03144, -7.59926, -6.10127, -7.4313, 4.83572, 0.994797, + -7.32695, -5.59909, 0.569683, 10.1339, 3.35957, 2.84563, -2.4122, 9.60944, + 5.00855, 1.57983, -2.57528, 7.80327, 3.96349, 3.77411, 4.59429, 2.21651, + -6.54765, -6.68961, 4.76798, 1.29212, -1.67351, 7.88458, 5.63615, 1.47941, + -2.5301, 9.13161, 4.26075, 1.76959, 4.67788, 2.0932, 4.39955, 1.59835, + 3.91274, 1.72565, -4.1786, 9.55765, -7.34566, -8.47481, 4.8364, 2.68217, + -7.36848, -7.99973, -5.84708, -5.7534, 5.37252, 1.89245, -2.1707, 8.599, + -1.3299, 9.0818, -6.79122, -5.40258, 5.56391, 1.78827, -0.194539, 7.14702, + 4.60489, 3.74397, 5.50995, 2.46885, -3.98772, 8.29444, -5.21837, -7.33721, + -1.63959, 10.3699, -5.92932, -5.1695, -5.88358, -7.6369, 4.11716, 3.02218, + -6.54114, -7.17551, 3.97179, 2.96521, -6.75325, -4.94118, 5.26169, 0.402945, + 3.25031, 0.327771, -0.44845, 10.7696, -2.15141, 9.57507, 7.04329, 1.91555, + -3.74615, 7.69383, -7.52318, -5.85015, -6.80419, -8.48208, -4.57664, 8.92517, + 4.57574, 2.30193, 4.84098, 3.02382, -9.43355, -5.94579, -3.52203, 9.32853, + 3.43018, 2.5731, -6.15725, -7.25294, -6.69861, -8.17694, -2.40955, 8.51081, + -4.82342, -7.98332, -7.10611, -6.51274, 5.86755, 0.763529, -6.56045, -5.53966, + -3.61553, 7.81808, 4.3825, 0.304586, -6.52818, -5.80996, 4.59972, 0.542395, + -6.90603, -6.59995, -6.3585, -6.23489, -6.01915, -7.46319, -5.38694, -7.15123, + -7.83475, -6.45651, 5.89564, 1.07856, -5.15266, -7.27975, -6.97978, -7.08378, + 5.83493, 0.449983, -2.62374, 10.2521, -7.34494, -6.98606, -6.79719, -8.33766, + 3.54757, 1.65676, -8.40528, -5.61753, -5.85556, -6.28758, 4.66862, 3.25162, + -6.26047, -4.82261, 4.61552, 4.11544, -1.36637, 9.76622, 4.2517, 2.14359, + -2.45099, 7.87132, -0.376164, 7.0622, 4.34493, 3.22091, 6.95921, 2.36649, + -6.70319, -7.24714, -5.56932, -5.48443, -7.43149, -4.32191, -3.23956, 9.23074, + -5.77255, -7.00049, 4.96601, 0.722056, -7.88617, -5.74023, 4.18757, -0.45071, + -7.12569, -7.72336, 5.27366, 2.38697, 3.93487, 1.9174, 3.19186, -0.225636, + -3.41722, 7.60198, -3.08286, 8.46743, -5.87905, -7.55073, -5.26425, -7.20243, + -2.97867, 9.55685, -1.23153, 8.42272, -2.33602, 9.3996, -3.33819, 8.45411, + -3.58009, 9.49676, 3.78152, 2.67348, -1.54582, 9.42707, -4.04331, 10.292, + 3.3452, 3.134, -2.75494, 8.74156, -3.26555, 7.59203, -7.27139, -7.80252, + 3.5293, 3.72544, 6.11642, 3.35326, 4.01611, 3.8872, 4.89591, 2.95586, + -7.06677, -5.89438, 4.19438, 3.42655, -6.11355, -5.65318, -7.59645, -8.74665, + -5.80362, -6.8588, 3.80453, 4.11832, 5.70655, 3.14247, -4.98084, 8.21739, + -1.87642, 11.285, 4.39864, 2.32523, -3.48388, 9.80137, 4.02836, 0.566509, + -2.41212, 9.98293, -5.40846, -7.08943, 4.01506, 1.99926, -3.43613, 8.95476, + -7.24458, -7.71932, 6.02204, 2.62188, -6.29999, -6.55431, 6.19038, 0.974816, + 3.55882, 3.02632, -7.06011, -3.687, -1.55877, 8.43738, -5.14711, -4.64881, + 4.7167, 0.690177, -7.90381, -5.02602, 4.17218, 2.31967, -0.643423, 9.48812, + -7.95237, -6.64086, -4.05986, 9.08285, -6.24158, -6.37927, -6.6105, -7.2233, + -6.21675, -5.70664, -3.29967, 9.48575, 3.41775, 2.68617, -2.24948, 8.10997, + -2.24931, 9.79611, -9.0523, -6.03269, -2.2587, 9.36073, 5.20965, 2.42088, + -3.10159, 8.1503, -6.67906, -5.73147, 4.0687, 2.54575, -1.24229, 8.30662, + -2.09627, 8.45056, -7.87801, -6.57832, 4.72216, 3.03865, -0.929985, 9.78172, + -8.56307, -7.68598, -7.05257, -5.1684, -7.09076, -7.86729, 4.61432, 3.1459, + -6.34133, -5.8076, -3.82943, 10.8457, -8.46082, -5.98507, 5.34763, 1.4107, + -1.68714, 10.9111, -1.67886, 8.1582, -0.623012, 9.18886, -4.21258, 8.95874, + -2.16744, 10.8905, -6.57158, -7.27176, 2.14047, 4.26411, -8.44217, -7.40916, + 5.29008, 1.87399, 4.31824, 4.04992, -3.77008, 9.93215, -2.72688, 10.1131, + -6.14278, -7.16144, -3.92457, 8.59364, -5.92649, -6.59299, 4.68369, 1.82617, + -6.89905, -7.18329, 3.95173, 4.22561, -7.66453, -6.23183, -2.44167, 7.58954, + -6.36603, -7.41281, -6.45081, -6.187, -6.6125, -6.37138, 5.46036, 2.48044, + -2.14756, 8.36917, -2.3889, 9.52872, 3.80752, 2.44459, -3.98778, 10.158, + -6.63887, -4.27843, -8.65266, -5.61819, -7.97003, -5.46918, -5.9604, -7.54825, + -0.916011, 8.50307, -3.69246, 6.97505, -7.98533, -7.09503, -2.30033, 7.05462, + 4.76218, 2.51647, -7.04981, -7.33334, 3.66401, 3.02681, -2.50408, 8.7797, + 7.19996, 1.87711, 4.01291, 3.78562, -0.356015, 8.24694, -0.958046, 9.12996, + 4.60675, 3.76773, 6.21945, 1.45031, 4.27744, 0.8535, -4.72232, -7.48582, + 6.03923, 2.8978, -3.26833, 9.16468, -7.97059, -7.29092, -2.3998, 9.74005, + -2.66721, 8.58741, -7.36269, -6.73332, -7.87893, -7.38488, 4.65023, 0.661333, + -4.8171, -7.94764, -4.11564, 9.21775, 4.80633, 2.46562, -2.72887, 9.3714, + -5.26735, -5.5652, 4.9826, 2.42992, -6.17018, -7.3156, 4.38084, 1.77682, + 5.35084, 2.41743, -2.61796, 9.416, 5.27229, 2.94572, -7.52315, -5.95227, + -1.45077, 7.25555, -3.79916, 7.71921, -2.23251, 9.84147, 3.70054, 1.82908, + -1.93831, 10.1499, -6.18324, -5.9248, -3.33142, 9.25797, -6.08536, -8.1344, + 5.95727, 2.17077, 4.87366, 0.417274, -6.529, -6.39092, -9.24256, -7.88984, + -6.36652, -7.13966, -3.90777, 9.57726, -7.06252, -5.50523, -2.26423, 8.50734, + -2.84498, 10.6833, 5.0391, 2.62037, -2.74815, 8.10672, 3.35945, 3.72796, + -4.11668, 9.19892, 5.66903, 2.44577, -1.63807, 8.68826, -7.42587, -6.48831, + 6.17063, 3.19193, -2.28511, 9.02688, -7.10088, -7.15692, 4.46293, 1.17487, + -5.91017, -6.45292, -2.26724, 7.10101, -2.43339, 8.33712, -4.63309, 8.48853, + -3.31769, 8.51253, -2.49078, 10.6907, -1.30798, 8.60621, 6.30535, 2.98754, + -5.79384, -6.78213, -1.93213, 8.81124, 4.55773, 3.09047, 6.37584, 2.17108, + 4.3927, 1.29119, -3.2245, 9.69388, -1.69634, 9.64392, 2.799, 0.693593, + -2.1426, 8.07441, -8.4505, -8.00688, 4.736, 1.51089, -2.5863, 9.35544, + -2.94924, 9.14503, 6.2054, 1.90742, 5.67172, 0.487609, -5.69071, -6.17181, + -8.24651, -7.10488, -7.34424, -6.67895, -6.71977, -7.90778, -1.82294, 7.40157, + -9.40991, -7.16611, -4.37999, 8.66277, -1.42615, 10.0681, -2.00828, 8.03673, + -7.50228, -6.6855, -5.65859, -6.29801, -8.02335, -6.77155, -3.40761, 9.50621, + -2.82447, 9.77326, -1.5938, 9.34304, -3.5213, 7.35943, -3.36961, 8.62973, + -7.01708, -5.92724, 5.20886, 3.60157, -1.71817, 8.1049, -2.46363, 8.36269, + -2.77809, 7.90776, -2.75459, 8.26055, -2.03596, 8.94146, -4.53434, 9.20074, + -7.44387, -6.69556, -6.90099, -7.62732, 3.29169, 2.71643, 6.08686, 2.16972, + -2.31111, 8.86993, -5.75046, 7.9899, 4.69951, 1.32623, 4.71851, -0.025031, + -6.42374, -4.71511, -8.04974, -8.68209, -3.16103, 9.06168, -6.18267, -7.21393, + -7.94202, -6.4518, -7.07697, -7.03138, 3.93554, 0.564708, -1.20372, 9.03529, + -7.10611, -7.83955, -7.47529, -5.50567, -6.15453, -6.36393, -2.98024, 9.24634, + -7.75761, -7.70699, -3.08597, 9.76968, -8.04954, -9.75237, 5.2534, 0.950377, + 5.63789, -0.923086, -5.7065, -6.51047, -8.02132, -7.07377, -8.28594, -6.96322, + -7.70722, -6.79397, -2.4962, 10.4678, 5.02846, 4.46617, 4.02648, 1.6707, + -0.319395, 8.20599, 4.74525, 0.639144, -1.0313, 8.49602, 4.08766, 2.6061, + 3.63826, 1.69207, 2.55795, 3.66963, 5.2826, 3.30232, -1.04355, 8.78851, + -6.84762, -7.63353, -4.70868, -7.056, 3.53651, -0.179721, -3.38482, 7.63149, + -5.9265, -6.36702, -0.986074, 9.5532, -2.42261, 8.85861, -7.42835, -6.78726, + -4.02857, 8.53005, -8.22675, -7.85172, -5.57529, -8.5426, 6.03009, 2.53098, + -7.10448, -7.53011, -3.4988, 8.8885, -2.62485, 8.71318, -6.39489, -7.72647, + 3.93789, 1.31027, 4.27627, 1.91622, -0.923181, 7.77647, -5.16017, 10.1058, + -6.44307, -5.97617, -7.24495, -6.69543, 6.27331, 0.826824, -6.55655, -7.13246, + 5.66245, 4.41292, -2.13805, 8.4103, 5.23463, 2.82659, -4.86624, -6.74357, + -6.14082, -6.26474, -2.67048, 9.41834, -1.26311, 6.9409, -7.20231, -7.13094, + -1.35109, 9.80595, 3.9906, 0.749229, -6.75696, -5.25543, 4.84826, -0.0685652, + -7.4914, -6.91715, 4.46725, 2.85683, -2.95571, 9.87068, 6.32381, 1.51429, + -6.81177, -6.02734, -2.57188, 9.96943, -4.28792, 10.5103, 3.65025, 2.91394, + -7.11856, -7.24693, -6.98693, -6.43239, 4.7651, 1.54376, 4.00092, 0.65008, + -7.14816, -7.7713, -7.58803, -8.39382, 4.3321, 2.19232, -7.89545, -6.81843, + -2.11475, 8.5933, -0.743743, 9.41927, 3.64849, -0.18022, -1.68665, 7.79344, + 4.00214, 1.44217, -6.96799, -7.25012, -1.58302, 10.9237, -6.68524, -7.23328, + 4.65831, 2.32075, 4.62024, 2.52566, -4.23412, 8.452, -0.822056, 9.89593, + -7.19868, -7.67614, -3.32742, 11.1067, 5.27861, 0.830165, 4.48982, 2.09875, + -6.58087, -7.6319, -0.880582, 7.63418, -7.01088, -6.80326, -7.31601, -6.98972, + -6.85883, -7.60811, 6.14328, 2.85053, -7.49206, -6.51861, -2.28174, 10.3214, + 4.81074, 1.78919, -5.58987, -6.20693, 4.08096, 2.35038, -1.5029, 8.43739, + 4.11536, 2.46254, -3.28299, 7.76963, 4.31953, 2.39734, 4.91146, 0.696421, + -1.4782, 9.94557, -3.34842, 8.70507, -6.97822, -6.86126, 4.10012, 1.19486, + -2.50395, 9.06127, 4.41891, 2.00006, -2.73266, 9.72829, 3.5436, 0.533119, + 5.78864, 0.233456, -6.62589, -6.41242, -2.21942, 11.0897, -6.76636, -8.31839, + -2.71732, 8.52129, -5.20972, -6.48544, 3.26056, 1.24224, 3.45228, 2.28299, + 4.72171, 1.87428, -7.52585, -5.1048, 5.0695, 2.18086, -6.55646, -7.02771, + 3.23727, 3.72275, 3.41411, 0.508795, -7.80698, -6.64174, -5.90443, -6.37902, + -0.387041, 10.0468, -1.3506, 8.1936, -6.08614, -8.62864, -5.91478, -5.26453, + -2.61623, 7.97904, 4.45459, 1.84335, -6.66643, -7.63208, 3.6729, 1.92546, + -1.32976, 8.54511, 6.31758, 1.41958, 4.63381, 2.81166, -7.01394, -6.0693, + -2.7786, 9.73183, -2.90131, 7.55077, -7.13842, -5.28146, 6.71514, 1.28398, + -6.98408, -7.04893, -3.03946, 8.22141, -2.76417, 10.5183, -7.35347, -6.89456, + 4.19345, 2.16726, -2.02819, 9.23817, 4.97076, 2.8067, -0.544473, 9.04955, + 4.90727, 2.29487, -6.31871, -7.17559, 3.71665, 0.621485, 4.7903, 2.33813, + -6.47994, -7.53147, -6.80958, -5.71823, -8.07326, -5.96096, 4.77342, 1.8207, + 5.71856, 1.93466, -2.70156, 9.31583, -2.1478, 10.5523, 4.78855, 1.63608, + 5.53507, 2.60834, -7.00058, -6.46058, 5.4738, 2.43235, -1.34603, 9.02452, + -7.5337, -8.71074, -7.30893, -7.57253, -5.33752, -4.87402, -7.01364, -6.86542, + -7.93331, -7.94791, -5.69392, -6.16116, -7.32291, -7.76491, -6.41965, -7.55783, + -7.87996, -7.55785, -6.69005, -5.87906, 3.92147, 2.86809, -1.5552, 9.66568, + 5.07989, 1.47112, -7.48524, -5.0541, -1.82724, 8.70402, -2.00421, 9.88004, + -2.62153, 8.79332, -7.52111, -6.44819, 4.06424, 2.09518, -6.65494, -5.94752, + 6.93878, 1.61033, -3.95728, 7.60682, 5.67016, 2.21196, -7.81507, -5.79413, + -2.41152, 8.24128, -3.83738, 9.21115, 4.5516, 4.55288, -5.75551, -5.93258, + 4.56545, 2.59384, -7.45614, -9.47115, -2.39568, 9.67642, 5.57816, 1.45712, + -7.48184, -6.41134, -1.99415, 12.867, -8.35854, -6.69675, -7.52559, -7.6793, + 5.7454, 3.1602, 2.94692, 1.87483, -8.77324, -6.66682, -3.21125, 8.68662, + -6.25806, -7.24972, 5.17639, 1.0747, -2.44897, 11.4775, -3.30172, 8.89955, + -2.85191, 8.21201, -8.85893, -6.1322, 4.08957, 1.30155, -5.88132, -7.31173, + -7.10309, -7.22943, -2.46068, 8.18334, -7.01226, -7.85464, 4.75411, 2.12347, + -3.42862, 10.5642, 7.16681, 1.4423, 5.42568, 2.39863, -6.00833, -8.22609, + -1.7619, 9.62466, -2.49527, 8.99016, -2.98837, 8.82863, -2.97262, 8.54856, + -1.34142, 9.26871, -5.99652, -6.95795, -1.87061, 7.35277, -8.68277, -8.46425, + -7.01808, -8.10441, -7.04269, -7.62501, -7.69783, -6.88348, -2.19829, 10.4896, + 4.67396, 1.2032, -5.58263, -6.90298, -5.69224, -4.29055, 4.77285, 1.27305, + -3.33469, 8.6929, -2.54195, 8.47086, 4.46492, 1.21742, 5.41158, -0.875373, + -8.68069, -7.42278, -3.88687, 8.07646, 4.6682, 2.00293, -8.29799, -8.64092, + -1.86382, 10.3829, -6.51234, -5.04193, 4.54458, 2.25219, -1.93264, 9.32554, + -3.06285, 7.81641, -6.90714, -5.10786, 4.69653, 2.50286, 6.43757, 2.61401, + -1.85483, 8.9587, 4.60224, 3.07647, 4.4492, 2.1906, 5.02181, 2.40321, + -2.22923, 7.8888, 5.68943, 1.43793, -6.71097, -6.43817, -5.00633, -5.80006, + -2.43763, 8.53663, 5.72577, 2.44787, -6.57079, -5.17789, -5.77867, -4.92176, + -6.57222, -6.06437, 3.96639, 2.25216, -7.95177, -9.80146, 4.92574, 2.30763, + -7.6221, -8.20013, -6.4132, -6.91575, 4.01432, 2.36897, 3.0833, 1.54505, + -1.99416, 9.52807, -7.85128, -8.25973, -0.86423, 8.76525, -6.31412, -8.64087, + -8.07355, -6.73717, -2.52821, 8.01176, -5.82357, -6.65687, -7.08865, -7.73063, + -5.56251, -6.99818, -2.12513, 8.98159, -6.89834, -7.26863, -7.92654, -6.34346, + 4.86201, 1.49442, 4.92905, 4.42847, -5.57789, -5.3186, 4.34232, 3.34888, + 2.64614, 2.34723, -4.10363, 8.41491, -2.18648, 8.18706, -3.39871, 8.19848, + -2.66098, 9.6026, -6.95927, -6.42774, -5.61392, -7.74628, 5.60376, 4.18369, + 5.28536, 4.13642, 4.8428, 0.457426, -6.33816, -6.12095, -2.4394, 8.62897, + 4.56938, 2.45967, 4.0582, 0.958413, 5.62164, 1.64834, 5.73119, 2.58231, + 4.66806, 1.96405, -6.71905, -6.87706, -2.18503, 8.88414, -6.03901, -6.33338, + -8.38435, -6.12005, 0.0641622, 9.0735, 5.19967, 3.05395, -5.48716, -7.13016, + -6.85541, -5.46789, -1.88353, 8.15713, 4.27891, 3.1325, -2.75816, 9.98586, + -2.03022, 9.34795, -7.66741, -7.50096, -3.39305, 9.16801, -8.49476, -5.71537, + -1.68378, 9.8278, -7.41559, -6.07205, -3.15577, 7.93274, 5.22381, 1.61388, + 3.65739, 1.74854, 4.94251, 1.21889, -7.12832, -5.27276, -9.58286, -6.20223, + -2.21613, 8.29993, 5.34799, 2.92987, 4.09496, 2.37231, -7.25183, -5.79136, + -6.46981, -7.12137, -6.28607, -9.8205, 4.52865, 1.06926, -3.10984, 8.72259, + 3.61865, 2.68153, -5.96604, -7.68329, 3.11435, 1.28126, -1.1064, 7.61243, + -2.17688, 8.2658, -3.27246, 7.2094, -5.55143, -6.32388, -1.69667, 10.3705, + -2.16558, 7.25125, -6.36572, -6.70053, 4.12259, 3.38252, -4.80554, -7.79949, + -5.23966, -6.13798, 4.21969, 1.69139, -1.98985, 10.547, -2.52269, 7.95658, + -6.75642, -6.32862, -3.51521, 7.8001, 4.70435, -0.00229688, 6.25359, 2.4267, + 5.82935, 0.745562, 5.24778, 2.15978, 5.48052, 1.32055, -3.05358, 9.12521, + -3.18922, 9.24654, 4.47276, 2.11988, 5.36751, 2.02512, -2.18511, 8.6292, + -2.48469, 9.51228, 5.57556, 3.24472, -2.58121, 10.0178, -6.12629, -6.49895, + -4.54732, 8.0062, -4.20166, 10.5438, -7.61422, -7.69036, -4.42797, 8.98777, + 4.45301, 1.53344, 4.59296, 2.45021, -6.81264, -6.36417, 4.62346, 3.16156, + -5.93007, -8.36501, -2.78425, 6.71237, -6.17141, -6.64689, -5.20608, 8.95999, + -7.30598, -5.73166, 4.39572, 2.93726, -1.89503, 9.77179, -5.683, -7.48989, + 4.80924, 0.559455, -2.17793, 9.98983, 5.23728, 2.67434, -7.03976, -6.20877, + 3.90435, 3.20926, -7.78536, -7.53388, -1.00684, 9.08838, -5.26741, -5.98327, + 3.28002, 2.71942, -1.47166, 8.50427, -2.32733, 9.26251, 5.16271, 1.39947, + -6.59093, -6.61979, -2.44492, 7.93654, -1.05805, 9.97356, -3.1109, 10.8666, + 3.38834, 3.41693, 4.83098, 2.01961, -2.74013, 9.71049, -3.34892, 8.41489, + 4.94768, 0.263001, 3.57477, 1.66795, 5.78915, 1.26999, -4.81812, -5.67174, + -1.88508, 9.64263, 3.69048, 4.60555, 4.03037, 1.7862, -7.4418, -7.08933}, + {0.127717, 0.211407, 0.195547, 0.21633, 0.39671, 0.229008, 0.20839, 0.169236, 0.314314, + 0.322473, 0.169506, 0.45499, 0.147819, 0.296502, 0.15198, 0.356444, 0.0992833, 0.220833, + 0.296206, 0.178067, 0.135359, 0.189725, 0.243099, 0.519986, 0.168105, 0.273465, 0.126033, + 0.18045, 0.282832, 0.193901, 0.213704, 0.425046, 0.203191, 0.228674, 0.209267, 0.355039, + 0.212918, 0.315495, 0.294112, 0.257576, 0.5786, 0.186019, 0.171919, 0.171919, 0.449151, + 1.34947, 0.171919, 0.16341, 0.641387, 0.342115, 0.267343, 0.246125, 0.277612, 0.181462, + 0.22944, 1.95598, 0.164897, 0.235803, 0.228273, 0.314629, 0.127403, 0.241241, 0.189362, + 0.151691, 0.130085, 0.526707, 0.217069, 0.282306, 0.531523, 0.177035, 0.169776, 0.20395, + 0.177165, 0.146628, 0.280013, 0.223033, 0.50947, 0.184133, 0.295329, 0.183219, 0.28166, + 0.179348, 0.276462, 1.00283, 0.248147, 0.214453, 0.231732, 0.170672, 0.256893, 0.133271, + 0.151137, 0.500823, 0.23678, 0.376983, 0.362061, 0.140013, 0.388863, 0.398552, 0.38015, + 0.190081, 0.167115, 0.206884, 0.473849, 1.05117, 0.435665, 0.323618, 0.326201, 0.32226, + 0.201787, 0.246496, 0.28325, 0.226596, 0.238153, 0.277268, 0.674629, 0.179433, 0.175651, + 0.154778, 0.178195, 0.192796, 0.103571, 0.227621, 0.201124, 0.160525, 0.160964, 0.240099, + 0.258027, 0.134127, 0.127717, 0.341378, 0.311595, 0.282306, 0.168988, 0.40775, 0.246125, + 0.583131, 0.236804, 0.238633, 0.194824, 0.169315, 0.244227, 0.249511, 0.189725, 0.305662, + 0.301415, 0.658641, 0.250944, 0.151792, 0.141383, 0.143843, 0.563347, 0.184216, 0.204155, + 0.221764, 0.314908, 0.144518, 0.228808, 0.255785, 0.163457, 0.424705, 0.170202, 0.312598, + 0.300629, 0.532614, 0.661392, 0.228273, 0.543432, 0.257175, 0.258994, 0.281413, 0.273897, + 0.246837, 0.293489, 0.25533, 0.260492, 0.213704, 0.3091, 0.17103, 0.172285, 0.241399, + 0.35999, 0.372243, 0.269191, 0.390239, 0.31761, 0.200593, 0.22197, 0.752914, 0.266571, + 0.13102, 0.268659, 0.293723, 0.356294, 0.296258, 0.264531, 0.15468, 0.358535, 0.243711, + 0.112147, 0.121659, 0.197101, 0.515292, 0.245628, 0.279863, 0.789807, 0.195156, 0.196073, + 0.149564, 0.118675, 0.389373, 0.233821, 0.176128, 0.481088, 0.360027, 0.553152, 0.208207, + 0.171608, 0.160489, 0.334298, 0.139426, 0.168603, 0.266199, 0.326458, 0.103571, 0.171208, + 0.130961, 0.190887, 0.177229, 0.241651, 0.115152, 0.196753, 0.481088, 0.230965, 0.354631, + 0.14591, 0.328543, 0.141544, 0.195888, 0.290379, 0.245954, 0.184547, 0.575214, 0.186929, + 0.28527, 0.292213, 1.20033, 0.281528, 0.15625, 0.211524, 0.186398, 0.298061, 0.147393, + 0.245349, 0.164527, 0.224771, 0.222382, 0.251643, 0.148835, 0.135359, 0.204967, 0.193024, + 0.486309, 0.389686, 0.211921, 0.307405, 0.38666, 0.26802, 0.16605, 0.323134, 0.268397, + 0.217894, 0.974118, 0.371618, 0.156201, 0.305787, 0.339305, 0.371032, 0.381765, 0.22747, + 0.24906, 0.100884, 0.253192, 0.314253, 0.388289, 0.580947, 1.00267, 0.241998, 0.489101, + 0.341501, 0.247423, 0.328311, 0.440281, 0.14927, 0.244469, 0.846828, 0.191725, 0.217429, + 0.123403, 0.322875, 0.145373, 0.757259, 0.190086, 0.316286, 0.268397, 0.296721, 0.440472, + 0.186848, 0.232134, 0.180239, 0.219724, 0.205886, 0.250975, 0.145636, 0.312476, 0.366418, + 0.128135, 0.315235, 0.264531, 0.161815, 0.31631, 0.296489, 0.37171, 0.197217, 0.195625, + 0.479579, 0.443037, 0.323347, 0.193616, 0.160251, 0.8952, 0.256291, 0.593345, 0.177165, + 0.409514, 0.847863, 0.111448, 0.210031, 0.251347, 0.351953, 0.705204, 0.117901, 0.182343, + 0.230179, 0.83632, 0.22104, 0.145163, 0.200326, 0.23431, 0.21868, 0.253575, 0.186562, + 0.192757, 0.172716, 0.27396, 0.258581, 0.327892, 0.376138, 0.223477, 0.302375, 0.145845, + 0.436902, 0.421794, 0.328543, 0.19246, 0.238889, 0.254866, 0.284674, 0.457849, 0.202937, + 0.392568, 0.453083, 0.782713, 0.465401, 0.178623, 0.304863, 0.190081, 0.228641, 0.255135, + 0.245037, 0.217526, 0.109584, 0.276462, 0.182301, 0.38582, 0.349942, 1.3889, 0.30235, + 0.796353, 0.160168, 0.643204, 0.153752, 0.410268, 0.186439, 0.256834, 0.185783, 0.0957629, + 0.226596, 0.197951, 0.17123, 0.192836, 0.18405, 0.575784, 0.228874, 0.201787, 0.241209, + 0.217386, 0.195751, 0.291585, 0.144531, 0.14176, 0.157635, 0.410268, 0.476338, 0.308148, + 0.148077, 0.152093, 0.196791, 0.568087, 0.414026, 0.250587, 0.473463, 0.293645, 0.396768, + 0.2766, 0.38664, 0.135034, 1.50827, 0.472527, 0.268418, 0.40383, 0.375914, 0.246496, + 0.176474, 0.340405, 0.220833, 0.138782, 0.159009, 0.444219, 0.259582, 0.33638, 0.195586, + 0.210974, 0.200288, 0.148129, 0.0974216, 0.211588, 0.280081, 0.44113, 0.773921, 0.553848, + 0.448079, 0.183136, 0.380854, 0.685021, 0.308767, 0.553276, 0.181578, 0.164759, 0.313889, + 0.137886, 0.545387, 0.278449, 0.736895, 0.360054, 0.358929, 0.457315, 0.343278, 0.507662, + 0.280829, 0.113886, 0.23146, 0.160584, 0.192796, 0.147561, 0.241272, 0.168988, 0.730511, + 0.27836, 0.179847, 0.22555, 0.418069, 0.158348, 0.128965, 0.179454, 0.126366, 0.164434, + 0.273633, 0.309556, 0.500823, 0.367852, 0.192875, 0.230262, 0.32724, 0.249969, 0.142618, + 0.494229, 0.36108, 0.227931, 0.23113, 0.742825, 0.190126, 0.33741, 0.280598, 0.145268, + 0.378423, 0.211921, 0.183594, 0.59201, 0.279563, 0.195683, 0.248101, 0.199754, 0.342494, + 0.174343, 0.14149, 0.28085, 0.175781, 0.518738, 0.17223, 0.489904, 0.181167, 0.354286, + 0.297824, 0.280829, 0.219412, 0.22814, 0.195625, 0.313949, 0.294708, 0.211551, 0.236255, + 0.666933, 0.204808, 0.52591, 0.180725, 0.186889, 0.246589, 0.410575, 0.338348, 0.206219, + 0.361766, 0.158143, 0.280816, 0.4149, 0.773082, 0.340046, 0.369672, 0.256923, 0.167195, + 0.197217, 0.252339, 0.172716, 0.191526, 0.263085, 0.345698, 0.168286, 0.243099, 0.434631, + 0.22944, 0.161862, 0.206589, 0.23457, 0.181924, 0.419063, 0.183427, 0.186152, 0.236352, + 0.306336, 0.149002, 1.50086, 0.188231, 0.442757, 0.485602, 0.466662, 0.17329, 0.141329, + 0.180619, 0.160061, 0.192569, 0.270999, 0.117901, 0.362693, 0.217561, 0.208975, 0.233658, + 0.175173, 1.10307, 0.14625, 1.31124, 0.237608, 0.286784, 0.325112, 0.2485, 0.259641, + 0.553152, 0.179039, 0.780781, 0.174758, 0.297824, 0.2558, 0.235949, 0.952186, 0.356744, + 0.312646, 0.189362, 0.574524, 0.705204, 0.213168, 0.225956, 0.424165, 0.169506, 0.137109, + 0.352451, 0.454554, 0.653302, 0.31261, 0.194412, 0.23719, 0.137886, 0.31498, 0.199085, + 0.203875, 0.597248, 1.10036, 0.196869, 0.22104, 0.451345, 0.105613, 0.683928, 0.135204, + 0.25533, 0.607871, 0.219724, 0.184464, 0.725001, 0.160061, 0.333407, 0.192569, 0.234147, + 0.47178, 0.161815, 0.242455, 0.215305, 0.410575, 0.242376, 0.211335, 0.462804, 0.275065, + 0.126878, 0.170404, 0.179433, 0.147244, 0.109584, 0.352905, 0.158215, 0.197604, 0.172407, + 0.407506, 0.645446, 0.313061, 0.165602, 0.136663, 0.55444, 0.15527, 0.133128, 0.125912, + 0.340405, 0.44521, 0.122783, 0.814526, 0.243773, 0.15743, 0.266743, 0.684458, 0.22221, + 0.181294, 0.193901, 0.258802, 0.167195, 0.292056, 0.132309, 0.227671, 0.117334, 0.271758, + 0.146185, 0.225042, 0.225964, 0.194863, 0.290274, 0.138438, 0.196714, 0.266012, 0.267771, + 0.162544, 0.244258, 0.358038, 0.522617, 0.192875, 0.45066, 0.330396, 0.223477, 0.42967, + 0.350884, 0.404655, 0.123155, 0.431583, 0.191675, 0.147354, 0.609034, 0.459487, 0.187337, + 0.215128, 0.604169, 0.330165, 0.494229, 0.40775, 0.167377, 0.192648, 0.234635, 0.275578, + 0.253094, 0.420063, 0.228299, 0.206478, 0.20395, 0.377656, 0.317393, 0.478623, 0.159009, + 0.217034, 0.300933, 0.139754, 0.153901, 0.261077, 0.22834, 0.449609, 0.157672, 0.176474, + 0.285704, 0.180186, 0.212738, 0.266428, 0.388313, 0.0954637, 0.298093, 0.251643, 0.330696, + 0.159572, 0.210666, 0.149411, 0.139618, 0.338472, 0.450304, 0.208793, 0.583609, 0.185865, + 0.400576, 0.21626, 0.174867, 0.239144, 0.249113, 0.200402, 0.275065, 0.238793, 0.205784, + 0.4475, 0.231262, 0.259082, 0.20934, 0.16806, 0.193616, 0.213811, 0.395632, 0.482465, + 0.274649, 0.307405, 0.165866, 0.334275, 0.683337, 0.368825, 0.14625, 0.780742, 0.163457, + 0.226596, 0.138713, 1.79155, 0.400443, 0.233658, 0.426399, 0.623024, 0.670955, 0.123588, + 0.110899, 0.173751, 0.651068, 0.199983, 0.190887, 0.541435, 0.21324, 0.266571, 0.134638, + 0.179348, 0.145636, 0.170929, 0.623252, 0.587738, 0.109688, 0.515314, 0.217666, 0.213311, + 0.249144, 0.187947, 0.270999, 0.268311, 0.469782, 0.763609, 0.32124, 0.146315, 0.265223, + 0.298694, 0.197623, 0.21349, 0.845778, 0.175466, 0.123588, 0.17223, 0.258603, 1.17119, + 0.538142, 0.407675, 0.120288, 0.587238, 0.244664, 0.333956, 0.132812, 0.21399, 0.302375, + 0.275882, 0.134284, 0.377555, 0.228541, 0.187307, 0.143804, 0.180545, 0.222451, 0.239638, + 0.188028, 0.46334, 0.175868, 0.242392, 0.314762, 0.44473, 0.21962, 0.175966, 1.12364, + 0.138837, 0.400576, 0.18184, 0.137706, 0.409763, 0.216894, 0.466662, 0.376604, 0.487155, + 0.283143, 0.118547, 0.221591, 0.122783, 0.179007, 0.16628, 0.180999, 0.239845, 0.169607, + 0.578402, 0.396537, 0.222288, 0.563237, 0.371238, 0.138658, 0.324336, 0.191526, 0.168603, + 0.357715, 0.640905, 0.460706, 0.220902, 0.240797, 0.164062, 0.157853, 0.34457, 0.196092, + 0.289353, 0.104597, 0.259641, 0.126878, 0.175781, 0.441458, 0.820108, 0.261864, 0.23431, + 0.254506, 0.271955, 0.227529, 0.22834, 0.196753, 0.224906, 0.193783, 0.419481, 0.236933, + 0.229706, 0.29785, 0.222947, 0.177606, 0.216911, 0.305188, 0.933438, 0.116666, 0.278483, + 0.0973824, 0.271224, 0.127717, 1.28139, 0.276283, 0.180704, 0.234554, 0.285984, 0.290172, + 0.49594, 0.135879, 0.436784, 0.206219, 0.342215, 0.374165, 0.182217, 0.274864, 0.625, + 0.356925, 0.194324, 0.342215, 0.113012, 0.155123, 0.254207, 0.438919, 0.262548, 0.302299, + 0.179528, 0.312744, 0.168513, 0.142618, 0.150543, 0.231361, 0.166004, 0.186725, 0.38848, + 0.179857, 0.182301, 0.629476, 0.44113, 0.289669, 0.328543, 0.279938, 0.14625, 0.187174, + 0.157635, 0.396749, 0.798931, 0.201541, 0.778619, 0.265883, 0.258027, 0.218576, 0.266571, + 0.160168, 0.230303, 0.273633, 0.233298, 0.30175, 0.217069, 0.345145, 0.397901, 0.224499, + 0.248101, 0.241335, 0.222947, 0.237094, 0.176518, 0.380032, 0.634775, 0.426193, 0.16362, + 0.231097, 0.219898, 0.343789, 0.275578, 0.282022, 0.628542, 0.232184, 0.848367, 0.200754, + 0.179177}, + {0, 0, 2, 3, 3, 0, 2, 2, 2, 2, 3, 0, 3, 2, 2, 2, 3, 3, 3, 3, 2, 0, 0, 0, 2, 3, 3, 3, 2, 2, 0, 0, + 2, 3, 3, 0, 0, 2, 0, 0, 3, 2, 3, 0, 3, 0, 3, 3, 0, 2, 0, 3, 2, 0, 3, 0, 3, 3, 3, 2, 2, 3, 0, 0, + 3, 3, 0, 2, 2, 3, 0, 3, 2, 2, 2, 0, 2, 3, 3, 3, 2, 3, 3, 3, 2, 0, 2, 0, 3, 3, 3, 3, 2, 2, 0, 2, + 0, 3, 2, 2, 2, 0, 0, 3, 0, 2, 2, 3, 2, 3, 0, 2, 2, 2, 3, 2, 0, 0, 2, 3, 3, 2, 0, 2, 0, 0, 2, 0, + 2, 2, 3, 2, 2, 0, 3, 0, 3, 2, 2, 2, 3, 3, 0, 0, 0, 3, 2, 3, 3, 3, 3, 0, 2, 0, 3, 2, 3, 2, 3, 0, + 2, 3, 3, 2, 3, 3, 2, 2, 0, 0, 2, 3, 3, 2, 3, 0, 2, 0, 2, 0, 3, 2, 3, 2, 3, 0, 3, 0, 3, 0, 2, 3, + 2, 2, 3, 0, 2, 2, 2, 0, 3, 2, 3, 3, 2, 3, 2, 3, 3, 2, 2, 0, 0, 2, 2, 3, 0, 3, 0, 2, 0, 0, 2, 3, + 0, 3, 3, 2, 0, 3, 3, 0, 3, 0, 2, 2, 0, 2, 0, 2, 0, 0, 0, 2, 0, 3, 2, 3, 2, 3, 2, 2, 0, 2, 3, 2, + 3, 2, 2, 2, 2, 3, 0, 2, 0, 0, 2, 3, 3, 0, 2, 3, 2, 2, 3, 0, 3, 0, 0, 2, 0, 2, 0, 2, 2, 3, 3, 2, + 3, 0, 0, 3, 2, 2, 0, 3, 2, 0, 0, 3, 0, 0, 2, 0, 3, 2, 0, 2, 0, 0, 0, 0, 0, 2, 0, 0, 2, 3, 0, 0, + 2, 0, 0, 2, 0, 2, 3, 2, 3, 3, 2, 2, 0, 0, 0, 3, 0, 2, 0, 2, 0, 2, 2, 2, 3, 3, 0, 0, 3, 3, 3, 3, + 3, 2, 3, 3, 2, 3, 3, 0, 2, 2, 2, 2, 0, 2, 0, 0, 0, 2, 2, 3, 3, 2, 3, 2, 3, 0, 2, 3, 0, 2, 0, 2, + 2, 0, 3, 0, 2, 0, 2, 3, 0, 3, 0, 0, 0, 3, 2, 3, 3, 0, 3, 2, 3, 0, 2, 3, 3, 0, 2, 3, 0, 0, 0, 2, + 0, 3, 0, 2, 3, 3, 3, 3, 3, 0, 2, 0, 2, 2, 3, 3, 0, 3, 0, 2, 0, 2, 0, 3, 0, 0, 0, 2, 3, 3, 2, 3, + 0, 0, 0, 0, 3, 3, 0, 3, 2, 0, 2, 3, 2, 2, 3, 3, 2, 2, 2, 0, 2, 3, 0, 3, 3, 0, 0, 2, 0, 3, 2, 3, + 0, 2, 0, 2, 2, 3, 2, 0, 3, 3, 3, 2, 3, 0, 3, 0, 2, 2, 0, 0, 0, 3, 0, 3, 3, 2, 3, 2, 3, 2, 3, 0, + 2, 3, 0, 2, 0, 3, 3, 3, 3, 3, 3, 2, 0, 3, 2, 2, 2, 3, 3, 2, 3, 0, 2, 3, 3, 2, 2, 0, 0, 0, 0, 3, + 0, 3, 3, 3, 0, 0, 0, 3, 3, 3, 3, 3, 0, 2, 3, 3, 3, 3, 3, 3, 0, 0, 2, 2, 3, 3, 2, 2, 0, 0, 3, 0, + 0, 0, 2, 3, 0, 0, 0, 3, 0, 3, 0, 2, 2, 0, 0, 0, 0, 3, 2, 2, 3, 2, 3, 2, 2, 2, 2, 3, 0, 0, 2, 3, + 0, 3, 3, 0, 3, 0, 0, 2, 0, 3, 3, 0, 2, 2, 3, 3, 0, 0, 2, 0, 2, 3, 2, 0, 0, 3, 3, 0, 3, 2, 0, 2, + 0, 2, 3, 2, 0, 3, 3, 2, 0, 0, 2, 2, 0, 0, 2, 0, 3, 3, 2, 3, 2, 0, 3, 0, 2, 2, 3, 3, 0, 3, 2, 2, + 0, 3, 0, 0, 0, 2, 0, 3, 2, 0, 2, 3, 2, 3, 2, 2, 3, 3, 0, 2, 3, 2, 3, 2, 2, 0, 3, 0, 3, 0, 2, 2, + 2, 0, 2, 0, 2, 2, 0, 0, 3, 3, 0, 0, 3, 2, 0, 2, 3, 2, 2, 0, 3, 3, 0, 2, 0, 3, 3, 0, 2, 3, 2, 3, + 2, 0, 2, 2, 0, 0, 0, 2, 2, 3, 3, 2, 2, 0, 2, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 2, 3, 2, 0, 3, 3, + 3, 0, 2, 0, 2, 3, 2, 0, 3, 3, 2, 0, 2, 0, 3, 2, 0, 3, 0, 0, 2, 2, 0, 3, 0, 2, 3, 3, 3, 0, 2, 0, + 0, 3, 0, 2, 3, 2, 2, 0, 3, 3, 3, 3, 3, 0, 3, 0, 0, 0, 0, 3, 2, 0, 0, 2, 3, 3, 2, 2, 0, 3, 2, 0, + 3, 0, 2, 3, 3, 0, 2, 2, 3, 2, 2, 2, 3, 2, 0, 0, 3, 2, 0, 0, 0, 2, 0, 2, 0, 0, 2, 2, 3, 0, 3, 0, + 0, 3, 0, 0, 0, 3, 0, 0, 2, 2, 0, 2, 2, 3, 3, 3, 3, 0, 0, 2, 2, 2, 0, 3, 2, 2, 2, 2, 2, 0, 3, 0, + 0, 3, 2, 0, 0, 3, 2, 3, 3, 0, 3, 0, 3, 0, 3, 2, 2, 2, 0, 0, 3, 2, 2, 0, 0, 0, 2, 3, 2, 0, 2, 3, + 3, 3, 0, 3, 3, 0, 2, 0, 0, 2, 3, 3, 0, 3, 2, 2, 2, 2, 2, 3, 3, 2, 2, 3, 3, 2, 3, 0, 3, 3, 0, 3, + 2, 2, 0, 2, 0, 3, 0, 3, 0, 2, 3, 0, 2, 3, 2, 0, 2, 0, 3, 0, 2, 3, 3, 2, 0, 3, 3, 3, 2, 2, 3, 3, + 2, 2, 2, 0, 3, 2, 2, 0}, + {271, 271, 329, 343, 387, 426, 426, 601}, + {426, 601, 426, 387, 343, 271, 329, 271}, + {3.70991, 4.43491, 3.76334, 9.43944, 9.43944, 3.70991, 3.76334, 4.43491}}}; + +typedef ConnectComponentsEdgesTest ConnectComponentsEdgesTestF_Int; +TEST_P(ConnectComponentsEdgesTestF_Int, Result) { EXPECT_TRUE(true); } + +INSTANTIATE_TEST_CASE_P(ConnectComponentsEdgesTest, + ConnectComponentsEdgesTestF_Int, + ::testing::ValuesIn(mr_fix_conn_inputsf2)); + +}; // namespace sparse +}; // end namespace cuvs diff --git a/docs/source/integrations.rst b/docs/source/integrations.rst index 46cbaa52c9..19d72fd90e 100644 --- a/docs/source/integrations.rst +++ b/docs/source/integrations.rst @@ -5,7 +5,7 @@ Aside from using cuVS directly, it can be consumed through a number of sdk and v - `FAISS`_ - `Milvus`_ -- `Lucene` _ +- `Lucene`_ - `Kinetica`_