Abstract:

NVIDIA is introducing two significant, behavior-changing updates to the NVIDIA CUDA Compiler Driver (NVCC) in the upcoming CUDA 13.0 release, intended to mitigate subtle, hard-to-debug runtime errors related to ELF visibility and linkage of device entities. The first change, primarily affecting shared libraries on non-Windows platforms, forces hidden ELF visibility for __global__ functions and device variables (__managed__, __device__, __constant__) to prevent conflicts arising from linking multiple versions of the CUDA Runtime Library (CUDART). The second change forces internal linkage for host template stub definitions of __global__ functions in the default whole program compilation mode (-rdc=false), resolving ambiguities when template definitions are merged by the host linker. While these changes enhance robustness, developers using existing workflows must be prepared to update their code or utilize new NVCC flags (-device-entity-has-hidden-visibility and -static-global-template-stub) to explicitly opt out of the new default behaviors.

CUDA C++ Compiler Updates Impacting ELF Visibility and Linkage

• Introduction to Changes (CUDA 13.0)

  • NVIDIA is implementing two significant changes to NVCC defaults starting in CUDA 13.0 to address subtle runtime bugs that are difficult to detect (no build-time warnings or CUDA runtime errors).
  • Change #1: Force hidden ELF visibility for __global__ functions and device variables (__managed__/__device__/__constant__).
  • Change #2: Force internal linkage for __global__ function host template stub definitions (when using whole program compilation mode, -rdc=false).

• NVCC Change #1: ELF Visibility

  • The Problem: Prior to CUDA 13.0, when device entities were compiled into a shared library (.so), they maintained external visibility. If the main program also statically links CUDART (the default), this results in two distinct CUDART libraries being linked, causing runtime failures (e.g., kernel launches fail silently, or __managed__ variable access causes segmentation faults).
  • Platforms Affected: All non-Windows platforms (Windows host compilers typically do not export symbols by default).
  • CUDA 13.0 Solution (New Default): NVCC now forces hidden visibility for these device symbols, causing programs that rely on external visibility to fail at link time (a preferable outcome to runtime crashes).
  • Opt-Out/Legacy Behavior: Developers can restore pre-CUDA 13.0 behavior using the flag -device-entity-has-hidden-visibility=false or by explicitly annotating symbols with __attribute__((visibility("default"))).

• NVCC Change #2: Forced Internal Linkage

  • The Context: This change relates specifically to __global__ template functions compiled in the default whole program mode (-rdc=false). NVCC generates host stub functions for these kernels.
  • The Problem: Host stub functions for __global__ templates traditionally had external (weak) linkage. If two different translation units (a.cu and b.cu) include and launch the same template, the host linker may merge the stubs, causing a launch in a.cu to incorrectly execute the kernel in the device program generated from b.cu.
  • CUDA 13.0 Solution (New Default): NVCC now forces internal linkage for these template stubs, preventing the host linker from merging them, ensuring that each launch correctly links to its own translation unit's device program.
  • Opt-Out (If Needed): This change breaks programs relying on explicit template instantiation across translation units. To opt out, use the flag -static-global-template-stub=false.

• Key Takeaway:

  • CUDA 13.0 introduces new, more robust default linkage behaviors. Existing code bases, especially those involving shared libraries or header-only template libraries, may experience new compile or link failures due to these changes. Developers should review and potentially implement the opt-out flags if their workflows relied on the legacy behavior.