Switch from FindCUDA to the FindCUDAToolkit

src/cmake/cuda_macros.cmake

@@ -30,6 +30,8 @@ function ( NVCC_COMPILE cuda_src extra_headers ptx_generated extra_nvcc_args ext
           OUTPUT_VARIABLE ALL_OPENEXR_INCLUDES)
     list (TRANSFORM OpenImageIO_INCLUDES PREPEND -I
           OUTPUT_VARIABLE ALL_OpenImageIO_INCLUDES)
+    set (ALL_CUDA_INCLUDES ${OSL_CUDA_INCLUDE_DIRS})
+    list (TRANSFORM ALL_CUDA_INCLUDES PREPEND -I)
 
     if (NOT CUDA_NO_FTZ)
         set (NVCC_FTZ_FLAG "--ftz=true")
@@ -44,17 +46,17 @@ function ( NVCC_COMPILE cuda_src extra_headers ptx_generated extra_nvcc_args ext
     endforeach()
 
     add_custom_command ( OUTPUT ${cuda_ptx}
-        COMMAND ${CUDA_NVCC_EXECUTABLE}
+        COMMAND ${CUDAToolkit_NVCC_EXECUTABLE}
             "-I${OPTIX_INCLUDES}"
-            "-I${CUDA_INCLUDES}"
+            ${ALL_CUDA_INCLUDES}
             "-I${CMAKE_CURRENT_SOURCE_DIR}"
             "-I${CMAKE_BINARY_DIR}/include"
             "-I${PROJECT_SOURCE_DIR}/src/include"
             "-I${PROJECT_SOURCE_DIR}/src/cuda_common"
             ${EXTRA_INCLUDES}
             ${ALL_OpenImageIO_INCLUDES}
             ${ALL_IMATH_INCLUDES}
-            "-DFMT_DEPRECATED=\"\""
+            -DFMT_DEPRECATED=
             ${LLVM_COMPILE_FLAGS}
             -DOSL_USE_FAST_MATH=1
             -m64 -arch ${CUDA_TARGET_ARCH} -ptx
@@ -66,7 +68,8 @@ function ( NVCC_COMPILE cuda_src extra_headers ptx_generated extra_nvcc_args ext
             ${cuda_src} -o ${cuda_ptx}
         MAIN_DEPENDENCY ${cuda_src}
         DEPENDS ${cuda_src} ${cuda_headers} ${extra_libs} oslexec
-        WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}" )
+        WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}"
+        VERBATIM )
 endfunction ()
 
 # Function to compile a C++ source file to CUDA-compatible LLVM bitcode
@@ -142,7 +145,7 @@ function ( MAKE_CUDA_BITCODE src suffix generated_bc extra_clang_args )
         set (CLANG_FTZ_FLAG "-fcuda-flush-denormals-to-zero")
     endif ()
 
-    if ("${CUDA_VERSION}" VERSION_GREATER_EQUAL "12.0")
+    if ("${CUDAToolkit_VERSION}" VERSION_GREATER_EQUAL "12.0")
         # The textureReference API was removed in CUDA 12.0, but it's still referenced
         # in the clang headers, so compilation will fail with 12.0 and newer toolkits
         # due to the missing definitions.
@@ -152,7 +155,7 @@ function ( MAKE_CUDA_BITCODE src suffix generated_bc extra_clang_args )
         set (CUDA_TEXREF_FIX "-D__CLANG_CUDA_TEXTURE_INTRINSICS_H__")
     endif()
 
-    if ("${CUDA_VERSION}" VERSION_GREATER_EQUAL "13.0")
+    if ("${CUDAToolkit_VERSION}" VERSION_GREATER_EQUAL "13.0")
         # CUDA 13 removed texture_fetch_functions.h, but older Clang CUDA
         # wrappers still include it. Search OSL's compatibility headers first.
         set (CUDA_COMPAT_INCLUDE
@@ -174,13 +177,15 @@ function ( MAKE_CUDA_BITCODE src suffix generated_bc extra_clang_args )
           OUTPUT_VARIABLE ALL_OPENEXR_INCLUDES)
     list (TRANSFORM OpenImageIO_INCLUDES PREPEND -I
           OUTPUT_VARIABLE ALL_OpenImageIO_INCLUDES)
+    set (ALL_CUDA_INCLUDES ${OSL_CUDA_INCLUDE_DIRS})
+    list (TRANSFORM ALL_CUDA_INCLUDES PREPEND -I)
 
     add_custom_command (OUTPUT ${bc_cuda}
         COMMAND ${LLVM_BC_GENERATOR}
             ${CUDA_COMPAT_INCLUDE}
             ${CLANG_CUDA_COMPAT_FLAGS}
             "-I${OPTIX_INCLUDES}"
-            "-I${CUDA_INCLUDES}"
+            ${ALL_CUDA_INCLUDES}
             "-I${CMAKE_CURRENT_SOURCE_DIR}"
             "-I${CMAKE_SOURCE_DIR}/src/liboslexec"
             "-I${CMAKE_BINARY_DIR}/include"
@@ -189,15 +194,16 @@ function ( MAKE_CUDA_BITCODE src suffix generated_bc extra_clang_args )
             ${ALL_OpenImageIO_INCLUDES}
             ${ALL_IMATH_INCLUDES}
             ${ALL_OPENEXR_INCLUDES}
-            ${LLVM_COMPILE_FLAGS} ${CUDA_LIB_FLAGS} ${CLANG_MSVC_FIX} ${CUDA_TEXREF_FIX}
+            ${LLVM_COMPILE_FLAGS} ${OSL_CUDA_LIB_FLAGS} ${CLANG_MSVC_FIX} ${CUDA_TEXREF_FIX}
             -D__CUDACC__ -DOSL_COMPILING_TO_BITCODE=1 -DNDEBUG -DOIIO_NO_SSE -D__CUDADEVRT_INTERNAL__
             --language=cuda --cuda-device-only --cuda-gpu-arch=${CUDA_TARGET_ARCH}
             -Wno-deprecated-register -Wno-format-security
             -fno-math-errno -ffast-math ${CUDA_OPT_FLAG_CLANG} ${CLANG_FTZ_FLAG} -S -emit-llvm ${extra_clang_args}
             ${src} -o ${asm_cuda}
         COMMAND ${LLVM_AS_TOOL} -f -o ${bc_cuda} ${asm_cuda}
         DEPENDS ${exec_headers} ${PROJECT_PUBLIC_HEADERS} ${src}
-        WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}" )
+        WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}"
+        VERBATIM )
 endfunction ()
 
 # Compile a CUDA source file to LLVM bitcode, and then serialize the bitcode to

src/cmake/externalpackages.cmake

@@ -119,49 +119,73 @@ endif ()
 option (CUDA_PREFER_STATIC_LIBS "Prefer static CUDA libraries" OFF)
 if (OSL_USE_OPTIX)
     if (USE_LLVM_BITCODE)
-        if (NOT CUDA_TOOLKIT_ROOT_DIR AND NOT $ENV{CUDA_TOOLKIT_ROOT_DIR} STREQUAL "")
-            set (CUDA_TOOLKIT_ROOT_DIR $ENV{CUDA_TOOLKIT_ROOT_DIR})
+        # FindCUDAToolkit uses CUDAToolkit_ROOT. Honor the old FindCUDA input
+        # as a fallback so existing OSL build scripts keep working.
+        if (NOT CUDAToolkit_ROOT
+            AND "$ENV{CUDAToolkit_ROOT}" STREQUAL "")
+            if (CUDA_TOOLKIT_ROOT_DIR)
+                set (CUDAToolkit_ROOT "${CUDA_TOOLKIT_ROOT_DIR}")
+            elseif (NOT "$ENV{CUDA_TOOLKIT_ROOT_DIR}" STREQUAL "")
+                set (CUDAToolkit_ROOT "$ENV{CUDA_TOOLKIT_ROOT_DIR}")
+            endif ()
         endif ()
 
-        if (NOT CUDA_FIND_QUIETLY OR NOT OptiX_FIND_QUIETLY)
-            message (STATUS "CUDA_TOOLKIT_ROOT_DIR = ${CUDA_TOOLKIT_ROOT_DIR}")
+        if (NOT CUDAToolkit_FIND_QUIETLY OR NOT OptiX_FIND_QUIETLY)
+            message (STATUS "CUDAToolkit_ROOT = ${CUDAToolkit_ROOT}")
         endif ()
 
-        checked_find_package (CUDA REQUIRED
-                             VERSION_MIN 9.0
-                             RECOMMEND_MIN 11.0
-                             RECOMMEND_MIN_REASON
-                                "We don't actively test CUDA older than 11"
-                             PRINT CUDA_INCLUDES)
-        set (CUDA_INCLUDES ${CUDA_TOOLKIT_ROOT_DIR}/include)
-        include_directories (BEFORE "${CUDA_INCLUDES}")
+        checked_find_package (CUDAToolkit REQUIRED
+                              VERSION_MIN 9.0
+                              RECOMMEND_MIN 11.0
+                              RECOMMEND_MIN_REASON
+                                  "We don't actively test CUDA older than 11"
+                              PRINT CUDAToolkit_INCLUDE_DIRS
+                                    CUDAToolkit_NVCC_EXECUTABLE)
+
+        # CUDA 13 moved its C++ Core Libraries (CCCL) under include/cccl.
+        # New CMake releases report that path, but older supported releases
+        # do not, so augment the discovered list when that directory exists.
+        set (OSL_CUDA_INCLUDE_DIRS ${CUDAToolkit_INCLUDE_DIRS})
+        foreach (cuda_include_dir IN LISTS CUDAToolkit_INCLUDE_DIRS)
+            if (IS_DIRECTORY "${cuda_include_dir}/cccl")
+                list (APPEND OSL_CUDA_INCLUDE_DIRS
+                      "${cuda_include_dir}/cccl")
+            endif ()
+        endforeach ()
+        list (FILTER OSL_CUDA_INCLUDE_DIRS EXCLUDE REGEX "^$")
+        list (REMOVE_DUPLICATES OSL_CUDA_INCLUDE_DIRS)
+        include_directories (BEFORE ${OSL_CUDA_INCLUDE_DIRS})
+
+        # Clang's --cuda-path needs the toolkit root, not the target-specific
+        # library directory. FindCUDAToolkit normally reports it directly.
+        if (CUDAToolkit_LIBRARY_ROOT)
+            set (OSL_CUDA_TOOLKIT_ROOT "${CUDAToolkit_LIBRARY_ROOT}")
+        elseif (CUDAToolkit_BIN_DIR)
+            get_filename_component (OSL_CUDA_TOOLKIT_ROOT
+                                    "${CUDAToolkit_BIN_DIR}" DIRECTORY)
+        else ()
+            message (FATAL_ERROR
+                     "Could not determine the CUDA toolkit root directory")
+        endif ()
 
         STRING (FIND ${LLVM_TARGETS} "NVPTX" nvptx_index)
         if (NOT ${nvptx_index} GREATER -1)
             message (FATAL_ERROR "NVPTX target is not available in the provided LLVM build")
         endif()
 
-        set (CUDA_LIB_FLAGS "--cuda-path=${CUDA_TOOLKIT_ROOT_DIR}")
+        set (OSL_CUDA_LIB_FLAGS "--cuda-path=${OSL_CUDA_TOOLKIT_ROOT}")
 
-        # If the user wants, try to use static libs here to putting static lib
-        # suffixes earlier in the suffix list. Don't forget to restore after
-        # so that this only applies to these library searches right here.
-        set (save_lib_path ${CMAKE_FIND_LIBRARY_SUFFIXES})
+        # Imported targets carry the platform-specific system dependencies
+        # needed by the selected CUDA runtime, including on Windows.
         if (CUDA_PREFER_STATIC_LIBS)
-            set (CMAKE_FIND_LIBRARY_SUFFIXES .lib .a ${CMAKE_FIND_LIBRARY_SUFFIXES})
-            find_library(cudart_lib REQUIRED
-                         NAMES cudart_static cudart
-                         PATHS "${CUDA_TOOLKIT_ROOT_DIR}/lib64" "${CUDA_TOOLKIT_ROOT_DIR}/x64" "${CUDA_TOOLKIT_ROOT_DIR}/lib/x64")
+            set (OSL_CUDA_RUNTIME_TARGET CUDA::cudart_static)
         else ()
-            find_library(cudart_lib REQUIRED
-                         NAMES cudart
-                         PATHS "${CUDA_TOOLKIT_ROOT_DIR}/lib64" "${CUDA_TOOLKIT_ROOT_DIR}/x64" "${CUDA_TOOLKIT_ROOT_DIR}/lib/x64")
+            set (OSL_CUDA_RUNTIME_TARGET CUDA::cudart)
+        endif ()
+        if (NOT TARGET ${OSL_CUDA_RUNTIME_TARGET})
+            message (FATAL_ERROR
+                     "CUDA runtime target ${OSL_CUDA_RUNTIME_TARGET} not found")
         endif ()
-        # Is it really a good idea to completely reset CUDA_LIBRARIES here?
-        set(CUDA_LIBRARIES ${cudart_lib})
-        set(CUDA_EXTRA_LIBS ${CUDA_EXTRA_LIBS} dl rt)
-        set (CMAKE_FIND_LIBRARY_SUFFIXES ${save_lib_path})
-        unset (save_lib_path)
     endif()
 
     # OptiX setup
@@ -175,10 +199,16 @@ if (OSL_USE_OPTIX)
     endif ()
 
     function (osl_optix_target TARGET)
-        target_include_directories (${TARGET} BEFORE PRIVATE ${CUDA_INCLUDES} ${OPTIX_INCLUDES})
+        target_include_directories (${TARGET} BEFORE PRIVATE
+                                    ${OSL_CUDA_INCLUDE_DIRS}
+                                    ${OPTIX_INCLUDES})
         ## XXX: Should -DPTX_PATH point to (or include) CMAKE_CURRENT_BINARY_DIR so tests can run before installation ?
         target_compile_definitions (${TARGET} PRIVATE PTX_PATH="${OSL_PTX_INSTALL_DIR}")
-        target_link_libraries (${TARGET} PRIVATE ${CUDA_LIBRARIES} ${CUDA_EXTRA_LIBS} ${OPTIX_LIBRARIES} ${OPTIX_EXTRA_LIBS})
+        target_link_libraries (${TARGET} PRIVATE
+                               ${OSL_CUDA_RUNTIME_TARGET}
+                               ${CUDA_EXTRA_LIBS}
+                               ${OPTIX_LIBRARIES}
+                               ${OPTIX_EXTRA_LIBS})
     endfunction()
 else ()
     message(STATUS "CUDA/OptiX support disabled")

src/cmake/testing.cmake

@@ -120,16 +120,31 @@ macro (add_one_testsuite testname testsrcdir)
         set_tests_properties (${testname} PROPERTIES LABELS noise
                               PROCESSORS 2 COST 4)
     endif ()
-    if (${testname} MATCHES "optix")
+    if ("${testname}" MATCHES "optix")
         set_tests_properties (${testname} PROPERTIES LABELS optix)
-        if ("${CUDA_VERSION}" VERSION_GREATER_EQUAL "10.0")
-            # Make sure libnvrtc-builtins.so is reachable
-            set_property (TEST ${testname} APPEND PROPERTY ENVIRONMENT LD_LIBRARY_PATH=${CUDA_TOOLKIT_ROOT_DIR}/lib64:$ENV{LD_LIBRARY_PATH})
-        endif()
         if (DEFINED ENV{OSL_TESTSUITE_OPTIX_THRESH_SCALE})
             set_property (TEST ${testname} APPEND PROPERTY ENVIRONMENT OSL_TESTSUITE_THRESH_SCALE=$ENV{OSL_TESTSUITE_OPTIX_THRESH_SCALE})
         endif ()
     endif ()
+    if ("${testname}" MATCHES "optix|example-cuda")
+        # Make sure the CUDA and NVRTC runtime libraries are reachable. CTest
+        # adds the host platform's path separator for this operation.
+        if (WIN32)
+            set_property (TEST ${testname} APPEND PROPERTY
+                          ENVIRONMENT_MODIFICATION
+                          "PATH=path_list_prepend:${CUDAToolkit_BIN_DIR}")
+            # CUDA 13 places Windows runtime DLLs in bin/x64.
+            if (IS_DIRECTORY "${CUDAToolkit_BIN_DIR}/x64")
+                set_property (TEST ${testname} APPEND PROPERTY
+                              ENVIRONMENT_MODIFICATION
+                              "PATH=path_list_prepend:${CUDAToolkit_BIN_DIR}/x64")
+            endif ()
+        else ()
+            set_property (TEST ${testname} APPEND PROPERTY
+                          ENVIRONMENT_MODIFICATION
+                          "LD_LIBRARY_PATH=path_list_prepend:${CUDAToolkit_LIBRARY_DIR}")
+        endif ()
+    endif ()
     if (${testname} MATCHES "-reg")
         # These are batched shading regression tests. Some are quite
         # long, so give them a higher cost and timeout.
@@ -556,7 +571,7 @@ macro (osl_add_all_tests)
     endif ()
 
     # Only run the OptiX tests if OptiX and CUDA are found
-    if (OPTIX_FOUND AND CUDA_FOUND)
+    if (OPTIX_FOUND AND CUDAToolkit_FOUND)
         TESTSUITE ( testoptix testoptix-noise example-cuda)
     endif ()
 

src/liboslexec/CMakeLists.txt

@@ -203,7 +203,7 @@ if (USE_LLVM_BITCODE)
 
     EMBED_LLVM_BITCODE_IN_CPP ( "${rs_dependent_ops_srcs}" "_host_rs" "osl_llvm_compiled_rs_dependent_ops" lib_src "" "${include_dirs}")
 
-    if (CUDA_FOUND)
+    if (CUDAToolkit_FOUND)
         add_compile_definitions (OSL_LLVM_CUDA_BITCODE)
 
         # The shadeops are compiled to both LLVM bitcode and PTX, and embedded
@@ -588,7 +588,7 @@ target_compile_definitions (${local_lib}
     PRIVATE
         OSL_EXPORTS
         CUDA_TARGET_ARCH="${CUDA_TARGET_ARCH}"
-        OSL_CUDA_VERSION="${CUDA_VERSION}"
+        OSL_CUDA_VERSION="${CUDAToolkit_VERSION}"
         OSL_OPTIX_VERSION="${OPTIX_VERSION}"
         OSL_CPP_COMPILER_DEFAULT="${CMAKE_CXX_COMPILER}"
         OSL_CPP_COMPILER_FLAGS_DEFAULT="${_cpp_base_flags} ${_cpp_inc_flags_str}"

src/testrender/cuda/rend_lib.cu

@@ -11,11 +11,29 @@
 #include <cuda_runtime.h>
 #include <optix_device.h>
 
+#include <cstdint>
+
 #include "rend_lib.h"
 
-#define MEMCPY_ALIGNED(dst, src, size, alignment)    \
-    memcpy(__builtin_assume_aligned(dst, alignment), \
-           __builtin_assume_aligned(src, alignment), size);
+// Older Clang CUDA builds on Windows can resolve __builtin_assume_aligned to
+// an MSVC host-only declaration. Express the same device-side alignment
+// promise with __builtin_assume instead.
+template<int Alignment, class T>
+__device__ inline T*
+device_assume_aligned(T* ptr)
+{
+#if defined(_WIN32)
+    __builtin_assume((reinterpret_cast<std::uintptr_t>(ptr) & (Alignment - 1))
+                     == 0);
+    return ptr;
+#else
+    return static_cast<T*>(__builtin_assume_aligned(ptr, Alignment));
+#endif
+}
+
+#define MEMCPY_ALIGNED(dst, src, size, alignment) \
+    memcpy(device_assume_aligned<alignment>(dst), \
+           device_assume_aligned<alignment>(src), size);
 
 OSL_NAMESPACE_BEGIN
 namespace pvt {
@@ -448,12 +466,12 @@ osl_range_check_err(int indexvalue, int length, OSL::ustringhash_pod symname,
 }
 
 
-#define MAT(m) (*(OSL::Matrix44*)__builtin_assume_aligned(m, alignof(float)))
+#define MAT(m) (*(OSL::Matrix44*)device_assume_aligned<alignof(float)>(m))
 
 __device__ int
 osl_get_matrix(void* sg_, void* r, OSL::ustringhash_pod from_)
 {
-    r                           = __builtin_assume_aligned(r, alignof(float));
+    r                           = device_assume_aligned<alignof(float)>(r);
     OSL::ustringhash from       = OSL::ustringhash_from(from_);
     OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
     if (from == OSL::Hashes::common) {
@@ -498,7 +516,7 @@ osl_get_matrix(void* sg_, void* r, OSL::ustringhash_pod from_)
 __device__ int
 osl_get_inverse_matrix(void* sg_, void* r, OSL::ustringhash_pod to_)
 {
-    r                           = __builtin_assume_aligned(r, alignof(float));
+    r                           = device_assume_aligned<alignof(float)>(r);
     OSL::ustringhash to         = OSL::ustringhash_from(to_);
     OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_;
     if (to == OSL::Hashes::common) {