Fixes needed by CUDA 13.0 - now code compiles also on CUDA 12.x

Author: krzysg

CMakeLists.txt

@@ -209,16 +209,29 @@ set_property(TARGET aprObjLib PROPERTY POSITION_INDEPENDENT_CODE ON)
 
 if(APR_USE_CUDA)
     message(STATUS "APR: Building CUDA for APR")
-    set(CMAKE_CUDA_COMPILER "/usr/local/cuda/bin/nvcc")
+#    set(CMAKE_CUDA_COMPILER "/usr/local/cuda/bin/nvcc")
     set(CMAKE_CUDA_STANDARD 17)
+    if(CMAKE_VERSION VERSION_GREATER_EQUAL "3.24.0")
+        set(CMAKE_CUDA_ARCHITECTURES native)
+    endif()
+    set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
+    set(CMAKE_POSITION_INDEPENDENT_CODE ON)
     set(CMAKE_CUDA_RUNTIME_LIBRARY "Static")
-    set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -rdc=true --fmad=false --default-stream per-thread -Wno-deprecated-gpu-targets -Xptxas -v -DAPR_USE_CUDA")
+    set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --fmad=false --default-stream per-thread -Wno-deprecated-gpu-targets -Xptxas -v -DAPR_USE_CUDA")
+
     set(CMAKE_CUDA_FLAGS_RELEASE "-O3") # -lineinfo for profiling
     set(CMAKE_CUDA_FLAGS_DEBUG "-O0 -g -G")
     if(APR_BENCHMARK)
         set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -DAPR_BENCHMARK")
     endif()
     enable_language(CUDA)
+
+    if (${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER 13.0.0)
+        # For CUDA versions > 13.0.0 keep compatibility with older CUDAs - new CUDA introduce changes that
+        #  impact ELF visibility and linkage for __global__ functions and device variables
+        set(CMAKE_CUDA_FLAGS "--device-entity-has-hidden-visibility=false -static-global-template-stub=false  ${CMAKE_CUDA_FLAGS}")
+    endif ()
+
     set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -DAPR_USE_CUDA")
     set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DAPR_USE_CUDA")
     set(APR_CUDA_SOURCE_FILES
@@ -243,8 +256,9 @@ if(APR_BUILD_STATIC_LIB)
     # generate static library used as a intermediate step in generating fat lib
     set(STATIC_TARGET_NAME staticLib)
     add_library(${STATIC_TARGET_NAME} STATIC $<TARGET_OBJECTS:aprObjLib> ${APR_CUDA_SOURCE_FILES})
+    set_property(TARGET ${STATIC_TARGET_NAME} PROPERTY CUDA_SEPARABLE_COMPILATION ON)
     set_property(TARGET ${STATIC_TARGET_NAME} PROPERTY CUDA_ARCHITECTURES OFF)
-    target_compile_features(${STATIC_TARGET_NAME} PUBLIC cxx_std_14)
+    target_compile_features(${STATIC_TARGET_NAME} PUBLIC cxx_std_17)
     set_target_properties(${STATIC_TARGET_NAME} PROPERTIES OUTPUT_NAME ${LIBRARY_NAME})
     set_target_properties(${STATIC_TARGET_NAME} PROPERTIES CUDA_SEPARABLE_COMPILATION OFF)
     target_include_directories(${STATIC_TARGET_NAME} PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/src> $<BUILD_INTERFACE:${PROJECT_BINARY_DIR}>)
@@ -266,6 +280,7 @@ if(APR_BUILD_SHARED_LIB)
     set(SHARED_TARGET_NAME sharedLib)
     add_library(${SHARED_TARGET_NAME} SHARED $<TARGET_OBJECTS:aprObjLib> ${APR_CUDA_SOURCE_FILES})
     set_property(TARGET ${SHARED_TARGET_NAME} PROPERTY CUDA_ARCHITECTURES OFF)
+    set_property(TARGET ${SHARED_TARGET_NAME} PROPERTY CUDA_SEPARABLE_COMPILATION ON)
     target_include_directories(${SHARED_TARGET_NAME} PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/src> $<BUILD_INTERFACE:${PROJECT_BINARY_DIR}>)
     set_target_properties(${SHARED_TARGET_NAME} PROPERTIES OUTPUT_NAME ${LIBRARY_NAME})
     set_target_properties(${SHARED_TARGET_NAME} PROPERTIES LIBRARY_OUTPUT_NAME ${LIBRARY_NAME})