Hiding symbols from .so file

Developer Tools CUDA Developer Tools
1

Hello,

We are about to release a library of functions. For the C/C++ code, I have been using the -fvisibility=hidden flag to hide our internal function names from the outside world, and explicitly setting the visibility of the functions that make up our public API. However, I have noticed that any __global__ or __device__ functions still appear using objdump -T even after running strip on the .so file, ex:

00000000000xxxxx g    DF .text	000000000000003c  Base        _Z33superSecretFunctionThatDoesThingsyP6uchar4mmm6float2S1_
00000000000xxxxx g    DF .text	00000000000000f4  Base        _Z75__device_stub__Z33superSecretFunctionThatDoesThingsyP6uchar4mmm6float2S1_yP6uchar4mmmR6float2S2_

I am concerned that this could reveal the internal working of our algorithms, or worse, would allow users to call those functions directly using a carefully crafted program.

Is there a way to hide these symbols?

Thanks.

2

Finally, I made all of the __global__ and __device__ functions static. Since they are generally co-located with a launcher function, only the launcher function needs to see them. The launcher function is then marked with a “hidden” visibility at compile time, effectively suppressing all of the symbols when they are merged into a .so file.

3

Unfortunately, my “launcher” functions are actually still visible I wonder if the -fvisibility=hidden flag is getting passed to the host. Investigation continues…

4

Problem solved. I am using “modern” CUDA support in CMAKE and I needed to set CMAKE_CUDA_VISIBILITY_PRESET correctly, in addition to CXX and C

set(CMAKE_CUDA_VISIBILITY_PRESET hidden)
set(CMAKE_CXX_VISIBILITY_PRESET hidden)
set(CMAKE_C_VISIBILITY_PRESET hidden)