Build fix for CUDA 12.2 (rapidsai#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: rapidsai#1870

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

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