Open3D (C++ API)  0.19.0
Loading...
Searching...
No Matches
open3d::core::SlabHashBackendImpl< Key, Hash, Eq > Class Template Reference

#include <SlabHashBackendImpl.h>

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.
__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.
__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.
__device__ void WarpSyncKey (const Key &key, uint32_t lane_id, Key &ret_key)
 Warp-synchronize a key in a slab.
__device__ int32_t WarpFindKey (const Key &src_key, uint32_t lane_id, uint32_t slab_entry)
 Warp-find a key in a slab.
__device__ int32_t WarpFindEmpty (uint32_t slab_entry)
 Warp-find the first empty slot in a slab.
__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)

Data Fields

Hash hash_fn_
Eq eq_fn_
int64_t bucket_count_
Slabbucket_list_head_
SlabNodeManagerImpl node_mgr_impl_
CUDAHashBackendBufferAccessor buffer_accessor_
int key_size_in_int_ = sizeof(Key) / sizeof(int)

Constructor & Destructor Documentation

◆ SlabHashBackendImpl()

template<typename Key, typename Hash, typename Eq>
open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::SlabHashBackendImpl ( )

Member Function Documentation

◆ AllocateSlab()

template<typename Key, typename Hash, typename Eq>
__device__ uint32_t open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::AllocateSlab ( uint32_t lane_id)

◆ ComputeBucket()

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

◆ Erase()

template<typename Key, typename Hash, typename Eq>
__device__ Pair< buf_index_t, bool > open3d::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.

◆ Find()

template<typename Key, typename Hash, typename Eq>
__device__ Pair< buf_index_t, bool > open3d::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.

◆ FreeSlab()

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

◆ Insert()

template<typename Key, typename Hash, typename Eq>
__device__ bool open3d::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.

◆ Setup()

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

◆ SlabEntryPtr()

template<typename Key, typename Hash, typename Eq>
__device__ uint32_t * open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::SlabEntryPtr ( uint32_t bucket_id,
uint32_t lane_id,
uint32_t slab_ptr )
inline

◆ SlabEntryPtrFromHead()

template<typename Key, typename Hash, typename Eq>
__device__ uint32_t * open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::SlabEntryPtrFromHead ( uint32_t bucket_id,
uint32_t lane_id )
inline

◆ SlabEntryPtrFromNodes()

template<typename Key, typename Hash, typename Eq>
__device__ uint32_t * open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::SlabEntryPtrFromNodes ( uint32_t slab_ptr,
uint32_t lane_id )
inline

◆ WarpFindEmpty()

template<typename Key, typename Hash, typename Eq>
__device__ int32_t open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::WarpFindEmpty ( uint32_t slab_entry)

Warp-find the first empty slot in a slab.

◆ WarpFindKey()

template<typename Key, typename Hash, typename Eq>
__device__ int32_t open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::WarpFindKey ( const Key & src_key,
uint32_t lane_id,
uint32_t slab_entry )

Warp-find a key in a slab.

◆ WarpSyncKey()

template<typename Key, typename Hash, typename Eq>
__device__ void open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::WarpSyncKey ( const Key & key,
uint32_t lane_id,
Key & ret_key )

Warp-synchronize a key in a slab.

Field Documentation

◆ bucket_count_

template<typename Key, typename Hash, typename Eq>
int64_t open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::bucket_count_

◆ bucket_list_head_

template<typename Key, typename Hash, typename Eq>
Slab* open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::bucket_list_head_

◆ buffer_accessor_

template<typename Key, typename Hash, typename Eq>
CUDAHashBackendBufferAccessor open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::buffer_accessor_

◆ eq_fn_

template<typename Key, typename Hash, typename Eq>
Eq open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::eq_fn_

◆ hash_fn_

template<typename Key, typename Hash, typename Eq>
Hash open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::hash_fn_

◆ key_size_in_int_

template<typename Key, typename Hash, typename Eq>
int open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::key_size_in_int_ = sizeof(Key) / sizeof(int)

◆ node_mgr_impl_

template<typename Key, typename Hash, typename Eq>
SlabNodeManagerImpl open3d::core::SlabHashBackendImpl< Key, Hash, Eq >::node_mgr_impl_

The documentation for this class was generated from the following file: