diff --git a/cpp/cmake/thirdparty/get_cuco.cmake b/cpp/cmake/thirdparty/get_cuco.cmake index 9b60839b81..299146420a 100644 --- a/cpp/cmake/thirdparty/get_cuco.cmake +++ b/cpp/cmake/thirdparty/get_cuco.cmake @@ -22,7 +22,7 @@ function(find_and_configure_cuco VERSION) INSTALL_EXPORT_SET raft-distance-lib-exports CPM_ARGS GIT_REPOSITORY https://github.com/NVIDIA/cuCollections.git - GIT_TAG 6ec8b6dcdeceea07ab4456d32461a05c18864411 + GIT_TAG 55029034c3f82bca36148c9be29941b37492394d OPTIONS "BUILD_TESTS OFF" "BUILD_BENCHMARKS OFF" "BUILD_EXAMPLES OFF" diff --git a/cpp/include/raft/sparse/distance/detail/coo_spmv_strategies/hash_strategy.cuh b/cpp/include/raft/sparse/distance/detail/coo_spmv_strategies/hash_strategy.cuh index 4f8637b425..0893a5bb27 100644 --- a/cpp/include/raft/sparse/distance/detail/coo_spmv_strategies/hash_strategy.cuh +++ b/cpp/include/raft/sparse/distance/detail/coo_spmv_strategies/hash_strategy.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021, NVIDIA CORPORATION. + * Copyright (c) 2021-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -226,8 +226,11 @@ class hash_strategy : public coo_spmv_strategy { __device__ inline insert_type init_insert(smem_type cache, const value_idx& cache_size) { - return insert_type::make_from_uninitialized_slots( - cooperative_groups::this_thread_block(), cache, cache_size, -1, 0); + return insert_type::make_from_uninitialized_slots(cooperative_groups::this_thread_block(), + cache, + cache_size, + cuco::sentinel::empty_key{value_idx{-1}}, + cuco::sentinel::empty_value{value_t{0}}); } __device__ inline void insert(insert_type cache, const value_idx& key, const value_t& value) @@ -237,7 +240,10 @@ class hash_strategy : public coo_spmv_strategy { __device__ inline find_type init_find(smem_type cache, const value_idx& cache_size) { - return find_type(cache, cache_size, -1, 0); + return find_type(cache, + cache_size, + cuco::sentinel::empty_key{value_idx{-1}}, + cuco::sentinel::empty_value{value_t{0}}); } __device__ inline value_t find(find_type cache, const value_idx& key)