Back to home page

EIC code displayed by LXR

 
 

    


File indexing completed on 2026-11-05 10:28:24

0001 //------------------------------- -*- C++ -*- -------------------------------//
0002 // Copyright Celeritas contributors: see top-level COPYRIGHT file for details
0003 // SPDX-License-Identifier: (Apache-2.0 OR MIT)
0004 //---------------------------------------------------------------------------//
0005 //! \file corecel/data/DeviceVector.hh
0006 //---------------------------------------------------------------------------//
0007 #pragma once
0008 
0009 #include "corecel/Types.hh"
0010 #include "corecel/cont/InitializedValue.hh"
0011 #include "corecel/cont/Span.hh"
0012 #include "corecel/sys/ThreadId.hh"
0013 
0014 #include "DeviceAllocation.hh"
0015 #include "ObserverPtr.hh"
0016 
0017 namespace celeritas
0018 {
0019 //---------------------------------------------------------------------------//
0020 /*!
0021  * Host vector for managing uninitialized device-storage data.
0022  *
0023  * This is a class used only in host memory (not passed to kernels) to manage
0024  * device allocation and host/device copies.  It does \em not perform
0025  * initialization on the data: the host code must define and copy over suitable
0026  * data.
0027  *
0028  * For more complex data usage (dynamic size increases, std vector-like access,
0029  * object initialization), use \c thrust::device_vector inside a \c .cu file.
0030  *
0031  * When a \c StreamId is passed as the last constructor argument,
0032  * all memory operations are asynchronous and ordered within that stream.
0033  *
0034  * \code
0035     DeviceVector<double> myvec(100);
0036     myvec.copy_to_device(make_span(hostvec));
0037     myvec.copy_to_host(make_span(hostvec));
0038    \endcode
0039  *
0040  * \todo move to detail since this is basically only a backend for Collection
0041  */
0042 template<class T>
0043 class DeviceVector
0044 {
0045     static_assert(is_trivially_copyable_v<T>,
0046                   "DeviceVector element is not trivially copyable");
0047 
0048   public:
0049     //!@{
0050     //! \name Type aliases
0051     using value_type = T;
0052     using SpanT = Span<T>;
0053     using SpanConstT = Span<T const>;
0054     //!@}
0055 
0056   public:
0057     // Construct with no elements
0058     DeviceVector() = default;
0059 
0060     // Construct with no elements
0061     explicit DeviceVector(StreamId stream);
0062 
0063     // Construct with a number of elements
0064     explicit DeviceVector(size_type count);
0065 
0066     // Construct with a number of elements
0067     DeviceVector(size_type count, StreamId stream);
0068 
0069     // Swap with another vector
0070     inline void swap(DeviceVector& other) noexcept;
0071 
0072     // Allocate and copy from host pointers
0073     void assign(T const* first, T const* last);
0074 
0075     //// ACCESSORS ////
0076 
0077     //! Get the number of elements
0078     size_type size() const { return size_; }
0079 
0080     //! Whether any elements are stored
0081     bool empty() const { return size_ == 0; }
0082 
0083     //// DEVICE ACCESSORS ////
0084 
0085     // Copy data to device
0086     inline void copy_to_device(SpanConstT host_data);
0087 
0088     // Copy data to host
0089     inline void copy_to_host(SpanT host_data) const;
0090 
0091     // Get a mutable view to device data
0092     SpanT device_ref() { return {this->data(), this->size()}; }
0093 
0094     // Get a const view to device data
0095     SpanConstT device_ref() const { return {this->data(), this->size()}; }
0096 
0097     // Raw pointer to device data (dangerous!)
0098     inline T* data();
0099 
0100     // Raw pointer to device data (dangerous!)
0101     inline T const* data() const;
0102 
0103   private:
0104     DeviceAllocation allocation_;
0105     InitializedValue<size_type> size_;
0106 };
0107 
0108 // Swap two vectors.
0109 template<class T>
0110 inline void swap(DeviceVector<T>& a, DeviceVector<T>& b) noexcept;
0111 
0112 //---------------------------------------------------------------------------//
0113 // INLINE DEFINITIONS
0114 //---------------------------------------------------------------------------//
0115 /*!
0116  * Construct with a stream.
0117  */
0118 template<class T>
0119 DeviceVector<T>::DeviceVector(StreamId stream) : allocation_{stream}, size_{0}
0120 {
0121 }
0122 
0123 //---------------------------------------------------------------------------//
0124 /*!
0125  * Construct with a number of allocated elements.
0126  */
0127 template<class T>
0128 DeviceVector<T>::DeviceVector(size_type count)
0129     : allocation_{count * sizeof(T)}, size_{count}
0130 {
0131 }
0132 
0133 //---------------------------------------------------------------------------//
0134 /*!
0135  * Construct with a number of allocated elements and a stream.
0136  *
0137  * To make resizing eaasier, the stream may be null.
0138  */
0139 template<class T>
0140 DeviceVector<T>::DeviceVector(size_type count, StreamId stream)
0141 {
0142     if (stream)
0143     {
0144         allocation_ = DeviceAllocation{count * sizeof(T), stream};
0145     }
0146     else
0147     {
0148         allocation_ = DeviceAllocation{count * sizeof(T)};
0149     }
0150     size_ = count;
0151 }
0152 
0153 //---------------------------------------------------------------------------//
0154 /*!
0155  * Get the device data pointer.
0156  */
0157 template<class T>
0158 void DeviceVector<T>::swap(DeviceVector& other) noexcept
0159 {
0160     using std::swap;
0161     swap(size_, other.size_);
0162     swap(allocation_, other.allocation_);
0163 }
0164 
0165 //---------------------------------------------------------------------------//
0166 /*!
0167  * Allocate and copy from \em host pointers.
0168  *
0169  * Not exception safe: if the copy fails, the original contents are lost.
0170  */
0171 template<class T>
0172 void DeviceVector<T>::assign(T const* first, T const* last)
0173 {
0174     auto const new_size = static_cast<size_type>(last - first);
0175     if (new_size > size_ && new_size * sizeof(T) > allocation_.size())
0176     {
0177         // Reallocate
0178         *this = DeviceVector<T>(new_size, allocation_.stream_id());
0179     }
0180     else
0181     {
0182         // Update size to fit capacity
0183         size_ = new_size;
0184     }
0185 
0186     this->copy_to_device({first, new_size});
0187 }
0188 
0189 //---------------------------------------------------------------------------//
0190 /*!
0191  * Copy data to device.
0192  */
0193 template<class T>
0194 void DeviceVector<T>::copy_to_device(SpanConstT data)
0195 {
0196     CELER_EXPECT(data.size() == this->size());
0197     // NOLINTNEXTLINE(cppcoreguidelines-pro-type-reinterpret-cast)
0198     allocation_.copy_to_device({reinterpret_cast<std::byte const*>(data.data()),
0199                                 data.size() * sizeof(T)});
0200 }
0201 
0202 //---------------------------------------------------------------------------//
0203 /*!
0204  * Copy data to host.
0205  */
0206 template<class T>
0207 void DeviceVector<T>::copy_to_host(SpanT data) const
0208 {
0209     CELER_EXPECT(data.size() == this->size());
0210     allocation_.copy_to_host(
0211         // NOLINTNEXTLINE(cppcoreguidelines-pro-type-reinterpret-cast)
0212         {reinterpret_cast<std::byte*>(data.data()), data.size() * sizeof(T)});
0213 }
0214 
0215 //---------------------------------------------------------------------------//
0216 /*!
0217  * Get a device data pointer.
0218  */
0219 template<class T>
0220 T* DeviceVector<T>::data()
0221 {
0222     // NOLINTNEXTLINE(cppcoreguidelines-pro-type-reinterpret-cast)
0223     return reinterpret_cast<T*>(allocation_.device_ref().data());
0224 }
0225 
0226 //---------------------------------------------------------------------------//
0227 /*!
0228  * Get a device data pointer.
0229  */
0230 template<class T>
0231 T const* DeviceVector<T>::data() const
0232 {
0233     // NOLINTNEXTLINE(cppcoreguidelines-pro-type-reinterpret-cast)
0234     return reinterpret_cast<T const*>(allocation_.device_ref().data());
0235 }
0236 
0237 //---------------------------------------------------------------------------//
0238 /*!
0239  * Swap two vectors.
0240  */
0241 template<class T>
0242 void swap(DeviceVector<T>& a, DeviceVector<T>& b) noexcept
0243 {
0244     return a.swap(b);
0245 }
0246 
0247 //---------------------------------------------------------------------------//
0248 /*!
0249  * Prevent accidental construction of Span from a device vector.
0250  *
0251  * Use \c dv.device_ref() to get a span.
0252  */
0253 template<class T>
0254 CELER_FUNCTION Span<T const> make_span(DeviceVector<T> const& dv)
0255 {
0256     static_assert(sizeof(T) == 0, "Cannot 'make_span' from a device vector");
0257     return {dv.data(), dv.size()};
0258 }
0259 
0260 //---------------------------------------------------------------------------//
0261 //! Prevent accidental construction of Span from a device vector.
0262 template<class T>
0263 CELER_FUNCTION Span<T> make_span(DeviceVector<T>& dv)
0264 {
0265     static_assert(sizeof(T) == 0, "Cannot 'make_span' from a device vector");
0266     return {dv.data(), dv.size()};
0267 }
0268 
0269 //---------------------------------------------------------------------------//
0270 //! Create an observer pointer from a device vector.
0271 template<class T>
0272 ObserverPtr<T, MemSpace::device> make_observer(DeviceVector<T>& vec) noexcept
0273 {
0274     return ObserverPtr<T, MemSpace::device>{vec.data()};
0275 }
0276 
0277 //---------------------------------------------------------------------------//
0278 //! Create an observer pointer from a pointer in the native memspace.
0279 template<class T>
0280 ObserverPtr<T const, MemSpace::device> make_observer(
0281     DeviceVector<T> const& vec) noexcept
0282 {
0283     return ObserverPtr<T const, MemSpace::device>{vec.data()};
0284 }
0285 
0286 //---------------------------------------------------------------------------//
0287 }  // namespace celeritas