# Modifications Copyright (c) 2025-2026 Advanced Micro Devices, Inc. # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal # in the Software without restriction, including without limitation the rights # to use, copy, modify, merge, publish, distribute, sublicense, and/or sell # copies of the Software, and to permit persons to whom the Software is # furnished to do so, subject to the following conditions: # The above copyright notice and this permission notice shall be included in # all copies or substantial portions of the Software. # THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR # IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, # FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE # AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER # LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. # For every private internal header build a TU that directly includes it # without anything else to verify whether its properly modular. # We also want to check whether all potential fallbacks are actually valid, # even if we never exercise them in the current support matrix add_custom_target(libcudacxx.test.internal_headers) if ("NVHPC" STREQUAL "${CMAKE_CXX_COMPILER_ID}") find_package(NVHPC) else() find_package(CUDAToolkit) endif() # Grep all internal headers file(GLOB_RECURSE internal_headers RELATIVE "${libhipcxx_SOURCE_DIR}/include/cuda/" CONFIGURE_DEPENDS ${libhipcxx_SOURCE_DIR}/include/cuda/__*/*.h ${libhipcxx_SOURCE_DIR}/include/cuda/std/__*/*.h ) # headers in `__cuda` are meant to come after the related "cuda" headers so they do not compile on their own list(FILTER internal_headers EXCLUDE REGEX "__cuda/*") # mdspan is currently not supported on msvc outside of C++20 if ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "MSVC" AND NOT "${CMAKE_CXX_STANDARD}" MATCHES "20") list(FILTER internal_headers EXCLUDE REGEX "mdspan") endif() # Exclude __limits/msvc_win32.h on non-MSVC compilers if (NOT "${CMAKE_CXX_COMPILER_ID}" STREQUAL "MSVC") list(FILTER internal_headers EXCLUDE REGEX "__limits/msvc_win32.h") endif() # generated cuda::ptx headers are not standalone list(FILTER internal_headers EXCLUDE REGEX "__ptx/instructions/generated") function(libcudacxx_create_internal_header_test header_name, headertest_src, fallback) if(fallback) set(header_name "${header_name}_fallback") endif() set(headertest_${header_name} verify_${header_name}) # cuda internal headers should not be tested against a host compiler # NOTE(HIP/AMD): PTX tests are skipped for AMD devices. if(header_name MATCHES "__ptx.*") return() endif() if(header_name MATCHES "__memory_resource.*") return() endif() # cuda internal headers should not be tested against a host compiler add_library(headertest_${header_name} SHARED "${headertest_src}.cpp") set_source_files_properties("${headertest_src}.cpp" PROPERTIES LANGUAGE HIP) target_include_directories(headertest_${header_name} PRIVATE "${libhipcxx_SOURCE_DIR}/include") target_compile_options(headertest_${header_name} PRIVATE $<$:${headertest_warning_levels_device}> $<$:${headertest_warning_levels_host}> -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE) target_compile_definitions(headertest_${header_name} PRIVATE CCCL_ENABLE_ASSERTIONS) target_compile_definitions(headertest_${header_name} PRIVATE CCCL_ENABLE_OPTIONAL_REF) # if ("NVHPC" STREQUAL "${CMAKE_CXX_COMPILER_ID}") # target_link_libraries(headertest_${header_name} NVHPC::CUDART) # else() # target_link_libraries(headertest_${header_name} CUDA::cudart) # endif() target_link_libraries(headertest_${header_name} hip::host) if(fallback) target_compile_definitions(headertest_${header_name} PRIVATE "-D${fallback}") endif() target_compile_definitions(headertest_${header_name} PRIVATE "-DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE") add_dependencies(libcudacxx.test.internal_headers headertest_${header_name}) endfunction() function(libcudacxx_add_internal_header_test header) # ${header} contains the "/" from the subfolder, replace by "_" for actual names string(REPLACE "/" "_" header_name "${header}") # NOTE(HIP/AMD): Skip unsupported headers for AMD devices (no support as of November 2025). string(REGEX MATCH "ptx|barrier|latch|semaphore|annotated_ptr|pipeline|memcpy_async" match "${header}") if(match) return() endif() # Create the source file for the header target from the template and add the file to the global project set(headertest_src "headers/${header_name}") configure_file("${CMAKE_CURRENT_SOURCE_DIR}/header_test.cpp.in" "${headertest_src}.cpp") # Create the default target for that file libcudacxx_create_internal_header_test(${header_name}, ${headertest_src}, "") # MSVC cannot handle some of the fallbacks if ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "MSVC") if("${header}" MATCHES "is_base_of" OR "${header}" MATCHES "is_nothrow_destructible" OR "${header}" MATCHES "is_polymorphic") return() endif() endif() # Search the file for a fallback definition file(READ ${libhipcxx_SOURCE_DIR}/include/cuda/${header} header_file) string(REGEX MATCH "_LIBCUDACXX_[A-Z_]*_FALLBACK" fallback "${header_file}") if(fallback) libcudacxx_create_internal_header_test(${header_name}, ${headertest_src}, ${fallback}) endif() endfunction() foreach(header IN LISTS internal_headers) libcudacxx_add_internal_header_test(${header}) endforeach()