Merge remote-tracking branch 'upstream/dev' into dev

Author: PointKernel

CMakeLists.txt

@@ -98,7 +98,10 @@ include(${CMAKE_CURRENT_LIST_DIR}/cmake/roaring_testdata.cmake)
 # - optionally build tests ------------------------------------------------------------------------
 
 if(BUILD_TESTS)
+    include(CTest)
     add_subdirectory(tests)
+    # Include header tests as part of the regular test suite
+    include(cmake/header_testing.cmake)
 endif(BUILD_TESTS)
 
 ###################################################################################################
@@ -119,6 +122,7 @@ if(BUILD_EXAMPLES)
     add_subdirectory(examples)
 endif(BUILD_EXAMPLES)
 
+
 ###################################################################################################
 # - Install targets -------------------------------------------------------------------------------
 

cmake/header_testing.cmake

@@ -0,0 +1,86 @@
+#=============================================================================
+# Copyright (c) 2025, NVIDIA CORPORATION.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+#=============================================================================
+
+# For every public header, build a translation unit containing `#include <header>`
+# to let the compiler try to figure out warnings in that header if it is not otherwise
+# included in tests, and also to verify if the headers are modular enough.
+# .inl files are not globbed for, because they are not supposed to be used as public
+# entrypoints.
+
+function(cuco_add_header_tests)
+  file(GLOB_RECURSE headers
+    RELATIVE "${CUCO_SOURCE_DIR}/include"
+    CONFIGURE_DEPENDS
+    "${CUCO_SOURCE_DIR}/include/cuco/*.cuh"
+    "${CUCO_SOURCE_DIR}/include/cuco/*.hpp"
+  )
+  
+  list(LENGTH headers headers_count)
+  message(STATUS "Found ${headers_count} headers for testing")
+
+  # List of headers that have known issues or are not meant to be included directly
+  set(excluded_headers
+    # Add any headers that should be excluded from testing here
+    # Example: cuco/internal_header.cuh
+  )
+  
+  # Remove excluded headers
+  if(excluded_headers)
+    list(REMOVE_ITEM headers ${excluded_headers})
+  endif()
+
+  foreach (header IN LISTS headers)
+    # Create a safe target name by replacing path separators and dots
+    string(REPLACE "/" "_" header_target_name "${header}")
+    string(REPLACE "." "_" header_target_name "${header_target_name}")
+    # Use a hash to ensure uniqueness in case of similar names
+    string(MD5 header_hash "${header}")
+    string(SUBSTRING "${header_hash}" 0 8 header_hash_short)
+    set(headertest_target "cuco_header_${header_target_name}_${header_hash_short}")
+    
+    set(header_src "${CMAKE_CURRENT_BINARY_DIR}/headers/${headertest_target}/${header}.cu")
+    
+    # Create the directory if it doesn't exist
+    get_filename_component(header_dir "${header_src}" DIRECTORY)
+    file(MAKE_DIRECTORY "${header_dir}")
+    
+    # Write simple test file that includes the header
+    file(WRITE "${header_src}" "#include <${header}>\nint main() { return 0; }\n")
+
+    # Create executable test for this specific header
+    add_executable(${headertest_target} ${header_src})
+    target_link_libraries(${headertest_target} PRIVATE cuco::cuco CUDA::cudart)
+
+    target_compile_options(${headertest_target} PRIVATE
+      $<$<COMPILE_LANGUAGE:CUDA>:--expt-extended-lambda>
+      --compiler-options=-Wall --compiler-options=-Wextra
+      --compiler-options=-Werror -Wno-deprecated-gpu-targets
+    )
+
+    if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
+      target_compile_options(${headertest_target} PRIVATE -Xcompiler -Wno-subobject-linkage)
+    endif()
+
+    set_target_properties(${headertest_target} PROPERTIES
+      RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/tests/headers"
+    )
+
+    # Add as a CTest test
+    add_test(NAME ${headertest_target} COMMAND ${headertest_target})
+  endforeach()
+endfunction()
+
+cuco_add_header_tests()

include/cuco/detail/bloom_filter/bloom_filter_impl.cuh

@@ -20,6 +20,7 @@
 #include <cuco/detail/error.hpp>
 #include <cuco/detail/utility/cuda.cuh>
 #include <cuco/detail/utility/cuda.hpp>
+#include <cuco/detail/utility/math.cuh>
 #include <cuco/detail/utils.hpp>
 #include <cuco/utility/cuda_thread_scope.cuh>
 

include/cuco/detail/bloom_filter/kernels.cuh

@@ -16,6 +16,7 @@
 #pragma once
 
 #include <cuco/detail/utility/cuda.cuh>
+#include <cuco/detail/utility/math.cuh>
 
 #include <cuda/std/iterator>
 

include/cuco/detail/hash_functions/utils.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,6 +17,7 @@
 #pragma once
 
 #include <cuda/std/cstddef>
+#include <cuda/std/cstdint>
 
 namespace cuco::detail {
 
@@ -29,12 +30,14 @@ constexpr __host__ __device__ T load_chunk(U const* const data, Extent index) no
   return chunk;
 }
 
-constexpr __host__ __device__ std::uint32_t rotl32(std::uint32_t x, std::int8_t r) noexcept
+constexpr __host__ __device__ cuda::std::uint32_t rotl32(cuda::std::uint32_t x,
+                                                         cuda::std::int8_t r) noexcept
 {
   return (x << r) | (x >> (32 - r));
 }
 
-constexpr __host__ __device__ std::uint64_t rotl64(std::uint64_t x, std::int8_t r) noexcept
+constexpr __host__ __device__ cuda::std::uint64_t rotl64(cuda::std::uint64_t x,
+                                                         cuda::std::int8_t r) noexcept
 {
   return (x << r) | (x >> (64 - r));
 }

include/cuco/detail/pair/tuple_helpers.inl

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,8 +26,8 @@ struct tuple_size<volatile cuco::pair<T1, T2>> : tuple_size<cuco::pair<T1, T2>>
 template <typename T1, typename T2>
 struct tuple_size<const volatile cuco::pair<T1, T2>> : tuple_size<cuco::pair<T1, T2>> {};
 
-template <std::size_t I, typename T1, typename T2>
-struct tuple_element<I, cuco::pair<T1, T2>> {
+template <std::size_t Index, typename T1, typename T2>
+struct tuple_element<Index, cuco::pair<T1, T2>> {
   using type = void;
 };
 
@@ -61,48 +61,48 @@ template <typename T1, typename T2>
 struct tuple_element<1, const volatile cuco::pair<T1, T2>> : tuple_element<1, cuco::pair<T1, T2>> {
 };
 
-template <std::size_t I, typename T1, typename T2>
+template <std::size_t Index, typename T1, typename T2>
 __host__ __device__ constexpr auto get(cuco::pair<T1, T2>& p) ->
-  typename tuple_element<I, cuco::pair<T1, T2>>::type&
+  typename tuple_element<Index, cuco::pair<T1, T2>>::type&
 {
-  static_assert(I < 2);
-  if constexpr (I == 0) {
+  static_assert(Index < 2);
+  if constexpr (Index == 0) {
     return p.first;
   } else {
     return p.second;
   }
 }
 
-template <std::size_t I, typename T1, typename T2>
+template <std::size_t Index, typename T1, typename T2>
 __host__ __device__ constexpr auto get(cuco::pair<T1, T2>&& p) ->
-  typename tuple_element<I, cuco::pair<T1, T2>>::type&&
+  typename tuple_element<Index, cuco::pair<T1, T2>>::type&&
 {
-  static_assert(I < 2);
-  if constexpr (I == 0) {
+  static_assert(Index < 2);
+  if constexpr (Index == 0) {
     return cuda::std::move(p.first);
   } else {
     return cuda::std::move(p.second);
   }
 }
 
-template <std::size_t I, typename T1, typename T2>
+template <std::size_t Index, typename T1, typename T2>
 __host__ __device__ constexpr auto get(cuco::pair<T1, T2> const& p) ->
-  typename tuple_element<I, cuco::pair<T1, T2>>::type const&
+  typename tuple_element<Index, cuco::pair<T1, T2>>::type const&
 {
-  static_assert(I < 2);
-  if constexpr (I == 0) {
+  static_assert(Index < 2);
+  if constexpr (Index == 0) {
     return p.first;
   } else {
     return p.second;
   }
 }
 
-template <std::size_t I, typename T1, typename T2>
+template <std::size_t Index, typename T1, typename T2>
 __host__ __device__ constexpr auto get(cuco::pair<T1, T2> const&& p) ->
-  typename tuple_element<I, cuco::pair<T1, T2>>::type const&&
+  typename tuple_element<Index, cuco::pair<T1, T2>>::type const&&
 {
-  static_assert(I < 2);
-  if constexpr (I == 0) {
+  static_assert(Index < 2);
+  if constexpr (Index == 0) {
     return cuda::std::move(p.first);
   } else {
     return cuda::std::move(p.second);

include/cuco/detail/roaring_bitmap/util.cuh

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <cuco/detail/error.hpp>
 #include <cuco/utility/traits.hpp>
 
 #include <cuda/std/cstddef>

include/cuco/detail/static_map/helpers.cuh

@@ -17,6 +17,9 @@
 
 #include <cuco/detail/bitwise_compare.cuh>
 #include <cuco/detail/utility/cuda.cuh>
+#include <cuco/operator.hpp>
+#include <cuco/pair.cuh>
+#include <cuco/types.cuh>
 
 #include <cub/block/block_reduce.cuh>
 #include <cuda/atomic>