Skip to content

Commit

Permalink
Inline get_cache_idx (#1492)
Browse files Browse the repository at this point in the history
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 01cf409 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: #1492
  • Loading branch information
ahendriksen authored May 5, 2023
1 parent 3978b32 commit 576b22f
Showing 1 changed file with 9 additions and 13 deletions.
22 changes: 9 additions & 13 deletions cpp/include/raft/util/cache_util.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -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.
*
Expand All @@ -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) {
Expand All @@ -363,6 +360,5 @@ __global__ void get_cache_idx(int* keys,
}
}
}
}; // end unnamed namespace
}; // namespace cache
}; // namespace raft

0 comments on commit 576b22f

Please sign in to comment.