#include <SlabHashBackendImpl.h>

Collaboration diagram for cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >:

[legend]

Public Member Functions
	SlabHashBackendImpl ()

__host__ void	Setup (int64_t init_buckets, const SlabNodeManagerImpl &node_mgr_impl, const CUDAHashBackendBufferAccessor &buffer_accessor)

__device__ bool	Insert (bool lane_active, uint32_t lane_id, uint32_t bucket_id, const Key &key, buf_index_t buf_index)
	Warp-insert a pre-allocated buf_index at key. More...

__device__ Pair< buf_index_t, bool >	Find (bool lane_active, uint32_t lane_id, uint32_t bucket_id, const Key &key)
	Warp-find a buf_index and its mask at key. More...

__device__ Pair< buf_index_t, bool >	Erase (bool lane_active, uint32_t lane_id, uint32_t bucket_id, const Key &key)
	Warp-erase an entry at key. More...

__device__ void	WarpSyncKey (const Key &key, uint32_t lane_id, Key &ret_key)
	Warp-synchronize a key in a slab. More...

__device__ int32_t	WarpFindKey (const Key &src_key, uint32_t lane_id, uint32_t slab_entry)
	Warp-find a key in a slab. More...

__device__ int32_t	WarpFindEmpty (uint32_t slab_entry)
	Warp-find the first empty slot in a slab. More...

__device__ int64_t	ComputeBucket (const Key &key) const

__device__ uint32_t	AllocateSlab (uint32_t lane_id)

__device__ void	FreeSlab (uint32_t slab_ptr)

__device__ uint32_t *	SlabEntryPtr (uint32_t bucket_id, uint32_t lane_id, uint32_t slab_ptr)

__device__ uint32_t *	SlabEntryPtrFromNodes (uint32_t slab_ptr, uint32_t lane_id)

__device__ uint32_t *	SlabEntryPtrFromHead (uint32_t bucket_id, uint32_t lane_id)

Public Attributes
Hash	hash_fn_

Eq	eq_fn_

int64_t	bucket_count_

Slab *	bucket_list_head_

SlabNodeManagerImpl	node_mgr_impl_

CUDAHashBackendBufferAccessor	buffer_accessor_

int	key_size_in_int_ = sizeof(Key) / sizeof(int)

Detailed Description

template<typename Key, typename Hash, typename Eq>
class cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >

Definition at line 45 of file SlabHashBackendImpl.h.

Constructor & Destructor Documentation

◆ SlabHashBackendImpl()

template<typename Key , typename Hash , typename Eq >

cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >::SlabHashBackendImpl

Definition at line 175 of file SlabHashBackendImpl.h.

Member Function Documentation

◆ AllocateSlab()

template<typename Key , typename Hash , typename Eq >

__device__ uint32_t cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >::AllocateSlab ( uint32_t lane_id )

Definition at line 463 of file SlabHashBackendImpl.h.

References cloudViewer::core::SlabNodeManagerImpl::WarpAllocate().

◆ ComputeBucket()

template<typename Key , typename Hash , typename Eq >

__device__ int64_t cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >::ComputeBucket ( const Key & key ) const

Definition at line 457 of file SlabHashBackendImpl.h.

Referenced by cloudViewer::core::EraseKernelPass0(), cloudViewer::core::FindKernel(), and cloudViewer::core::InsertKernelPass1().

◆ Erase()

template<typename Key , typename Hash , typename Eq >

__device__ Pair< buf_index_t, bool > cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >::Erase	(	bool	lane_active,
		uint32_t	lane_id,
		uint32_t	bucket_id,
		const Key &	key
	)

Warp-erase an entry at key.

Definition at line 361 of file SlabHashBackendImpl.h.

References cloudViewer::core::kEmptyNodeAddr, cloudViewer::core::kEmptySlabAddr, cloudViewer::core::kHeadSlabAddr, cloudViewer::core::kNextSlabPtrLaneId, cloudViewer::core::kSyncLanesMask, cloudViewer::core::kWarpSize, and cloudViewer::core::make_pair().

Referenced by cloudViewer::core::EraseKernelPass0().

◆ Find()

template<typename Key , typename Hash , typename Eq >

__device__ Pair< buf_index_t, bool > cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >::Find	(	bool	lane_active,
		uint32_t	lane_id,
		uint32_t	bucket_id,
		const Key &	key
	)

Warp-find a buf_index and its mask at key.

Definition at line 294 of file SlabHashBackendImpl.h.

References cloudViewer::core::kEmptySlabAddr, cloudViewer::core::kHeadSlabAddr, cloudViewer::core::kNextSlabPtrLaneId, cloudViewer::core::kNullAddr, cloudViewer::core::kSyncLanesMask, cloudViewer::core::kWarpSize, and cloudViewer::core::make_pair().

Referenced by cloudViewer::core::FindKernel().

◆ FreeSlab()

template<typename Key , typename Hash , typename Eq >

__device__ __forceinline__ void cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >::FreeSlab ( uint32_t slab_ptr )

Definition at line 468 of file SlabHashBackendImpl.h.

References cloudViewer::core::SlabNodeManagerImpl::FreeUntouched().

◆ Insert()

template<typename Key , typename Hash , typename Eq >

__device__ bool cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >::Insert	(	bool	lane_active,
		uint32_t	lane_id,
		uint32_t	bucket_id,
		const Key &	key,
		buf_index_t	buf_index
	)

Warp-insert a pre-allocated buf_index at key.

Definition at line 189 of file SlabHashBackendImpl.h.

References cloudViewer::core::kEmptyNodeAddr, cloudViewer::core::kEmptySlabAddr, cloudViewer::core::kHeadSlabAddr, cloudViewer::core::kNextSlabPtrLaneId, cloudViewer::core::kSyncLanesMask, and cloudViewer::core::kWarpSize.

Referenced by cloudViewer::core::InsertKernelPass1().

◆ Setup()

template<typename Key , typename Hash , typename Eq >

void cloudViewer::core::SlabHashBackendImpl< Key, Hash, Eq >::Setup	(	int64_t	init_buckets,
		const SlabNodeManagerImpl &	node_mgr_impl,
		const CUDAHashBackendBufferAccessor &	buffer_accessor
	)