diff --git a/CMakeLists.txt b/CMakeLists.txt index f3ca85a8a..25fb514d9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -15,6 +15,9 @@ #============================================================================= cmake_minimum_required(VERSION 3.23.1 FATAL_ERROR) +set(rapids-cmake-repo bdice/rapids-cmake) +set(rapids-cmake-branch cccl-update-2.1.0) + if(NOT EXISTS ${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake) file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-23.10/RAPIDS.cmake ${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake) diff --git a/examples/static_map/custom_type_example.cu b/examples/static_map/custom_type_example.cu index e150a858e..7430fda84 100644 --- a/examples/static_map/custom_type_example.cu +++ b/examples/static_map/custom_type_example.cu @@ -22,6 +22,8 @@ #include #include +#include + // User-defined key type #if !defined(CUCO_HAS_INDEPENDENT_THREADS) struct custom_key_type { @@ -88,7 +90,8 @@ int main(void) // Create an iterator of input key/value pairs auto pairs_begin = thrust::make_transform_iterator( thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::make_pair(custom_key_type{i}, custom_value_type{i}); }); + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::make_pair(custom_key_type{i}, custom_value_type{i}); })); // Construct a map with 100,000 slots using the given empty key/value sentinels. Note the // capacity is chosen knowing we will insert 80,000 keys, for an load factor of 80%. @@ -101,7 +104,8 @@ int main(void) // Reproduce inserted keys auto insert_keys = thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [] __device__(auto i) { return custom_key_type{i}; }); + cuda::proclaim_return_type( + [] __device__(auto i) { return custom_key_type{i}; })); thrust::device_vector contained(num_pairs); @@ -114,7 +118,9 @@ int main(void) // All inserted keys are contained assert( - thrust::all_of(contained.begin(), contained.end(), [] __device__(auto const& b) { return b; })); + thrust::all_of(contained.begin(), + contained.end(), + cuda::proclaim_return_type([] __device__(auto const& b) { return b; }))); return 0; } diff --git a/tests/dynamic_map/unique_sequence_test.cu b/tests/dynamic_map/unique_sequence_test.cu index aa01ca51a..d9364f341 100644 --- a/tests/dynamic_map/unique_sequence_test.cu +++ b/tests/dynamic_map/unique_sequence_test.cu @@ -27,6 +27,8 @@ #include #include +#include + #include TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", @@ -48,9 +50,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); thrust::sequence(thrust::device, d_values.begin(), d_values.end()); - auto pairs_begin = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i, i); }); + auto pairs_begin = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); thrust::device_vector d_results(num_keys); thrust::device_vector d_contained(num_keys); @@ -62,9 +65,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", map.find(d_keys.begin(), d_keys.end(), d_results.begin()); auto zip = thrust::make_zip_iterator(thrust::make_tuple(d_results.begin(), d_values.begin())); - REQUIRE(cuco::test::all_of(zip, zip + num_keys, [] __device__(auto const& p) { - return thrust::get<0>(p) == thrust::get<1>(p); - })); + REQUIRE(cuco::test::all_of( + zip, zip + num_keys, cuda::proclaim_return_type([] __device__(auto const& p) { + return thrust::get<0>(p) == thrust::get<1>(p); + }))); } SECTION("All non-inserted keys-value pairs should have the empty sentinel value recovered") @@ -72,7 +76,9 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", map.find(d_keys.begin(), d_keys.end(), d_results.begin()); REQUIRE(cuco::test::all_of( - d_results.begin(), d_results.end(), [] __device__(auto const& p) { return p == -1; })); + d_results.begin(), + d_results.end(), + cuda::proclaim_return_type([] __device__(auto const& p) { return p == -1; }))); } SECTION("All inserted keys-value pairs should be contained") diff --git a/tests/static_map/custom_type_test.cu b/tests/static_map/custom_type_test.cu index e23216ca3..fbba8ec3b 100644 --- a/tests/static_map/custom_type_test.cu +++ b/tests/static_map/custom_type_test.cu @@ -25,6 +25,8 @@ #include #include +#include + #include #include @@ -123,17 +125,18 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", thrust::counting_iterator(0), thrust::counting_iterator(num), insert_keys.begin(), - [] __device__(auto i) { return Key{i}; }); + cuda::proclaim_return_type([] __device__(auto i) { return Key{i}; })); thrust::transform(thrust::device, thrust::counting_iterator(0), thrust::counting_iterator(num), insert_values.begin(), - [] __device__(auto i) { return Value{i}; }); + cuda::proclaim_return_type([] __device__(auto i) { return Value{i}; })); - auto insert_pairs = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i, i); }); + auto insert_pairs = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); SECTION("All inserted keys-value pairs should be correctly recovered during find") { @@ -151,9 +154,9 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", REQUIRE(cuco::test::equal(insert_values.begin(), insert_values.end(), found_values.begin(), - [] __device__(Value lhs, Value rhs) { + cuda::proclaim_return_type([] __device__(Value lhs, Value rhs) { return std::tie(lhs.f, lhs.s) == std::tie(rhs.f, rhs.s); - })); + }))); } SECTION("All inserted keys-value pairs should be contained") @@ -175,7 +178,7 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", insert_pairs, insert_pairs + num, thrust::counting_iterator(0), - [] __device__(auto const& key) { return (key % 2) == 0; }, + cuda::proclaim_return_type([] __device__(auto const& key) { return (key % 2) == 0; }), hash_custom_key{}, custom_key_equals{}); @@ -187,12 +190,13 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", hash_custom_key{}, custom_key_equals{}); - REQUIRE(cuco::test::equal(contained.begin(), - contained.end(), - thrust::counting_iterator(0), - [] __device__(auto const& idx_contained, auto const& idx) { - return ((idx % 2) == 0) == idx_contained; - })); + REQUIRE(cuco::test::equal( + contained.begin(), + contained.end(), + thrust::counting_iterator(0), + cuda::proclaim_return_type([] __device__(auto const& idx_contained, auto const& idx) { + return ((idx % 2) == 0) == idx_contained; + }))); } SECTION("Non-inserted keys-value pairs should not be contained") @@ -212,9 +216,11 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", map.insert(insert_pairs, insert_pairs + num, hash_custom_key{}, custom_key_equals{}); auto view = map.get_device_view(); REQUIRE(cuco::test::all_of( - insert_pairs, insert_pairs + num, [view] __device__(cuco::pair const& pair) { + insert_pairs, + insert_pairs + num, + cuda::proclaim_return_type([view] __device__(cuco::pair const& pair) { return view.contains(pair.first, hash_custom_key{}, custom_key_equals{}); - })); + }))); } SECTION("Inserting unique keys should return insert success.") @@ -222,9 +228,11 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", auto m_view = map.get_device_mutable_view(); REQUIRE(cuco::test::all_of(insert_pairs, insert_pairs + num, - [m_view] __device__(cuco::pair const& pair) mutable { - return m_view.insert(pair, hash_custom_key{}, custom_key_equals{}); - })); + cuda::proclaim_return_type( + [m_view] __device__(cuco::pair const& pair) mutable { + return m_view.insert( + pair, hash_custom_key{}, custom_key_equals{}); + }))); } SECTION("Cannot find any key in an empty hash map") @@ -235,18 +243,21 @@ TEMPLATE_TEST_CASE_SIG("User defined key and value type", REQUIRE(cuco::test::all_of( insert_pairs, insert_pairs + num, - [view] __device__(cuco::pair const& pair) mutable { - return view.find(pair.first, hash_custom_key{}, custom_key_equals{}) == view.end(); - })); + cuda::proclaim_return_type( + [view] __device__(cuco::pair const& pair) mutable { + return view.find(pair.first, hash_custom_key{}, custom_key_equals{}) == view.end(); + }))); } SECTION("const view") { auto const view = map.get_device_view(); REQUIRE(cuco::test::all_of( - insert_pairs, insert_pairs + num, [view] __device__(cuco::pair const& pair) { + insert_pairs, + insert_pairs + num, + cuda::proclaim_return_type([view] __device__(cuco::pair const& pair) { return view.find(pair.first, hash_custom_key{}, custom_key_equals{}) == view.end(); - })); + }))); } } } diff --git a/tests/static_map/duplicate_keys_test.cu b/tests/static_map/duplicate_keys_test.cu index 5620fa4e9..64f9d35cf 100644 --- a/tests/static_map/duplicate_keys_test.cu +++ b/tests/static_map/duplicate_keys_test.cu @@ -27,6 +27,8 @@ #include #include +#include + #include TEMPLATE_TEST_CASE_SIG("Duplicate keys", @@ -49,7 +51,8 @@ TEMPLATE_TEST_CASE_SIG("Duplicate keys", auto pairs_begin = thrust::make_transform_iterator( thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i / 2, i / 2); }); + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i / 2, i / 2); })); thrust::device_vector d_results(num_keys); thrust::device_vector d_contained(num_keys); diff --git a/tests/static_map/heterogeneous_lookup_test.cu b/tests/static_map/heterogeneous_lookup_test.cu index e842612b1..260dc34cb 100644 --- a/tests/static_map/heterogeneous_lookup_test.cu +++ b/tests/static_map/heterogeneous_lookup_test.cu @@ -25,6 +25,8 @@ #include #include +#include + #include #include @@ -99,11 +101,13 @@ TEMPLATE_TEST_CASE("Heterogeneous lookup", cuco::static_map map{ capacity, cuco::empty_key{sentinel_key}, cuco::empty_value{sentinel_value}}; - auto insert_pairs = - thrust::make_transform_iterator(thrust::counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i, i); }); - auto probe_keys = thrust::make_transform_iterator(thrust::counting_iterator(0), - [] __device__(auto i) { return ProbeKey(i); }); + auto insert_pairs = thrust::make_transform_iterator( + thrust::counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); + auto probe_keys = thrust::make_transform_iterator( + thrust::counting_iterator(0), + cuda::proclaim_return_type([] __device__(auto i) { return ProbeKey(i); })); SECTION("All inserted keys-value pairs should be contained") { diff --git a/tests/static_map/key_sentinel_test.cu b/tests/static_map/key_sentinel_test.cu index 74a1badd1..349df62dd 100644 --- a/tests/static_map/key_sentinel_test.cu +++ b/tests/static_map/key_sentinel_test.cu @@ -22,6 +22,8 @@ #include #include +#include + #include #define SIZE 10 @@ -51,19 +53,21 @@ TEMPLATE_TEST_CASE_SIG( } CUCO_CUDA_TRY(cudaMemcpyToSymbol(A, h_A, SIZE * sizeof(int))); - auto pairs_begin = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i, i); }); + auto pairs_begin = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); SECTION( "Tests of non-CG insert: The custom `key_equal` can never be used to compare against sentinel") { - REQUIRE(cuco::test::all_of(pairs_begin, - pairs_begin + num_keys, - [m_view] __device__(cuco::pair const& pair) mutable { - return m_view.insert( - pair, cuco::default_hash_function{}, custom_equals{}); - })); + REQUIRE(cuco::test::all_of( + pairs_begin, + pairs_begin + num_keys, + cuda::proclaim_return_type( + [m_view] __device__(cuco::pair const& pair) mutable { + return m_view.insert(pair, cuco::default_hash_function{}, custom_equals{}); + }))); } SECTION( @@ -75,10 +79,12 @@ TEMPLATE_TEST_CASE_SIG( custom_equals{}); // All keys inserted via custom `key_equal` should be found REQUIRE(cuco::test::all_of( - pairs_begin, pairs_begin + num_keys, [view] __device__(cuco::pair const& pair) { + pairs_begin, + pairs_begin + num_keys, + cuda::proclaim_return_type([view] __device__(cuco::pair const& pair) { auto const found = view.find(pair.first); return (found != view.end()) and (found->first.load() == pair.first and found->second.load() == pair.second); - })); + }))); } } diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu index 444f1c7e7..351ef9f6a 100644 --- a/tests/static_map/shared_memory_test.cu +++ b/tests/static_map/shared_memory_test.cu @@ -25,6 +25,8 @@ #include #include +#include + #include #include @@ -127,9 +129,11 @@ TEMPLATE_TEST_CASE_SIG("Shared memory static map", auto zip = thrust::make_zip_iterator( thrust::make_tuple(d_keys_exist.begin(), d_keys_and_values_correct.begin())); - REQUIRE(cuco::test::all_of(zip, zip + d_keys_exist.size(), [] __device__(auto const& z) { - return thrust::get<0>(z) and thrust::get<1>(z); - })); + REQUIRE(cuco::test::all_of(zip, + zip + d_keys_exist.size(), + cuda::proclaim_return_type([] __device__(auto const& z) { + return thrust::get<0>(z) and thrust::get<1>(z); + }))); } SECTION("No key is found before insertion.") diff --git a/tests/static_map/stream_test.cu b/tests/static_map/stream_test.cu index 6121cbd62..7344ab785 100644 --- a/tests/static_map/stream_test.cu +++ b/tests/static_map/stream_test.cu @@ -27,6 +27,8 @@ #include #include +#include + #include TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream", @@ -53,9 +55,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream", thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); thrust::sequence(thrust::device, d_values.begin(), d_values.end()); - auto pairs_begin = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i, i); }); + auto pairs_begin = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); auto hash_fn = cuco::default_hash_function{}; auto equal_fn = thrust::equal_to{}; @@ -69,11 +72,12 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream", map.find(d_keys.begin(), d_keys.end(), d_results.begin(), hash_fn, equal_fn, stream); auto zip = thrust::make_zip_iterator(thrust::make_tuple(d_results.begin(), d_values.begin())); - REQUIRE(cuco::test::all_of( - zip, - zip + num_keys, - [] __device__(auto const& p) { return thrust::get<0>(p) == thrust::get<1>(p); }, - stream)); + REQUIRE(cuco::test::all_of(zip, + zip + num_keys, + cuda::proclaim_return_type([] __device__(auto const& p) { + return thrust::get<0>(p) == thrust::get<1>(p); + }), + stream)); } SECTION("All inserted keys-value pairs should be contained") diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/unique_sequence_test.cu index 6a0165cc2..0cc8c8fc5 100644 --- a/tests/static_map/unique_sequence_test.cu +++ b/tests/static_map/unique_sequence_test.cu @@ -29,6 +29,8 @@ #include #include +#include + #include TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", @@ -52,9 +54,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); thrust::sequence(thrust::device, d_values.begin(), d_values.end()); - auto pairs_begin = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i, i); }); + auto pairs_begin = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); thrust::device_vector d_results(num_keys); thrust::device_vector d_contained(num_keys); @@ -66,9 +69,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", map.find(d_keys.begin(), d_keys.end(), d_results.begin()); auto zip = thrust::make_zip_iterator(thrust::make_tuple(d_results.begin(), d_values.begin())); - REQUIRE(cuco::test::all_of(zip, zip + num_keys, [] __device__(auto const& p) { - return thrust::get<0>(p) == thrust::get<1>(p); - })); + REQUIRE(cuco::test::all_of( + zip, zip + num_keys, cuda::proclaim_return_type([] __device__(auto const& p) { + return thrust::get<0>(p) == thrust::get<1>(p); + }))); } SECTION("All inserted keys-value pairs should be contained") @@ -90,9 +94,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", { REQUIRE(cuco::test::all_of(pairs_begin, pairs_begin + num_keys, - [m_view] __device__(cuco::pair const& pair) mutable { - return m_view.insert(pair); - })); + cuda::proclaim_return_type( + [m_view] __device__(cuco::pair const& pair) mutable { + return m_view.insert(pair); + }))); } SECTION("Cannot find any key in an empty hash map with non-const view") @@ -101,16 +106,19 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", { REQUIRE(cuco::test::all_of(pairs_begin, pairs_begin + num_keys, - [view] __device__(cuco::pair const& pair) mutable { - return view.find(pair.first) == view.end(); - })); + cuda::proclaim_return_type( + [view] __device__(cuco::pair const& pair) mutable { + return view.find(pair.first) == view.end(); + }))); } SECTION("const view") { REQUIRE(cuco::test::all_of( - pairs_begin, pairs_begin + num_keys, [view] __device__(cuco::pair const& pair) { + pairs_begin, + pairs_begin + num_keys, + cuda::proclaim_return_type([view] __device__(cuco::pair const& pair) { return view.find(pair.first) == view.end(); - })); + }))); } } @@ -121,29 +129,33 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys", thrust::device, pairs_begin, pairs_begin + num_keys, - [m_view] __device__(cuco::pair const& pair) mutable { m_view.insert(pair); }); + cuda::proclaim_return_type( + [m_view] __device__(cuco::pair const& pair) mutable { m_view.insert(pair); })); SECTION("non-const view") { // All keys should be found REQUIRE(cuco::test::all_of(pairs_begin, pairs_begin + num_keys, - [view] __device__(cuco::pair const& pair) mutable { - auto const found = view.find(pair.first); - return (found != view.end()) and - (found->first.load() == pair.first and - found->second.load() == pair.second); - })); + cuda::proclaim_return_type( + [view] __device__(cuco::pair const& pair) mutable { + auto const found = view.find(pair.first); + return (found != view.end()) and + (found->first.load() == pair.first and + found->second.load() == pair.second); + }))); } SECTION("const view") { // All keys should be found REQUIRE(cuco::test::all_of( - pairs_begin, pairs_begin + num_keys, [view] __device__(cuco::pair const& pair) { + pairs_begin, + pairs_begin + num_keys, + cuda::proclaim_return_type([view] __device__(cuco::pair const& pair) { auto const found = view.find(pair.first); return (found != view.end()) and (found->first.load() == pair.first and found->second.load() == pair.second); - })); + }))); } } } @@ -160,14 +172,17 @@ __inline__ void test_unique_sequence(Map& map, size_type num_keys) thrust::sequence(thrust::device, d_keys.begin(), d_keys.end()); - auto keys_begin = d_keys.begin(); - auto pairs_begin = - thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i, i); }); + auto keys_begin = d_keys.begin(); + auto pairs_begin = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); thrust::device_vector d_contained(num_keys); - auto zip_equal = [] __device__(auto const& p) { return thrust::get<0>(p) == thrust::get<1>(p); }; - auto is_even = [] __device__(auto const& i) { return i % 2 == 0; }; + auto zip_equal = cuda::proclaim_return_type( + [] __device__(auto const& p) { return thrust::get<0>(p) == thrust::get<1>(p); }); + auto is_even = + cuda::proclaim_return_type([] __device__(auto const& i) { return i % 2 == 0; }); SECTION("Non-inserted keys should not be contained.") { @@ -196,12 +211,13 @@ __inline__ void test_unique_sequence(Map& map, size_type num_keys) REQUIRE(map.size() == num_keys / 2); map.contains(keys_begin, keys_begin + num_keys, d_contained.begin()); - REQUIRE(cuco::test::equal(d_contained.begin(), - d_contained.end(), - thrust::counting_iterator(0), - [] __device__(auto const& idx_contained, auto const& idx) { - return ((idx % 2) == 0) == idx_contained; - })); + REQUIRE(cuco::test::equal( + d_contained.begin(), + d_contained.end(), + thrust::counting_iterator(0), + cuda::proclaim_return_type([] __device__(auto const& idx_contained, auto const& idx) { + return ((idx % 2) == 0) == idx_contained; + }))); } map.insert(pairs_begin, pairs_begin + num_keys); diff --git a/tests/static_multimap/custom_pair_retrieve_test.cu b/tests/static_multimap/custom_pair_retrieve_test.cu index 7856b9e20..f6f1b33ae 100644 --- a/tests/static_multimap/custom_pair_retrieve_test.cu +++ b/tests/static_multimap/custom_pair_retrieve_test.cu @@ -28,6 +28,8 @@ #include #include +#include + #include #include @@ -93,9 +95,9 @@ void test_non_shmem_pair_retrieve(Map& map, std::size_t const num_pairs) thrust::counting_iterator(0), thrust::counting_iterator(num_pairs), d_pairs.begin(), - [] __device__(auto i) { + cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i / 2, i}; - }); + })); auto pair_begin = d_pairs.begin(); @@ -106,15 +108,17 @@ void test_non_shmem_pair_retrieve(Map& map, std::size_t const num_pairs) thrust::counting_iterator(0), thrust::counting_iterator(num_pairs), pair_begin, - [] __device__(auto i) { + cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i, i}; - }); + })); // create an array of prefix sum thrust::device_vector d_scan(num_pairs); - auto count_begin = thrust::make_transform_iterator( - thrust::make_counting_iterator(0), - [num_pairs] __device__(auto i) { return i < (num_pairs / 2) ? 2 : 1; }); + auto count_begin = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + cuda::proclaim_return_type([num_pairs] __device__(auto i) { + return i < (num_pairs / 2) ? 2 : 1; + })); thrust::exclusive_scan(thrust::device, count_begin, count_begin + num_pairs, d_scan.begin(), 0); auto constexpr gold_size = 300; @@ -151,21 +155,24 @@ void test_non_shmem_pair_retrieve(Map& map, std::size_t const num_pairs) thrust::sort(thrust::device, contained_vals.begin(), contained_vals.end()); // set gold references - auto gold_probe = thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [num_pairs] __device__(auto i) { - if (i < num_pairs) { return i / 2; } - return i - (int(num_pairs) / 2); - }); - auto gold_contained_key = thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [num_pairs] __device__(auto i) { - if (i < num_pairs / 2) { return -1; } - return (i - (int(num_pairs) / 2)) / 2; - }); - auto gold_contained_val = thrust::make_transform_iterator(thrust::make_counting_iterator(0), - [num_pairs] __device__(auto i) { - if (i < num_pairs / 2) { return -1; } - return i - (int(num_pairs) / 2); - }); + auto gold_probe = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + cuda::proclaim_return_type([num_pairs] __device__(auto i) { + if (i < num_pairs) { return i / 2; } + return i - (int(num_pairs) / 2); + })); + auto gold_contained_key = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + cuda::proclaim_return_type([num_pairs] __device__(auto i) { + if (i < num_pairs / 2) { return -1; } + return (i - (int(num_pairs) / 2)) / 2; + })); + auto gold_contained_val = + thrust::make_transform_iterator(thrust::make_counting_iterator(0), + cuda::proclaim_return_type([num_pairs] __device__(auto i) { + if (i < num_pairs / 2) { return -1; } + return i - (int(num_pairs) / 2); + })); auto key_equal = thrust::equal_to{}; auto value_equal = thrust::equal_to{}; diff --git a/tests/static_multimap/custom_type_test.cu b/tests/static_multimap/custom_type_test.cu index f53719205..58d653f76 100644 --- a/tests/static_multimap/custom_type_test.cu +++ b/tests/static_multimap/custom_type_test.cu @@ -27,6 +27,8 @@ #include #include +#include + #include #include @@ -72,17 +74,17 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) thrust::counting_iterator(0), thrust::counting_iterator(num_pairs), insert_keys.begin(), - [] __device__(auto i) { + cuda::proclaim_return_type([] __device__(auto i) { return Key{i, i}; - }); + })); thrust::transform(thrust::device, thrust::counting_iterator(0), thrust::counting_iterator(num_pairs), insert_values.begin(), - [] __device__(auto i) { + cuda::proclaim_return_type([] __device__(auto i) { return Value{i, i}; - }); + })); auto pair_begin = thrust::make_zip_iterator(thrust::make_tuple(insert_keys.begin(), insert_values.begin())); @@ -106,21 +108,22 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) REQUIRE(size == num_pairs); // sort before compare - thrust::sort( - thrust::device, - found_pairs.begin(), - found_pairs.end(), - [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { - return lhs.first.a < rhs.first.a; - }); + thrust::sort(thrust::device, + found_pairs.begin(), + found_pairs.end(), + cuda::proclaim_return_type([] __device__(const cuco::pair& lhs, + const cuco::pair& rhs) { + return lhs.first.a < rhs.first.a; + })); REQUIRE( cuco::test::equal(pair_begin, pair_begin + num_pairs, found_pairs.begin(), - [] __device__(cuco::pair lhs, cuco::pair rhs) { - return lhs.first.a == rhs.first.a; - })); + cuda::proclaim_return_type( + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first.a == rhs.first.a; + }))); } SECTION("Non-matches are not included in the output") @@ -135,9 +138,9 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) thrust::counting_iterator(0), thrust::counting_iterator(num), query_key_begin, - [] __device__(auto i) { + cuda::proclaim_return_type([] __device__(auto i) { return Key{i, i}; - }); + })); auto count = map.count(query_key_begin, query_key_begin + num, stream, key_pair_equals{}); REQUIRE(count == num_pairs); @@ -150,20 +153,21 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) REQUIRE(size == num_pairs); // sort before compare - thrust::sort( - thrust::device, - found_pairs.begin(), - found_pairs.end(), - [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { - return lhs.first.a < rhs.first.a; - }); + thrust::sort(thrust::device, + found_pairs.begin(), + found_pairs.end(), + cuda::proclaim_return_type([] __device__(const cuco::pair& lhs, + const cuco::pair& rhs) { + return lhs.first.a < rhs.first.a; + })); REQUIRE( cuco::test::equal(pair_begin, pair_begin + num_pairs, found_pairs.begin(), - [] __device__(cuco::pair lhs, cuco::pair rhs) { - return lhs.first.a == rhs.first.a; - })); + cuda::proclaim_return_type( + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first.a == rhs.first.a; + }))); } SECTION("Outer functions include non-matches in the output") @@ -177,9 +181,9 @@ __inline__ void test_custom_key_value_type(Map& map, std::size_t num_pairs) thrust::counting_iterator(0), thrust::counting_iterator(num), query_key_begin, - [] __device__(auto i) { + cuda::proclaim_return_type([] __device__(auto i) { return Key{i, i}; - }); + })); auto count_outer = map.count_outer(query_key_begin, query_key_begin + num, stream, key_pair_equals{}); diff --git a/tests/static_multimap/heterogeneous_lookup_test.cu b/tests/static_multimap/heterogeneous_lookup_test.cu index 5a5b8b242..6f8b2c2b8 100644 --- a/tests/static_multimap/heterogeneous_lookup_test.cu +++ b/tests/static_multimap/heterogeneous_lookup_test.cu @@ -25,6 +25,8 @@ #include #include +#include + #include #include @@ -103,11 +105,13 @@ TEMPLATE_TEST_CASE("Heterogeneous lookup", cuco::linear_probing<1, custom_hasher>> map{capacity, cuco::empty_key{sentinel_key}, cuco::empty_value{sentinel_value}}; - auto insert_pairs = - thrust::make_transform_iterator(thrust::counting_iterator(0), - [] __device__(auto i) { return cuco::pair(i, i); }); - auto probe_keys = thrust::make_transform_iterator(thrust::counting_iterator(0), - [] __device__(auto i) { return ProbeKey(i); }); + auto insert_pairs = thrust::make_transform_iterator( + thrust::counting_iterator(0), + cuda::proclaim_return_type>( + [] __device__(auto i) { return cuco::pair(i, i); })); + auto probe_keys = thrust::make_transform_iterator( + thrust::counting_iterator(0), + cuda::proclaim_return_type([] __device__(auto i) { return ProbeKey(i); })); SECTION("All inserted keys-value pairs should be contained") { diff --git a/tests/static_multimap/insert_if_test.cu b/tests/static_multimap/insert_if_test.cu index 5d5648e71..53e150730 100644 --- a/tests/static_multimap/insert_if_test.cu +++ b/tests/static_multimap/insert_if_test.cu @@ -24,13 +24,15 @@ #include #include +#include + #include template __inline__ void test_insert_if(Map& map, PairIt pair_begin, KeyIt key_begin, std::size_t size) { // 50% insertion - auto pred_lambda = [] __device__(Key k) { return k % 2 == 0; }; + auto pred_lambda = cuda::proclaim_return_type([] __device__(Key k) { return k % 2 == 0; }); map.insert_if(pair_begin, pair_begin + size, key_begin, pred_lambda); @@ -63,9 +65,9 @@ TEMPLATE_TEST_CASE_SIG( thrust::counting_iterator(0), thrust::counting_iterator(num_keys), d_pairs.begin(), - [] __device__(auto i) { + cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i, i}; - }); + })); using probe = std::conditional_t>, diff --git a/tests/static_multimap/multiplicity_test.cu b/tests/static_multimap/multiplicity_test.cu index 5de83a042..b1c4cd1ce 100644 --- a/tests/static_multimap/multiplicity_test.cu +++ b/tests/static_multimap/multiplicity_test.cu @@ -27,6 +27,8 @@ #include #include +#include + #include template @@ -44,9 +46,9 @@ __inline__ void test_multiplicity_two(Map& map, std::size_t num_items) thrust::counting_iterator(0), thrust::counting_iterator(num_items), d_pairs.begin(), - [] __device__(auto i) { + cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i / 2, i}; - }); + })); thrust::device_vector> d_results(num_items); @@ -91,22 +93,23 @@ __inline__ void test_multiplicity_two(Map& map, std::size_t num_items) REQUIRE(size == num_items); // sort before compare - thrust::sort( - thrust::device, - d_results.begin(), - d_results.end(), - [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { - if (lhs.first != rhs.first) { return lhs.first < rhs.first; } - return lhs.second < rhs.second; - }); + thrust::sort(thrust::device, + d_results.begin(), + d_results.end(), + cuda::proclaim_return_type([] __device__(const cuco::pair& lhs, + const cuco::pair& rhs) { + if (lhs.first != rhs.first) { return lhs.first < rhs.first; } + return lhs.second < rhs.second; + })); REQUIRE( cuco::test::equal(pair_begin, pair_begin + num_items, output_begin, - [] __device__(cuco::pair lhs, cuco::pair rhs) { - return lhs.first == rhs.first and lhs.second == rhs.second; - })); + cuda::proclaim_return_type( + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first == rhs.first and lhs.second == rhs.second; + }))); } SECTION("count and count_outer should return the same value.") @@ -129,22 +132,23 @@ __inline__ void test_multiplicity_two(Map& map, std::size_t num_items) REQUIRE(size == size_outer); // sort before compare - thrust::sort( - thrust::device, - d_results.begin(), - d_results.end(), - [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { - if (lhs.first != rhs.first) { return lhs.first < rhs.first; } - return lhs.second < rhs.second; - }); + thrust::sort(thrust::device, + d_results.begin(), + d_results.end(), + cuda::proclaim_return_type([] __device__(const cuco::pair& lhs, + const cuco::pair& rhs) { + if (lhs.first != rhs.first) { return lhs.first < rhs.first; } + return lhs.second < rhs.second; + })); REQUIRE( cuco::test::equal(pair_begin, pair_begin + num_items, output_begin, - [] __device__(cuco::pair lhs, cuco::pair rhs) { - return lhs.first == rhs.first and lhs.second == rhs.second; - })); + cuda::proclaim_return_type( + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first == rhs.first and lhs.second == rhs.second; + }))); } } diff --git a/tests/static_multimap/non_match_test.cu b/tests/static_multimap/non_match_test.cu index 94023af56..1be809b63 100644 --- a/tests/static_multimap/non_match_test.cu +++ b/tests/static_multimap/non_match_test.cu @@ -26,6 +26,8 @@ #include #include +#include + #include template @@ -50,22 +52,23 @@ __inline__ void test_non_matches(Map& map, PairIt pair_begin, KeyIt key_begin, s REQUIRE(size == num_keys); // sort before compare - thrust::sort( - thrust::device, - output_begin, - output_end, - [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { - if (lhs.first != rhs.first) { return lhs.first < rhs.first; } - return lhs.second < rhs.second; - }); + thrust::sort(thrust::device, + output_begin, + output_end, + cuda::proclaim_return_type([] __device__(const cuco::pair& lhs, + const cuco::pair& rhs) { + if (lhs.first != rhs.first) { return lhs.first < rhs.first; } + return lhs.second < rhs.second; + })); REQUIRE( cuco::test::equal(pair_begin, pair_begin + num_keys, output_begin, - [] __device__(cuco::pair lhs, cuco::pair rhs) { - return lhs.first == rhs.first and lhs.second == rhs.second; - })); + cuda::proclaim_return_type( + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first == rhs.first and lhs.second == rhs.second; + }))); } SECTION("Output of count_outer and retrieve_outer should be coherent.") @@ -82,34 +85,36 @@ __inline__ void test_non_matches(Map& map, PairIt pair_begin, KeyIt key_begin, s REQUIRE(size == (num_keys + num_keys / 2)); // sort before compare - thrust::sort( - thrust::device, - output_begin, - output_end, - [] __device__(const cuco::pair& lhs, const cuco::pair& rhs) { - if (lhs.first != rhs.first) { return lhs.first < rhs.first; } - return lhs.second < rhs.second; - }); + thrust::sort(thrust::device, + output_begin, + output_end, + cuda::proclaim_return_type([] __device__(const cuco::pair& lhs, + const cuco::pair& rhs) { + if (lhs.first != rhs.first) { return lhs.first < rhs.first; } + return lhs.second < rhs.second; + })); // create gold reference thrust::device_vector> gold(size); auto gold_begin = gold.begin(); - thrust::transform(thrust::device, - thrust::counting_iterator(0), - thrust::counting_iterator(size), - gold_begin, - [num_keys] __device__(auto i) { - if (i < num_keys) { return cuco::pair{i / 2, i}; } - return cuco::pair{i - num_keys / 2, -1}; - }); + thrust::transform( + thrust::device, + thrust::counting_iterator(0), + thrust::counting_iterator(size), + gold_begin, + cuda::proclaim_return_type>([num_keys] __device__(auto i) { + if (i < num_keys) { return cuco::pair{i / 2, i}; } + return cuco::pair{i - num_keys / 2, -1}; + })); REQUIRE( cuco::test::equal(gold_begin, gold_begin + size, output_begin, - [] __device__(cuco::pair lhs, cuco::pair rhs) { - return lhs.first == rhs.first and lhs.second == rhs.second; - })); + cuda::proclaim_return_type( + [] __device__(cuco::pair lhs, cuco::pair rhs) { + return lhs.first == rhs.first and lhs.second == rhs.second; + }))); } } @@ -135,9 +140,9 @@ TEMPLATE_TEST_CASE_SIG( thrust::counting_iterator(0), thrust::counting_iterator(num_keys), d_pairs.begin(), - [] __device__(auto i) { + cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i / 2, i}; - }); + })); using probe = std::conditional_t>, diff --git a/tests/static_multimap/pair_function_test.cu b/tests/static_multimap/pair_function_test.cu index 3ef49377d..8da11521a 100644 --- a/tests/static_multimap/pair_function_test.cu +++ b/tests/static_multimap/pair_function_test.cu @@ -27,6 +27,8 @@ #include #include +#include + #include // Custom pair equal @@ -53,9 +55,9 @@ __inline__ void test_pair_functions(Map& map, PairIt pair_begin, std::size_t num thrust::counting_iterator(0), thrust::counting_iterator(num_pairs), pair_begin, - [] __device__(auto i) { + cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i, i}; - }); + })); SECTION("pair_contains returns true for all inserted pairs and false for non-inserted ones.") { @@ -128,9 +130,9 @@ TEMPLATE_TEST_CASE_SIG( thrust::counting_iterator(0), thrust::counting_iterator(num_pairs), d_pairs.begin(), - [] __device__(auto i) { + cuda::proclaim_return_type>([] __device__(auto i) { return cuco::pair{i / 2, i}; - }); + })); using probe = std::conditional_t>, diff --git a/tests/static_set/heterogeneous_lookup_test.cu b/tests/static_set/heterogeneous_lookup_test.cu index cbc0efac3..3bd631a40 100644 --- a/tests/static_set/heterogeneous_lookup_test.cu +++ b/tests/static_set/heterogeneous_lookup_test.cu @@ -25,6 +25,8 @@ #include #include +#include + #include #include @@ -98,10 +100,12 @@ TEMPLATE_TEST_CASE_SIG( probe_type>{ capacity, cuco::empty_key{sentinel_key}, custom_key_equal{}, probe}; - auto insert_pairs = thrust::make_transform_iterator(thrust::counting_iterator(0), - [] __device__(auto i) { return Key{i}; }); - auto probe_keys = thrust::make_transform_iterator(thrust::counting_iterator(0), - [] __device__(auto i) { return ProbeKey(i); }); + auto insert_pairs = thrust::make_transform_iterator( + thrust::counting_iterator(0), + cuda::proclaim_return_type([] __device__(auto i) { return Key{i}; })); + auto probe_keys = thrust::make_transform_iterator( + thrust::counting_iterator(0), + cuda::proclaim_return_type([] __device__(auto i) { return ProbeKey(i); })); SECTION("All inserted keys should be contained") { diff --git a/tests/static_set/insert_and_find_test.cu b/tests/static_set/insert_and_find_test.cu index 278510e08..55f293898 100644 --- a/tests/static_set/insert_and_find_test.cu +++ b/tests/static_set/insert_and_find_test.cu @@ -22,6 +22,8 @@ #include #include +#include + #include template @@ -34,8 +36,9 @@ __inline__ void test_insert_and_find(Set& set, std::size_t num_keys) if constexpr (cg_size == 1) { return thrust::counting_iterator(0); } else { - return thrust::make_transform_iterator(thrust::counting_iterator(0), - [] __device__(auto i) { return i / cg_size; }); + return thrust::make_transform_iterator( + thrust::counting_iterator(0), + cuda::proclaim_return_type([] __device__(auto i) { return i / cg_size; })); } }(); auto const keys_end = [&]() { diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu index 53ede7524..663560807 100644 --- a/tests/static_set/unique_sequence_test.cu +++ b/tests/static_set/unique_sequence_test.cu @@ -27,6 +27,8 @@ #include #include +#include + #include using size_type = int32_t; @@ -43,8 +45,10 @@ __inline__ void test_unique_sequence(Set& set, size_type num_keys) auto keys_begin = d_keys.begin(); thrust::device_vector d_contained(num_keys); - auto zip_equal = [] __device__(auto const& p) { return thrust::get<0>(p) == thrust::get<1>(p); }; - auto is_even = [] __device__(auto const& i) { return i % 2 == 0; }; + auto zip_equal = cuda::proclaim_return_type( + [] __device__(auto const& p) { return thrust::get<0>(p) == thrust::get<1>(p); }); + auto is_even = + cuda::proclaim_return_type([] __device__(auto const& i) { return i % 2 == 0; }); SECTION("Non-inserted keys should not be contained.") { @@ -73,12 +77,13 @@ __inline__ void test_unique_sequence(Set& set, size_type num_keys) REQUIRE(set.size() == num_keys / 2); set.contains(keys_begin, keys_begin + num_keys, d_contained.begin()); - REQUIRE(cuco::test::equal(d_contained.begin(), - d_contained.end(), - thrust::counting_iterator(0), - [] __device__(auto const& idx_contained, auto const& idx) { - return ((idx % 2) == 0) == idx_contained; - })); + REQUIRE(cuco::test::equal( + d_contained.begin(), + d_contained.end(), + thrust::counting_iterator(0), + cuda::proclaim_return_type([] __device__(auto const& idx_contained, auto const& idx) { + return ((idx % 2) == 0) == idx_contained; + }))); } set.insert(keys_begin, keys_begin + num_keys);