diff --git a/benchmarks/hash_table/dynamic_map_bench.cu b/benchmarks/hash_table/dynamic_map_bench.cu index 1b6b3c784..3034efd30 100644 --- a/benchmarks/hash_table/dynamic_map_bench.cu +++ b/benchmarks/hash_table/dynamic_map_bench.cu @@ -254,7 +254,7 @@ static void BM_dynamic_erase_none(::benchmark::State& state) state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) * int64_t(state.range(0))); } -/* + BENCHMARK_TEMPLATE(BM_dynamic_insert, int32_t, int32_t, dist_type::UNIQUE) ->Unit(benchmark::kMillisecond) ->Apply(gen_final_size) @@ -264,84 +264,8 @@ BENCHMARK_TEMPLATE(BM_dynamic_search_all, int32_t, int32_t, dist_type::UNIQUE) ->Unit(benchmark::kMillisecond) ->Apply(gen_final_size) ->UseManualTime(); -*/ -BENCHMARK_TEMPLATE(BM_dynamic_erase_all, int32_t, int32_t, dist_type::UNIQUE) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); -/* -BENCHMARK_TEMPLATE(BM_dynamic_insert, int32_t, int32_t, dist_type::UNIFORM) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -BENCHMARK_TEMPLATE(BM_dynamic_search_all, int32_t, int32_t, dist_type::UNIFORM) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -BENCHMARK_TEMPLATE(BM_dynamic_insert, int32_t, int32_t, dist_type::GAUSSIAN) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -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_erase_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) - ->UseManualTime(); -BENCHMARK_TEMPLATE(BM_dynamic_search_all, int64_t, int64_t, dist_type::UNIQUE) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); -*/ -BENCHMARK_TEMPLATE(BM_dynamic_erase_none, int32_t, int32_t, dist_type::UNIQUE) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); -/* -BENCHMARK_TEMPLATE(BM_dynamic_erase_none, int32_t, int32_t, dist_type::GAUSSIAN) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -BENCHMARK_TEMPLATE(BM_dynamic_search_none, int32_t, int32_t, dist_type::GAUSSIAN) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -BENCHMARK_TEMPLATE(BM_dynamic_erase_all, int64_t, int64_t, dist_type::UNIQUE) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -BENCHMARK_TEMPLATE(BM_dynamic_insert, int64_t, int64_t, dist_type::UNIFORM) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -BENCHMARK_TEMPLATE(BM_dynamic_search_all, int64_t, int64_t, dist_type::UNIFORM) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -BENCHMARK_TEMPLATE(BM_dynamic_insert, int64_t, int64_t, dist_type::GAUSSIAN) - ->Unit(benchmark::kMillisecond) - ->Apply(gen_final_size) - ->UseManualTime(); - -BENCHMARK_TEMPLATE(BM_dynamic_search_all, int64_t, int64_t, dist_type::GAUSSIAN) +BENCHMARK_TEMPLATE(BM_dynamic_erase_all, int32_t, int32_t, dist_type::UNIQUE) ->Unit(benchmark::kMillisecond) ->Apply(gen_final_size) - ->UseManualTime(); -*/ \ No newline at end of file + ->UseManualTime(); \ No newline at end of file diff --git a/include/cuco/detail/dynamic_map.inl b/include/cuco/detail/dynamic_map.inl index 2d61ae685..e35e27f62 100644 --- a/include/cuco/detail/dynamic_map.inl +++ b/include/cuco/detail/dynamic_map.inl @@ -203,7 +203,6 @@ void dynamic_map::erase(InputIt first, <<>>( first, first + num_keys, - submap_views_.data().get(), submap_mutable_views_.data().get(), num_successes_, d_submap_num_successes_.data().get(), diff --git a/include/cuco/detail/dynamic_map_kernels.cuh b/include/cuco/detail/dynamic_map_kernels.cuh index 00423701b..9250eb86a 100644 --- a/include/cuco/detail/dynamic_map_kernels.cuh +++ b/include/cuco/detail/dynamic_map_kernels.cuh @@ -179,17 +179,40 @@ __global__ void insert(InputIt first, if (threadIdx.x == 0) { *num_successes += block_num_successes; } } +/** + * @brief Erases the key/value pairs corresponding to all keys in the range `[first, last)`. + * + * If the key `*(first + i)` exists in the map, its slot is erased and made available for future + insertions. + * Else, no effect. + * @tparam block_size The size of the thread block + * @tparam pair_type Type of the pairs contained in the map + * @tparam InputIt Device accessible input iterator whose `value_type` is + * convertible to the map's `key_type` + * @tparam mutableViewT Type of device view allowing modification of hash map storage + * @tparam atomicT Type of atomic storage + * @tparam Hash Unary callable type + * @tparam KeyEqual Binary callable type + * @param first Beginning of the sequence of keys + * @param last End of the sequence of keys + * @param submap_mutable_views Array of `static_map::mutable_device_view` objects used to + * perform `erase` operations on each underlying `static_map` + * @param num_successes The number of successfully erased key/value pairs + * @param submap_num_successes The number of successfully erased key/value pairs + * in each submap + * @param num_submaps The number of submaps in the map + * @param hash The unary function to apply to hash each key + * @param key_equal The binary function to compare two keys for equality + */ template __global__ void erase(InputIt first, InputIt last, - viewT* submap_views, mutableViewT* submap_mutable_views, atomicT* num_successes, atomicT** submap_num_successes, @@ -245,18 +268,42 @@ __global__ void erase(InputIt first, } } +/** + * @brief Erases the key/value pairs corresponding to all keys in the range `[first, last)`. + * + * If the key `*(first + i)` exists in the map, its slot is erased and made available for future + insertions. + * Else, no effect. + * @tparam block_size The size of the thread block + * @tparam tile_size The number of threads in the Cooperative Groups used to perform erase + * @tparam pair_type Type of the pairs contained in the map + * @tparam InputIt Device accessible input iterator whose `value_type` is + * convertible to the map's `key_type` + * @tparam mutableViewT Type of device view allowing modification of hash map storage + * @tparam atomicT Type of atomic storage + * @tparam Hash Unary callable type + * @tparam KeyEqual Binary callable type + * @param first Beginning of the sequence of keys + * @param last End of the sequence of keys + * @param submap_mutable_views Array of `static_map::mutable_device_view` objects used to + * perform `erase` operations on each underlying `static_map` + * @param num_successes The number of successfully erased key/value pairs + * @param submap_num_successes The number of successfully erased key/value pairs + * in each submap + * @param num_submaps The number of submaps in the map + * @param hash The unary function to apply to hash each key + * @param key_equal The binary function to compare two keys for equality + */ template __global__ void erase(InputIt first, InputIt last, - viewT* submap_views, mutableViewT* submap_mutable_views, atomicT* num_successes, atomicT** submap_num_successes, diff --git a/include/cuco/detail/static_map_kernels.cuh b/include/cuco/detail/static_map_kernels.cuh index d1c2ac5c1..9fe4f1cd6 100644 --- a/include/cuco/detail/static_map_kernels.cuh +++ b/include/cuco/detail/static_map_kernels.cuh @@ -156,6 +156,26 @@ __global__ void insert( if (threadIdx.x == 0) { *num_successes += block_num_successes; } } +/** + * @brief Erases the key/value pairs corresponding to all keys in the range `[first, last)`. + * + * If the key `*(first + i)` exists in the map, its slot is erased and made available for future + insertions. + * Else, no effect. + * @tparam block_size The size of the thread block + * @tparam InputIt Device accessible input iterator whose `value_type` is + * convertible to the map's `key_type` + * @tparam atomicT Type of atomic storage + * @tparam viewT Type of device view allowing access of hash map storage + * @tparam Hash Unary callable type + * @tparam KeyEqual Binary callable type + * @param first Beginning of the sequence of keys + * @param last End of the sequence of keys + * @param num_successes The number of successfully erased key/value pairs + * @param view Device view used to access the hash map's slot storage + * @param hash The unary function to apply to hash each key + * @param key_equal The binary function to compare two keys for equality + */ template empty_key_sentinel, sentinel::empty_value empty_value_sentinel, Allocator const& alloc = Allocator{}); - + + /** + * @brief Construct a dynamically-sized map with erase capability. + * + * The capacity of the map will automatically increase as the user adds key/value pairs using + * `insert`. + * + * Capacity increases by a factor of growth_factor each time the size of the map exceeds a + * threshold occupancy. The performance of `find` and `contains` decreases somewhat each time the + * map's capacity grows. + * + * The `empty_key_sentinel` and `empty_value_sentinel` values are reserved and + * undefined behavior results from attempting to insert any key/value pair + * that contains either. + * + * @param initial_capacity The initial number of slots in the map + * @param growth_factor The factor by which the capacity increases when resizing + * @param empty_key_sentinel The reserved key value for empty slots + * @param empty_value_sentinel The reserved mapped value for empty slots + * @param erased_key_sentinel The reserved key value for erased slots + * @param alloc Allocator used to allocate submap device storage + * + * @throw std::runtime error if the empty key sentinel and erased key sentinel + * are the same value + */ dynamic_map(std::size_t initial_capacity, sentinel::empty_key empty_key_sentinel, sentinel::empty_value empty_value_sentinel, @@ -178,7 +202,36 @@ class dynamic_map { typename Hash = cuco::detail::MurmurHash3_32, typename KeyEqual = thrust::equal_to> void insert(InputIt first, InputIt last, Hash hash = Hash{}, KeyEqual key_equal = KeyEqual{}); - + + /** + * @brief Erases keys in the range `[first, last)`. + * + * For each key `k` in `[first, last)`, if `contains(k) == true), removes `k` and it's + * associated value from the map. Else, no effect. + * + * Side-effects: + * - `contains(k) == false` + * - `find(k) == end()` + * - `insert({k,v}) == true` + * - `get_size()` is reduced by the total number of erased keys + * + * This function synchronizes `stream`. + * + * Keep in mind that `erase` does not cause the map to shrink its memory allocation. + * + * @tparam InputIt Device accessible input iterator whose `value_type` is + * convertible to the map's `value_type` + * @tparam Hash Unary callable type + * @tparam KeyEqual Binary callable type + * @param first Beginning of the sequence of keys + * @param last End of the sequence of keys + * @param hash The unary function to apply to hash each key + * @param key_equal The binary function to compare two keys for equality + * @param stream Stream used for executing the kernels + * + * @throw std::runtime_error if a unique erased key sentinel value was not + * provided at construction + */ template , typename KeyEqual = thrust::equal_to> @@ -263,7 +316,7 @@ class dynamic_map { private: key_type empty_key_sentinel_{}; ///< Key value that represents an empty slot mapped_type empty_value_sentinel_{}; ///< Initial value of empty slot - key_type erased_key_sentinel_{}; + key_type erased_key_sentinel_{}; ///< Key value that represents an erased slot // TODO: initialize this std::size_t size_{}; ///< Number of keys in the map @@ -278,7 +331,7 @@ class dynamic_map { std::size_t min_insert_size_{}; ///< min remaining capacity of submap for insert atomic_ctr_type* num_successes_; ///< number of successfully inserted keys on insert std::vector submap_num_successes_; ///< number of succesfully erased keys for each submap - thrust::device_vector d_submap_num_successes_; + thrust::device_vector d_submap_num_successes_; ///< device-side number of successfully erased keys for each submap Allocator alloc_{}; ///< Allocator passed to submaps to allocate their device storage counter_allocator_type counter_allocator_{}; ///< Allocator used to allocate `num_successes_` };