Hiding symbols from shared object libraries

[cjnolet]

Partially addresses #1722.

[ahendriksen]

We should check the binary size of libraft after this change. We have a lot of duplicate kernels. When they are static they will be duplicated in each TU and linked into libraft.so. The impact might also be smaller than expected because the linker seems not to deduplicate kernels with weak linkage anyway(?)

[jrhemstad]

Strictly speaking, this is only half of the solution. The other half is to annotate all non-kernel template functions that aren't part of the API of any pre-built DLLs with __attribute__((visibility("hidden"))).

[cjnolet]

Strictly speaking, this is only half of the solution

@jrhemstad yep, that's right. The first half of the solution was a simple search and replace. I realized the second half is much more involved and is going to require scraping through the entire codebase and adding a macro before every host template function that launches a kernel.

[jrhemstad]

After discussing internally with @ahendriksen, allow me to contradict myself (again) and suggest that the macro you want is actually:

#if defined(__CUDACC__)
#   if !defined(__CUDACC_RDC__)
#       define _RAFT_KERNEL \
            static __global__
#   else
#       define _RAFT_KERNEL \
            __attribute__ ((visibility ("hidden"))) __global__
#   endif
#else
#   define _RAFT_KERNEL
#endif

This is because when compiling with -rdc=true then the device linker will deduplicate all the identical kernels in the same shared object library and you likely want to preserve that behavior by not marking the kernels as static. You still want to hide the symbols from the resulting .so though, so instead we annotate it with the hidden attribute.

For more info, see: NVIDIA/cccl#166 (comment)

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[cjnolet]

/ok to test

[cjnolet]

/ok to test

[jrhemstad]

fyi, here's the solution we ended up using in Thrust/CUB: NVIDIA/cccl#443