From 3241addf2f80cac2913a6d5b2c75d139da61da0c Mon Sep 17 00:00:00 2001 From: Ben Frederickson Date: Thu, 5 Oct 2023 19:57:58 -0700 Subject: [PATCH 1/2] Build fix for CUDA 12.2 (#1870) Building on cuda 12.2 shows errors like ``` /code/raft/cpp/include/raft/spatial/knn/detail/ball_cover/registers-inl.cuh(177): error #20054-D: dynamic initialization is not supported for a function-scope static __shared__ variable within a __device__/__global__ function KeyValuePair shared_memV[kNumWarps * warp_q]; ``` Fix by using default constructors for structures in shared memory, even trivial constructors will cause this issue Authors: - Ben Frederickson (https://github.com/benfred) - Corey J. Nolet (https://github.com/cjnolet) Approvers: - Corey J. Nolet (https://github.com/cjnolet) URL: https://github.com/rapidsai/raft/pull/1870 --- cpp/cmake/patches/ggnn.patch | 23 +++++++++++++++++++++++ cpp/include/raft/core/kvp.hpp | 2 +- 2 files changed, 24 insertions(+), 1 deletion(-) diff --git a/cpp/cmake/patches/ggnn.patch b/cpp/cmake/patches/ggnn.patch index 95e1aaff4b..21df3bd738 100644 --- a/cpp/cmake/patches/ggnn.patch +++ b/cpp/cmake/patches/ggnn.patch @@ -1,3 +1,26 @@ +diff --git a/include/ggnn/cache/cuda_simple_knn_sym_cache.cuh b/include/ggnn/cache/cuda_simple_knn_sym_cache.cuh +index 890420e..d792903 100644 +--- a/include/ggnn/cache/cuda_simple_knn_sym_cache.cuh ++++ b/include/ggnn/cache/cuda_simple_knn_sym_cache.cuh +@@ -62,7 +62,7 @@ struct SimpleKNNSymCache { + const ValueT dist_half) + : dist_query(dist_query), dist_half(dist_half) {} + +- __device__ __forceinline__ DistQueryAndHalf() {} ++ DistQueryAndHalf() = default; + }; + + struct DistanceAndNorm { +@@ -98,8 +98,7 @@ struct SimpleKNNSymCache { + KeyT cache; + DistQueryAndHalf dist; + bool flag; +- +- __device__ __forceinline__ SyncTempStorage() {} ++ SyncTempStorage() = default; + }; + + public: diff --git a/include/ggnn/cuda_knn_ggnn_gpu_instance.cuh b/include/ggnn/cuda_knn_ggnn_gpu_instance.cuh index 8cbaf0d..6eb72ac 100644 --- a/include/ggnn/cuda_knn_ggnn_gpu_instance.cuh diff --git a/cpp/include/raft/core/kvp.hpp b/cpp/include/raft/core/kvp.hpp index 192d160d45..610945a76c 100644 --- a/cpp/include/raft/core/kvp.hpp +++ b/cpp/include/raft/core/kvp.hpp @@ -36,7 +36,7 @@ struct KeyValuePair { Value value; ///< Item value /// Constructor - RAFT_INLINE_FUNCTION KeyValuePair() {} + KeyValuePair() = default; #ifdef _RAFT_HAS_CUDA /// Conversion Constructor to allow integration w/ cub From d9fde97fd147f1452d472cebc75433d381e78d6c Mon Sep 17 00:00:00 2001 From: tsuki <12711693+enp1s0@users.noreply.github.com> Date: Fri, 6 Oct 2023 10:59:19 +0800 Subject: [PATCH 2/2] [BUG] Fix a bug in NN descent (#1869) I got errors when compiling a program using raft NN-descent. This PR fixes the bug. ## Error e.g. ``` /home/.../include/raft/neighbors/detail/nn_descent.cuh(1158): error: invalid narrowing conversion from "unsigned long" to "int" h_rev_graph_old_{static_cast(nrow_ * NUM_SAMPLES)}, ``` - nvcc ``` Built on Tue_Aug_15_22:02:13_PDT_2023 Cuda compilation tools, release 12.2, V12.2.140 Build cuda_12.2.r12.2/compiler.33191640_0 ``` - gcc ``` gcc (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0 ``` - thrust : 2.2 (This is the cause of this error [[detail](https://github.com/rapidsai/raft/pull/1869#issuecomment-1746637944)]) # Change Use `Class(...)` instead of `Class{...}`. # Cause The NN-descent code calls constructors of `thrust::host_vector` as shown below: ```cpp graph_host_buffer_{static_cast(nrow_ * DEGREE_ON_DEVICE)}, ``` However, this constructor is regarded as a list initialization. This is the same as the following code outputting 1 instead of 2. ```cpp #include #include int main() { std::vector list{2}; std::cout << list.size() << std::endl; } ``` [detail](https://en.cppreference.com/w/cpp/language/list_initialization) Authors: - tsuki (https://github.com/enp1s0) - Corey J. Nolet (https://github.com/cjnolet) Approvers: - Corey J. Nolet (https://github.com/cjnolet) - Divye Gala (https://github.com/divyegala) URL: https://github.com/rapidsai/raft/pull/1869 --- cpp/include/raft/neighbors/detail/nn_descent.cuh | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/cpp/include/raft/neighbors/detail/nn_descent.cuh b/cpp/include/raft/neighbors/detail/nn_descent.cuh index 009ffd4684..1fb568a934 100644 --- a/cpp/include/raft/neighbors/detail/nn_descent.cuh +++ b/cpp/include/raft/neighbors/detail/nn_descent.cuh @@ -980,9 +980,9 @@ GnndGraph::GnndGraph(const size_t nrow, num_samples(num_samples), bloom_filter(nrow, internal_node_degree / segment_size, 3), h_dists{raft::make_host_matrix(nrow, node_degree)}, - h_graph_new{nrow * num_samples}, - h_list_sizes_new{nrow}, - h_graph_old{nrow * num_samples}, + h_graph_new(nrow * num_samples), + h_list_sizes_new(nrow), + h_graph_old(nrow * num_samples), h_list_sizes_old{nrow} { // node_degree must be a multiple of segment_size; @@ -1150,12 +1150,12 @@ GNND::GNND(raft::resources const& res, const BuildConfig& build raft::make_device_matrix(res, nrow_, DEGREE_ON_DEVICE)}, dists_buffer_{ raft::make_device_matrix(res, nrow_, DEGREE_ON_DEVICE)}, - graph_host_buffer_{static_cast(nrow_ * DEGREE_ON_DEVICE)}, - dists_host_buffer_{static_cast(nrow_ * DEGREE_ON_DEVICE)}, + graph_host_buffer_(nrow_ * DEGREE_ON_DEVICE), + dists_host_buffer_(nrow_ * DEGREE_ON_DEVICE), d_locks_{raft::make_device_vector(res, nrow_)}, - h_rev_graph_new_{static_cast(nrow_ * NUM_SAMPLES)}, - h_graph_old_{static_cast(nrow_ * NUM_SAMPLES)}, - h_rev_graph_old_{static_cast(nrow_ * NUM_SAMPLES)}, + h_rev_graph_new_(nrow_ * NUM_SAMPLES), + h_graph_old_(nrow_ * NUM_SAMPLES), + h_rev_graph_old_(nrow_ * NUM_SAMPLES), d_list_sizes_new_{raft::make_device_vector(res, nrow_)}, d_list_sizes_old_{raft::make_device_vector(res, nrow_)} {