diff --git a/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh b/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh index 31f2acd0b7..b927f184e2 100644 --- a/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh @@ -180,10 +180,8 @@ void launch_vpq_search_main_core( DatasetT, 8 /*PQ bit*/, 2 /* Subspace dimension*/, - 0, DistanceT, - InternalIdxT, - 0>; + InternalIdxT>; dataset_desc_t dataset_desc(vpq_dset->data.data_handle(), vpq_dset->encoded_row_length(), vpq_dset->pq_dim(), @@ -200,10 +198,8 @@ void launch_vpq_search_main_core( DatasetT, 8 /*PQ bit*/, 4 /* Subspace dimension*/, - 0, DistanceT, - InternalIdxT, - 0>; + InternalIdxT>; dataset_desc_t dataset_desc(vpq_dset->data.data_handle(), vpq_dset->encoded_row_length(), vpq_dset->pq_dim(), @@ -266,7 +262,7 @@ void search_main(raft::resources const& res, strided_dset != nullptr) { // Set TEAM_SIZE and DATASET_BLOCK_SIZE to zero tentatively since these parameters cannot be // determined here. They are set just before kernel launch. - using dataset_desc_t = standard_dataset_descriptor_t; + using dataset_desc_t = standard_dataset_descriptor_t; // Search using a plain (strided) row-major dataset const dataset_desc_t dataset_desc(strided_dset->view().data_handle(), strided_dset->n_rows(), diff --git a/cpp/include/raft/neighbors/detail/cagra/compute_distance.hpp b/cpp/include/raft/neighbors/detail/cagra/compute_distance.hpp index 7fffef2497..2436d0a3ca 100644 --- a/cpp/include/raft/neighbors/detail/cagra/compute_distance.hpp +++ b/cpp/include/raft/neighbors/detail/cagra/compute_distance.hpp @@ -78,7 +78,8 @@ _RAFT_DEVICE void compute_distance_to_random_nodes( } } - const auto norm2 = dataset_desc.compute_similarity(query_buffer, seed_index, valid_i); + const auto norm2 = dataset_desc.template compute_similarity( + query_buffer, seed_index, valid_i); if (valid_i && (norm2 < best_norm2_team_local)) { best_norm2_team_local = norm2; @@ -152,8 +153,8 @@ _RAFT_DEVICE void compute_distance_to_child_nodes( INDEX_T child_id = invalid_index; if (valid_i) { child_id = result_child_indices_ptr[i]; } - const auto norm2 = - dataset_desc.compute_similarity(query_buffer, child_id, child_id != invalid_index); + const auto norm2 = dataset_desc.template compute_similarity( + query_buffer, child_id, child_id != invalid_index); // Store the distance const unsigned lane_id = threadIdx.x % TEAM_SIZE; @@ -181,18 +182,12 @@ struct dataset_descriptor_base_t { dataset_descriptor_base_t(const INDEX_T size, const std::uint32_t dim) : size(size), dim(dim) {} }; -template +template struct standard_dataset_descriptor_t : public dataset_descriptor_base_t { using LOAD_T = device::LOAD_128BIT_T; using DATA_T = DATA_T_; using QUERY_T = typename dataset_descriptor_base_t::QUERY_T; - static const std::uint32_t DATASET_BLOCK_DIM = DATASET_BLOCK_DIM_; - static const std::uint32_t TEAM_SIZE = TEAM_SIZE_; const DATA_T* const ptr; const std::size_t ld; @@ -210,6 +205,7 @@ struct standard_dataset_descriptor_t static const std::uint32_t smem_buffer_size_in_byte = 0; __device__ void set_smem_ptr(void* const){}; + template __device__ DISTANCE_T compute_similarity(const QUERY_T* const query_ptr, const INDEX_T dataset_i, const bool valid) const @@ -255,24 +251,4 @@ struct standard_dataset_descriptor_t } }; -template -standard_dataset_descriptor_t -set_compute_template_params( - standard_dataset_descriptor_t& - desc_in) -{ - return standard_dataset_descriptor_t( - desc_in.ptr, desc_in.size, desc_in.dim, desc_in.ld); -} - } // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/compute_distance_vpq.cuh b/cpp/include/raft/neighbors/detail/cagra/compute_distance_vpq.cuh index a29f22f506..998cf40f77 100644 --- a/cpp/include/raft/neighbors/detail/cagra/compute_distance_vpq.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/compute_distance_vpq.cuh @@ -24,18 +24,14 @@ namespace raft::neighbors::cagra::detail { template + class INDEX_T> struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t { using LOAD_T = device::LOAD_128BIT_T; using DATA_T = DATA_T_; using CODE_BOOK_T = CODE_BOOK_T_; using QUERY_T = typename dataset_descriptor_base_t::QUERY_T; - static const std::uint32_t DATASET_BLOCK_DIM = DATASET_BLOCK_DIM_; - static const std::uint32_t TEAM_SIZE = TEAM_SIZE_; const std::uint8_t* encoded_dataset_ptr; const std::uint32_t encoded_dataset_dim; @@ -50,7 +46,7 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t(); + (1 << PQ_BITS) * PQ_LEN * utils::size_of(); __device__ void set_smem_ptr(void* const smem_ptr) { @@ -58,15 +54,14 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t::value) { - for (unsigned i = threadIdx.x * 2; i < (1 << PQ_BITS) * PQ_CODE_BOOK_DIM; - i += blockDim.x * 2) { + for (unsigned i = threadIdx.x * 2; i < (1 << PQ_BITS) * PQ_LEN; i += blockDim.x * 2) { half2 buf2; buf2.x = pq_code_book_ptr[i]; buf2.y = pq_code_book_ptr[i + 1]; (reinterpret_cast(smem_pq_code_book_ptr + i))[0] = buf2; } } else { - for (unsigned i = threadIdx.x; i < (1 << PQ_BITS) * PQ_CODE_BOOK_DIM; i += blockDim.x) { + for (unsigned i = threadIdx.x; i < (1 << PQ_BITS) * PQ_LEN; i += blockDim.x) { // TODO: vectorize smem_pq_code_book_ptr[i] = pq_code_book_ptr[i]; } @@ -93,6 +88,7 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t __device__ DISTANCE_T compute_similarity(const QUERY_T* const query_ptr, const INDEX_T node_id, const bool valid) const @@ -104,9 +100,9 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t(encoded_dataset_dim) * node_id))); if (PQ_BITS == 8) { for (uint32_t elem_offset = 0; elem_offset < dim; elem_offset += DATASET_BLOCK_DIM) { - constexpr unsigned vlen = 4; // **** DO NOT CHANGE **** - constexpr unsigned nelem = raft::div_rounding_up_unsafe( - DATASET_BLOCK_DIM / PQ_CODE_BOOK_DIM, TEAM_SIZE * vlen); + constexpr unsigned vlen = 4; // **** DO NOT CHANGE **** + constexpr unsigned nelem = + raft::div_rounding_up_unsafe(DATASET_BLOCK_DIM / PQ_LEN, TEAM_SIZE * vlen); // Loading PQ codes uint32_t pq_codes[nelem]; #pragma unroll @@ -119,7 +115,7 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t::value) && (PQ_CODE_BOOK_DIM % 2 == 0)) { + if constexpr ((std::is_same::value) && (PQ_LEN % 2 == 0)) { // **** Use half2 for distance computation **** half2 norm2{0, 0}; #pragma unroll @@ -127,10 +123,10 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t= n_subspace) break; // Loading VQ code-book - raft::TxN_t vq_vals[PQ_CODE_BOOK_DIM]; + raft::TxN_t vq_vals[PQ_LEN]; #pragma unroll - for (std::uint32_t m = 0; m < PQ_CODE_BOOK_DIM; m += 1) { - const uint32_t d = (vlen * m) + (PQ_CODE_BOOK_DIM * k) + elem_offset; + for (std::uint32_t m = 0; m < PQ_LEN; m += 1) { + const uint32_t d = (vlen * m) + (PQ_LEN * k) + elem_offset; if (d >= dim) break; vq_vals[m].load( reinterpret_cast(vq_code_book_ptr + d + (dim * vq_code)), 0); @@ -139,11 +135,11 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t= dim) break; + if (PQ_LEN * (v + k) >= dim) break; #pragma unroll - for (std::uint32_t m = 0; m < PQ_CODE_BOOK_DIM; m += 2) { - const std::uint32_t d1 = m + (PQ_CODE_BOOK_DIM * v); - const std::uint32_t d = d1 + (PQ_CODE_BOOK_DIM * k); + for (std::uint32_t m = 0; m < PQ_LEN; m += 2) { + const std::uint32_t d1 = m + (PQ_LEN * v); + const std::uint32_t d = d1 + (PQ_LEN * k); // Loading query vector in smem half2 diff2 = (reinterpret_cast( query_ptr))[device::swizzling(d / 2)]; @@ -164,10 +160,10 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t= n_subspace) break; // Loading VQ code-book - raft::TxN_t vq_vals[PQ_CODE_BOOK_DIM]; + raft::TxN_t vq_vals[PQ_LEN]; #pragma unroll - for (std::uint32_t m = 0; m < PQ_CODE_BOOK_DIM; m++) { - const std::uint32_t d = (vlen * m) + (PQ_CODE_BOOK_DIM * k) + elem_offset; + for (std::uint32_t m = 0; m < PQ_LEN; m++) { + const std::uint32_t d = (vlen * m) + (PQ_LEN * k) + elem_offset; if (d >= dim) break; // Loading 4 x 8/16-bit VQ-values using 32/64-bit load ops (from L2$ or device // memory) @@ -178,15 +174,15 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t= dim) break; - raft::TxN_t pq_vals; - pq_vals.load(reinterpret_cast(smem_pq_code_book_ptr + - PQ_CODE_BOOK_DIM * (pq_code & 0xff)), - 0); // (from L1$ or smem) + if (PQ_LEN * (v + k) >= dim) break; + raft::TxN_t pq_vals; + pq_vals.load( + reinterpret_cast(smem_pq_code_book_ptr + PQ_LEN * (pq_code & 0xff)), + 0); // (from L1$ or smem) #pragma unroll - for (std::uint32_t m = 0; m < PQ_CODE_BOOK_DIM; m++) { - const std::uint32_t d1 = m + (PQ_CODE_BOOK_DIM * v); - const std::uint32_t d = d1 + (PQ_CODE_BOOK_DIM * k); + for (std::uint32_t m = 0; m < PQ_LEN; m++) { + const std::uint32_t d1 = m + (PQ_LEN * v); + const std::uint32_t d = d1 + (PQ_LEN * k); // if (d >= dataset_dim) break; DISTANCE_T diff = query_ptr[d]; // (from smem) diff -= pq_scale * static_cast(pq_vals.data[m]); @@ -207,48 +203,4 @@ struct cagra_q_dataset_descriptor_t : public dataset_descriptor_base_t -cagra_q_dataset_descriptor_t -set_compute_template_params(cagra_q_dataset_descriptor_t& desc_in) -{ - return cagra_q_dataset_descriptor_t(desc_in.encoded_dataset_ptr, - desc_in.encoded_dataset_dim, - desc_in.n_subspace, - desc_in.vq_code_book_ptr, - desc_in.vq_scale, - desc_in.pq_code_book_ptr, - desc_in.pq_scale, - desc_in.size, - desc_in.dim); -} } // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh index 1e3095771f..50f9e69593 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh @@ -57,37 +57,37 @@ void select_and_run( cudaStream_t stream) RAFT_EXPLICIT; #endif // RAFT_EXPLICIT_INSTANTIATE_ONLY -#define instantiate_kernel_selection( \ - TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_T) \ - extern template void \ - select_and_run, \ - SAMPLE_FILTER_T>( \ - raft::neighbors::cagra::detail::standard_dataset_descriptor_t dataset_desc, \ - raft::device_matrix_view graph, \ - INDEX_T* const topk_indices_ptr, \ - DISTANCE_T* const topk_distances_ptr, \ - const DATA_T* const queries_ptr, \ - const uint32_t num_queries, \ - const INDEX_T* dev_seed_ptr, \ - uint32_t* const num_executed_iterations, \ - uint32_t topk, \ - uint32_t block_size, \ - uint32_t result_buffer_size, \ - uint32_t smem_size, \ - int64_t hash_bitlen, \ - INDEX_T* hashmap_ptr, \ - uint32_t num_cta_per_query, \ - uint32_t num_random_samplings, \ - uint64_t rand_xor_mask, \ - uint32_t num_seeds, \ - size_t itopk_size, \ - size_t search_width, \ - size_t min_iterations, \ - size_t max_iterations, \ - SAMPLE_FILTER_T sample_filter, \ +#define instantiate_kernel_selection( \ + TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_T) \ + extern template void select_and_run< \ + TEAM_SIZE, \ + MAX_DATASET_DIM, \ + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, \ + SAMPLE_FILTER_T>( \ + raft::neighbors::cagra::detail::standard_dataset_descriptor_t \ + dataset_desc, \ + raft::device_matrix_view graph, \ + INDEX_T* const topk_indices_ptr, \ + DISTANCE_T* const topk_distances_ptr, \ + const DATA_T* const queries_ptr, \ + const uint32_t num_queries, \ + const INDEX_T* dev_seed_ptr, \ + uint32_t* const num_executed_iterations, \ + uint32_t topk, \ + uint32_t block_size, \ + uint32_t result_buffer_size, \ + uint32_t smem_size, \ + int64_t hash_bitlen, \ + INDEX_T* hashmap_ptr, \ + uint32_t num_cta_per_query, \ + uint32_t num_random_samplings, \ + uint64_t rand_xor_mask, \ + uint32_t num_seeds, \ + size_t itopk_size, \ + size_t search_width, \ + size_t min_iterations, \ + size_t max_iterations, \ + SAMPLE_FILTER_T sample_filter, \ cudaStream_t stream); instantiate_kernel_selection( @@ -141,19 +141,15 @@ instantiate_kernel_selection( CODE_BOOK_T, \ PQ_BITS, \ PQ_CODE_BOOK_DIM, \ - 0, \ DISTANCE_T, \ - INDEX_T, \ - 0>, \ + INDEX_T>, \ SAMPLE_FILTER_T>( \ raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t dataset_desc, \ + INDEX_T> dataset_desc, \ raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh index eb771c9325..001b007000 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh @@ -123,7 +123,11 @@ __device__ inline void topk_by_bitonic_sort(float* distances, // [num_elements] // // multiple CTAs per single query // -template +template __launch_bounds__(1024, 1) RAFT_KERNEL search_kernel( typename DATASET_DESCRIPTOR_T::INDEX_T* const result_indices_ptr, // [num_queries, num_cta_per_query, itopk_size] @@ -147,12 +151,10 @@ __launch_bounds__(1024, 1) RAFT_KERNEL search_kernel( uint32_t* const num_executed_iterations, /* stats */ SAMPLE_FILTER_T sample_filter) { - constexpr std::uint32_t TEAM_SIZE = DATASET_DESCRIPTOR_T::TEAM_SIZE; - constexpr std::uint32_t DATASET_BLOCK_DIM = DATASET_DESCRIPTOR_T::DATASET_BLOCK_DIM; - using DATA_T = typename DATASET_DESCRIPTOR_T::DATA_T; - using INDEX_T = typename DATASET_DESCRIPTOR_T::INDEX_T; - using DISTANCE_T = typename DATASET_DESCRIPTOR_T::DISTANCE_T; - using QUERY_T = typename DATASET_DESCRIPTOR_T::QUERY_T; + using DATA_T = typename DATASET_DESCRIPTOR_T::DATA_T; + using INDEX_T = typename DATASET_DESCRIPTOR_T::INDEX_T; + using DISTANCE_T = typename DATASET_DESCRIPTOR_T::DISTANCE_T; + using QUERY_T = typename DATASET_DESCRIPTOR_T::QUERY_T; const auto num_queries = gridDim.y; const auto query_id = blockIdx.y; @@ -398,21 +400,36 @@ void set_value_batch(T* const dev_ptr, <<>>(dev_ptr, ld, val, count, batch_size); } -template +template struct search_kernel_config { // Search kernel function type. Note that the actual values for the template value // parameters do not matter, because they are not part of the function signature. The // second to fourth value parameters will be selected by the choose_* functions below. - using kernel_t = decltype(&search_kernel<128, DATASET_DESCRIPTOR_T, SAMPLE_FILTER_T>); + using kernel_t = decltype(&search_kernel); static auto choose_buffer_size(unsigned result_buffer_size, unsigned block_size) -> kernel_t { if (result_buffer_size <= 64) { - return search_kernel<64, DATASET_DESCRIPTOR_T, SAMPLE_FILTER_T>; + return search_kernel; } else if (result_buffer_size <= 128) { - return search_kernel<128, DATASET_DESCRIPTOR_T, SAMPLE_FILTER_T>; + return search_kernel; } else if (result_buffer_size <= 256) { - return search_kernel<256, DATASET_DESCRIPTOR_T, SAMPLE_FILTER_T>; + return search_kernel; } THROW("Result buffer size %u larger than max buffer size %u", result_buffer_size, 256); } @@ -449,16 +466,13 @@ void select_and_run( SAMPLE_FILTER_T sample_filter, cudaStream_t stream) { - const auto dataset_desc_ = - set_compute_template_params(dataset_desc); - - using dataset_desc_t = typename std::remove_const::type; - auto kernel = search_kernel_config::choose_buffer_size( - result_buffer_size, block_size); + auto kernel = + search_kernel_config:: + choose_buffer_size(result_buffer_size, block_size); RAFT_CUDA_TRY(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, - smem_size + dataset_desc_t::smem_buffer_size_in_byte)); + smem_size + DATASET_DESCRIPTOR_T::smem_buffer_size_in_byte)); // Initialize hash table const uint32_t hash_size = hashmap::get_size(hash_bitlen); set_value_batch(hashmap_ptr, @@ -477,7 +491,7 @@ void select_and_run( smem_size); kernel<<>>(topk_indices_ptr, topk_distances_ptr, - dataset_desc_, + dataset_desc, queries_ptr, graph.data_handle(), graph.extent(1), diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh index 05a02f7384..10788da432 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh @@ -87,7 +87,7 @@ void get_value(T* const host_ptr, const T* const dev_ptr, cudaStream_t cuda_stre } // MAX_DATASET_DIM : must equal to or greater than dataset_dim -template +template RAFT_KERNEL random_pickup_kernel( const DATASET_DESCRIPTOR_T dataset_desc, const typename DATASET_DESCRIPTOR_T::DATA_T* const queries_ptr, // [num_queries, dataset_dim] @@ -102,11 +102,9 @@ RAFT_KERNEL random_pickup_kernel( typename DATASET_DESCRIPTOR_T::INDEX_T* const visited_hashmap_ptr, // [num_queries, 1 << bitlen] const std::uint32_t hash_bitlen) { - constexpr std::uint32_t TEAM_SIZE = DATASET_DESCRIPTOR_T::TEAM_SIZE; - constexpr std::uint32_t DATASET_BLOCK_DIM = DATASET_DESCRIPTOR_T::DATASET_BLOCK_DIM; - using DATA_T = typename DATASET_DESCRIPTOR_T::DATA_T; - using INDEX_T = typename DATASET_DESCRIPTOR_T::INDEX_T; - using DISTANCE_T = typename DATASET_DESCRIPTOR_T::DISTANCE_T; + using DATA_T = typename DATASET_DESCRIPTOR_T::DATA_T; + using INDEX_T = typename DATASET_DESCRIPTOR_T::INDEX_T; + using DISTANCE_T = typename DATASET_DESCRIPTOR_T::DISTANCE_T; const auto ldb = hashmap::get_size(hash_bitlen); const auto global_team_index = (blockIdx.x * blockDim.x + threadIdx.x) / TEAM_SIZE; @@ -139,7 +137,8 @@ RAFT_KERNEL random_pickup_kernel( device::xorshift64((global_team_index ^ rand_xor_mask) * (i + 1)) % dataset_desc.size; } - const auto norm2 = dataset_desc.compute_similarity(query_buffer, seed_index, true); + const auto norm2 = dataset_desc.template compute_similarity( + query_buffer, seed_index, true); if (norm2 < best_norm2_team_local) { best_norm2_team_local = norm2; @@ -161,7 +160,7 @@ RAFT_KERNEL random_pickup_kernel( } // MAX_DATASET_DIM : must be equal to or greater than dataset_dim -template +template void random_pickup( const DATASET_DESCRIPTOR_T dataset_desc, const typename DATASET_DESCRIPTOR_T::DATA_T* const queries_ptr, // [num_queries, dataset_dim] @@ -178,10 +177,8 @@ void random_pickup( const std::uint32_t hash_bitlen, cudaStream_t const cuda_stream = 0) { - constexpr std::uint32_t TEAM_SIZE = DATASET_DESCRIPTOR_T::TEAM_SIZE; - constexpr std::uint32_t DATASET_BLOCK_DIM = DATASET_DESCRIPTOR_T::DATASET_BLOCK_DIM; - const auto block_size = 256u; - const auto num_teams_per_threadblock = block_size / TEAM_SIZE; + const auto block_size = 256u; + const auto num_teams_per_threadblock = block_size / TEAM_SIZE; const dim3 grid_size((num_pickup + num_teams_per_threadblock - 1) / num_teams_per_threadblock, num_queries); @@ -189,7 +186,7 @@ void random_pickup( raft::ceildiv(dataset_desc.dim, DATASET_BLOCK_DIM) * DATASET_BLOCK_DIM; const auto smem_size = query_smem_buffer_length * sizeof(float); - random_pickup_kernel + random_pickup_kernel <<>>(dataset_desc, queries_ptr, num_pickup, @@ -304,7 +301,9 @@ void pickup_next_parents(INDEX_T* const parent_candidates_ptr, // [num_queries, terminate_flag); } -template RAFT_KERNEL compute_distance_to_child_nodes_kernel( const typename DATASET_DESCRIPTOR_T::INDEX_T* const @@ -331,9 +330,6 @@ RAFT_KERNEL compute_distance_to_child_nodes_kernel( using INDEX_T = typename DATASET_DESCRIPTOR_T::INDEX_T; using DISTANCE_T = typename DATASET_DESCRIPTOR_T::DISTANCE_T; - constexpr std::uint32_t TEAM_SIZE = DATASET_DESCRIPTOR_T::TEAM_SIZE; - constexpr std::uint32_t DATASET_BLOCK_DIM = DATASET_DESCRIPTOR_T::DATASET_BLOCK_DIM; - const uint32_t ldb = hashmap::get_size(hash_bitlen); const auto tid = threadIdx.x + blockDim.x * blockIdx.x; const auto global_team_id = tid / TEAM_SIZE; @@ -375,7 +371,8 @@ RAFT_KERNEL compute_distance_to_child_nodes_kernel( const auto compute_distance_flag = hashmap::insert( visited_hashmap_ptr + (ldb * blockIdx.y), hash_bitlen, child_id); - const auto norm2 = dataset_desc.compute_similarity(query_buffer, child_id, compute_distance_flag); + const auto norm2 = dataset_desc.template compute_similarity( + query_buffer, child_id, compute_distance_flag); if (compute_distance_flag) { if (threadIdx.x % TEAM_SIZE == 0) { @@ -398,7 +395,9 @@ RAFT_KERNEL compute_distance_to_child_nodes_kernel( } } -template void compute_distance_to_child_nodes( const typename DATASET_DESCRIPTOR_T::INDEX_T* const @@ -424,9 +423,6 @@ void compute_distance_to_child_nodes( SAMPLE_FILTER_T sample_filter, cudaStream_t cuda_stream = 0) { - constexpr std::uint32_t TEAM_SIZE = DATASET_DESCRIPTOR_T::TEAM_SIZE; - constexpr std::uint32_t DATASET_BLOCK_DIM = DATASET_DESCRIPTOR_T::DATASET_BLOCK_DIM; - const auto block_size = 128; const dim3 grid_size( (search_width * graph_degree + (block_size / TEAM_SIZE) - 1) / (block_size / TEAM_SIZE), @@ -438,7 +434,10 @@ void compute_distance_to_child_nodes( const auto smem_size = query_smem_buffer_length * sizeof(float) + DATASET_DESCRIPTOR_T::smem_buffer_size_in_byte; - compute_distance_to_child_nodes_kernel + compute_distance_to_child_nodes_kernel <<>>(parent_node_list, parent_candidates_ptr, parent_distance_ptr, @@ -795,7 +794,7 @@ struct search : search_plan_impl { } void operator()(raft::resources const& res, - DATASET_DESCRIPTOR_T dataset_desc_, + DATASET_DESCRIPTOR_T dataset_desc, raft::device_matrix_view graph, INDEX_T* const topk_indices_ptr, // [num_queries, topk] DISTANCE_T* const topk_distances_ptr, // [num_queries, topk] @@ -806,8 +805,6 @@ struct search : search_plan_impl { uint32_t topk, SAMPLE_FILTER_T sample_filter) { - const auto dataset_desc = - set_compute_template_params(dataset_desc_); // Init hashmap cudaStream_t stream = resource::get_cuda_stream(res); const uint32_t hash_size = hashmap::get_size(hash_bitlen); @@ -825,20 +822,20 @@ struct search : search_plan_impl { } // Choose initial entry point candidates at random - random_pickup(dataset_desc, - queries_ptr, - num_queries, - result_buffer_size, - num_random_samplings, - rand_xor_mask, - dev_seed_ptr, - num_seeds, - result_indices.data(), - result_distances.data(), - result_buffer_allocation_size, - hashmap.data(), - hash_bitlen, - stream); + random_pickup(dataset_desc, + queries_ptr, + num_queries, + result_buffer_size, + num_random_samplings, + rand_xor_mask, + dev_seed_ptr, + num_seeds, + result_indices.data(), + result_distances.data(), + result_buffer_allocation_size, + hashmap.data(), + hash_bitlen, + stream); unsigned iter = 0; while (1) { @@ -890,7 +887,7 @@ struct search : search_plan_impl { } // Compute distance to child nodes that are adjacent to the parent node - compute_distance_to_child_nodes( + compute_distance_to_child_nodes( parent_node_list.data(), result_indices.data() + (1 - (iter & 0x1)) * result_buffer_size, result_distances.data() + (1 - (iter & 0x1)) * result_buffer_size, @@ -998,28 +995,24 @@ struct search, + INDEX_T_>, SAMPLE_FILTER_T> : public search_plan_impl, + INDEX_T_>, SAMPLE_FILTER_T> { using DATASET_DESCRIPTOR_T = cagra_q_dataset_descriptor_t; + INDEX_T_>; using INDEX_T = typename DATASET_DESCRIPTOR_T::INDEX_T; using DISTANCE_T = typename DATASET_DESCRIPTOR_T::DISTANCE_T; @@ -1036,7 +1029,7 @@ struct search graph, INDEX_T* const topk_indices_ptr, // [num_queries, topk] DISTANCE_T* const topk_distances_ptr, // [num_queries, topk] diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh index 44cc575b52..a836334667 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh @@ -58,39 +58,38 @@ void select_and_run( // raft::resources const& res, #endif // RAFT_EXPLICIT_INSTANTIATE_ONLY -#define instantiate_single_cta_select_and_run( \ - TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_T) \ - extern template void \ - select_and_run, \ - SAMPLE_FILTER_T>( \ - raft::neighbors::cagra::detail:: \ - standard_dataset_descriptor_t dataset, \ - raft::device_matrix_view graph, \ - INDEX_T* const topk_indices_ptr, \ - DISTANCE_T* const topk_distances_ptr, \ - const DATA_T* const queries_ptr, \ - const uint32_t num_queries, \ - const INDEX_T* dev_seed_ptr, \ - uint32_t* const num_executed_iterations, \ - uint32_t topk, \ - uint32_t num_itopk_candidates, \ - uint32_t block_size, \ - uint32_t smem_size, \ - int64_t hash_bitlen, \ - INDEX_T* hashmap_ptr, \ - size_t small_hash_bitlen, \ - size_t small_hash_reset_interval, \ - uint32_t num_random_samplings, \ - uint64_t rand_xor_mask, \ - uint32_t num_seeds, \ - size_t itopk_size, \ - size_t search_width, \ - size_t min_iterations, \ - size_t max_iterations, \ - SAMPLE_FILTER_T sample_filter, \ +#define instantiate_single_cta_select_and_run( \ + TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_T) \ + extern template void select_and_run< \ + TEAM_SIZE, \ + MAX_DATASET_DIM, \ + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, \ + SAMPLE_FILTER_T>( \ + raft::neighbors::cagra::detail::standard_dataset_descriptor_t \ + dataset, \ + raft::device_matrix_view graph, \ + INDEX_T* const topk_indices_ptr, \ + DISTANCE_T* const topk_distances_ptr, \ + const DATA_T* const queries_ptr, \ + const uint32_t num_queries, \ + const INDEX_T* dev_seed_ptr, \ + uint32_t* const num_executed_iterations, \ + uint32_t topk, \ + uint32_t num_itopk_candidates, \ + uint32_t block_size, \ + uint32_t smem_size, \ + int64_t hash_bitlen, \ + INDEX_T* hashmap_ptr, \ + size_t small_hash_bitlen, \ + size_t small_hash_reset_interval, \ + uint32_t num_random_samplings, \ + uint64_t rand_xor_mask, \ + uint32_t num_seeds, \ + size_t itopk_size, \ + size_t search_width, \ + size_t min_iterations, \ + size_t max_iterations, \ + SAMPLE_FILTER_T sample_filter, \ cudaStream_t stream); instantiate_single_cta_select_and_run( @@ -144,19 +143,15 @@ instantiate_single_cta_select_and_run( CODE_BOOK_T, \ PQ_BITS, \ PQ_CODE_BOOK_DIM, \ - 0, \ DISTANCE_T, \ - INDEX_T, \ - 0>, \ + INDEX_T>, \ SAMPLE_FILTER_T>( \ raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t dataset, \ + INDEX_T> dataset, \ raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ 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 52b7e549d8..8d8dad6708 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 @@ -456,7 +456,9 @@ __device__ inline void set_value_device(T* const ptr, const T fill, const std::u } // One query one thread block -template (dataset_desc.dim, DATASET_DESCRIPTOR_T::DATASET_BLOCK_DIM) * - DATASET_DESCRIPTOR_T::DATASET_BLOCK_DIM; + raft::ceildiv(dataset_desc.dim, DATASET_BLOCK_DIM) * DATASET_BLOCK_DIM; auto query_buffer = reinterpret_cast(smem); auto result_indices_buffer = reinterpret_cast(query_buffer + query_smem_buffer_length); auto result_distances_buffer = @@ -577,18 +576,17 @@ __launch_bounds__(1024, 1) RAFT_KERNEL search_kernel( // compute distance to randomly selecting nodes _CLK_START(); const INDEX_T* const local_seed_ptr = seed_ptr ? seed_ptr + (num_seeds * query_id) : nullptr; - device::compute_distance_to_random_nodes( - result_indices_buffer, - result_distances_buffer, - query_buffer, - dataset_desc, - result_buffer_size, - num_distilation, - rand_xor_mask, - local_seed_ptr, - num_seeds, - local_visited_hashmap_ptr, - hash_bitlen); + device::compute_distance_to_random_nodes(result_indices_buffer, + result_distances_buffer, + query_buffer, + dataset_desc, + result_buffer_size, + num_distilation, + rand_xor_mask, + local_seed_ptr, + num_seeds, + local_visited_hashmap_ptr, + hash_bitlen); __syncthreads(); _CLK_REC(clk_compute_1st_distance); @@ -818,33 +816,50 @@ __launch_bounds__(1024, 1) RAFT_KERNEL search_kernel( #endif } -template +template struct search_kernel_config { - using kernel_t = decltype(&search_kernel<64, 64, 0, DATASET_DESCRIPTOR_T, SAMPLE_FILTER_T>); + using kernel_t = decltype(&search_kernel); template static auto choose_search_kernel(unsigned itopk_size) -> kernel_t { if (itopk_size <= 64) { - return search_kernel<64, + return search_kernel; } else if (itopk_size <= 128) { - return search_kernel<128, + return search_kernel; } else if (itopk_size <= 256) { - return search_kernel<256, + return search_kernel; } else if (itopk_size <= 512) { - return search_kernel<512, + return search_kernel; + return search_kernel; } else if (itopk_size <= 512) { - return search_kernel<512, max_candidates, 0, DATASET_DESCRIPTOR_T, SAMPLE_FILTER_T>; + return search_kernel; } } THROW("No kernel for parametels itopk_size %u, num_itopk_candidates %u", @@ -910,15 +937,12 @@ void select_and_run( SAMPLE_FILTER_T sample_filter, cudaStream_t stream) { - const auto dataset_desc_ = - set_compute_template_params(dataset_desc); - using dataset_desc_t = typename std::remove_const::type; auto kernel = - search_kernel_config::choose_itopk_and_mx_candidates( - itopk_size, num_itopk_candidates, block_size); + search_kernel_config:: + choose_itopk_and_mx_candidates(itopk_size, num_itopk_candidates, block_size); RAFT_CUDA_TRY(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, - smem_size + dataset_desc_t::smem_buffer_size_in_byte)); + smem_size + DATASET_DESCRIPTOR_T::smem_buffer_size_in_byte)); dim3 thread_dims(block_size, 1, 1); dim3 block_dims(1, num_queries, 1); RAFT_LOG_DEBUG( @@ -926,7 +950,7 @@ void select_and_run( kernel<<>>(topk_indices_ptr, topk_distances_ptr, topk, - dataset_desc_, + dataset_desc, queries_ptr, graph.data_handle(), graph.extent(1), diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_00_generate.py b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_00_generate.py index bd5f6b278f..e827c06be5 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_00_generate.py +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_00_generate.py @@ -77,7 +77,7 @@ with open(path, "w") as f: f.write(header) f.write( - f"instantiate_kernel_selection(\n {team}, {mxdim}, raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t<{data_t} COMMA {code_book_t} COMMA {pq_bit} COMMA {subspace_dim} COMMA 0 COMMA {distance_t} COMMA {idx_t} COMMA 0>, raft::neighbors::filtering::none_cagra_sample_filter);\n" + f"instantiate_kernel_selection(\n {team}, {mxdim}, raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t<{data_t} COMMA {code_book_t} COMMA {pq_bit} COMMA {subspace_dim} COMMA {distance_t} COMMA {idx_t}>, raft::neighbors::filtering::none_cagra_sample_filter);\n" ) f.write(trailer) # For pasting into CMakeLists.txt diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim1024_t32_8pq_2subd_half.cu index 5b174ddaee..0bd386144c 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim1024_t32_8pq_4subd_half.cu index 11503927fb..cd891b8e97 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim128_t8_8pq_2subd_half.cu index 5ad1f942e6..66e8357498 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim128_t8_8pq_4subd_half.cu index fdab2893b2..eb84983f9e 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim256_t16_8pq_2subd_half.cu index 01e7dcba4c..c66f8a0ae3 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim256_t16_8pq_4subd_half.cu index 61afc4731b..2a1783944c 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim512_t32_8pq_2subd_half.cu index 8524a653c6..9fa74f1134 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim512_t32_8pq_4subd_half.cu index 28f75dcbff..8fc91b5a10 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint32_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim1024_t32_8pq_2subd_half.cu index 9c5b161911..4e68c00525 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim1024_t32_8pq_4subd_half.cu index ab13f43868..5fe526ae47 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim128_t8_8pq_2subd_half.cu index c03a57372d..64c89a880a 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim128_t8_8pq_4subd_half.cu index f4f7148580..c3e2427f57 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim256_t16_8pq_2subd_half.cu index 15c65830e9..0a8826df1c 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim256_t16_8pq_4subd_half.cu index ba00a5cf7e..8019bec3e3 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim512_t32_8pq_2subd_half.cu index ad101deec9..1a2a364037 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim512_t32_8pq_4subd_half.cu index 185fadae2e..2f661538e6 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_float_uint64_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim1024_t32_8pq_2subd_half.cu index d67903cea2..aec486769f 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim1024_t32_8pq_4subd_half.cu index 2dc54bcb59..03f27085d8 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim128_t8_8pq_2subd_half.cu index a9761e947a..119d1f2921 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim128_t8_8pq_4subd_half.cu index fd23a5a5b9..666c676e87 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim256_t16_8pq_2subd_half.cu index b350696bf0..e53b456a54 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim256_t16_8pq_4subd_half.cu index c6ecb67efe..2aee739141 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim512_t32_8pq_2subd_half.cu index ac4ffa356a..daa442b514 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim512_t32_8pq_4subd_half.cu index e7efd7e305..a19346d19b 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint32_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim1024_t32_8pq_2subd_half.cu index f78f7ae508..1c1d5381c9 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim1024_t32_8pq_4subd_half.cu index 38cf9c85a5..b7402a3c38 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim128_t8_8pq_2subd_half.cu index c280585101..f493b83bee 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim128_t8_8pq_4subd_half.cu index bddf03f42a..8efcbe0650 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim256_t16_8pq_2subd_half.cu index ddb1304325..cb770f44ba 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim256_t16_8pq_4subd_half.cu index 61e4305ffa..0fd8ab809c 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim512_t32_8pq_2subd_half.cu index 0967bbb039..50cf198883 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim512_t32_8pq_4subd_half.cu index 1b4db0b1f6..1548ed831e 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_half_uint64_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim1024_t32_8pq_2subd_half.cu index b0d33811f0..c60ea7c87d 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim1024_t32_8pq_4subd_half.cu index 84ee6dc773..4a68e1e43c 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim128_t8_8pq_2subd_half.cu index 887c152c0a..df9fabd6a5 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim128_t8_8pq_4subd_half.cu index 2af58c6211..77075b0a44 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim256_t16_8pq_2subd_half.cu index 753a3d8f49..374af8b56b 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim256_t16_8pq_4subd_half.cu index b2528e4446..ddb80458fd 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim512_t32_8pq_2subd_half.cu index e3b691f10d..14e5c5d3dc 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim512_t32_8pq_4subd_half.cu index aa070a2763..3c1776760a 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_int8_uint32_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim1024_t32_8pq_2subd_half.cu index c28b846130..e5a0a8882c 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim1024_t32_8pq_4subd_half.cu index 93b3036aee..cee80390e8 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim128_t8_8pq_2subd_half.cu index 3d0553f043..88678bf4ff 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim128_t8_8pq_4subd_half.cu index 35a09411ce..baa7ee358a 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim256_t16_8pq_2subd_half.cu index 96aec1b2a2..5c44f052f2 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim256_t16_8pq_4subd_half.cu index 0f19ae504f..127a065fb5 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim512_t32_8pq_2subd_half.cu index 9083453c7c..fcf6985f97 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim512_t32_8pq_4subd_half.cu index 0d8ec8dde6..f361e771b5 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_multi_cta_uint8_uint32_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_00_generate.py b/cpp/src/neighbors/detail/cagra/q_search_single_cta_00_generate.py index 2cbc423b9e..418d528a82 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_00_generate.py +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_00_generate.py @@ -81,7 +81,7 @@ with open(path, "w") as f: f.write(header) f.write( - f"instantiate_kernel_selection(\n {team}, {mxdim}, raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t<{data_t} COMMA {code_book_t} COMMA {pq_bit} COMMA {subspace_dim} COMMA 0 COMMA {distance_t} COMMA {idx_t} COMMA 0>, raft::neighbors::filtering::none_cagra_sample_filter);\n" + f"instantiate_kernel_selection(\n {team}, {mxdim}, raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t<{data_t} COMMA {code_book_t} COMMA {pq_bit} COMMA {subspace_dim} COMMA {distance_t} COMMA {idx_t}>, raft::neighbors::filtering::none_cagra_sample_filter);\n" ) f.write(trailer) diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim1024_t32_8pq_2subd_half.cu index 9d7d64e544..d61ad0ce15 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim1024_t32_8pq_4subd_half.cu index 5a4ccf7682..410d2377ec 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim128_t8_8pq_2subd_half.cu index 4bcdc42008..60cd58bab9 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim128_t8_8pq_4subd_half.cu index 9e64faa3c8..dfe5e6f14e 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim256_t16_8pq_2subd_half.cu index c942fb075a..9a5d862276 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim256_t16_8pq_4subd_half.cu index dd1a9ac9e1..d92ab50a58 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim512_t32_8pq_2subd_half.cu index dd50c4b063..aac197d590 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim512_t32_8pq_4subd_half.cu index b69ffaa988..f38a10e6d0 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint32_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim1024_t32_8pq_2subd_half.cu index e5cbe01c51..5523e63038 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim1024_t32_8pq_4subd_half.cu index c8b5cfe3f6..b06ef3d4fd 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim128_t8_8pq_2subd_half.cu index 69551690e2..1fddee0e06 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim128_t8_8pq_4subd_half.cu index ba0095892c..2aee442186 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim256_t16_8pq_2subd_half.cu index 960d2e1f28..7a15e85280 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim256_t16_8pq_4subd_half.cu index b589c1af8c..efba46c248 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim512_t32_8pq_2subd_half.cu index 11f21a24a5..990582f18b 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim512_t32_8pq_4subd_half.cu index 0f06167fc2..a55907c66f 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_float_uint64_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - float COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + float COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim1024_t32_8pq_2subd_half.cu index 308684a452..55fd749720 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim1024_t32_8pq_4subd_half.cu index 5a88bae588..4b4063652a 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim128_t8_8pq_2subd_half.cu index a3b58d44cf..bae83dc0fa 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim128_t8_8pq_4subd_half.cu index bb19e6bd54..99492db344 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim256_t16_8pq_2subd_half.cu index 78a2ef1cd1..797142e317 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim256_t16_8pq_4subd_half.cu index 54b277fdb5..9a36c35ae0 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim512_t32_8pq_2subd_half.cu index 7df4430968..e0a01e84cc 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim512_t32_8pq_4subd_half.cu index 5bb3ede62b..14de1b8941 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint32_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim1024_t32_8pq_2subd_half.cu index 11f81782ae..b1d50fb445 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim1024_t32_8pq_4subd_half.cu index 03ad10016f..c189a91764 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim128_t8_8pq_2subd_half.cu index 1b1b1551ab..8693ee3716 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim128_t8_8pq_4subd_half.cu index d3ca5fa44d..216ffd1ec5 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim256_t16_8pq_2subd_half.cu index 357f4e9ca4..36985d218b 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim256_t16_8pq_4subd_half.cu index f50a789bf9..8d55fe2b09 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim512_t32_8pq_2subd_half.cu index 6a7e80dd92..2fdb1cbc20 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim512_t32_8pq_4subd_half.cu index 2a1817d9f8..6dc3dc2ca8 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_half_uint64_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - half COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint64_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + half COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint64_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim1024_t32_8pq_2subd_half.cu index 4f1f51b147..21f8633033 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim1024_t32_8pq_4subd_half.cu index eea96f3d15..1a3867e06f 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim128_t8_8pq_2subd_half.cu index 14ef0c7a33..9cbb16188a 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim128_t8_8pq_4subd_half.cu index a88d71554b..305a1754bc 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim256_t16_8pq_2subd_half.cu index 96dde7ab99..900e1b69d9 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim256_t16_8pq_4subd_half.cu index 0dcb6fbaf9..a0bb2259f0 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim512_t32_8pq_2subd_half.cu index 3facd5f91e..09d36a39a0 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim512_t32_8pq_4subd_half.cu index fc4d49e207..dc9cbb2b56 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_int8_uint32_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - int8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + int8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim1024_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim1024_t32_8pq_2subd_half.cu index ecd7c8c918..c5508a38e2 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim1024_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim1024_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim1024_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim1024_t32_8pq_4subd_half.cu index 9f94847864..7024425155 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim1024_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim1024_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 1024, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 1024, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim128_t8_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim128_t8_8pq_2subd_half.cu index 3474092311..68687bc9cf 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim128_t8_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim128_t8_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim128_t8_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim128_t8_8pq_4subd_half.cu index 49addc2bc4..60efc55a30 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim128_t8_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim128_t8_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 8, - 128, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(8, + 128, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim256_t16_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim256_t16_8pq_2subd_half.cu index 25f2a2ba25..b2dfaac5fe 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim256_t16_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim256_t16_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim256_t16_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim256_t16_8pq_4subd_half.cu index 1c1ae0e845..891e9ef7cc 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim256_t16_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim256_t16_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 16, - 256, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(16, + 256, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim512_t32_8pq_2subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim512_t32_8pq_2subd_half.cu index 8ec11f13fa..91e617204c 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim512_t32_8pq_2subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim512_t32_8pq_2subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 2 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 2 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim512_t32_8pq_4subd_half.cu b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim512_t32_8pq_4subd_half.cu index dc6b1f8ce2..a01d497676 100644 --- a/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim512_t32_8pq_4subd_half.cu +++ b/cpp/src/neighbors/detail/cagra/q_search_single_cta_uint8_uint32_dim512_t32_8pq_4subd_half.cu @@ -28,11 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection( - 32, - 512, - raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< - uint8_t COMMA half COMMA 8 COMMA 4 COMMA 0 COMMA float COMMA uint32_t COMMA 0>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection(32, + 512, + raft::neighbors::cagra::detail::cagra_q_dataset_descriptor_t< + uint8_t COMMA half COMMA 8 COMMA 4 COMMA float COMMA uint32_t>, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py b/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py index 472886729b..6f023c39f1 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py @@ -71,7 +71,7 @@ with open(path, "w") as f: f.write(header) f.write( - f"instantiate_kernel_selection(\n {team}, {mxdim}, raft::neighbors::cagra::detail::standard_dataset_descriptor_t<{data_t} COMMA {idx_t} COMMA 0 COMMA 0 COMMA {distance_t}>, raft::neighbors::filtering::none_cagra_sample_filter);\n" + f"instantiate_kernel_selection(\n {team}, {mxdim}, raft::neighbors::cagra::detail::standard_dataset_descriptor_t<{data_t} COMMA {idx_t} COMMA {distance_t}>, raft::neighbors::filtering::none_cagra_sample_filter);\n" ) f.write(trailer) # For pasting into CMakeLists.txt diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu index 777dab40c6..0e28d7a876 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu index ca69079ebd..5e5e80a5de 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu index 956bd57f78..9039496968 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu index 5016ca4e29..fe1c7e77e5 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu index f71447a84e..7ef36baf7d 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu index 3da7a01d2b..da51c16314 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu index a91883c1a3..99a4f7feb7 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu index 979bf614dc..50cdc97dd7 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim1024_t32.cu index 3b4e0dedd0..b2d9cdb600 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim128_t8.cu index 490fabe5cb..d756b295b7 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim256_t16.cu index 7095a2bcbd..b1e998762c 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim512_t32.cu index c4d7e89656..e712de6390 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint32_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim1024_t32.cu index 8976924220..282de4a851 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim128_t8.cu index 22e120b374..71ef968575 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim256_t16.cu index a0e52734f7..7c88406d71 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim512_t32.cu index 1637af3973..360635dddb 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_half_uint64_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu index bfd55f004f..3f129bd7cf 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - int8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu index a4c338e6db..053b73275e 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - int8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu index 597e862bdd..a1bb20369a 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - int8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu index bd2e584e75..dbbc8bdd21 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - int8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu index a936c424c0..125499e319 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - uint8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu index 8a707ccaf2..f2117c4f80 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - uint8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu index 4ebe2e3b77..8e5ba0f98f 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - uint8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu index 459219c8d1..bea7d25392 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::multi_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - uint8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::multi_cta_search 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 6fedae0fa3..0e809e4dc3 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 @@ -74,7 +74,7 @@ with open(path, "w") as f: f.write(header) f.write( - f"instantiate_kernel_selection(\n {team}, {mxdim}, raft::neighbors::cagra::detail::standard_dataset_descriptor_t<{data_t} COMMA {idx_t} COMMA 0 COMMA 0 COMMA {distance_t}>, raft::neighbors::filtering::none_cagra_sample_filter);\n" + f"instantiate_kernel_selection(\n {team}, {mxdim}, raft::neighbors::cagra::detail::standard_dataset_descriptor_t<{data_t} COMMA {idx_t} COMMA {distance_t}>, raft::neighbors::filtering::none_cagra_sample_filter);\n" ) f.write(trailer) diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu index 81d242d37f..8a9fc408ee 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu index f26ca71b8e..c6f7c90c69 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu index 048d740dd3..2766286673 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu index 99ea62960c..98ee189766 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu index 27f06caa8c..c3ea39a729 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu index 573a5688a8..a53457656c 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu index cbe53429e9..52318efb85 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu index e297abde56..6451fdc7f3 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - float COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim1024_t32.cu index bde6407312..e927fd0878 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim128_t8.cu index b256de71e5..3f3d22ee08 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim256_t16.cu index 019d64b7cb..a84e5b8bd7 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim512_t32.cu index 2da2263380..af4248865b 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint32_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim1024_t32.cu index 20dd3857b1..16bd0cb647 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim128_t8.cu index a4f2d10f23..afc59c8a59 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim256_t16.cu index 7460bd54af..147d31cf85 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim512_t32.cu index a1983aa9ad..5624a71c3c 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_half_uint64_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - half COMMA uint64_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu index cd5a061c2e..761fb705ba 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - int8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu index d5845f1cc7..84b76cba53 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - int8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu index e370b547dd..598fff9cdf 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - int8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu index c5b2841027..e7a1a9d9c6 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - int8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu index a16d3cac38..d40b9285fc 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 1024, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - uint8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 1024, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu index ceaa625f05..073bb350da 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(8, - 128, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - uint8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 8, + 128, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu index cb8f2c16cf..29b0224b4d 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(16, - 256, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - uint8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 16, + 256, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu index 72307f5c39..d9601de2ad 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu @@ -28,10 +28,10 @@ #include namespace raft::neighbors::cagra::detail::single_cta_search { -instantiate_kernel_selection(32, - 512, - raft::neighbors::cagra::detail::standard_dataset_descriptor_t< - uint8_t COMMA uint32_t COMMA 0 COMMA 0 COMMA float>, - raft::neighbors::filtering::none_cagra_sample_filter); +instantiate_kernel_selection( + 32, + 512, + raft::neighbors::cagra::detail::standard_dataset_descriptor_t, + raft::neighbors::filtering::none_cagra_sample_filter); } // namespace raft::neighbors::cagra::detail::single_cta_search