How to traverse all values for given key in kernel function of static_multimap?

[djie-0329]

Hi,

I want to traverse all values for given key in each thread. And I'm not sure how to use static_multimap to implement it.
Can I use find(key) like std::multimap?
for (auto it = mapRef.find(key); it != mapRef.end(); it++)
{
.....
}
or
for (auto it = mapRef.find(key); it != mapRef.end() && it.first() == key; it++) // since all key-value are stored continutely?
{

}

[PointKernel]

The easiest way is to use the device for_each operator

cuCollections/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh

Lines 1482 to 1537 in 420d502

	/**
	* @brief For a given key, applies the function object `callback_op` to the copy of all
	* corresponding matches found in the container.
	*
	* @note This function uses cooperative group semantics, meaning that any thread may call the
	* callback if it finds a matching element. If multiple elements are found within the same group,
	* each thread with a match will call the callback with its associated element.
	*
	* @note The return value of `callback_op`, if any, is ignored.
	*
	* @note Synchronizing `group` within `callback_op` is undefined behavior.
	*
	* @tparam ProbeKey Probe key type
	* @tparam CallbackOp Type of unary callback function object
	* @tparam ParentCG Type of parent Cooperative Group
	*
	* @param group The Cooperative Group used to perform this operation
	* @param key The key to search for
	* @param callback_op Function to apply to every matched slot
	*/
	template <class ProbeKey, class CallbackOp, typename ParentCG>
	__device__ void for_each(cooperative_groups::thread_block_tile<cg_size, ParentCG> group,
	ProbeKey key,
	CallbackOp&& callback_op) const noexcept
	{
	auto probing_iter =
	probing_scheme_.template make_iterator<bucket_size>(group, key, storage_ref_.extent());
	auto const init_idx = *probing_iter;
	bool empty = false;
	while (true) {
	// TODO atomic_ref::load if insert operator is present
	auto const bucket_slots = this->storage_ref_[*probing_iter];
	for (cuda::std::int32_t i = 0; i < bucket_size and !empty; ++i) {
	switch (this->predicate_.template operator()<is_insert::NO>(
	key, this->extract_key(bucket_slots[i]))) {
	case detail::equal_result::EMPTY: {
	empty = true;
	continue;
	}
	case detail::equal_result::EQUAL: {
	callback_op(bucket_slots[i]);
	continue;
	}
	default: {
	continue;
	}
	}
	}
	if (group.any(empty)) { return; }
	++probing_iter;
	if (*probing_iter == init_idx) { return; }
	}
	}

.

https://godbolt.org/z/oYrKacTdd demonstrates inserting the same key 1 ten times with values [0, 1, …, 9]. It also shows how to use for_each to print all matching keys and their corresponding values. Note that the default cooperative group (CG) size is 8, meaning 8 threads are used to query each key. Therefore, the kernel launch must specify at least 8 threads per block to retrieve all 10 matches.

since all key-value are stored continutely?

By default, static_multimap uses a double hashing probing scheme:

cuCollections/include/cuco/static_multimap.cuh

Lines 96 to 97 in 420d502

	class ProbingScheme = cuco::double_hashing<8, // CG size
	cuco::default_hash_function<Key>>,

Unlike linear probing, values associated with the same key are not stored in contiguous memory. Therefore, it++ will not work unless you are using linear probing with a cooperative group (CG) size of 1.

[djie-0329]

Thanks your reply. Looks like all threads parallel-find values for one key simultaneously. How to make different thread separately find all values for given key? I mean, for example {1-2, 1-3, 1-4, 5-6, 5-7}, then thread 0 find value: 2 3 4 for key: 1, thread 1 find values: 6 7 for key: 5.

[PointKernel]

As I mentioned above, the CG size controls how many threads are used simultaneously to query a given key. We use a default CG size of 8 for the multimap based on empirical results, with the goal of improving performance.

If you want each thread to handle a single query key independently, you can set the CG size to 1. However, this will result in reduced runtime performance.