From 7c3738e052863ce25195d2a358cb3070e6d7a756 Mon Sep 17 00:00:00 2001
From: Ben Frederickson <ben@benfrederickson.com>
Date: Thu, 5 Oct 2023 19:57:58 -0700
Subject: [PATCH] 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<value_t, value_idx> 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