From be8cf268d1e6781eafb4b9c6e082aa99ae6ff082 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 16 Jun 2022 09:46:35 -0700 Subject: [PATCH 1/3] Update cuco hash. --- cpp/cmake/thirdparty/get_cuco.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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" From 903a47530cf6f879fcb6cec967f0f916e410fd4e Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 16 Jun 2022 11:22:09 -0700 Subject: [PATCH 2/3] Update sentinel types. --- .../distance/detail/coo_spmv_strategies/hash_strategy.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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..8d5c229316 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 @@ -227,7 +227,7 @@ 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); + 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 +237,7 @@ 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) From e70c061d24d090262c1f92471c21918e515e3546 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 16 Jun 2022 12:04:52 -0700 Subject: [PATCH 3/3] Fix formatting. --- .../detail/coo_spmv_strategies/hash_strategy.cuh | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) 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 8d5c229316..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, cuco::sentinel::empty_key{value_idx{-1}}, cuco::sentinel::empty_value{value_t{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, cuco::sentinel::empty_key{value_idx{-1}}, cuco::sentinel::empty_value{value_t{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)