Use static_for for fixed-size loops (#771)

PointKernel · web-flow · commit b8ef0bf1265b · 2025-12-05T10:41:08.000-08:00
Closes #770 This PR replaces naive for loops with `cuda::static_for` in cases where the loop size is known at compile time, enabling guaranteed compile-time unrolling and improving runtime performance. It also updates CMake to treat device compiler warnings as errors and addresses warnings for variables like `gen` being declared but unused, which appears to be a compiler bug in CUDA 12.0.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -94,6 +94,30 @@ target_compile_features(cuco INTERFACE cxx_std_17 cuda_std_17)
 option(CUCO_DOWNLOAD_ROARING_TESTDATA "Download RoaringFormatSpec test data" ON)
 include(${CMAKE_CURRENT_LIST_DIR}/cmake/roaring_testdata.cmake)
 
+###################################################################################################
+# - common compile options function ---------------------------------------------------------------
+
+function(cuco_set_common_compile_options target_name)
+    # Parse optional arguments
+    cmake_parse_arguments(CUCO_OPTS "ADD_LINEINFO" "" "" ${ARGN})
+    
+    # Base compile options common to all targets
+    target_compile_options(${target_name} PRIVATE 
+        --compiler-options=-Wall --compiler-options=-Wextra --compiler-options=-Werror 
+        -Wno-deprecated-gpu-targets --expt-extended-lambda -Werror=all-warnings
+    )
+    
+    # Add lineinfo option if requested (typically for benchmarks)
+    if(CUCO_OPTS_ADD_LINEINFO)
+        target_compile_options(${target_name} PRIVATE -lineinfo)
+    endif()
+    
+    # Add GCC-specific warning suppression only for GCC
+    if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
+        target_compile_options(${target_name} PRIVATE -Xcompiler -Wno-subobject-linkage)
+    endif()
+endfunction()
+
 ###################################################################################################
 # - optionally build tests ------------------------------------------------------------------------
 
diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt
@@ -27,12 +27,7 @@ function(ConfigureBench BENCH_NAME)
                                         RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/benchmarks")
     target_include_directories(${BENCH_NAME} PRIVATE
                                              "${CMAKE_CURRENT_SOURCE_DIR}")
-    target_compile_options(${BENCH_NAME} PRIVATE --compiler-options=-Wall --compiler-options=-Wextra
-      --compiler-options=-Werror -Wno-deprecated-gpu-targets --expt-extended-lambda -lineinfo)
-    # Add GCC-specific warning suppression only for GCC
-    if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
-      target_compile_options(${BENCH_NAME} PRIVATE -Xcompiler -Wno-subobject-linkage)
-    endif()
+    cuco_set_common_compile_options(${BENCH_NAME} ADD_LINEINFO)
     target_link_libraries(${BENCH_NAME} PRIVATE
                                         nvbench::main
                                         pthread
diff --git a/benchmarks/dynamic_map/contains_bench.cu b/benchmarks/dynamic_map/contains_bench.cu
@@ -43,7 +43,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_contains(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/dynamic_map/erase_bench.cu b/benchmarks/dynamic_map/erase_bench.cu
@@ -43,7 +43,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_erase(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/dynamic_map/find_bench.cu b/benchmarks/dynamic_map/find_bench.cu
@@ -43,7 +43,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_find(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/dynamic_map/insert_bench.cu b/benchmarks/dynamic_map/insert_bench.cu
@@ -45,7 +45,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_insert(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/dynamic_map/retrieve_all_bench.cu b/benchmarks/dynamic_map/retrieve_all_bench.cu
@@ -42,7 +42,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_retrieve_all(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/hyperloglog/hyperloglog_bench.cu b/benchmarks/hyperloglog/hyperloglog_bench.cu
@@ -58,7 +58,7 @@ template <class Estimator, class Dist>
 
   thrust::device_vector<T> items(num_items);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   Estimator estimator{cuco::sketch_size_kb(sketch_size_kb)};
   double error_sum = 0;
   for (std::size_t i = 0; i < num_samples; ++i) {
@@ -97,7 +97,7 @@ void hyperloglog_e2e(nvbench::state& state, nvbench::type_list<T, Dist>)
 
   thrust::device_vector<T> items(num_items);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), items.begin(), items.end());
 
   estimator_type estimator{cuco::sketch_size_kb(sketch_size_kb)};
@@ -126,7 +126,7 @@ void hyperloglog_add(nvbench::state& state, nvbench::type_list<T, Dist>)
 
   thrust::device_vector<T> items(num_items);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), items.begin(), items.end());
 
   state.add_element_count(num_items);
diff --git a/benchmarks/roaring_bitmap/contains_bench.cu b/benchmarks/roaring_bitmap/contains_bench.cu
@@ -55,7 +55,7 @@ void roaring_bitmap_contains(nvbench::state& state, nvbench::type_list<T>)
 
   thrust::device_vector<T> items(num_items);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(distribution::unique{}, items.begin(), items.end());
 
   thrust::device_vector<bool> contained(items.size(), false);
diff --git a/benchmarks/static_map/contains_bench.cu b/benchmarks/static_map/contains_bench.cu
@@ -45,7 +45,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_contains(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_map/erase_bench.cu b/benchmarks/static_map/erase_bench.cu
@@ -45,7 +45,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_erase(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_map/find_bench.cu b/benchmarks/static_map/find_bench.cu
@@ -45,7 +45,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_find(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_map/insert_bench.cu b/benchmarks/static_map/insert_bench.cu
@@ -44,7 +44,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_insert(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_map/insert_or_apply_bench.cu b/benchmarks/static_map/insert_or_apply_bench.cu
@@ -46,7 +46,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_insert_or_appl
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_multimap/count_bench.cu b/benchmarks/static_multimap/count_bench.cu
@@ -45,7 +45,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_count(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_multimap/insert_bench.cu b/benchmarks/static_multimap/insert_bench.cu
@@ -44,7 +44,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_insert(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_multimap/query_bench.cu b/benchmarks/static_multimap/query_bench.cu
@@ -45,7 +45,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_query(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_multimap/retrieve_bench.cu b/benchmarks/static_multimap/retrieve_bench.cu
@@ -45,7 +45,7 @@ std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_retrieve(
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   thrust::device_vector<pair_type> pairs(num_keys);
diff --git a/benchmarks/static_multiset/contains_bench.cu b/benchmarks/static_multiset/contains_bench.cu
@@ -41,7 +41,7 @@ void static_multiset_contains(nvbench::state& state, nvbench::type_list<Key, Dis
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   cuco::static_multiset<Key> set{size, cuco::empty_key<Key>{-1}};
diff --git a/benchmarks/static_multiset/count_bench.cu b/benchmarks/static_multiset/count_bench.cu
@@ -42,7 +42,7 @@ void static_multiset_count(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   auto map = cuco::static_multiset{size, cuco::empty_key<Key>{-1}};
diff --git a/benchmarks/static_multiset/find_bench.cu b/benchmarks/static_multiset/find_bench.cu
@@ -42,7 +42,7 @@ void static_multiset_find(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   cuco::static_multiset<Key> set{size, cuco::empty_key<Key>{-1}};
diff --git a/benchmarks/static_multiset/insert_bench.cu b/benchmarks/static_multiset/insert_bench.cu
@@ -40,7 +40,7 @@ void static_multiset_insert(nvbench::state& state, nvbench::type_list<Key, Dist>
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   state.add_element_count(num_keys);
diff --git a/benchmarks/static_multiset/retrieve_bench.cu b/benchmarks/static_multiset/retrieve_bench.cu
@@ -42,7 +42,7 @@ void static_multiset_retrieve(nvbench::state& state, nvbench::type_list<Key, Dis
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   gen.dropout(keys.begin(), keys.end(), matching_rate);
diff --git a/benchmarks/static_set/contains_bench.cu b/benchmarks/static_set/contains_bench.cu
@@ -41,7 +41,7 @@ void static_set_contains(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   cuco::static_set<Key> set{size, cuco::empty_key<Key>{-1}};
diff --git a/benchmarks/static_set/find_bench.cu b/benchmarks/static_set/find_bench.cu
@@ -42,7 +42,7 @@ void static_set_find(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   cuco::static_set<Key> set{size, cuco::empty_key<Key>{-1}};
diff --git a/benchmarks/static_set/insert_bench.cu b/benchmarks/static_set/insert_bench.cu
@@ -40,7 +40,7 @@ void static_set_insert(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   state.add_element_count(num_keys);
diff --git a/benchmarks/static_set/rehash_bench.cu b/benchmarks/static_set/rehash_bench.cu
@@ -40,7 +40,7 @@ void static_set_rehash(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   thrust::device_vector<Key> keys(num_keys);  // slots per second
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   state.add_element_count(capacity);
diff --git a/benchmarks/static_set/retrieve_all_bench.cu b/benchmarks/static_set/retrieve_all_bench.cu
@@ -40,7 +40,7 @@ void static_set_retrieve_all(nvbench::state& state, nvbench::type_list<Key, Dist
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   cuco::static_set<Key> set{size, cuco::empty_key<Key>{-1}};
diff --git a/benchmarks/static_set/retrieve_bench.cu b/benchmarks/static_set/retrieve_bench.cu
@@ -42,7 +42,7 @@ void static_set_retrieve(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   gen.dropout(keys.begin(), keys.end(), matching_rate);
diff --git a/benchmarks/static_set/size_bench.cu b/benchmarks/static_set/size_bench.cu
@@ -40,7 +40,7 @@ void static_set_size(nvbench::state& state, nvbench::type_list<Key, Dist>)
 
   thrust::device_vector<Key> keys(num_keys);
 
-  key_generator gen{};
+  [[maybe_unused]] key_generator gen{};
   gen.generate(dist_from_state<Dist>(state), keys.begin(), keys.end());
 
   state.add_element_count(num_keys);
diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
@@ -24,12 +24,7 @@ function(ConfigureExample EXAMPLE_NAME EXAMPLE_SRC)
                                           RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/examples")
     target_include_directories(${EXAMPLE_NAME} PRIVATE
                                              "${CMAKE_CURRENT_SOURCE_DIR}")
-    target_compile_options(${EXAMPLE_NAME} PRIVATE --compiler-options=-Wall --compiler-options=-Wextra
-      --compiler-options=-Werror -Wno-deprecated-gpu-targets --expt-extended-lambda)
-    # Add GCC-specific warning suppression only for GCC
-    if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
-      target_compile_options(${EXAMPLE_NAME} PRIVATE -Xcompiler -Wno-subobject-linkage)
-    endif()
+    cuco_set_common_compile_options(${EXAMPLE_NAME})
     target_link_libraries(${EXAMPLE_NAME} PRIVATE cuco CUDA::cudart)
 endfunction(ConfigureExample)
 
diff --git a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh
diff --git a/include/cuco/detail/hyperloglog/kernels.cuh b/include/cuco/detail/hyperloglog/kernels.cuh
diff --git a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
diff --git a/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh b/include/cuco/detail/roaring_bitmap/roaring_bitmap_impl.cuh
diff --git a/include/cuco/detail/static_map/static_map_ref.inl b/include/cuco/detail/static_map/static_map_ref.inl
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
