[FEA] RAFT should ensure all its symbols are hidden from shared object libraries

[jrhemstad]

Is your feature request related to a problem? Please describe.

As a user of RAFT, I would like to build a shared object library, libA.so, that internally uses RAFT function templates, including __global__ function templates.

Today, RAFT does nothing to hide the visibility of its __global__ function templates or any other host template functions, and by default these symbols have weak visibility. In short, this means if I link two dynamic libraries A.so and B.so into my application that both contain identical instantiations of a RAFT template, then the linker will discard one of the two instantiations and use only one of them. This can lead to disastrous and insidious issues like spurious silent failures.

This issue is true of any header-only, C++ template library, but is particularly bad for CUDA C++ libraries that ship __global__ function templates. Consider this trivial example of one of many ways things can go wrong

The following code has two TUs:

• volta.cu​ compiled with sm_70​
• pascal.cu​ compiled with sm_60​

Each TU has a single function ( volta()​ or pascal()​ respectively) and this function queries and prints the ptxVersion​ of a kernel<void>​ using cudaFuncGetAttributes​.

These TUs are linked into a program that determines the compute capability of device 0 and invokes volta()​ or pascal​() accordingly.

One would expect that invoking volta​() would always print 70​ and invoking pascal()​ would print 60​.

However, this is not the case. As described above, the kernel​ template has weak linkage, and so when linking the volta.o​ and pascal.o​ TUs together, the linker selects one of the instantiations of kernel<void>​ and discards the other.

The end result is that the program will randomly print 60 or 70 depending on which instantiation the linker picked.

// kernel.cuh
template <typename T>
__global__ void kernel(){}

// volta.cu
#include "kernel.cuh"
void volta(){
   cudaFuncAttributes attrs;
   cudaFuncGetAttributes(&attrs, (void *)(kernel<void>));
   printf("%d\n, attrs.ptxVersion*10);
}

// pascal.cu
#include "kernel.cuh"
void pascal(){

   cudaFuncAttributes attrs;
   cudaFuncGetAttributes(&attrs, (void *)(kernel<void>));
   printf("%d\n, attrs.ptxVersion*10);

}

// main.cpp
void volta();
void pascal();
int main(){
   int compute_capability;
   cudaDeviceGetAttribute(&compute_capability, cudaDevAttrComputeCapabilityMajor, 0);
   if(compute_capability >= 70)
      volta();
   else
      pascal();
}

nvcc -c -arch=sm_70 volta.cu
nvcc -c -arch=sm_60 pascal.cu
nvcc -c main.cpp
nvcc -o test volta.o pascal.o main.o
./test

TL;DR:

• Given a __global__​ function, kernel​, with weak linkage (like a template)
• Instantiate kernel​ in separate TUs compiled with different PTX architectures
• Link the separate TUs into a single program
• The linker will see the two instantiations of kernel​ as identical and discard one
• Invoking kernel()​ results in potentially executing code you did not expect

Describe the solution you'd like

Luckily the solution is quite simple.

• Every __global__ function should be annotated as static
• Every host template function should be annotated as __attribute__((visibility("hidden"))).

This makes the symbol hidden in any resulting dynamic library.

For convenience, you'd likely want to use a macro to wrap these like this:

#if defined(__CUDACC__)
#       define RAFT_KERNEL \
            static __global__
#else
#   define CUB_KERNEL
#endif

Additional Context

We've been bitten by this in Thrust/CUB several times over the years.

Thrust/CUB also have the ability to allow users to customize the namespace in order to differentiate the symbols and avoid this problem. However, this solution is not robust. First of all, it requires every user to remember to customize the namespace. Secondly, it's possible for users to properly customize the namespace and still run afoul of the issues that can result.

See:

• Dispatch mechanism may break when any two libraries that use CUB and/thrust have been compiled for different set of GPU architectures NVIDIA/cub#545
• [EPIC]: Hide the visibility of Thrust/CUB symbols in shared object libraries NVIDIA/cccl#166
• [FEA] CUTLASS should ensure all its symbols are hidden from shared object libraries  NVIDIA/cutlass#1027
• https://confluence.nvidia.com/x/esayW (internal only)

[jrhemstad]

@cjnolet note that there are a few alternatives to explicitly annotating every function with __attribute__((visibility("hidden"))).

gcc also provides a pragma based solution that allows modifying the visibility of many symbols at once:

#pragma GCC visibility push(hidden)
extern void someprivatefunct(int);
#pragma GCC visibility pop

You can also do the inverse of annotating non-binary APIs with __attribute__((visibility("hidden"))), you could instead annotate your public binary APIs with __attribute__((visibility("default"))) and instead compile with fvisibility=hidden to implicitly make the visibility of every function that is not otherwise annotated be hidden.

There's pros and cons among all these options, so it'll come down to what makes the most sense for RAFT.

See https://gcc.gnu.org/wiki/Visibility

[ahendriksen]

Hi Jake, thanks for the excellent write up. It's really clear.

What potential ramifications could it have if "Invoking kernel()​ results in potentially executing code you did not expect". What I can think of is:

• Two libraries use different versions of RAFT.
• They are linked into the same program.
• The host code of one version of RAFT calls the device code of another version.

Is this a concern?

The failure case you mentioned where within one version of RAFT the cudaPTXversion of a kernel could be different is something that I already expected could happen. I think our dispatch mechanism should continue to work in this case (because it calls cudaFuncGetAttributes on the kernel that it is going to run, not on another kernel).

[jrhemstad]

Is this a concern?

Yes. Getting the different PTX versions is just one example. In general, I think "running code that you did not except" is not good.

The most trivial example would be if there was a bug in the older version of RAFT that was fixed in the newer version, and the library using the newer version inadvertently picks up the kernel with the bug.

Another other example I've seen where this causes problems is if in one library invokes runtime APIs to configure attributes of a kernel, but instead the kernel from the other library gets invoked and the kernel configuration is ignored.