1#ifndef AMREX_TABLE_DATA_H_
2#define AMREX_TABLE_DATA_H_
3#include <AMReX_Config.H>
18template <
typename T,
typename IDX =
int>
28 requires (std::is_const_v<U>)
37 constexpr Table1D (T* a_p, IDX a_begin, IDX a_end) noexcept
44 explicit operator bool () const noexcept {
return p !=
nullptr; }
48 requires (!std::is_void_v<T>)
50#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
56#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
58 void index_assert (IDX i)
const
60 if (i < begin || i >=
end) {
61 if constexpr (std::is_same_v<IDX,int>) {
67 }
else if constexpr (std::is_same_v<IDX,long>) {
73 }
else if constexpr (std::is_same_v<IDX,long long>) {
84 ss <<
" (" << i <<
") is out of bound ("
93template <
typename T, Order ORDER = Order::F>
104 requires (std::is_const_v<U>)
124 explicit operator bool () const noexcept {
return p !=
nullptr; }
128 requires (!std::is_void_v<T>)
130#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
140#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
142 void index_assert (
int i,
int j)
const
152 std::stringstream ss;
153 ss << " (" << i << "," << j << ") is out of bound ("
155 << "," <<
begin[1] << ":" <<
end[1]-1 << ")";
168 return a_end[0] - a_begin[0];
170 return a_end[1] - a_begin[1];
175template <
typename T, Order ORDER = Order::F>
187 requires (std::is_const_v<U>)
191 stride1(rhs.stride1),
192 stride2(rhs.stride2),
202 stride1( len0(a_begin,a_end)),
203 stride2(stride1*len1(a_begin,a_end)),
209 explicit operator bool () const noexcept {
return p !=
nullptr; }
212 T& operator() (
int i,
int j,
int k)
const noexcept
213 requires (!std::is_void_v<T>)
215#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
218 if constexpr (ORDER == Order::F) {
225#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
227 void index_assert (
int i,
int j,
int k)
const
229 if (i < begin[0] || i >= end[0] ||
230 j < begin[1] || j >= end[1] ||
231 k < begin[2] || k >= end[2]) {
234 i, j, k, begin[0], end[0]-1, begin[1], end[1]-1,
239 std::stringstream ss;
240 ss << " (" << i << "," << j << "," << k << ") is out of bound ("
241 << begin[0] << ":" << end[0]-1
242 << "," << begin[1] << ":" << end[1]-1
243 << "," << begin[2] << ":" << end[2]-1 << ")";
244 amrex::Abort(ss.str());
255 if constexpr (ORDER == Order::F) {
256 return a_end[0] - a_begin[0];
258 return a_end[2] - a_begin[2];
262 static constexpr int len1 (GpuArray<int,3>
const& a_begin,
263 GpuArray<int,3>
const& a_end)
noexcept
265 return a_end[1] - a_begin[1];
269template <
typename T, Order ORDER = Order::F>
282 requires (std::is_const_v<U>)
286 stride1(rhs.stride1),
287 stride2(rhs.stride2),
288 stride3(rhs.stride3),
298 stride1( len0(a_begin,a_end)),
299 stride2(stride1*len1(a_begin,a_end)),
300 stride3(stride2*len2(a_begin,a_end)),
306 explicit operator bool () const noexcept {
return p !=
nullptr; }
309 T& operator() (
int i,
int j,
int k,
int n)
const noexcept
310 requires (!std::is_void_v<T>)
312#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
313 index_assert(i,j,k,n);
315 if constexpr (ORDER == Order::F) {
322#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
324 void index_assert (
int i,
int j,
int k,
int n)
const
326 if (i < begin[0] || i >= end[0] ||
327 j < begin[1] || j >= end[1] ||
328 k < begin[2] || k >= end[2] ||
329 n < begin[3] || n >= end[3]) {
332 i, j, k, n, begin[0], end[0]-1, begin[1], end[1]-1,
333 begin[2], end[2]-1, begin[3], end[3]-1);
337 std::stringstream ss;
338 ss << " (" << i << "," << j << "," << k << "," << n << ") is out of bound ("
339 << begin[0] << ":" << end[0]-1
340 << "," << begin[1] << ":" << end[1]-1
341 << "," << begin[2] << ":" << end[2]-1
342 << "," << begin[3] << ":" << end[3]-1 << ")";
343 amrex::Abort(ss.str());
354 if constexpr (ORDER == Order::F) {
355 return a_end[0] - a_begin[0];
357 return a_end[3] - a_begin[3];
361 static constexpr int len1 (GpuArray<int,4>
const& a_begin,
362 GpuArray<int,4>
const& a_end)
noexcept
364 if constexpr (ORDER == Order::F) {
365 return a_end[1] - a_begin[1];
367 return a_end[2] - a_begin[2];
371 static constexpr int len2 (GpuArray<int,4>
const& a_begin,
372 GpuArray<int,4>
const& a_end)
noexcept
374 if constexpr (ORDER == Order::F) {
375 return a_end[2] - a_begin[2];
377 return a_end[1] - a_begin[1];
420template <
typename T,
int N, Order ORDER = Order::F>
426 template <
class U,
int M, Order O>
friend class TableData;
429 std::conditional_t<N==2, Table2D<T, ORDER>,
430 std::conditional_t<N==3, Table3D<T, ORDER>,
433 std::conditional_t<N==2, Table2D<T const, ORDER>,
434 std::conditional_t<N==3, Table3D<T const, ORDER>,
451 [[nodiscard]] constexpr
int dim () const noexcept {
return N; }
455 [[nodiscard]]
Long size () const noexcept;
457 Array<
int,N> const& lo () const noexcept {
return m_lo; }
461 void clear () noexcept;
463 void copy (
TableData<T,N,ORDER> const& rhs) noexcept;
465 table_type table () noexcept;
466 const_table_type table () const noexcept;
467 const_table_type const_table () const noexcept;
476 Long m_truesize = 0L;
477 bool m_ptr_owner = false;
480template <typename T,
int N,
Order ORDER>
488template <
typename T,
int N, Order ORDER>
494 m_truesize(rhs.m_truesize),
495 m_ptr_owner(rhs.m_ptr_owner)
497 rhs.m_dptr =
nullptr;
498 rhs.m_ptr_owner =
false;
501template <
typename T,
int N, Order ORDER>
507 m_arena = rhs.m_arena;
511 m_truesize = rhs.m_truesize;
512 m_ptr_owner = rhs.m_ptr_owner;
513 rhs.m_dptr =
nullptr;
514 rhs.m_ptr_owner =
false;
519template <
typename T,
int N, Order ORDER>
522 static_assert(std::is_trivially_copyable<T>() &&
523 std::is_trivially_destructible<T>(),
524 "TableData<T,N,ORDER>: T must be trivially copyable and trivially destructible");
525 static_assert(N>=1 && N <=4,
"TableData<T,N,ORDER>: N must be in the range of [1,4]");
529template <
typename T,
int N, Order ORDER>
544 }
else if (m_dptr ==
nullptr || !m_ptr_owner) {
547 }
else if (size() > m_truesize) {
553template <
typename T,
int N, Order ORDER>
558 for (
int i = 0; i < N; ++i) {
559 r *= m_hi[i] - m_lo[i] + 1;
564template <
typename T,
int N, Order ORDER>
577template <
typename T,
int N, Order ORDER>
583 if (m_truesize == 0) {
587 m_dptr =
static_cast<T*
>(this->alloc(m_truesize*
sizeof(T)));
593 template <
typename T, Order>
594 Table1D<T> make_table (T* p, Array<int,1>
const& lo, Array<int,1>
const& hi) {
595 return Table1D<T>(p, lo[0], hi[0]+1);
597 template <
typename T, Order ORDER>
598 Table2D<T,ORDER> make_table (T* p, Array<int,2>
const& lo, Array<int,2>
const& hi) {
599 return Table2D<T,ORDER>(p, {lo[0],lo[1]}, {hi[0]+1,hi[1]+1});
601 template <
typename T, Order ORDER>
602 Table3D<T,ORDER> make_table (T* p, Array<int,3>
const& lo, Array<int,3>
const& hi) {
603 return Table3D<T,ORDER>(p, {lo[0],lo[1],lo[2]}, {hi[0]+1,hi[1]+1,hi[2]+1});
605 template <
typename T, Order ORDER>
606 Table4D<T,ORDER> make_table (T* p, Array<int,4>
const& lo, Array<int,4>
const& hi) {
607 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});
612template <
typename T,
int N, Order ORDER>
613typename TableData<T,N,ORDER>::table_type
616 return detail::make_table<T,ORDER>(m_dptr, m_lo, m_hi);
619template <
typename T,
int N, Order ORDER>
623 return detail::make_table<T const, ORDER>(m_dptr, m_lo, m_hi);
626template <
typename T,
int N, Order ORDER>
630 return detail::make_table<T const, ORDER>(m_dptr, m_lo, m_hi);
633template <
typename T,
int N, Order ORDER>
637 std::size_t count =
sizeof(T)*size();
638 if (count == 0) {
return; }
641 bool this_on_device = arena()->isManaged() || arena()->isDevice();
642 bool rhs_on_device = rhs.arena()->isManaged() || rhs.arena()->isDevice();
643 if (this_on_device && rhs_on_device) {
644 Gpu::dtod_memcpy_async(m_dptr, rhs.m_dptr, count);
645 }
else if (this_on_device && !rhs_on_device) {
646 Gpu::htod_memcpy_async(m_dptr, rhs.m_dptr, count);
647 }
else if (!this_on_device && rhs_on_device) {
648 Gpu::dtoh_memcpy_async(m_dptr, rhs.m_dptr, count);
652 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:423
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:431
TableData() noexcept=default
Array< int, N > const & hi() const noexcept
Definition AMReX_TableData.H:459
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:435
T value_type
Definition AMReX_TableData.H:427
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:47
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:37
Definition AMReX_TableData.H:95
__host__ __device__ T & operator()(int i, int j) const noexcept
Definition AMReX_TableData.H:127
T *__restrict__ p
Definition AMReX_TableData.H:96
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:114
GpuArray< int, 2 > end
Definition AMReX_TableData.H:99
GpuArray< int, 2 > begin
Definition AMReX_TableData.H:98
Long stride1
Definition AMReX_TableData.H:97
Definition AMReX_TableData.H:177
__host__ __device__ constexpr Table3D(T *a_p, GpuArray< int, 3 > const &a_begin, GpuArray< int, 3 > const &a_end) noexcept
Definition AMReX_TableData.H:198
constexpr Table3D() noexcept=default
Definition AMReX_TableData.H:271
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:294