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

Code formatting #125

Merged
merged 1 commit into from
Dec 9, 2021
Merged

Code formatting #125

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
100 changes: 49 additions & 51 deletions benchmarks/hash_table/dynamic_map_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,120 +15,118 @@
*/

#include <benchmark/benchmark.h>
#include <synchronization.hpp>
#include <cuco/dynamic_map.cuh>
#include <iostream>
#include <random>
#include <synchronization.hpp>

enum class dist_type {
UNIQUE,
UNIFORM,
GAUSSIAN
};
enum class dist_type { UNIQUE, UNIFORM, GAUSSIAN };

template<dist_type Dist, typename Key, typename OutputIt>
static void generate_keys(OutputIt output_begin, OutputIt output_end) {
template <dist_type Dist, typename Key, typename OutputIt>
static void generate_keys(OutputIt output_begin, OutputIt output_end)
{
auto num_keys = std::distance(output_begin, output_end);

std::random_device rd;
std::mt19937 gen{rd()};

switch(Dist) {
switch (Dist) {
case dist_type::UNIQUE:
for(auto i = 0; i < num_keys; ++i) {
for (auto i = 0; i < num_keys; ++i) {
output_begin[i] = i;
}
break;
case dist_type::UNIFORM:
for(auto i = 0; i < num_keys; ++i) {
for (auto i = 0; i < num_keys; ++i) {
output_begin[i] = std::abs(static_cast<Key>(gen()));
}
break;
case dist_type::GAUSSIAN:
std::normal_distribution<> dg{1e9, 1e7};
for(auto i = 0; i < num_keys; ++i) {
for (auto i = 0; i < num_keys; ++i) {
output_begin[i] = std::abs(static_cast<Key>(dg(gen)));
}
break;
}
}

static void gen_final_size(benchmark::internal::Benchmark* b) {
for(auto size = 10'000'000; size <= 150'000'000; size += 20'000'000) {
static void gen_final_size(benchmark::internal::Benchmark* b)
{
for (auto size = 10'000'000; size <= 150'000'000; size += 20'000'000) {
b->Args({size});
}
}

template <typename Key, typename Value, dist_type Dist>
static void BM_dynamic_insert(::benchmark::State& state) {
static void BM_dynamic_insert(::benchmark::State& state)
{
using map_type = cuco::dynamic_map<Key, Value>;
std::size_t num_keys = state.range(0);
std::size_t initial_size = 1<<27;
std::vector<Key> h_keys( num_keys );
std::vector<cuco::pair_type<Key, Value>> h_pairs ( num_keys );

std::size_t num_keys = state.range(0);
std::size_t initial_size = 1 << 27;

std::vector<Key> h_keys(num_keys);
std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);

generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());

for(auto i = 0; i < num_keys; ++i) {
Key key = h_keys[i];
Value val = h_keys[i];
h_pairs[i].first = key;
for (auto i = 0; i < num_keys; ++i) {
Key key = h_keys[i];
Value val = h_keys[i];
h_pairs[i].first = key;
h_pairs[i].second = val;
}

thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs( h_pairs );
thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);

std::size_t batch_size = 1E6;
for(auto _ : state) {
for (auto _ : state) {
map_type map{initial_size, -1, -1};
{
cuda_event_timer raii{state};
for(auto i = 0; i < num_keys; i += batch_size) {
cuda_event_timer raii{state};
for (auto i = 0; i < num_keys; i += batch_size) {
map.insert(d_pairs.begin() + i, d_pairs.begin() + i + batch_size);
}
}
}

state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) *
int64_t(state.iterations()) *
state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
int64_t(state.range(0)));
}

template <typename Key, typename Value, dist_type Dist>
static void BM_dynamic_search_all(::benchmark::State& state) {
static void BM_dynamic_search_all(::benchmark::State& state)
{
using map_type = cuco::dynamic_map<Key, Value>;

std::size_t num_keys = state.range(0);
std::size_t initial_size = 1<<27;

std::vector<Key> h_keys( num_keys );
std::vector<cuco::pair_type<Key, Value>> h_pairs ( num_keys );
std::size_t num_keys = state.range(0);
std::size_t initial_size = 1 << 27;

std::vector<Key> h_keys(num_keys);
std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);

generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
for(auto i = 0; i < num_keys; ++i) {
Key key = h_keys[i];
Value val = h_keys[i];
h_pairs[i].first = key;

for (auto i = 0; i < num_keys; ++i) {
Key key = h_keys[i];
Value val = h_keys[i];
h_pairs[i].first = key;
h_pairs[i].second = val;
}

thrust::device_vector<Key> d_keys( h_keys );
thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs( h_pairs );
thrust::device_vector<Value> d_results( num_keys );
thrust::device_vector<Key> d_keys(h_keys);
thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
thrust::device_vector<Value> d_results(num_keys);

map_type map{initial_size, -1, -1};
map.insert(d_pairs.begin(), d_pairs.end());

for(auto _ : state) {
for (auto _ : state) {
cuda_event_timer raii{state};
map.find(d_keys.begin(), d_keys.end(), d_results.begin());
}

state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) *
int64_t(state.iterations()) *
state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
int64_t(state.range(0)));
}

Expand Down Expand Up @@ -161,7 +159,7 @@ BENCHMARK_TEMPLATE(BM_dynamic_search_all, int32_t, int32_t, dist_type::GAUSSIAN)
->Unit(benchmark::kMillisecond)
->Apply(gen_final_size)
->UseManualTime();

BENCHMARK_TEMPLATE(BM_dynamic_insert, int64_t, int64_t, dist_type::UNIQUE)
->Unit(benchmark::kMillisecond)
->Apply(gen_final_size)
Expand Down
105 changes: 49 additions & 56 deletions benchmarks/hash_table/static_map_bench.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,40 +15,37 @@
*/

#include <benchmark/benchmark.h>
#include "cuco/static_map.cuh"
#include <thrust/for_each.h>
#include <thrust/device_vector.h>
#include <iostream>
#include <thrust/for_each.h>
#include <fstream>
#include <iostream>
#include <random>
#include "cuco/static_map.cuh"

enum class dist_type {
UNIQUE,
UNIFORM,
GAUSSIAN
};
enum class dist_type { UNIQUE, UNIFORM, GAUSSIAN };

template<dist_type Dist, typename Key, typename OutputIt>
static void generate_keys(OutputIt output_begin, OutputIt output_end) {
template <dist_type Dist, typename Key, typename OutputIt>
static void generate_keys(OutputIt output_begin, OutputIt output_end)
{
auto num_keys = std::distance(output_begin, output_end);

std::random_device rd;
std::mt19937 gen{rd()};

switch(Dist) {
switch (Dist) {
case dist_type::UNIQUE:
for(auto i = 0; i < num_keys; ++i) {
for (auto i = 0; i < num_keys; ++i) {
output_begin[i] = i;
}
break;
case dist_type::UNIFORM:
for(auto i = 0; i < num_keys; ++i) {
for (auto i = 0; i < num_keys; ++i) {
output_begin[i] = std::abs(static_cast<Key>(gen()));
}
break;
case dist_type::GAUSSIAN:
std::normal_distribution<> dg{1e9, 1e7};
for(auto i = 0; i < num_keys; ++i) {
for (auto i = 0; i < num_keys; ++i) {
output_begin[i] = std::abs(static_cast<Key>(dg(gen)));
}
break;
Expand All @@ -59,39 +56,39 @@ static void generate_keys(OutputIt output_begin, OutputIt output_end) {
* @brief Generates input sizes and hash table occupancies
*
*/
static void generate_size_and_occupancy(benchmark::internal::Benchmark* b) {
static void generate_size_and_occupancy(benchmark::internal::Benchmark* b)
{
for (auto size = 100'000'000; size <= 100'000'000; size *= 10) {
for (auto occupancy = 10; occupancy <= 90; occupancy += 10) {
b->Args({size, occupancy});
}
}
}



template <typename Key, typename Value, dist_type Dist>
static void BM_static_map_insert(::benchmark::State& state) {
static void BM_static_map_insert(::benchmark::State& state)
{
using map_type = cuco::static_map<Key, Value>;

std::size_t num_keys = state.range(0);
float occupancy = state.range(1) / float{100};
std::size_t size = num_keys / occupancy;
float occupancy = state.range(1) / float{100};
std::size_t size = num_keys / occupancy;

std::vector<Key> h_keys(num_keys);
std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);

std::vector<Key> h_keys( num_keys );
std::vector<cuco::pair_type<Key, Value>> h_pairs( num_keys );

generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
for(auto i = 0; i < num_keys; ++i) {
Key key = h_keys[i];
Value val = h_keys[i];
h_pairs[i].first = key;

for (auto i = 0; i < num_keys; ++i) {
Key key = h_keys[i];
Value val = h_keys[i];
h_pairs[i].first = key;
h_pairs[i].second = val;
}

thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs( h_pairs );
thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);

for(auto _ : state) {
for (auto _ : state) {
state.ResumeTiming();
state.PauseTiming();
map_type map{size, -1, -1};
Expand All @@ -102,54 +99,50 @@ static void BM_static_map_insert(::benchmark::State& state) {
state.PauseTiming();
}

state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) *
int64_t(state.iterations()) *
state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
int64_t(state.range(0)));
}



template <typename Key, typename Value, dist_type Dist>
static void BM_static_map_search_all(::benchmark::State& state) {
static void BM_static_map_search_all(::benchmark::State& state)
{
using map_type = cuco::static_map<Key, Value>;

std::size_t num_keys = state.range(0);
float occupancy = state.range(1) / float{100};
std::size_t size = num_keys / occupancy;
float occupancy = state.range(1) / float{100};
std::size_t size = num_keys / occupancy;

map_type map{size, -1, -1};
auto view = map.get_device_mutable_view();

std::vector<Key> h_keys( num_keys );
std::vector<Value> h_values( num_keys );
std::vector<cuco::pair_type<Key, Value>> h_pairs ( num_keys );
std::vector<Value> h_results (num_keys);
std::vector<Key> h_keys(num_keys);
std::vector<Value> h_values(num_keys);
std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);
std::vector<Value> h_results(num_keys);

generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
for(auto i = 0; i < num_keys; ++i) {
Key key = h_keys[i];
Value val = h_keys[i];
h_pairs[i].first = key;

for (auto i = 0; i < num_keys; ++i) {
Key key = h_keys[i];
Value val = h_keys[i];
h_pairs[i].first = key;
h_pairs[i].second = val;
}

thrust::device_vector<Key> d_keys( h_keys );
thrust::device_vector<Value> d_results( num_keys);
thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs( h_pairs );
thrust::device_vector<Key> d_keys(h_keys);
thrust::device_vector<Value> d_results(num_keys);
thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);

map.insert(d_pairs.begin(), d_pairs.end());
for(auto _ : state) {

for (auto _ : state) {
map.find(d_keys.begin(), d_keys.end(), d_results.begin());
}

state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
int64_t(state.range(0)));
}



BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIQUE)
->Unit(benchmark::kMillisecond)
->Apply(generate_size_and_occupancy);
Expand Down
15 changes: 10 additions & 5 deletions examples/static_map/static_map_example.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,8 @@ int main(void)
// for an load factor of 50%.
cudaStream_t str;
cudaStreamCreate(&str);
cuco::static_map<int, int> map{100'000, empty_key_sentinel, empty_value_sentinel, cuco::cuda_allocator<char>{}, str};
cuco::static_map<int, int> map{
100'000, empty_key_sentinel, empty_value_sentinel, cuco::cuda_allocator<char>{}, str};

thrust::device_vector<thrust::pair<int, int>> pairs(50'000);

Expand All @@ -43,8 +44,8 @@ int main(void)
[] __device__(auto i) { return thrust::make_pair(i, i); });

// Inserts all pairs into the map
map.insert(pairs.begin(), pairs.end(), cuco::detail::MurmurHash3_32<int>{},
thrust::equal_to<int>{}, str);
map.insert(
pairs.begin(), pairs.end(), cuco::detail::MurmurHash3_32<int>{}, thrust::equal_to<int>{}, str);

// Sequence of keys {0, 1, 2, ...}
thrust::device_vector<int> keys_to_find(50'000);
Expand All @@ -53,8 +54,12 @@ int main(void)

// Finds all keys {0, 1, 2, ...} and stores associated values into `found_values`
// If a key `keys_to_find[i]` doesn't exist, `found_values[i] == empty_value_sentinel`
map.find(keys_to_find.begin(), keys_to_find.end(), found_values.begin(),
cuco::detail::MurmurHash3_32<int>{}, thrust::equal_to<int>{}, str);
map.find(keys_to_find.begin(),
keys_to_find.end(),
found_values.begin(),
cuco::detail::MurmurHash3_32<int>{},
thrust::equal_to<int>{},
str);
cudaStreamSynchronize(str);

return 0;
Expand Down
Loading