Merge pull request #49230 from fwyzard/from-CMSSW_16_0_X_2025-10-27-2300

Explicitly export CUDA kernels in shared libraries

Author: cmsbuild

HeterogeneousTest/CUDAKernel/README.md

@@ -14,8 +14,8 @@ kernels that call the device functions defined in the `HeterogeneousTest/CUDADev
 ```c++
 namespace cms::cudatest {
 
-  __global__ void kernel_add_vectors_f(...);
-  __global__ void kernel_add_vectors_d(...);
+  __global__ __attribute__((visibility("default"))) void kernel_add_vectors_f(...);
+  __global__ __attribute__((visibility("default"))) void kernel_add_vectors_d(...);
 
 }  // namespace cms::cudatest
 ```

HeterogeneousTest/CUDAKernel/interface/DeviceAdditionKernel.h

@@ -7,15 +7,17 @@
 
 namespace cms::cudatest {
 
-  __global__ void kernel_add_vectors_f(const float* __restrict__ in1,
-                                       const float* __restrict__ in2,
-                                       float* __restrict__ out,
-                                       size_t size);
+  // Mark the kernel with default visibility to export it as a public symbol for CUDA 12.8 and later
+  __global__ __attribute__((visibility("default"))) void kernel_add_vectors_f(const float* __restrict__ in1,
+                                                                              const float* __restrict__ in2,
+                                                                              float* __restrict__ out,
+                                                                              size_t size);
 
-  __global__ void kernel_add_vectors_d(const double* __restrict__ in1,
-                                       const double* __restrict__ in2,
-                                       double* __restrict__ out,
-                                       size_t size);
+  // Mark the kernel with default visibility to export it as a public symbol for CUDA 12.8 and later
+  __global__ __attribute__((visibility("default"))) void kernel_add_vectors_d(const double* __restrict__ in1,
+                                                                              const double* __restrict__ in2,
+                                                                              double* __restrict__ out,
+                                                                              size_t size);
 
 }  // namespace cms::cudatest
 

HeterogeneousTest/CUDAKernel/src/DeviceAdditionKernel.cu

@@ -7,17 +7,17 @@
 
 namespace cms::cudatest {
 
-  __global__ void kernel_add_vectors_f(const float* __restrict__ in1,
-                                       const float* __restrict__ in2,
-                                       float* __restrict__ out,
-                                       size_t size) {
+  __global__ __attribute__((visibility("default"))) void kernel_add_vectors_f(const float* __restrict__ in1,
+                                                                              const float* __restrict__ in2,
+                                                                              float* __restrict__ out,
+                                                                              size_t size) {
     add_vectors_f(in1, in2, out, size);
   }
 
-  __global__ void kernel_add_vectors_d(const double* __restrict__ in1,
-                                       const double* __restrict__ in2,
-                                       double* __restrict__ out,
-                                       size_t size) {
+  __global__ __attribute__((visibility("default"))) void kernel_add_vectors_d(const double* __restrict__ in1,
+                                                                              const double* __restrict__ in2,
+                                                                              double* __restrict__ out,
+                                                                              size_t size) {
     add_vectors_d(in1, in2, out, size);
   }