Skip to content

Use static_for for fixed-size loops #771

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 3 commits into
base: dev
Choose a base branch
from
Open

Use static_for for fixed-size loops #771

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
fabc460
Use static_for for fixed-size loops
PointKernel Oct 23, 2025
634c671
Use ftors instead of lambda to fix 12.0 warnings
PointKernel Oct 24, 2025
17ec7ec
CMake cleanups
PointKernel Oct 24, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 24 additions & 0 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 ------------------------------------------------------------------------

Expand Down
7 changes: 1 addition & 6 deletions benchmarks/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/dynamic_map/contains_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/dynamic_map/erase_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/dynamic_map/find_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/dynamic_map/insert_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/dynamic_map/retrieve_all_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
6 changes: 3 additions & 3 deletions benchmarks/hyperloglog/hyperloglog_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down Expand Up @@ -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)};
Expand Down Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/roaring_bitmap/contains_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_map/contains_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_map/erase_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_map/find_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_map/insert_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_map/insert_or_apply_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multimap/count_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multimap/insert_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multimap/query_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multimap/retrieve_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multiset/contains_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}};
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multiset/count_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}};
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multiset/find_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}};
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multiset/insert_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_multiset/retrieve_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_set/contains_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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());
Comment on lines +44 to 45
Copy link
Collaborator

@sleeepyjack sleeepyjack Oct 30, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

How can gen be maybe_unused when it is used immediately in the following line? What am I missing?

Copy link
Member Author

@PointKernel PointKernel Nov 3, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That’s a compiler bug in CUDA 12.0 where host-side use of host-device APIs isn’t correctly recognized.


cuco::static_set<Key> set{size, cuco::empty_key<Key>{-1}};
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_set/find_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}};
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_set/insert_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_set/rehash_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_set/retrieve_all_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}};
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_set/retrieve_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/static_set/size_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
7 changes: 1 addition & 6 deletions examples/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)

Expand Down
Loading
Loading