43 #include <thrust/device_vector.h>
73 memory_block_index_(0),
74 super_block_index_(0) {}
78 return super_blocks_ + addressDecoder(next_slab_ptr) + lane_id;
82 return super_blocks_ + super_block_idx * kUIntsPerSuperBlock +
89 createMemBlockIndex(tid >> 5);
92 memory_block_bitmap_ =
94 memory_block_index_ * kSlabsPerBlock + lane_id];
107 uint32_t read_bitmap = memory_block_bitmap_;
108 uint32_t allocated_result = kNotFoundFlag;
113 while (allocated_result == kNotFoundFlag) {
114 empty_lane = __ffs(~memory_block_bitmap_) - 1;
115 free_lane = __ballot_sync(kSyncLanesMask, empty_lane >= 0);
116 if (free_lane == 0) {
118 updateMemBlockIndex((threadIdx.x + blockIdx.x * blockDim.x) >>
120 read_bitmap = memory_block_bitmap_;
123 uint32_t src_lane = __ffs(free_lane) - 1;
124 if (src_lane == lane_id) {
125 read_bitmap = atomicCAS(
127 super_block_index_ * kUIntsPerSuperBlock +
128 memory_block_index_ * kSlabsPerBlock + lane_id,
129 memory_block_bitmap_,
130 memory_block_bitmap_ | (1 << empty_lane));
131 if (read_bitmap == memory_block_bitmap_) {
133 memory_block_bitmap_ |= (1 << empty_lane);
135 (super_block_index_ << kSuperBlockMaskBits) |
136 (memory_block_index_ << kBlockMaskBits) |
137 (lane_id << kSlabMaskBits) | empty_lane;
140 memory_block_bitmap_ = read_bitmap;
145 __shfl_sync(kSyncLanesMask, allocated_result, src_lane);
147 return allocated_result;
155 getSuperBlockIndex(ptr) * kUIntsPerSuperBlock +
156 getMemBlockIndex(ptr) * kSlabsPerBlock +
157 (getMemUnitIndex(ptr) >> 5),
158 ~(1 << (getMemUnitIndex(ptr) & 0x1F)));
162 __device__ __host__ __forceinline__
uint32_t
164 return address >> kSuperBlockMaskBits;
166 __device__ __host__ __forceinline__
uint32_t
168 return ((address >> kBlockMaskBits) & 0x1FFFF);
172 return (kBitmapsPerSuperBlock +
173 getMemBlockIndex(address) * kUIntsPerBlock);
175 __device__ __host__ __forceinline__
uint32_t
177 return address & 0x3FF;
181 return getMemUnitIndex(address) * kWarpSize;
185 __device__
void createMemBlockIndex(
uint32_t global_warp_id) {
186 super_block_index_ = global_warp_id % kSuperBlocks;
187 memory_block_index_ = (
hash_coef_ * global_warp_id) >>
188 (32 - kBlocksPerSuperBlockInBits);
192 __device__
void updateMemBlockIndex(
uint32_t global_warp_id) {
194 super_block_index_++;
196 (super_block_index_ == kSuperBlocks) ? 0 : super_block_index_;
197 memory_block_index_ = (
hash_coef_ * (global_warp_id + num_attempts_)) >>
198 (32 - kBlocksPerSuperBlockInBits);
200 memory_block_bitmap_ =
201 *((
super_blocks_ + super_block_index_ * kUIntsPerSuperBlock) +
202 memory_block_index_ * kSlabsPerBlock + (threadIdx.x & 0x1f));
207 return getSuperBlockIndex(address_ptr_index) * kUIntsPerSuperBlock +
208 getMemBlockAddress(address_ptr_index) +
209 getMemUnitIndex(address_ptr_index) * kWarpSize;
212 __host__ __device__
void print_address(
buf_index_t address_ptr_index) {
213 printf(
"Super block Index: %d, Memory block index: %d, Memory unit "
216 getSuperBlockIndex(address_ptr_index),
217 getMemBlockIndex(address_ptr_index),
218 getMemUnitIndex(address_ptr_index));
242 std::mt19937 rng(time(0));
248 kUIntsPerSuperBlock * kSuperBlocks *
sizeof(
uint32_t),
258 kUIntsPerSuperBlock * kSuperBlocks *
sizeof(
uint32_t)));
260 for (
uint32_t i = 0; i < kSuperBlocks; i++) {
264 kBlocksPerSuperBlock * kSlabsPerBlock *
sizeof(
uint32_t)));
271 const uint32_t num_super_blocks = kSuperBlocks;
273 thrust::device_vector<uint32_t> slabs_per_superblock(kSuperBlocks);
274 thrust::fill(slabs_per_superblock.begin(), slabs_per_superblock.end(),
278 int num_mem_units = kBlocksPerSuperBlock * 32;
279 int num_cuda_blocks =
280 (num_mem_units + kThreadsPerBlock - 1) / kThreadsPerBlock;
282 core::cuda::GetStream()>>>(
283 impl_, thrust::raw_pointer_cast(slabs_per_superblock.data()));
287 std::vector<int> result(num_super_blocks);
288 thrust::copy(slabs_per_superblock.begin(), slabs_per_superblock.end(),
#define OPEN3D_CUDA_CHECK(err)
Definition: CUDAUtils.h:66
static void * Malloc(size_t byte_size, const Device &device)
Definition: MemoryManager.cpp:41
static void Free(void *ptr, const Device &device)
Frees previously allocated memory at address ptr on device device.
Definition: MemoryManager.cpp:47
Definition: SlabNodeManager.h:58
buf_index_t kv_pair_ptrs[kWarpSize - 1]
Definition: SlabNodeManager.h:62
buf_index_t next_slab_ptr
An internal ptr managed by InternalNodeManager.
Definition: SlabNodeManager.h:64
Definition: SlabNodeManager.h:238
SlabNodeManager(const Device &device)
Definition: SlabNodeManager.h:240
void Reset()
Definition: SlabNodeManager.h:255
~SlabNodeManager()
Definition: SlabNodeManager.h:253
SlabNodeManagerImpl impl_
Definition: SlabNodeManager.h:295
Device device_
Definition: SlabNodeManager.h:296
std::vector< int > CountSlabsPerSuperblock()
Definition: SlabNodeManager.h:270
Definition: SlabNodeManager.h:67
uint32_t * super_blocks_
A pointer to each super-block.
Definition: SlabNodeManager.h:223
__device__ __forceinline__ uint32_t * get_ptr_for_bitmap(const uint32_t super_block_idx, const uint32_t bitmap_idx)
Definition: SlabNodeManager.h:80
__device__ __forceinline__ uint32_t * get_unit_ptr_from_slab(const buf_index_t &next_slab_ptr, const uint32_t &lane_id)
Definition: SlabNodeManager.h:76
__device__ void Init(uint32_t &tid, uint32_t &lane_id)
Definition: SlabNodeManager.h:87
uint32_t hash_coef_
hash_coef (register): used as (16 bits, 16 bits) for hashing.
Definition: SlabNodeManager.h:225
__device__ void FreeUntouched(buf_index_t ptr)
Definition: SlabNodeManager.h:153
__device__ uint32_t WarpAllocate(const uint32_t &lane_id)
Definition: SlabNodeManager.h:97
SlabNodeManagerImpl()
Definition: SlabNodeManager.h:69
void Synchronize()
Definition: CUDAUtils.cpp:78
uint32_t buf_index_t
Definition: HashBackendBuffer.h:63
__global__ void CountSlabsPerSuperblockKernel(SlabNodeManagerImpl impl, uint32_t *slabs_per_superblock)
const char const char value recording_handle imu_sample recording_handle uint8_t size_t data_size k4a_record_configuration_t config target_format k4a_capture_t capture_handle k4a_imu_sample_t imu_sample playback_handle k4a_logging_message_cb_t void min_level device_handle k4a_imu_sample_t timeout_in_ms capture_handle capture_handle capture_handle image_handle temperature_c k4a_image_t image_handle uint8_t image_handle image_handle image_handle image_handle uint32_t
Definition: K4aPlugin.cpp:567
Definition: PinholeCameraIntrinsic.cpp:35