Open3D (C++ API)  0.15.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 
34 #include "open3d/core/CUDAUtils.h"
39 
40 namespace open3d {
41 namespace core {
42 
44 public:
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 
98  __device__ buf_index_t DeviceAllocate() {
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 
114 public:
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 
126  int64_t *value_dsizes_;
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
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 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
__device__ void * GetKeyPtr(buf_index_t ptr)
Definition: CUDAHashBackendBufferAccessor.h:107
__device__ void * GetValuePtr(buf_index_t ptr, int value_idx=0)
Definition: CUDAHashBackendBufferAccessor.h:110
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:39
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:86
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:1108
void Synchronize()
Definition: CUDAUtils.cpp:78
uint32_t buf_index_t
Definition: HashBackendBuffer.h:63
Definition: PinholeCameraIntrinsic.cpp:35
Tensor cuda
Definition: HashBackendBuffer.h:68