diff --git a/src/cmake/cuda_macros.cmake b/src/cmake/cuda_macros.cmake index 9e44ed1d6e..bbee515c4b 100644 --- a/src/cmake/cuda_macros.cmake +++ b/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,9 +46,9 @@ 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" @@ -54,7 +56,7 @@ function ( NVCC_COMPILE cuda_src extra_headers ptx_generated extra_nvcc_args ext ${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,7 +194,7 @@ 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 @@ -197,7 +202,8 @@ function ( MAKE_CUDA_BITCODE src suffix generated_bc 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 diff --git a/src/cmake/externalpackages.cmake b/src/cmake/externalpackages.cmake index 11343a0228..2e55e0ebae 100644 --- a/src/cmake/externalpackages.cmake +++ b/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") diff --git a/src/cmake/testing.cmake b/src/cmake/testing.cmake index 46fab0e0e7..b9577a6f67 100644 --- a/src/cmake/testing.cmake +++ b/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 () diff --git a/src/liboslexec/CMakeLists.txt b/src/liboslexec/CMakeLists.txt index a72a50aef5..2ea8e34cee 100644 --- a/src/liboslexec/CMakeLists.txt +++ b/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}" diff --git a/src/testrender/cuda/rend_lib.cu b/src/testrender/cuda/rend_lib.cu index 681fae9579..02fe2491ae 100644 --- a/src/testrender/cuda/rend_lib.cu +++ b/src/testrender/cuda/rend_lib.cu @@ -11,11 +11,29 @@ #include #include +#include + #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 +__device__ inline T* +device_assume_aligned(T* ptr) +{ +#if defined(_WIN32) + __builtin_assume((reinterpret_cast(ptr) & (Alignment - 1)) + == 0); + return ptr; +#else + return static_cast(__builtin_assume_aligned(ptr, Alignment)); +#endif +} + +#define MEMCPY_ALIGNED(dst, src, size, alignment) \ + memcpy(device_assume_aligned(dst), \ + device_assume_aligned(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(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(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(r); OSL::ustringhash to = OSL::ustringhash_from(to_); OSL_CUDA::ShaderGlobals* sg = (OSL_CUDA::ShaderGlobals*)sg_; if (to == OSL::Hashes::common) { diff --git a/testsuite/example-cuda/CMakeLists.txt b/testsuite/example-cuda/CMakeLists.txt index 6797802870..5d95ffebfc 100644 --- a/testsuite/example-cuda/CMakeLists.txt +++ b/testsuite/example-cuda/CMakeLists.txt @@ -2,28 +2,25 @@ # SPDX-License-Identifier: BSD-3-Clause # https://github.com/AcademySoftwareFoundation/OpenShadingLanguage -cmake_minimum_required(VERSION 3.15) -project(examplecuda LANGUAGES CXX) +# Keep this in sync with the top-level project. +cmake_minimum_required (VERSION 3.23) +project (examplecuda LANGUAGES CXX) if (NOT CMAKE_BUILD_TYPE) set (CMAKE_BUILD_TYPE "Release") endif () # Is there a better way to get the location of the main OSL CMakeLists.txt? -set(OSL_LIST_DIR "${CMAKE_CURRENT_LIST_DIR}/../../../../") +set (OSL_LIST_DIR "${CMAKE_CURRENT_LIST_DIR}/../../../../") # We have to make the path absolute or CMake will fail to find the file... -get_filename_component(OSL_LIST_DIR ${OSL_LIST_DIR} ABSOLUTE) -set(CMAKE_MODULE_PATH "${OSL_LIST_DIR}/src/cmake;${OSL_LIST_DIR}/src/cmake/modules") +get_filename_component (OSL_LIST_DIR "${OSL_LIST_DIR}" ABSOLUTE) +list (APPEND CMAKE_MODULE_PATH + "${OSL_LIST_DIR}/src/cmake" + "${OSL_LIST_DIR}/src/cmake/modules") -include(check_is_enabled) -include(checked_find_package) - -find_package(OSL REQUIRED) -find_package(CUDA REQUIRED) -checked_find_package(LLVM 7.0 REQUIRED) -checked_find_package(Imath 3.1 REQUIRED) -checked_find_package(OpenImageIO 2.4 REQUIRED) -checked_find_package(OptiX REQUIRED) +find_package (OSL REQUIRED) +find_package (CUDAToolkit 9.0 REQUIRED) +find_package (OptiX 7.0 REQUIRED) # Make the build area layout look like we expect @@ -31,59 +28,81 @@ set (CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib) set (CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib) set (CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin) -# TODO: what are the path suffixes on other platforms? -find_library(CUDA_nvrtc_LIBRARY nvrtc HINTS ${CUDA_TOOLKIT_ROOT_DIR} PATH_SUFFIXES lib lib64) -find_library(CUDA_cuda_LIBRARY cuda HINTS ${CUDA_TOOLKIT_ROOT_DIR} PATH_SUFFIXES lib/stubs lib64/stubs) - -# TODO: move to sm_60? -set(CUDA_TARGET_ARCH sm_35) +if (CUDAToolkit_VERSION VERSION_GREATER_EQUAL 13.0) + set (cuda_default_arch "sm_75") +else () + set (cuda_default_arch "sm_60") +endif () +set (CUDA_TARGET_ARCH "${cuda_default_arch}" CACHE STRING + "CUDA GPU architecture (e.g. sm_60)") set (CMAKE_CXX_STANDARD 17 CACHE STRING "C++ standard to build with (17, 20, etc.)") -# Compile our "renderer" to PTX -cuda_compile_ptx(CUDA_PTX_FILES cuda_grid_renderer.cu - OPTIONS --gpu-architecture=${CUDA_TARGET_ARCH} --use_fast_math -dc - --std=c++${CMAKE_CXX_STANDARD} - --expt-relaxed-constexpr - -I${OSL_INCLUDES} - -I${IMATH_INCLUDES} - -I${OpenImageIO_INCLUDES} - -I${OPTIX_INCLUDES} -) -add_custom_command(OUTPUT ${CMAKE_BINARY_DIR}/cuda_grid_renderer.ptx - COMMAND ${CMAKE_COMMAND} -E rename ${CUDA_PTX_FILES} ${CMAKE_BINARY_DIR}/cuda_grid_renderer.ptx - DEPENDS ${CUDA_PTX_FILES} -) -add_custom_target(cuda_grid_renderer_ptx ALL - DEPENDS ${CMAKE_BINARY_DIR}/cuda_grid_renderer.ptx cuda_grid_renderer.cu - SOURCES cuda_grid_renderer.cu -) - -# Compile the rend_lib shadeops to PTX -cuda_compile_ptx(CUDA_PTX_FILES rend_lib.cu - OPTIONS --gpu-architecture=${CUDA_TARGET_ARCH} --use_fast_math -dc - --std=c++${CMAKE_CXX_STANDARD} - --expt-relaxed-constexpr - -I${OSL_INCLUDES} - -I${IMATH_INCLUDES} - -I${OpenImageIO_INCLUDES} - -I${OPTIX_INCLUDES} -) -add_custom_command(OUTPUT ${CMAKE_BINARY_DIR}/rend_lib.ptx - COMMAND ${CMAKE_COMMAND} -E rename ${CUDA_PTX_FILES} ${CMAKE_BINARY_DIR}/rend_lib.ptx - DEPENDS ${CUDA_PTX_FILES} -) -add_custom_target(rend_lib_ptx ALL - DEPENDS ${CMAKE_BINARY_DIR}/rend_lib.ptx rend_lib.cu - SOURCES rend_lib.cu -) +get_target_property (imath_include_dirs Imath::Imath + INTERFACE_INCLUDE_DIRECTORIES) +get_target_property (oiio_include_dirs OpenImageIO::OpenImageIO + INTERFACE_INCLUDE_DIRECTORIES) +set (cuda_toolkit_include_dirs ${CUDAToolkit_INCLUDE_DIRS}) +foreach (cuda_include_dir IN LISTS CUDAToolkit_INCLUDE_DIRS) + if (IS_DIRECTORY "${cuda_include_dir}/cccl") + list (APPEND cuda_toolkit_include_dirs "${cuda_include_dir}/cccl") + endif () +endforeach () +list (FILTER cuda_toolkit_include_dirs EXCLUDE REGEX "^$") +list (REMOVE_DUPLICATES cuda_toolkit_include_dirs) +set (cuda_ptx_include_dirs ${OSL_INCLUDES} + ${imath_include_dirs} + ${oiio_include_dirs} + ${OPTIX_INCLUDES} + ${cuda_toolkit_include_dirs}) +list (TRANSFORM cuda_ptx_include_dirs PREPEND -I) + +# Compile a CUDA source file to PTX with NVCC. This replaces the helper that +# was supplied by the deprecated FindCUDA module. +function (compile_cuda_to_ptx cuda_src ptx_generated) + get_filename_component (cuda_src_we "${cuda_src}" NAME_WE) + set (cuda_ptx "${CMAKE_CURRENT_BINARY_DIR}/${cuda_src_we}.ptx") + file (GLOB cuda_headers "${CMAKE_CURRENT_SOURCE_DIR}/*.h") + + add_custom_command (OUTPUT "${cuda_ptx}" + COMMAND "${CUDAToolkit_NVCC_EXECUTABLE}" + --gpu-architecture=${CUDA_TARGET_ARCH} + --use_fast_math + --std=c++${CMAKE_CXX_STANDARD} + --expt-relaxed-constexpr + -dc -ptx + ${cuda_ptx_include_dirs} + "${CMAKE_CURRENT_SOURCE_DIR}/${cuda_src}" + -o "${cuda_ptx}" + MAIN_DEPENDENCY "${cuda_src}" + DEPENDS ${cuda_headers} + VERBATIM) + + set (${ptx_generated} "${cuda_ptx}" PARENT_SCOPE) +endfunction () + +compile_cuda_to_ptx (cuda_grid_renderer.cu cuda_grid_renderer_ptx_file) +add_custom_target (cuda_grid_renderer_ptx ALL + DEPENDS "${cuda_grid_renderer_ptx_file}" + SOURCES cuda_grid_renderer.cu) + +compile_cuda_to_ptx (rend_lib.cu rend_lib_ptx_file) +add_custom_target (rend_lib_ptx ALL + DEPENDS "${rend_lib_ptx_file}" + SOURCES rend_lib.cu) # Compile and link the main executable -add_executable(example-cuda example-cuda.cpp cuda_grid_renderer.cpp) +add_executable (example-cuda example-cuda.cpp cuda_grid_renderer.cpp) +add_dependencies (example-cuda cuda_grid_renderer_ptx rend_lib_ptx) -target_link_libraries(example-cuda PRIVATE OSL::oslexec OSL::oslquery ${CUDA_LIBRARIES} ${CUDA_nvrtc_LIBRARY} ${CUDA_cuda_LIBRARY}) -target_include_directories(example-cuda PRIVATE ${CUDA_INCLUDE_DIRS}) -set_property(TARGET example-cuda PROPERTY CXX_STANDARD 17) +target_link_libraries (example-cuda PRIVATE + OSL::oslexec + OSL::oslquery + CUDA::cudart + CUDA::nvrtc + CUDA::cuda_driver) +target_include_directories (example-cuda PRIVATE ${cuda_toolkit_include_dirs}) +set_property (TARGET example-cuda PROPERTY CXX_STANDARD 17) -install(TARGETS example-cuda DESTINATION ${CMAKE_CURRENT_SOURCE_DIR}) +install (TARGETS example-cuda DESTINATION ${CMAKE_CURRENT_SOURCE_DIR})