[BUG] linker complains about duplicate symbol get_cache_idx

[alexanderguzhva]

Describe the bug
Many errors of the following kind:

ld.lld: error: duplicate symbol: raft::cache::(anonymous namespace)::get_cache_idx(int*, int, int*, int, int, int*, int*, bool*, int)

Expected behavior
No errors

Environment details (please complete the following information):
upstream hash is "a98295b516ef58bc855177077860bab2a2a76d77" (Apr 12 ?)
CUDA 11.4.2

Additional context
The error can be fixed in a hacky way by turning

__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)

into

template<typename T>
__global__ void get_cache_idx(T* keys,
                              T n,
                              T* cached_keys,
                              T n_cache_sets,
                              T associativity,
                              T* cache_time,
                              T* cache_idx,
                              bool* is_cached,
                              T time)

I haven't tried moving it to a source file.

[cjnolet]

@alexanderguzhva this is indeed strange because the anonymous namespace should be protecting us from this compile error. Does adding __inline__ to the definition help overcome this error for you? If so, we can remove the anonymous namespace and just use that.

[alexanderguzhva]

@cjnolet A kernel cannot be inlined, can it be?

[cjnolet]

@alexanderguzhva oddly, I'm not able to reproduce the error when building FAISS. Is there a specific branch I can check out? What version of gcc are you using? I'm on CUDA 11.8 but I can also try 11.4.

I verified I'm able to build raft successfully by removing the offending kernel from the anonymous namespace and adding __forceinline__ to the definition. If I create a RAFT PR for this, can you give it a try and see if the error goes away?

[ahendriksen]

A kernel cannot be inlined, can it be?

A kernel can be inlined, i.e, this works: inline __global__ void kernel() {}. It doesn't result in the kernel being literally "inlined" into a function, but it does change the linking behavior (symbol is defined as weak and may be omitted from symbol table by compiler).

In general, I think we should prefer inline over an anonymous namespace. With an anonymous namespace, a copy of the kernel is included in each translation unit and also in the final linked object. With inline, duplicate definitions should be deduplicated when linked.

@cjnolet : note that the linker in the bug report, ld.lld, is LLVM's linker. It might have slightly different behavior than GCC's linker that we use by default. In any case, clang also behaves slightly differently than gcc wrt private symbols. See for example the section "symbols" in this recent guide to linking: https://fabiensanglard.net/dc/cc.php

[ahendriksen]

@alexanderguzhva, I pushed PR #1492. Could you test if that PR fixes the issue for you? I was unfortunately not able to reproduce your issue, but I think that inlining the kernel should fix any linker errors.

[alexanderguzhva]

@ahendriksen #1492 seems to fix the problem. Thanks!