Skip to content
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

Migrate tests to excercise the new map #393

Merged
merged 8 commits into from
Nov 20, 2023
Merged

Migrate tests to excercise the new map #393

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
a6202c2
Migrate duplicate keys tests
PointKernel Nov 3, 2023
372bf1e
Migrate key sentinel tests
PointKernel Nov 15, 2023
a05d3f7
Migrate stream tests
PointKernel Nov 15, 2023
bca4832
Migrate unique sequence tests
PointKernel Nov 15, 2023
ffcf9b0
Remove unused type alias
PointKernel Nov 15, 2023
e3deceb
Merge remote-tracking branch 'upstream/dev' into migrate-map-tests
PointKernel Nov 16, 2023
47a8fff
Merge remote-tracking branch 'upstream/dev' into migrate-map-tests
PointKernel Nov 20, 2023
6fa4e56
Guard device lambda with proclaim_return_type
PointKernel Nov 20, 2023
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
54 changes: 44 additions & 10 deletions tests/static_map/duplicate_keys_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,16 +29,50 @@

#include <catch2/catch_template_test_macros.hpp>

TEMPLATE_TEST_CASE_SIG("Duplicate keys",
"",
((typename Key, typename Value), Key, Value),
(int32_t, int32_t),
(int32_t, int64_t),
(int64_t, int32_t),
(int64_t, int64_t))
using size_type = std::size_t;

TEMPLATE_TEST_CASE_SIG(
"static_map duplicate keys",
"",
((typename Key, typename Value, cuco::test::probe_sequence Probe, int CGSize),
Key,
Value,
Probe,
CGSize),
(int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 1),
(int32_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
(int32_t, int32_t, cuco::test::probe_sequence::double_hashing, 2),
(int32_t, int64_t, cuco::test::probe_sequence::double_hashing, 2),
(int64_t, int32_t, cuco::test::probe_sequence::double_hashing, 1),
(int64_t, int64_t, cuco::test::probe_sequence::double_hashing, 1),
(int64_t, int32_t, cuco::test::probe_sequence::double_hashing, 2),
(int64_t, int64_t, cuco::test::probe_sequence::double_hashing, 2),
(int32_t, int32_t, cuco::test::probe_sequence::linear_probing, 1),
(int32_t, int64_t, cuco::test::probe_sequence::linear_probing, 1),
(int32_t, int32_t, cuco::test::probe_sequence::linear_probing, 2),
(int32_t, int64_t, cuco::test::probe_sequence::linear_probing, 2),
(int64_t, int32_t, cuco::test::probe_sequence::linear_probing, 1),
(int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 1),
(int64_t, int32_t, cuco::test::probe_sequence::linear_probing, 2),
(int64_t, int64_t, cuco::test::probe_sequence::linear_probing, 2))
{
constexpr std::size_t num_keys{500'000};
cuco::static_map<Key, Value> map{
constexpr size_type num_keys{500'000};

using probe =
std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
cuco::experimental::linear_probing<CGSize, cuco::murmurhash3_32<Key>>,
cuco::experimental::double_hashing<CGSize,
cuco::murmurhash3_32<Key>,
cuco::murmurhash3_32<Key>>>;

auto map = cuco::experimental::static_map<Key,
Value,
cuco::experimental::extent<size_type>,
cuda::thread_scope_device,
thrust::equal_to<Key>,
probe,
cuco::cuda_allocator<std::byte>,
cuco::experimental::storage<2>>{
num_keys * 2, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};

thrust::device_vector<Key> d_keys(num_keys);
Expand Down Expand Up @@ -68,7 +102,7 @@ TEMPLATE_TEST_CASE_SIG("Duplicate keys",

map.insert(pairs_begin, pairs_begin + num_keys);

auto const num_entries = map.get_size();
auto const num_entries = map.size();
REQUIRE(num_entries == gold);

auto [key_out_end, value_out_end] =
Expand Down
38 changes: 20 additions & 18 deletions tests/static_map/key_sentinel_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ __device__ int A[SIZE];

template <typename T>
struct custom_equals {
__device__ bool operator()(T lhs, T rhs) { return A[lhs] == A[rhs]; }
__device__ bool operator()(T lhs, T rhs) const { return A[lhs] == A[rhs]; }
};

TEMPLATE_TEST_CASE_SIG(
Expand All @@ -39,11 +39,15 @@ TEMPLATE_TEST_CASE_SIG(
using Value = T;

constexpr std::size_t num_keys{SIZE};
cuco::static_map<Key, Value> map{
SIZE * 2, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
auto map = cuco::experimental::static_map{
SIZE * 2,
cuco::empty_key<Key>{-1},
cuco::empty_value<Value>{-1},
custom_equals<Key>{},
cuco::experimental::linear_probing<1, cuco::default_hash_function<Key>>{}};

auto m_view = map.get_device_mutable_view();
auto view = map.get_device_view();
auto insert_ref = map.ref(cuco::experimental::op::insert);
auto find_ref = map.ref(cuco::experimental::op::find);

int h_A[SIZE];
for (int i = 0; i < SIZE; i++) {
Expand All @@ -60,25 +64,23 @@ TEMPLATE_TEST_CASE_SIG(
{
REQUIRE(cuco::test::all_of(pairs_begin,
pairs_begin + num_keys,
[m_view] __device__(cuco::pair<Key, Value> const& pair) mutable {
return m_view.insert(
pair, cuco::default_hash_function<Key>{}, custom_equals<Key>{});
[insert_ref] __device__(cuco::pair<Key, Value> const& pair) mutable {
return insert_ref.insert(pair);
}));
}

SECTION(
"Tests of CG insert: The custom `key_equal` can never be used to compare against sentinel")
{
map.insert(pairs_begin,
pairs_begin + num_keys,
cuco::default_hash_function<Key>{},
custom_equals<Key>{});
map.insert(pairs_begin, pairs_begin + num_keys);
// 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<Key, Value> 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);
}));
REQUIRE(cuco::test::all_of(pairs_begin,
pairs_begin + num_keys,
[find_ref] __device__(cuco::pair<Key, Value> const& pair) {
Copy link
Collaborator

Choose a reason for hiding this comment

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

Do we need to specify the return type here via cuda::proclaim_return_type?

Copy link
Member Author

Choose a reason for hiding this comment

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

good point. Let me do this to make #343 easier.

auto const found = find_ref.find(pair.first);
return (found != find_ref.end()) and
(found->first == pair.first and
found->second == pair.second);
}));
}
}
5 changes: 2 additions & 3 deletions tests/static_map/shared_memory_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -74,9 +74,8 @@ TEMPLATE_TEST_CASE_SIG("Shared memory static map",
(int64_t, int32_t),
(int64_t, int64_t))
{
using MapType = cuco::static_map<Key, Value>;
using DeviceViewType = typename MapType::device_view;
using DeviceViewIteratorType = typename DeviceViewType::iterator;
using MapType = cuco::static_map<Key, Value>;
using DeviceViewType = typename MapType::device_view;

constexpr std::size_t number_of_maps = 1000;
constexpr std::size_t elements_in_map = 500;
Expand Down
26 changes: 13 additions & 13 deletions tests/static_map/stream_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@

#include <catch2/catch_template_test_macros.hpp>

TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
TEMPLATE_TEST_CASE_SIG("static_map: unique sequence of keys on given stream",
"",
((typename Key, typename Value), Key, Value),
(int32_t, int32_t),
Expand All @@ -41,11 +41,14 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
CUCO_CUDA_TRY(cudaStreamCreate(&stream));

constexpr std::size_t num_keys{500'000};
cuco::static_map<Key, Value> map{1'000'000,
cuco::empty_key<Key>{-1},
cuco::empty_value<Value>{-1},
cuco::cuda_allocator<char>{},
stream};
auto map = cuco::experimental::static_map{
num_keys * 2,
cuco::empty_key<Key>{-1},
cuco::empty_value<Value>{-1},
thrust::equal_to<Key>{},
cuco::experimental::linear_probing<1, cuco::default_hash_function<Key>>{},
cuco::cuda_allocator<char>{},
stream};

thrust::device_vector<Key> d_keys(num_keys);
thrust::device_vector<Value> d_values(num_keys);
Expand All @@ -57,16 +60,13 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
thrust::make_transform_iterator(thrust::make_counting_iterator<int>(0),
[] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });

auto hash_fn = cuco::default_hash_function<Key>{};
auto equal_fn = thrust::equal_to<Value>{};

// bulk function test cases
SECTION("All inserted keys-value pairs should be correctly recovered during find")
{
thrust::device_vector<Value> d_results(num_keys);

map.insert(pairs_begin, pairs_begin + num_keys, hash_fn, equal_fn, stream);
map.find(d_keys.begin(), d_keys.end(), d_results.begin(), hash_fn, equal_fn, stream);
map.insert(pairs_begin, pairs_begin + num_keys, stream);
map.find(d_keys.begin(), d_keys.end(), d_results.begin(), stream);
auto zip = thrust::make_zip_iterator(thrust::make_tuple(d_results.begin(), d_values.begin()));

REQUIRE(cuco::test::all_of(
Expand All @@ -80,8 +80,8 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys on given stream",
{
thrust::device_vector<bool> d_contained(num_keys);

map.insert(pairs_begin, pairs_begin + num_keys, hash_fn, equal_fn, stream);
map.contains(d_keys.begin(), d_keys.end(), d_contained.begin(), hash_fn, equal_fn, stream);
map.insert(pairs_begin, pairs_begin + num_keys, stream);
map.contains(d_keys.begin(), d_keys.end(), d_contained.begin(), stream);

REQUIRE(cuco::test::all_of(d_contained.begin(), d_contained.end(), thrust::identity{}, stream));
}
Expand Down
119 changes: 1 addition & 118 deletions tests/static_map/unique_sequence_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,123 +31,6 @@

#include <catch2/catch_template_test_macros.hpp>

TEMPLATE_TEST_CASE_SIG("Unique sequence of keys",
"",
((typename Key, typename Value), Key, Value),
(int32_t, int32_t),
(int32_t, int64_t),
(int64_t, int32_t),
(int64_t, int64_t))
{
constexpr std::size_t num_keys{500'000};
cuco::static_map<Key, Value> map{
1'000'000, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};

auto m_view = map.get_device_mutable_view();
auto view = map.get_device_view();

thrust::device_vector<Key> d_keys(num_keys);
thrust::device_vector<Value> d_values(num_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<int>(0),
[] __device__(auto i) { return cuco::pair<Key, Value>(i, i); });

thrust::device_vector<Value> d_results(num_keys);
thrust::device_vector<bool> d_contained(num_keys);

// bulk function test cases
SECTION("All inserted keys-value pairs should be correctly recovered during find")
{
map.insert(pairs_begin, pairs_begin + num_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);
}));
}

SECTION("All inserted keys-value pairs should be contained")
{
map.insert(pairs_begin, pairs_begin + num_keys);
map.contains(d_keys.begin(), d_keys.end(), d_contained.begin());

REQUIRE(cuco::test::all_of(d_contained.begin(), d_contained.end(), thrust::identity{}));
}

SECTION("Non-inserted keys-value pairs should not be contained")
{
map.contains(d_keys.begin(), d_keys.end(), d_contained.begin());

REQUIRE(cuco::test::none_of(d_contained.begin(), d_contained.end(), thrust::identity{}));
}

SECTION("Inserting unique keys should return insert success.")
{
REQUIRE(cuco::test::all_of(pairs_begin,
pairs_begin + num_keys,
[m_view] __device__(cuco::pair<Key, Value> const& pair) mutable {
return m_view.insert(pair);
}));
}

SECTION("Cannot find any key in an empty hash map with non-const view")
{
SECTION("non-const view")
{
REQUIRE(cuco::test::all_of(pairs_begin,
pairs_begin + num_keys,
[view] __device__(cuco::pair<Key, Value> 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<Key, Value> const& pair) {
return view.find(pair.first) == view.end();
}));
}
}

SECTION("Keys are all found after inserting many keys.")
{
// Bulk insert keys
thrust::for_each(
thrust::device,
pairs_begin,
pairs_begin + num_keys,
[m_view] __device__(cuco::pair<Key, Value> 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<Key, Value> 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<Key, Value> 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);
}));
}
}
}

using size_type = int32_t;

template <typename Map>
Expand Down Expand Up @@ -253,7 +136,7 @@ __inline__ void test_unique_sequence(Map& map, size_type num_keys)
}

TEMPLATE_TEST_CASE_SIG(
"Unique sequence",
"static_map: unique sequence",
"",
((typename Key, typename Value, cuco::test::probe_sequence Probe, int CGSize),
Key,
Expand Down
Loading