Open3D (C++ API)  0.16.1
CUDAHashBackendBufferAccessor.h
Go to the documentation of this file.
1// ----------------------------------------------------------------------------
2// - Open3D: www.open3d.org -
3// ----------------------------------------------------------------------------
4// The MIT License (MIT)
5//
6// Copyright (c) 2018-2021 www.open3d.org
7//
8// Permission is hereby granted, free of charge, to any person obtaining a copy
9// of this software and associated documentation files (the "Software"), to deal
10// in the Software without restriction, including without limitation the rights
11// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
12// copies of the Software, and to permit persons to whom the Software is
13// furnished to do so, subject to the following conditions:
14//
15// The above copyright notice and this permission notice shall be included in
16// all copies or substantial portions of the Software.
17//
18// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
19// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
20// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
21// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
22// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
23// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
24// IN THE SOFTWARE.
25// ----------------------------------------------------------------------------
26
27#pragma once
28
29#include <assert.h>
30
31#include <memory>
32#include <vector>
33
39
40namespace open3d {
41namespace core {
42
44public:
45 __host__ void Setup(HashBackendBuffer &hashmap_buffer) {
46 Device device = hashmap_buffer.GetDevice();
47
48 // Properties
49 capacity_ = hashmap_buffer.GetCapacity();
50 key_dsize_ = hashmap_buffer.GetKeyDsize();
51
52 std::vector<int64_t> value_dsizes_host =
53 hashmap_buffer.GetValueDsizes();
54 std::vector<int64_t> value_blocks_per_element_host =
55 hashmap_buffer.GetValueBlocksPerElement();
56 n_values_ = value_blocks_per_element_host.size();
57
58 value_dsizes_ = static_cast<int64_t *>(
59 MemoryManager::Malloc(n_values_ * sizeof(int64_t), device));
60 value_blocks_per_element_ = static_cast<int64_t *>(
61 MemoryManager::Malloc(n_values_ * sizeof(int64_t), device));
62
64 value_dsizes_host.data(),
65 n_values_ * sizeof(int64_t));
67 value_blocks_per_element_host.data(),
68 n_values_ * sizeof(int64_t));
69
70 common_block_size_ = hashmap_buffer.GetCommonBlockSize();
71
72 // Pointers
73 heap_ = hashmap_buffer.GetIndexHeap().GetDataPtr<buf_index_t>();
74 keys_ = hashmap_buffer.GetKeyBuffer().GetDataPtr<uint8_t>();
75
76 std::vector<Tensor> value_buffers = hashmap_buffer.GetValueBuffers();
77 std::vector<uint8_t *> value_ptrs(n_values_);
78 for (size_t i = 0; i < n_values_; ++i) {
79 value_ptrs[i] = value_buffers[i].GetDataPtr<uint8_t>();
80 cudaMemset(value_ptrs[i], 0, capacity_ * value_dsizes_host[i]);
81 }
82 values_ = static_cast<uint8_t **>(
83 MemoryManager::Malloc(n_values_ * sizeof(uint8_t *), device));
84 MemoryManager::MemcpyFromHost(values_, device, value_ptrs.data(),
85 n_values_ * sizeof(uint8_t *));
86
87 heap_top_ = hashmap_buffer.GetHeapTop().cuda.GetDataPtr<int>();
89 OPEN3D_CUDA_CHECK(cudaGetLastError());
90 }
91
92 __host__ void Shutdown(const Device &device) {
96 }
97
99 int index = atomicAdd(heap_top_, 1);
100 return heap_[index];
101 }
102 __device__ void DeviceFree(buf_index_t ptr) {
103 int index = atomicSub(heap_top_, 1);
104 heap_[index - 1] = ptr;
105 }
106
107 __device__ void *GetKeyPtr(buf_index_t ptr) {
108 return keys_ + ptr * key_dsize_;
109 }
110 __device__ void *GetValuePtr(buf_index_t ptr, int value_idx = 0) {
111 return values_[value_idx] + ptr * value_dsizes_[value_idx];
112 }
113
114public:
115 buf_index_t *heap_; /* [N] */
116 int *heap_top_ = nullptr; /* [1] */
117
118 uint8_t *keys_; /* [N] * sizeof(Key) */
119 int64_t key_dsize_;
120
121 size_t n_values_;
122 uint8_t **values_;
123
125
128
129 int64_t capacity_;
130};
131
132} // namespace core
133} // namespace open3d
Common CUDA utilities.
#define OPEN3D_CUDA_CHECK(err)
Definition: CUDAUtils.h:66
Definition: CUDAHashBackendBufferAccessor.h:43
int64_t * value_blocks_per_element_
Definition: CUDAHashBackendBufferAccessor.h:127
__host__ void Setup(HashBackendBuffer &hashmap_buffer)
Definition: CUDAHashBackendBufferAccessor.h:45
int64_t * value_dsizes_
Definition: CUDAHashBackendBufferAccessor.h:126
__device__ void * GetValuePtr(buf_index_t ptr, int value_idx=0)
Definition: CUDAHashBackendBufferAccessor.h:110
uint8_t ** values_
Definition: CUDAHashBackendBufferAccessor.h:122
__device__ buf_index_t DeviceAllocate()
Definition: CUDAHashBackendBufferAccessor.h:98
int * heap_top_
Definition: CUDAHashBackendBufferAccessor.h:116
__device__ void * GetKeyPtr(buf_index_t ptr)
Definition: CUDAHashBackendBufferAccessor.h:107
__device__ void DeviceFree(buf_index_t ptr)
Definition: CUDAHashBackendBufferAccessor.h:102
int64_t capacity_
Definition: CUDAHashBackendBufferAccessor.h:129
__host__ void Shutdown(const Device &device)
Definition: CUDAHashBackendBufferAccessor.h:92
uint8_t * keys_
Definition: CUDAHashBackendBufferAccessor.h:118
size_t n_values_
Definition: CUDAHashBackendBufferAccessor.h:121
buf_index_t * heap_
Definition: CUDAHashBackendBufferAccessor.h:115
int64_t key_dsize_
Definition: CUDAHashBackendBufferAccessor.h:119
int64_t common_block_size_
Definition: CUDAHashBackendBufferAccessor.h:124
Definition: Device.h:37
Definition: HashBackendBuffer.h:65
int64_t GetKeyDsize() const
Return key's data size in bytes.
Definition: HashBackendBuffer.cpp:96
Device GetDevice() const
Return device of the buffer.
Definition: HashBackendBuffer.cpp:92
Tensor GetIndexHeap() const
Return the index heap tensor.
Definition: HashBackendBuffer.cpp:116
int64_t GetCapacity() const
Return capacity of the buffer.
Definition: HashBackendBuffer.cpp:94
HeapTop & GetHeapTop()
Definition: HashBackendBuffer.cpp:118
Tensor GetKeyBuffer() const
Return the key buffer tensor.
Definition: HashBackendBuffer.cpp:129
std::vector< int64_t > GetValueDsizes() const
Return value's data sizes in bytes.
Definition: HashBackendBuffer.cpp:100
int64_t GetCommonBlockSize() const
Get the common block size divisor of all values types.
Definition: HashBackendBuffer.cpp:108
std::vector< Tensor > GetValueBuffers() const
Return the value buffer tensors.
Definition: HashBackendBuffer.cpp:131
std::vector< int64_t > GetValueBlocksPerElement() const
Return value's data sizes in the unit of common block size divisor.
Definition: HashBackendBuffer.cpp:112
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:96
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
T * GetDataPtr()
Definition: Tensor.h:1149
void Synchronize()
Definition: CUDAUtils.cpp:77
uint32_t buf_index_t
Definition: HashBackendBuffer.h:63
Definition: PinholeCameraIntrinsic.cpp:35
Tensor cuda
Definition: HashBackendBuffer.h:68