How to "deref" a static_map_ref

[Gibyeng]

When using bulk allocation storage for multiple maps, I have no access to the underlying static_map, which has member functions like .size() and 'retrieve_all' to get keys and values. I only have map_refs of it. Is there any way to dereference it so that I can access the keys and values?

[PointKernel]

Similar to other operations like insert or find, there is static_map_ref::for_each #565 available for use. Otherwise, you can always access map contents via probing_scheme and storage_ref:

cuCollections/include/cuco/static_map_ref.cuh

Lines 215 to 227 in 5602381

	/**
	* @brief Gets the non-owning storage ref.
	*
	* @return The non-owning storage ref of the container
	*/
	[[nodiscard]] __host__ __device__ constexpr auto storage_ref() const noexcept;
	/**
	* @brief Gets the probing scheme.
	*
	* @return The probing scheme used for the container
	*/
	[[nodiscard]] __host__ __device__ constexpr auto probing_scheme() const noexcept;

[PointKernel]

voila https://godbolt.org/z/vqPErncPc

[Gibyeng]

@PointKernel Thanks for your example, I write a kernel link which copies the stored pairs into a device vector. But I still have a few questions:

1. How to get the actual size of the map in device by map_ref ? (I know there is a size function for map but it seems not work for ref types)
2. In the code, thread tid get a item from *(iter+tid), and write it into output[tid]. Are they coalesced reads? it seems that *(iter) and *(iter+1) may be not adjacent in the memory.
3. Is it possible to copy/move the content of the map into a device vector directly, like thrust::copy or std::move?

[PointKernel]

1. Instead of printing or storing, you can check if the slot key equals the empty key sentinel. If it does, do nothing; otherwise, increment the counter. Depending on your case, the counter might need to be atomic.
2. "tid reads *(iter+tid)" is the definition of coalesced reads. :)
3. Not sure what you mean by "copy/move the content of the map into a device vector directly", your write_ftor is already doing exactly that.

[Gibyeng]

@PointKernel Thank you for your response.
I think using 'atomicadd' to count the unempty slot for each hashtable solution absolutely works. But does cuco static_map already do it for me when I insert items?
I mean using a solution like cudaMemcpyDeviceToDevice when I say "copy/move the content of the map into a device vector directly". It is because the cudaMemcpy solution always performs better than the write_ftor kernel, is it right?

[PointKernel]

does cuco static_map already do it for me when I insert items?

For host APIs, yes. When users insert via device APIs, they are responsible for counting the number of successful insertions on their own.

It is because the cudaMemcpy solution always performs better than the write_ftor kernel

It's two operations. You need write_ftor to store the output to the device storage first then use cudaMemcpyDeviceToDevice to copy the device storage back to the host.