1#ifndef AMREX_TABLE_DATA_H_
2#define AMREX_TABLE_DATA_H_
3#include <AMReX_Config.H>
18template <
typename T,
typename IDX =
int>
30 requires (std::is_const_v<U>)
39 constexpr Table1D (T* a_p, IDX a_begin, IDX a_end) noexcept
46 explicit operator bool () const noexcept {
return p !=
nullptr; }
50 requires (!std::is_void_v<T>)
52#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
58#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
60 void index_assert (IDX i)
const
62 if (i < begin || i >=
end) {
63 if constexpr (std::is_same_v<IDX,int>) {
69 }
else if constexpr (std::is_same_v<IDX,long>) {
75 }
else if constexpr (std::is_same_v<IDX,long long>) {
86 ss <<
" (" << i <<
") is out of bound ("
95template <
typename T, Order ORDER = Order::F>
108 requires (std::is_const_v<U>)
128 explicit operator bool () const noexcept {
return p !=
nullptr; }
132 requires (!std::is_void_v<T>)
134#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
144#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
146 void index_assert (
int i,
int j)
const
156 std::stringstream ss;
157 ss << " (" << i << "," << j << ") is out of bound ("
159 << "," <<
begin[1] << ":" <<
end[1]-1 << ")";
172 return a_end[0] - a_begin[0];
174 return a_end[1] - a_begin[1];
179template <
typename T, Order ORDER = Order::F>
193 requires (std::is_const_v<U>)
197 stride1(rhs.stride1),
198 stride2(rhs.stride2),
208 stride1( len0(a_begin,a_end)),
209 stride2(stride1*len1(a_begin,a_end)),
215 explicit operator bool () const noexcept {
return p !=
nullptr; }
218 T& operator() (
int i,
int j,
int k)
const noexcept
219 requires (!std::is_void_v<T>)
221#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
224 if constexpr (ORDER == Order::F) {
231#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
233 void index_assert (
int i,
int j,
int k)
const
235 if (i < begin[0] || i >= end[0] ||
236 j < begin[1] || j >= end[1] ||
237 k < begin[2] || k >= end[2]) {
240 i, j, k, begin[0], end[0]-1, begin[1], end[1]-1,
245 std::stringstream ss;
246 ss << " (" << i << "," << j << "," << k << ") is out of bound ("
247 << begin[0] << ":" << end[0]-1
248 << "," << begin[1] << ":" << end[1]-1
249 << "," << begin[2] << ":" << end[2]-1 << ")";
250 amrex::Abort(ss.str());
261 if constexpr (ORDER == Order::F) {
262 return a_end[0] - a_begin[0];
264 return a_end[2] - a_begin[2];
268 static constexpr int len1 (GpuArray<int,3>
const& a_begin,
269 GpuArray<int,3>
const& a_end)
noexcept
271 return a_end[1] - a_begin[1];
275template <
typename T, Order ORDER = Order::F>
290 requires (std::is_const_v<U>)
294 stride1(rhs.stride1),
295 stride2(rhs.stride2),
296 stride3(rhs.stride3),
306 stride1( len0(a_begin,a_end)),
307 stride2(stride1*len1(a_begin,a_end)),
308 stride3(stride2*len2(a_begin,a_end)),
314 explicit operator bool () const noexcept {
return p !=
nullptr; }
317 T& operator() (
int i,
int j,
int k,
int n)
const noexcept
318 requires (!std::is_void_v<T>)
320#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
321 index_assert(i,j,k,n);
323 if constexpr (ORDER == Order::F) {
330#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
332 void index_assert (
int i,
int j,
int k,
int n)
const
334 if (i < begin[0] || i >= end[0] ||
335 j < begin[1] || j >= end[1] ||
336 k < begin[2] || k >= end[2] ||
337 n < begin[3] || n >= end[3]) {
340 i, j, k, n, begin[0], end[0]-1, begin[1], end[1]-1,
341 begin[2], end[2]-1, begin[3], end[3]-1);
345 std::stringstream ss;
346 ss << " (" << i << "," << j << "," << k << "," << n << ") is out of bound ("
347 << begin[0] << ":" << end[0]-1
348 << "," << begin[1] << ":" << end[1]-1
349 << "," << begin[2] << ":" << end[2]-1
350 << "," << begin[3] << ":" << end[3]-1 << ")";
351 amrex::Abort(ss.str());
362 if constexpr (ORDER == Order::F) {
363 return a_end[0] - a_begin[0];
365 return a_end[3] - a_begin[3];
369 static constexpr int len1 (GpuArray<int,4>
const& a_begin,
370 GpuArray<int,4>
const& a_end)
noexcept
372 if constexpr (ORDER == Order::F) {
373 return a_end[1] - a_begin[1];
375 return a_end[2] - a_begin[2];
379 static constexpr int len2 (GpuArray<int,4>
const& a_begin,
380 GpuArray<int,4>
const& a_end)
noexcept
382 if constexpr (ORDER == Order::F) {
383 return a_end[2] - a_begin[2];
385 return a_end[1] - a_begin[1];
428template <
typename T,
int N, Order ORDER = Order::F>
434 template <
class U,
int M, Order O>
friend class TableData;
437 std::conditional_t<N==2, Table2D<T, ORDER>,
438 std::conditional_t<N==3, Table3D<T, ORDER>,
441 std::conditional_t<N==2, Table2D<T const, ORDER>,
442 std::conditional_t<N==3, Table3D<T const, ORDER>,
459 [[nodiscard]] constexpr
int dim () const noexcept {
return N; }
463 [[nodiscard]]
Long size () const noexcept;
465 Array<
int,N> const& lo () const noexcept {
return m_lo; }
469 void clear () noexcept;
471 void copy (
TableData<T,N,ORDER> const& rhs) noexcept;
473 table_type table () noexcept;
474 const_table_type table () const noexcept;
475 const_table_type const_table () const noexcept;
484 Long m_truesize = 0L;
485 bool m_ptr_owner = false;
488template <typename T,
int N,
Order ORDER>
496template <
typename T,
int N, Order ORDER>
502 m_truesize(rhs.m_truesize),
503 m_ptr_owner(rhs.m_ptr_owner)
505 rhs.m_dptr =
nullptr;
506 rhs.m_ptr_owner =
false;
509template <
typename T,
int N, Order ORDER>
515 m_arena = rhs.m_arena;
519 m_truesize = rhs.m_truesize;
520 m_ptr_owner = rhs.m_ptr_owner;
521 rhs.m_dptr =
nullptr;
522 rhs.m_ptr_owner =
false;
527template <
typename T,
int N, Order ORDER>
530 static_assert(std::is_trivially_copyable<T>() &&
531 std::is_trivially_destructible<T>(),
532 "TableData<T,N,ORDER>: T must be trivially copyable and trivially destructible");
533 static_assert(N>=1 && N <=4,
"TableData<T,N,ORDER>: N must be in the range of [1,4]");
537template <
typename T,
int N, Order ORDER>
552 }
else if (m_dptr ==
nullptr || !m_ptr_owner) {
555 }
else if (size() > m_truesize) {
561template <
typename T,
int N, Order ORDER>
566 for (
int i = 0; i < N; ++i) {
567 r *= m_hi[i] - m_lo[i] + 1;
572template <
typename T,
int N, Order ORDER>
585template <
typename T,
int N, Order ORDER>
591 if (m_truesize == 0) {
595 m_dptr =
static_cast<T*
>(this->alloc(m_truesize*
sizeof(T)));
601 template <
typename T, Order>
602 Table1D<T> make_table (T* p, Array<int,1>
const& lo, Array<int,1>
const& hi) {
603 return Table1D<T>(p, lo[0], hi[0]+1);
605 template <
typename T, Order ORDER>
606 Table2D<T,ORDER> make_table (T* p, Array<int,2>
const& lo, Array<int,2>
const& hi) {
607 return Table2D<T,ORDER>(p, {lo[0],lo[1]}, {hi[0]+1,hi[1]+1});
609 template <
typename T, Order ORDER>
610 Table3D<T,ORDER> make_table (T* p, Array<int,3>
const& lo, Array<int,3>
const& hi) {
611 return Table3D<T,ORDER>(p, {lo[0],lo[1],lo[2]}, {hi[0]+1,hi[1]+1,hi[2]+1});
613 template <
typename T, Order ORDER>
614 Table4D<T,ORDER> make_table (T* p, Array<int,4>
const& lo, Array<int,4>
const& hi) {
615 return Table4D<T,ORDER>(p, {lo[0],lo[1],lo[2],lo[3]}, {hi[0]+1,hi[1]+1,hi[2]+1,hi[3]+1});
620template <
typename T,
int N, Order ORDER>
621typename TableData<T,N,ORDER>::table_type
624 return detail::make_table<T,ORDER>(m_dptr, m_lo, m_hi);
627template <
typename T,
int N, Order ORDER>
631 return detail::make_table<T const, ORDER>(m_dptr, m_lo, m_hi);
634template <
typename T,
int N, Order ORDER>
638 return detail::make_table<T const, ORDER>(m_dptr, m_lo, m_hi);
641template <
typename T,
int N, Order ORDER>
645 std::size_t count =
sizeof(T)*size();
646 if (count == 0) {
return; }
649 bool this_on_device = arena()->isManaged() || arena()->isDevice();
650 bool rhs_on_device = rhs.arena()->isManaged() || rhs.arena()->isDevice();
651 if (this_on_device && rhs_on_device) {
652 Gpu::dtod_memcpy_async(m_dptr, rhs.m_dptr, count);
653 }
else if (this_on_device && !rhs_on_device) {
654 Gpu::htod_memcpy_async(m_dptr, rhs.m_dptr, count);
655 }
else if (!this_on_device && rhs_on_device) {
656 Gpu::dtoh_memcpy_async(m_dptr, rhs.m_dptr, count);
660 std::memcpy(m_dptr, rhs.m_dptr, count);
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_RESTRICT
Definition AMReX_Extension.H:32
#define AMREX_DEVICE_PRINTF(...)
Definition AMReX_GpuPrint.H:21
#define AMREX_IF_ON_DEVICE(CODE)
Definition AMReX_GpuQualifiers.H:56
#define AMREX_IF_ON_HOST(CODE)
Definition AMReX_GpuQualifiers.H:58
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
A virtual base class for objects that manage their own dynamic memory allocation.
Definition AMReX_Arena.H:132
Multi-dimensional array class.
Definition AMReX_TableData.H:431
std::conditional_t< N==1, Table1D< T >, std::conditional_t< N==2, Table2D< T, ORDER >, std::conditional_t< N==3, Table3D< T, ORDER >, Table4D< T, ORDER > > > > table_type
Definition AMReX_TableData.H:439
TableData() noexcept=default
Array< int, N > const & hi() const noexcept
Definition AMReX_TableData.H:467
std::conditional_t< N==1, Table1D< T const >, std::conditional_t< N==2, Table2D< T const, ORDER >, std::conditional_t< N==3, Table3D< T const, ORDER >, Table4D< T const, ORDER > > > > const_table_type
Definition AMReX_TableData.H:443
T value_type
Definition AMReX_TableData.H:435
amrex_long Long
Definition AMReX_INT.H:30
std::array< T, N > Array
Definition AMReX_Array.H:26
Definition AMReX_Amr.cpp:50
Order
Definition AMReX_Order.H:7
__host__ __device__ Dim3 begin(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2018
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition AMReX.cpp:241
__host__ __device__ Dim3 end(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2028
Definition AMReX_DataAllocator.H:9
Arena * arena() const noexcept
Definition AMReX_DataAllocator.H:24
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:43
Definition AMReX_TableData.H:20
IDX begin
Definition AMReX_TableData.H:22
T *__restrict__ p
Definition AMReX_TableData.H:21
__host__ __device__ T & operator()(IDX i) const noexcept
Definition AMReX_TableData.H:49
IDX end
Definition AMReX_TableData.H:23
constexpr Table1D() noexcept=default
__host__ __device__ constexpr Table1D(T *a_p, IDX a_begin, IDX a_end) noexcept
Definition AMReX_TableData.H:39
Definition AMReX_TableData.H:97
__host__ __device__ T & operator()(int i, int j) const noexcept
Definition AMReX_TableData.H:131
T *__restrict__ p
Definition AMReX_TableData.H:98
constexpr Table2D() noexcept=default
__host__ __device__ constexpr Table2D(T *a_p, GpuArray< int, 2 > const &a_begin, GpuArray< int, 2 > const &a_end) noexcept
Definition AMReX_TableData.H:118
GpuArray< int, 2 > end
Definition AMReX_TableData.H:101
GpuArray< int, 2 > begin
Definition AMReX_TableData.H:100
Long stride1
Definition AMReX_TableData.H:99
Definition AMReX_TableData.H:181
__host__ __device__ constexpr Table3D(T *a_p, GpuArray< int, 3 > const &a_begin, GpuArray< int, 3 > const &a_end) noexcept
Definition AMReX_TableData.H:204
constexpr Table3D() noexcept=default
Definition AMReX_TableData.H:277
constexpr Table4D() noexcept=default
__host__ __device__ constexpr Table4D(T *a_p, GpuArray< int, 4 > const &a_begin, GpuArray< int, 4 > const &a_end) noexcept
Definition AMReX_TableData.H:302