From 576b22fccdd19f430c7f02bd29b4bc69dd2a0bf7 Mon Sep 17 00:00:00 2001
From: Allard Hendriksen <ahendriksen@nvidia.com>
Date: Fri, 5 May 2023 23:19:57 +0200
Subject: [PATCH] Inline get_cache_idx (#1492)

Related to issue #1490.

The `get_cache_idx` kernel is currently defined in an anonymous namespace. In #1490, it was reported that this could lead to linker errors in CUDA 11.4 in combination with (presumably?) clang.

I tried to reproduce the linker error in commit 01cf4099b206b6b37dcdc84e852507306b018af5 but could not, as described in the commit message.

This PR attempts to fix the linker errors by marking `get_cache_idx` as inline and by removing the anonymous namespace.

Authors:
  - Allard Hendriksen (https://github.com/ahendriksen)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: https://github.com/rapidsai/raft/pull/1492
---
 cpp/include/raft/util/cache_util.cuh | 22 +++++++++-------------
 1 file changed, 9 insertions(+), 13 deletions(-)

diff --git a/cpp/include/raft/util/cache_util.cuh b/cpp/include/raft/util/cache_util.cuh
index b20982473f..bbd84d8bf2 100644
--- a/cpp/include/raft/util/cache_util.cuh
+++ b/cpp/include/raft/util/cache_util.cuh
@@ -303,9 +303,6 @@ __global__ void assign_cache_idx(const int* keys,
   }
 }
 
-/* Unnamed namespace is used to avoid multiple definition error for the
-  following non-template function */
-namespace {
 /**
  * @brief Get the cache indices for keys stored in the cache.
  *
@@ -331,15 +328,15 @@ namespace {
  * @param [out] is_cached whether the element is cached size[n]
  * @param [in] time iteration counter (used for time stamping)
  */
-__global__ void get_cache_idx(int* keys,
-                              int n,
-                              int* cached_keys,
-                              int n_cache_sets,
-                              int associativity,
-                              int* cache_time,
-                              int* cache_idx,
-                              bool* is_cached,
-                              int time)
+__global__ inline void get_cache_idx(int* keys,
+                                     int n,
+                                     int* cached_keys,
+                                     int n_cache_sets,
+                                     int associativity,
+                                     int* cache_time,
+                                     int* cache_idx,
+                                     bool* is_cached,
+                                     int time)
 {
   int tid = threadIdx.x + blockIdx.x * blockDim.x;
   if (tid < n) {
@@ -363,6 +360,5 @@ __global__ void get_cache_idx(int* keys,
     }
   }
 }
-};  // end unnamed namespace
 };  // namespace cache
 };  // namespace raft