From ec5689b1ac8170474e2df243a093e383b7a30a68 Mon Sep 17 00:00:00 2001 From: Tamas Bela Feher Date: Mon, 3 Jul 2023 23:29:46 +0200 Subject: [PATCH] Fix instantiations --- .../neighbors/detail/cagra/graph_core.cuh | 16 ++---- .../cagra/search_single_cta_kernel-inl.cuh | 1 + .../cagra/search_single_cta_00_generate.py | 1 + cpp/test/CMakeLists.txt | 50 +++++++------------ cpp/test/neighbors/ann_cagra.cuh | 4 +- .../neighbors/ann_cagra/test_float_int64_t.cu | 4 +- .../ann_cagra/test_float_uint32_t.cu | 2 + .../ann_cagra/test_int8_t_uint32_t.cu | 2 + .../ann_cagra/test_uint8_t_uint32_t.cu | 2 + 9 files changed, 36 insertions(+), 46 deletions(-) diff --git a/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh b/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh index 949dcfda8b..d915634df9 100644 --- a/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh @@ -72,7 +72,6 @@ template __global__ void kern_sort(const DATA_T* const dataset, // [dataset_chunk_size, dataset_dim] const IdxT dataset_size, const uint32_t dataset_dim, - const uint32_t dataset_ld, IdxT* const knn_graph, // [graph_chunk_size, graph_degree] const uint32_t graph_size, const uint32_t graph_degree) @@ -91,9 +90,9 @@ __global__ void kern_sort(const DATA_T* const dataset, // [dataset_chunk_size, float dist = 0.0; for (int d = lane_id; d < dataset_dim; d += raft::WarpSize) { float diff = spatial::knn::detail::utils::mapping{}( - dataset[d + static_cast(dataset_ld) * srcNode]) - + dataset[d + static_cast(dataset_dim) * srcNode]) - spatial::knn::detail::utils::mapping{}( - dataset[d + static_cast(dataset_ld) * dstNode]); + dataset[d + static_cast(dataset_dim) * dstNode]); dist += diff * diff; } dist += __shfl_xor_sync(0xffffffff, dist, 1); @@ -239,7 +238,6 @@ void sort_knn_graph(raft::resources const& res, "dataset size is expected to have the same number of graph index size"); const uint32_t dataset_size = dataset.extent(0); const uint32_t dataset_dim = dataset.extent(1); - const uint32_t dataset_ld = dataset.stride(0); const DataT* dataset_ptr = dataset.data_handle(); const IdxT graph_size = dataset_size; @@ -265,13 +263,8 @@ void sort_knn_graph(raft::resources const& res, graph_size * input_graph_degree, resource::get_cuda_stream(res)); - void (*kernel_sort)(const DataT* const, - const IdxT, - const uint32_t, - const uint32_t, - IdxT* const, - const uint32_t, - const uint32_t); + void (*kernel_sort)( + const DataT* const, const IdxT, const uint32_t, IdxT* const, const uint32_t, const uint32_t); if (input_graph_degree <= 32) { constexpr int numElementsPerThread = 1; kernel_sort = kern_sort; @@ -306,7 +299,6 @@ void sort_knn_graph(raft::resources const& res, d_dataset.data_handle(), dataset_size, dataset_dim, - dataset_ld, d_input_graph.data_handle(), graph_size, input_graph_degree); diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh index 8100db9dcb..db3bcc3e9c 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh @@ -17,6 +17,7 @@ #include #include +#include #include #include #include diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py b/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py index 44da773157..bb5e7d6838 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py @@ -111,6 +111,7 @@ float_uint32=("float", "uint32_t", "float"), # data_t, idx_t, distance_t int8_uint32=("int8_t", "uint32_t", "float"), uint8_uint32=("uint8_t", "uint32_t", "float"), + float_uint64=("float", "uint64_t", "float"), ) # knn diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt index 413cd89ce7..01f2380bfd 100644 --- a/cpp/test/CMakeLists.txt +++ b/cpp/test/CMakeLists.txt @@ -316,37 +316,25 @@ if(BUILD_TESTS) NEIGHBORS_TEST PATH test/neighbors/ann_cagra/test_float_uint32_t.cu - # test/neighbors/ann_cagra/test_int8_t_uint32_t.cu - # test/neighbors/ann_cagra/test_uint8_t_uint32_t.cu - # test/neighbors/ann_cagra/test_float_int64_t.cu - # test/neighbors/ann_ivf_flat/test_float_int64_t.cu - # test/neighbors/ann_ivf_flat/test_int8_t_int64_t.cu - # test/neighbors/ann_ivf_flat/test_uint8_t_int64_t.cu - # test/neighbors/ann_ivf_pq/test_float_int64_t.cu - # test/neighbors/ann_ivf_pq/test_float_uint32_t.cu - # test/neighbors/ann_ivf_pq/test_float_int64_t.cu - # test/neighbors/ann_ivf_pq/test_int8_t_int64_t.cu - # test/neighbors/ann_ivf_pq/test_uint8_t_int64_t.cu - # test/neighbors/knn.cu - # test/neighbors/fused_l2_knn.cu - # test/neighbors/tiled_knn.cu - # test/neighbors/haversine.cu - # test/neighbors/ball_cover.cu - # test/neighbors/epsilon_neighborhood.cu - # test/neighbors/refine.cu - # test/neighbors/selection.cu - # src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu + test/neighbors/ann_cagra/test_int8_t_uint32_t.cu + test/neighbors/ann_cagra/test_uint8_t_uint32_t.cu + test/neighbors/ann_cagra/test_float_int64_t.cu + test/neighbors/ann_ivf_flat/test_float_int64_t.cu + test/neighbors/ann_ivf_flat/test_int8_t_int64_t.cu + test/neighbors/ann_ivf_flat/test_uint8_t_int64_t.cu + test/neighbors/ann_ivf_pq/test_float_int64_t.cu + test/neighbors/ann_ivf_pq/test_float_uint32_t.cu + test/neighbors/ann_ivf_pq/test_float_int64_t.cu + test/neighbors/ann_ivf_pq/test_int8_t_int64_t.cu + test/neighbors/ann_ivf_pq/test_uint8_t_int64_t.cu + test/neighbors/knn.cu + test/neighbors/fused_l2_knn.cu + test/neighbors/tiled_knn.cu + test/neighbors/haversine.cu + test/neighbors/ball_cover.cu + test/neighbors/epsilon_neighborhood.cu + test/neighbors/refine.cu + test/neighbors/selection.cu src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu diff --git a/cpp/test/neighbors/ann_cagra.cuh b/cpp/test/neighbors/ann_cagra.cuh index 10d6d4b679..ed8d85dd7b 100644 --- a/cpp/test/neighbors/ann_cagra.cuh +++ b/cpp/test/neighbors/ann_cagra.cuh @@ -15,8 +15,8 @@ */ #pragma once -#define RAFT_EXPLICIT_INSTANTIATE_ONLY_CAGRA -#define RAFT_COMPILED_CAGRA +// #define RAFT_EXPLICIT_INSTANTIATE_ONLY_CAGRA +// #define RAFT_COMPILED_CAGRA #include "../test_utils.cuh" #include "ann_utils.cuh" #include diff --git a/cpp/test/neighbors/ann_cagra/test_float_int64_t.cu b/cpp/test/neighbors/ann_cagra/test_float_int64_t.cu index e473a72b2b..b76ff478db 100644 --- a/cpp/test/neighbors/ann_cagra/test_float_int64_t.cu +++ b/cpp/test/neighbors/ann_cagra/test_float_int64_t.cu @@ -16,7 +16,9 @@ #include -#undef RAFT_EXPLICIT_INSTANTIATE_ONLY +#undef RAFT_EXPLICIT_INSTANTIATE_ONLY_CAGRA +#undef RAFT_COMPILED_CAGRA + #include "../ann_cagra.cuh" namespace raft::neighbors::experimental::cagra { diff --git a/cpp/test/neighbors/ann_cagra/test_float_uint32_t.cu b/cpp/test/neighbors/ann_cagra/test_float_uint32_t.cu index dbaf4dedd9..28d28d473d 100644 --- a/cpp/test/neighbors/ann_cagra/test_float_uint32_t.cu +++ b/cpp/test/neighbors/ann_cagra/test_float_uint32_t.cu @@ -14,6 +14,8 @@ * limitations under the License. */ +#define RAFT_EXPLICIT_INSTANTIATE_ONLY_CAGRA +#define RAFT_COMPILED_CAGRA #include #include "../ann_cagra.cuh" diff --git a/cpp/test/neighbors/ann_cagra/test_int8_t_uint32_t.cu b/cpp/test/neighbors/ann_cagra/test_int8_t_uint32_t.cu index ba60131677..30e76643e2 100644 --- a/cpp/test/neighbors/ann_cagra/test_int8_t_uint32_t.cu +++ b/cpp/test/neighbors/ann_cagra/test_int8_t_uint32_t.cu @@ -14,6 +14,8 @@ * limitations under the License. */ +#define RAFT_EXPLICIT_INSTANTIATE_ONLY_CAGRA +#define RAFT_COMPILED_CAGRA #include #include "../ann_cagra.cuh" diff --git a/cpp/test/neighbors/ann_cagra/test_uint8_t_uint32_t.cu b/cpp/test/neighbors/ann_cagra/test_uint8_t_uint32_t.cu index cc172e4833..21af8ae14c 100644 --- a/cpp/test/neighbors/ann_cagra/test_uint8_t_uint32_t.cu +++ b/cpp/test/neighbors/ann_cagra/test_uint8_t_uint32_t.cu @@ -13,6 +13,8 @@ * See the License for the specific language governing permissions and * limitations under the License. */ +#define RAFT_EXPLICIT_INSTANTIATE_ONLY_CAGRA +#define RAFT_COMPILED_CAGRA #include