device_uvector.hpp
1 /*
2  * Copyright (c) 2020-2022, NVIDIA CORPORATION.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
17 #pragma once
18 
19 #include <rmm/cuda_stream_view.hpp>
20 #include <rmm/detail/error.hpp>
21 #include <rmm/detail/exec_check_disable.hpp>
22 #include <rmm/device_buffer.hpp>
23 #include <rmm/mr/device/device_memory_resource.hpp>
25 
26 #include <cstddef>
27 #include <vector>
28 
29 namespace rmm {
30 
68 template <typename T>
70  static_assert(std::is_trivially_copyable<T>::value,
71  "device_uvector only supports types that are trivially copyable.");
72 
73  public:
74  using value_type = T;
75  using size_type = std::size_t;
76  using reference = value_type&;
77  using const_reference = value_type const&;
78  using pointer = value_type*;
79  using const_pointer = value_type const*;
80  using iterator = pointer;
81  using const_iterator = const_pointer;
82 
83  RMM_EXEC_CHECK_DISABLE
84  ~device_uvector() = default;
85 
86  RMM_EXEC_CHECK_DISABLE
87  device_uvector(device_uvector&&) noexcept = default;
88 
89  device_uvector& operator=(device_uvector&&) noexcept = default;
90 
94  device_uvector(device_uvector const&) = delete;
95 
100 
104  device_uvector() = delete;
105 
117  explicit device_uvector(
118  std::size_t size,
120  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
121  : _storage{elements_to_bytes(size), stream, mr}
122  {
123  }
124 
134  explicit device_uvector(
135  device_uvector const& other,
137  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
138  : _storage{other._storage, stream, mr}
139  {
140  }
141 
150  [[nodiscard]] pointer element_ptr(std::size_t element_index) noexcept
151  {
152  assert(element_index < size());
153  return data() + element_index;
154  }
155 
164  [[nodiscard]] const_pointer element_ptr(std::size_t element_index) const noexcept
165  {
166  assert(element_index < size());
167  return data() + element_index;
168  }
169 
206  void set_element_async(std::size_t element_index,
207  value_type const& value,
209  {
210  RMM_EXPECTS(
211  element_index < size(), rmm::out_of_range, "Attempt to access out of bounds element.");
212 
213  if constexpr (std::is_same<value_type, bool>::value) {
214  RMM_CUDA_TRY(
215  cudaMemsetAsync(element_ptr(element_index), value, sizeof(value), stream.value()));
216  return;
217  }
218 
219  if constexpr (std::is_fundamental<value_type>::value) {
220  if (value == value_type{0}) {
221  set_element_to_zero_async(element_index, stream);
222  return;
223  }
224  }
225 
226  RMM_CUDA_TRY(cudaMemcpyAsync(
227  element_ptr(element_index), &value, sizeof(value), cudaMemcpyDefault, stream.value()));
228  }
229 
230  // We delete the r-value reference overload to prevent asynchronously copying from a literal or
231  // implicit temporary value after it is deleted or goes out of scope.
232  void set_element_async(std::size_t, value_type const&&, cuda_stream_view) = delete;
233 
256  void set_element_to_zero_async(std::size_t element_index, cuda_stream_view stream)
257  {
258  RMM_EXPECTS(
259  element_index < size(), rmm::out_of_range, "Attempt to access out of bounds element.");
260  RMM_CUDA_TRY(
261  cudaMemsetAsync(element_ptr(element_index), 0, sizeof(value_type), stream.value()));
262  }
263 
293  void set_element(std::size_t element_index, T const& value, cuda_stream_view stream)
294  {
295  set_element_async(element_index, value, stream);
297  }
298 
311  [[nodiscard]] value_type element(std::size_t element_index, cuda_stream_view stream) const
312  {
313  RMM_EXPECTS(
314  element_index < size(), rmm::out_of_range, "Attempt to access out of bounds element.");
315  value_type value;
316  RMM_CUDA_TRY(cudaMemcpyAsync(
317  &value, element_ptr(element_index), sizeof(value), cudaMemcpyDefault, stream.value()));
319  return value;
320  }
321 
333  [[nodiscard]] value_type front_element(cuda_stream_view stream) const
334  {
335  return element(0, stream);
336  }
337 
349  [[nodiscard]] value_type back_element(cuda_stream_view stream) const
350  {
351  return element(size() - 1, stream);
352  }
353 
370  void resize(std::size_t new_size, cuda_stream_view stream)
371  {
372  _storage.resize(elements_to_bytes(new_size), stream);
373  }
374 
383 
389  device_buffer release() noexcept { return std::move(_storage); }
390 
397  [[nodiscard]] std::size_t capacity() const noexcept
398  {
399  return bytes_to_elements(_storage.capacity());
400  }
401 
410  [[nodiscard]] pointer data() noexcept { return static_cast<pointer>(_storage.data()); }
411 
420  [[nodiscard]] const_pointer data() const noexcept
421  {
422  return static_cast<const_pointer>(_storage.data());
423  }
424 
432  [[nodiscard]] iterator begin() noexcept { return data(); }
433 
441  [[nodiscard]] const_iterator cbegin() const noexcept { return data(); }
442 
450  [[nodiscard]] const_iterator begin() const noexcept { return cbegin(); }
451 
460  [[nodiscard]] iterator end() noexcept { return data() + size(); }
461 
470  [[nodiscard]] const_iterator cend() const noexcept { return data() + size(); }
471 
480  [[nodiscard]] const_iterator end() const noexcept { return cend(); }
481 
485  [[nodiscard]] std::size_t size() const noexcept { return bytes_to_elements(_storage.size()); }
486 
490  [[nodiscard]] std::int64_t ssize() const noexcept
491  {
492  assert(size() < static_cast<std::size_t>(std::numeric_limits<int64_t>::max()) &&
493  "Size overflows signed integer");
494  return static_cast<int64_t>(size());
495  }
496 
503  [[nodiscard]] bool is_empty() const noexcept { return size() == 0; }
504 
510  [[nodiscard]] mr::device_memory_resource* memory_resource() const noexcept
511  {
512  return _storage.memory_resource();
513  }
514 
518  [[nodiscard]] cuda_stream_view stream() const noexcept { return _storage.stream(); }
519 
529  void set_stream(cuda_stream_view stream) noexcept { _storage.set_stream(stream); }
530 
531  private:
532  device_buffer _storage{};
533 
534  [[nodiscard]] std::size_t constexpr elements_to_bytes(std::size_t num_elements) const noexcept
535  {
536  return num_elements * sizeof(value_type);
537  }
538 
539  [[nodiscard]] std::size_t constexpr bytes_to_elements(std::size_t num_bytes) const noexcept
540  {
541  return num_bytes / sizeof(value_type);
542  }
543 };
544 } // namespace rmm
rmm::device_uvector::data
const_pointer data() const noexcept
Returns const pointer to underlying device storage.
Definition: device_uvector.hpp:420
rmm::device_uvector::set_element_async
void set_element_async(std::size_t element_index, value_type const &value, cuda_stream_view stream)
Performs an asynchronous copy of v to the specified element in device memory.
Definition: device_uvector.hpp:206
rmm::device_uvector::end
iterator end() noexcept
Returns an iterator to the element following the last element of the vector.
Definition: device_uvector.hpp:460
rmm::device_uvector::end
const_iterator end() const noexcept
Returns an iterator to the element following the last element of the vector.
Definition: device_uvector.hpp:480
rmm::device_uvector::element
value_type element(std::size_t element_index, cuda_stream_view stream) const
Returns the specified element from device memory.
Definition: device_uvector.hpp:311
rmm::device_uvector::set_stream
void set_stream(cuda_stream_view stream) noexcept
Sets the stream to be used for deallocation.
Definition: device_uvector.hpp:529
device_buffer.hpp
RAII construct for device memory allocation.
rmm::device_uvector::operator=
device_uvector & operator=(device_uvector const &)=delete
Copy assignment is deleted as it doesn't allow a stream argument.
rmm::device_uvector::begin
const_iterator begin() const noexcept
Returns a const_iterator to the first element.
Definition: device_uvector.hpp:450
rmm::device_buffer::memory_resource
mr::device_memory_resource * memory_resource() const noexcept
Returns pointer to the memory resource used to allocate and deallocate the device memory.
Definition: device_buffer.hpp:368
rmm::device_buffer::shrink_to_fit
void shrink_to_fit(cuda_stream_view stream)
Forces the deallocation of unused memory.
Definition: device_buffer.hpp:295
per_device_resource.hpp
Management of per-device device_memory_resources.
rmm::device_uvector::stream
cuda_stream_view stream() const noexcept
Returns stream most recently specified for allocation/deallocation.
Definition: device_uvector.hpp:518
rmm::device_uvector::device_uvector
device_uvector(std::size_t size, cuda_stream_view stream, rmm::mr::device_memory_resource *mr=rmm::mr::get_current_device_resource())
Construct a new device_uvector with sufficient uninitialized storage for size elements.
Definition: device_uvector.hpp:117
rmm::cuda_stream_view
Strongly-typed non-owning wrapper for CUDA streams with default constructor.
Definition: cuda_stream_view.hpp:34
rmm::device_buffer::size
std::size_t size() const noexcept
Returns the number of bytes.
Definition: device_buffer.hpp:320
rmm::device_uvector::memory_resource
mr::device_memory_resource * memory_resource() const noexcept
Returns pointer to the resource used to allocate and deallocate the device storage.
Definition: device_uvector.hpp:510
rmm::device_uvector::cbegin
const_iterator cbegin() const noexcept
Returns a const_iterator to the first element.
Definition: device_uvector.hpp:441
rmm::device_uvector::resize
void resize(std::size_t new_size, cuda_stream_view stream)
Resizes the vector to contain new_size elements.
Definition: device_uvector.hpp:370
rmm::device_uvector::device_uvector
device_uvector(device_uvector const &)=delete
Copy ctor is deleted as it doesn't allow a stream argument.
rmm::device_uvector::is_empty
bool is_empty() const noexcept
Returns true if the vector contains no elements, i.e., size() == 0.
Definition: device_uvector.hpp:503
rmm::device_uvector
An uninitialized vector of elements in device memory.
Definition: device_uvector.hpp:69
rmm::device_uvector::size
std::size_t size() const noexcept
Returns the number of elements.
Definition: device_uvector.hpp:485
rmm::device_buffer
Definition: device_buffer.hpp:77
rmm::device_uvector::capacity
std::size_t capacity() const noexcept
Returns the number of elements that can be held in currently allocated storage.
Definition: device_uvector.hpp:397
rmm::device_buffer::set_stream
void set_stream(cuda_stream_view stream) noexcept
Sets the stream to be used for deallocation.
Definition: device_buffer.hpp:362
rmm::device_uvector::cend
const_iterator cend() const noexcept
Returns a const_iterator to the element following the last element of the vector.
Definition: device_uvector.hpp:470
rmm::device_uvector::device_uvector
device_uvector(device_uvector const &other, cuda_stream_view stream, rmm::mr::device_memory_resource *mr=rmm::mr::get_current_device_resource())
Construct a new device_uvector by deep copying the contents of another device_uvector.
Definition: device_uvector.hpp:134
rmm::device_uvector::ssize
std::int64_t ssize() const noexcept
Returns the signed number of elements.
Definition: device_uvector.hpp:490
rmm::device_buffer::stream
cuda_stream_view stream() const noexcept
Returns stream most recently specified for allocation/deallocation.
Definition: device_buffer.hpp:351
rmm::device_uvector::device_uvector
device_uvector()=delete
Default constructor is deleted as it doesn't allow a stream argument.
rmm::device_uvector::back_element
value_type back_element(cuda_stream_view stream) const
Returns the last element.
Definition: device_uvector.hpp:349
rmm::device_buffer::resize
void resize(std::size_t new_size, cuda_stream_view stream)
Resize the device memory allocation.
Definition: device_buffer.hpp:264
rmm::device_uvector::release
device_buffer release() noexcept
Release ownership of device memory storage.
Definition: device_uvector.hpp:389
rmm::device_uvector::begin
iterator begin() noexcept
Returns an iterator to the first element.
Definition: device_uvector.hpp:432
rmm::cuda_stream_view::synchronize
void synchronize() const
Synchronize the viewed CUDA stream.
Definition: cuda_stream_view.hpp:81
rmm::device_uvector::shrink_to_fit
void shrink_to_fit(cuda_stream_view stream)
Forces deallocation of unused device memory.
Definition: device_uvector.hpp:382
rmm::device_buffer::data
void const * data() const noexcept
Returns raw pointer to underlying device memory allocation.
Definition: device_buffer.hpp:310
rmm::device_uvector::set_element_to_zero_async
void set_element_to_zero_async(std::size_t element_index, cuda_stream_view stream)
Asynchronously sets the specified element to zero in device memory.
Definition: device_uvector.hpp:256
rmm::device_uvector::element_ptr
const_pointer element_ptr(std::size_t element_index) const noexcept
Returns pointer to the specified element.
Definition: device_uvector.hpp:164
rmm::device_uvector::element_ptr
pointer element_ptr(std::size_t element_index) noexcept
Returns pointer to the specified element.
Definition: device_uvector.hpp:150
rmm::cuda_stream_view::value
constexpr cudaStream_t value() const noexcept
Get the wrapped stream.
Definition: cuda_stream_view.hpp:57
rmm::device_uvector::front_element
value_type front_element(cuda_stream_view stream) const
Returns the first element.
Definition: device_uvector.hpp:333
rmm::cuda_stream_view::synchronize_no_throw
void synchronize_no_throw() const noexcept
Synchronize the viewed CUDA stream. Does not throw if there is an error.
Definition: cuda_stream_view.hpp:88
rmm::mr::device_memory_resource
Base class for all libcudf device memory allocation.
Definition: device_memory_resource.hpp:82
rmm::device_buffer::capacity
std::size_t capacity() const noexcept
Returns actual size in bytes of device memory allocation.
Definition: device_buffer.hpp:346
rmm::out_of_range
Exception thrown when attempting to access outside of a defined range.
Definition: error.hpp:78
rmm::device_uvector::data
pointer data() noexcept
Returns pointer to underlying device storage.
Definition: device_uvector.hpp:410
rmm::device_uvector::set_element
void set_element(std::size_t element_index, T const &value, cuda_stream_view stream)
Performs a synchronous copy of v to the specified element in device memory.
Definition: device_uvector.hpp:293