Open3D (C++ API)  0.19.0
Loading...
Searching...
No Matches
CUDAHashBackendBufferAccessor.h
Go to the documentation of this file.
1// ----------------------------------------------------------------------------
2// - Open3D: www.open3d.org -
3// ----------------------------------------------------------------------------
4// Copyright (c) 2018-2024 www.open3d.org
5// SPDX-License-Identifier: MIT
6// ----------------------------------------------------------------------------
7
8#pragma once
9
10#include <assert.h>
11
12#include <memory>
13#include <vector>
14
20
21namespace open3d {
22namespace core {
23
25public:
26 __host__ void Setup(HashBackendBuffer &hashmap_buffer) {
27 Device device = hashmap_buffer.GetDevice();
28
29 // Properties
30 capacity_ = hashmap_buffer.GetCapacity();
31 key_dsize_ = hashmap_buffer.GetKeyDsize();
32
33 std::vector<int64_t> value_dsizes_host =
34 hashmap_buffer.GetValueDsizes();
35 std::vector<int64_t> value_blocks_per_element_host =
36 hashmap_buffer.GetValueBlocksPerElement();
37 n_values_ = value_blocks_per_element_host.size();
38
39 value_dsizes_ = static_cast<int64_t *>(
40 MemoryManager::Malloc(n_values_ * sizeof(int64_t), device));
41 value_blocks_per_element_ = static_cast<int64_t *>(
42 MemoryManager::Malloc(n_values_ * sizeof(int64_t), device));
43
45 value_dsizes_host.data(),
46 n_values_ * sizeof(int64_t));
48 value_blocks_per_element_host.data(),
49 n_values_ * sizeof(int64_t));
50
51 common_block_size_ = hashmap_buffer.GetCommonBlockSize();
52
53 // Pointers
54 heap_ = hashmap_buffer.GetIndexHeap().GetDataPtr<buf_index_t>();
55 keys_ = hashmap_buffer.GetKeyBuffer().GetDataPtr<uint8_t>();
56
57 std::vector<Tensor> value_buffers = hashmap_buffer.GetValueBuffers();
58 std::vector<uint8_t *> value_ptrs(n_values_);
59 for (size_t i = 0; i < n_values_; ++i) {
60 value_ptrs[i] = value_buffers[i].GetDataPtr<uint8_t>();
61 cudaMemset(value_ptrs[i], 0, capacity_ * value_dsizes_host[i]);
62 }
63 values_ = static_cast<uint8_t **>(
64 MemoryManager::Malloc(n_values_ * sizeof(uint8_t *), device));
65 MemoryManager::MemcpyFromHost(values_, device, value_ptrs.data(),
66 n_values_ * sizeof(uint8_t *));
67
68 heap_top_ = hashmap_buffer.GetHeapTop().cuda.GetDataPtr<int>();
70 OPEN3D_CUDA_CHECK(cudaGetLastError());
71 }
72
73 __host__ void Shutdown(const Device &device) {
77 }
78
80 int index = atomicAdd(heap_top_, 1);
81 return heap_[index];
82 }
83 __device__ void DeviceFree(buf_index_t ptr) {
84 int index = atomicSub(heap_top_, 1);
85 heap_[index - 1] = ptr;
86 }
87
88 __device__ void *GetKeyPtr(buf_index_t ptr) {
89 return keys_ + ptr * key_dsize_;
90 }
91 __device__ void *GetValuePtr(buf_index_t ptr, int value_idx = 0) {
92 return values_[value_idx] + ptr * value_dsizes_[value_idx];
93 }
94
95public:
96 buf_index_t *heap_; /* [N] */
97 int *heap_top_ = nullptr; /* [1] */
98
99 uint8_t *keys_; /* [N] * sizeof(Key) */
100 int64_t key_dsize_;
101
102 size_t n_values_;
103 uint8_t **values_;
104
106
109
110 int64_t capacity_;
111};
112
113} // namespace core
114} // namespace open3d
Common CUDA utilities.
#define OPEN3D_CUDA_CHECK(err)
Definition CUDAUtils.h:47
Definition CUDAHashBackendBufferAccessor.h:24
int64_t * value_blocks_per_element_
Definition CUDAHashBackendBufferAccessor.h:108
__host__ void Setup(HashBackendBuffer &hashmap_buffer)
Definition CUDAHashBackendBufferAccessor.h:26
int64_t * value_dsizes_
Definition CUDAHashBackendBufferAccessor.h:107
__device__ void * GetValuePtr(buf_index_t ptr, int value_idx=0)
Definition CUDAHashBackendBufferAccessor.h:91
uint8_t ** values_
Definition CUDAHashBackendBufferAccessor.h:103
__device__ buf_index_t DeviceAllocate()
Definition CUDAHashBackendBufferAccessor.h:79
int * heap_top_
Definition CUDAHashBackendBufferAccessor.h:97
__device__ void * GetKeyPtr(buf_index_t ptr)
Definition CUDAHashBackendBufferAccessor.h:88
__device__ void DeviceFree(buf_index_t ptr)
Definition CUDAHashBackendBufferAccessor.h:83
int64_t capacity_
Definition CUDAHashBackendBufferAccessor.h:110
__host__ void Shutdown(const Device &device)
Definition CUDAHashBackendBufferAccessor.h:73
uint8_t * keys_
Definition CUDAHashBackendBufferAccessor.h:99
size_t n_values_
Definition CUDAHashBackendBufferAccessor.h:102
buf_index_t * heap_
Definition CUDAHashBackendBufferAccessor.h:96
int64_t key_dsize_
Definition CUDAHashBackendBufferAccessor.h:100
int64_t common_block_size_
Definition CUDAHashBackendBufferAccessor.h:105
Definition Device.h:18
Definition HashBackendBuffer.h:46
int64_t GetKeyDsize() const
Return key's data size in bytes.
Definition HashBackendBuffer.cpp:77
Device GetDevice() const
Return device of the buffer.
Definition HashBackendBuffer.cpp:73
Tensor GetIndexHeap() const
Return the index heap tensor.
Definition HashBackendBuffer.cpp:97
int64_t GetCapacity() const
Return capacity of the buffer.
Definition HashBackendBuffer.cpp:75
HeapTop & GetHeapTop()
Definition HashBackendBuffer.cpp:99
Tensor GetKeyBuffer() const
Return the key buffer tensor.
Definition HashBackendBuffer.cpp:110
std::vector< int64_t > GetValueDsizes() const
Return value's data sizes in bytes.
Definition HashBackendBuffer.cpp:81
int64_t GetCommonBlockSize() const
Get the common block size divisor of all values types.
Definition HashBackendBuffer.cpp:89
std::vector< Tensor > GetValueBuffers() const
Return the value buffer tensors.
Definition HashBackendBuffer.cpp:112
std::vector< int64_t > GetValueBlocksPerElement() const
Return value's data sizes in the unit of common block size divisor.
Definition HashBackendBuffer.cpp:93
static void MemcpyFromHost(void *dst_ptr, const Device &dst_device, const void *host_ptr, size_t num_bytes)
Same as Memcpy, but with host (CPU:0) as default src_device.
Definition MemoryManager.cpp:77
static void * Malloc(size_t byte_size, const Device &device)
Definition MemoryManager.cpp:22
static void Free(void *ptr, const Device &device)
Frees previously allocated memory at address ptr on device device.
Definition MemoryManager.cpp:28
T * GetDataPtr()
Definition Tensor.h:1143
void Synchronize()
Definition CUDAUtils.cpp:58
uint32_t buf_index_t
Definition HashBackendBuffer.h:44
Definition PinholeCameraIntrinsic.cpp:16
Tensor cuda
Definition HashBackendBuffer.h:49