From 87eff461fe62547b0922f2c6ddba86138c021b88 Mon Sep 17 00:00:00 2001 From: Tim Grant Date: Thu, 15 Oct 2026 13:44:00 -0600 Subject: [PATCH 1/3] build(cuda): Migrate core build to FindCUDAToolkit Replace the deprecated FindCUDA module with CMake's supported FindCUDAToolkit module and imported CUDA runtime targets. Preserve CUDA_TOOLKIT_ROOT_DIR as a fallback input while using canonical CUDAToolkit results internally. Carry the complete toolkit include list through CUDA compilation and account for the CUDA 13 CCCL layout with older supported CMake releases. Keep OSL's CUDA compatibility headers ahead of toolkit headers. Use platform-aware CTest runtime paths so CUDA tests can locate shared libraries and Windows DLLs. CUDA 13 installs cudart and NVRTC DLLs in bin/x64, so prepend that directory when present while retaining bin for older toolkit layouts. Validated generated CTest paths for CUDA 12.4 and 13.3 on Windows. A clean-PATH loader probe using the generated CUDA 13.3 environment loads both cudart and NVRTC. Assisted-by: OpenAI Codex / GPT-6 Signed-off-by: Tim Grant --- src/cmake/cuda_macros.cmake | 24 +++++---- src/cmake/externalpackages.cmake | 92 +++++++++++++++++++++----------- src/cmake/testing.cmake | 27 +++++++--- src/liboslexec/CMakeLists.txt | 4 +- 4 files changed, 99 insertions(+), 48 deletions(-) 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}" From 9b15b73f3d0983ebe366ada21f5e7d07ecb802ee Mon Sep 17 00:00:00 2001 From: Tim Grant Date: Thu, 15 Oct 2026 13:44:10 -0600 Subject: [PATCH 2/3] build(example-cuda): Replace deprecated FindCUDA helpers Make the standalone CUDA example use FindCUDAToolkit and its imported runtime, NVRTC, and driver targets. Replace cuda_compile_ptx with a local NVCC custom command so the example continues to configure on CMake releases where FindCUDA has been removed. Propagate dependency include directories, account for the CUDA 13 CCCL layout, and choose a toolkit-compatible default GPU architecture while retaining a cache override. Assisted-by: OpenAI Codex / GPT-6 Signed-off-by: Tim Grant --- testsuite/example-cuda/CMakeLists.txt | 143 +++++++++++++++----------- 1 file changed, 81 insertions(+), 62 deletions(-) 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}) From dd127b4c8555022622bff7a90e092313377b9896 Mon Sep 17 00:00:00 2001 From: Tim Grant Date: Thu, 15 Oct 2026 13:44:18 -0600 Subject: [PATCH 3/3] fix(cuda): Make device alignment hints portable Older Clang CUDA builds on Windows can resolve __builtin_assume_aligned to an MSVC host-only declaration. Express the alignment promise with __builtin_assume in a device helper instead. The helper preserves the pointer type and evaluates each argument only once. Assisted-by: OpenAI Codex / GPT-6 Signed-off-by: Tim Grant --- src/testrender/cuda/rend_lib.cu | 30 ++++++++++++++++++++++++------ 1 file changed, 24 insertions(+), 6 deletions(-) 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) {