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 <<
")";
168 static constexpr int len0 (GpuArray<int,2>
const& a_begin,
169 GpuArray<int,2>
const& a_end)
noexcept
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>)
215 explicit operator bool () const noexcept {
return p !=
nullptr; }
219 requires (!std::is_void_v<T>)
221#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
231#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
233 void index_assert (
int i,
int j,
int k)
const
245 std::stringstream ss;
246 ss <<
" (" << i <<
"," << j <<
"," << k <<
") is out of bound ("
248 <<
"," <<
begin[1] <<
":" <<
end[1]-1
249 <<
"," <<
begin[2] <<
":" <<
end[2]-1 <<
")";
258 static constexpr int len0 (GpuArray<int,3>
const& a_begin,
259 GpuArray<int,3>
const& a_end)
noexcept
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>)
314 explicit operator bool () const noexcept {
return p !=
nullptr; }
318 requires (!std::is_void_v<T>)
320#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
321 index_assert(i,j,k,n);
330#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
332 void index_assert (
int i,
int j,
int k,
int n)
const
345 std::stringstream ss;
346 ss <<
" (" << i <<
"," << j <<
"," << k <<
"," << n <<
") is out of bound ("
348 <<
"," <<
begin[1] <<
":" <<
end[1]-1
349 <<
"," <<
begin[2] <<
":" <<
end[2]-1
350 <<
"," <<
begin[3] <<
":" <<
end[3]-1 <<
")";
359 static constexpr int len0 (GpuArray<int,4>
const& a_begin,
360 GpuArray<int,4>
const& a_end)
noexcept
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
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
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;
484 Long m_truesize = 0L;
485 bool m_ptr_owner = false;
488template <typename T,
int N,
Order ORDER>
493template <
typename T,
int N, Order ORDER>
501template <
typename T,
int N, Order ORDER>
507 m_truesize(rhs.m_truesize),
508 m_ptr_owner(rhs.m_ptr_owner)
510 rhs.m_dptr =
nullptr;
511 rhs.m_ptr_owner =
false;
514template <
typename T,
int N, Order ORDER>
524 m_truesize = rhs.m_truesize;
525 m_ptr_owner = rhs.m_ptr_owner;
526 rhs.m_dptr =
nullptr;
527 rhs.m_ptr_owner =
false;
532template <
typename T,
int N, Order ORDER>
535 static_assert(std::is_trivially_copyable<T>() &&
536 std::is_trivially_destructible<T>(),
537 "TableData<T,N,ORDER>: T must be trivially copyable and trivially destructible");
538 static_assert(N>=1 && N <=4,
"TableData<T,N,ORDER>: N must be in the range of [1,4]");
542template <
typename T,
int N, Order ORDER>
557 }
else if (m_dptr ==
nullptr || !m_ptr_owner) {
560 }
else if (
size() > m_truesize) {
566template <
typename T,
int N, Order ORDER>
571 for (
int i = 0; i < N; ++i) {
572 r *= m_hi[i] - m_lo[i] + 1;
577template <
typename T,
int N, Order ORDER>
590template <
typename T,
int N, Order ORDER>
592TableData<T,N,ORDER>::define ()
594 const Long n = size();
601 m_dptr =
static_cast<T*
>(this->alloc(n*
sizeof(T)));
608 template <
typename T, Order>
610 return Table1D<T>(
p, lo[0], hi[0]+1);
612 template <
typename T, Order ORDER>
614 return Table2D<T,ORDER>(
p, {lo[0],lo[1]}, {hi[0]+1,hi[1]+1});
616 template <
typename T, Order ORDER>
618 return Table3D<T,ORDER>(
p, {lo[0],lo[1],lo[2]}, {hi[0]+1,hi[1]+1,hi[2]+1});
620 template <
typename T, Order ORDER>
622 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});
627template <
typename T,
int N, Order ORDER>
631 return detail::make_table<T,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 return detail::make_table<T const, ORDER>(m_dptr, m_lo, m_hi);
648template <
typename T,
int N, Order ORDER>
652 std::size_t count =
sizeof(T)*
size();
653 if (count == 0) {
return; }
656 bool this_on_device =
arena()->isManaged() ||
arena()->isDevice();
657 bool rhs_on_device = rhs.arena()->isManaged() || rhs.arena()->isDevice();
658 if (this_on_device && rhs_on_device) {
660 }
else if (this_on_device && !rhs_on_device) {
662 }
else if (!this_on_device && rhs_on_device) {
667 std::memcpy(m_dptr, rhs.m_dptr, count);
Runtime initialization/finalization helpers and global diagnostics.
Fixed-size array types for use on GPU and CPU.
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
Lightweight Arena-backed allocator and deleter helpers.
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
#define AMREX_RESTRICT
Definition AMReX_Extension.H:37
#define AMREX_DEVICE_PRINTF(...)
Definition AMReX_GpuPrint.H:15
#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
while parser_exe_t p
Definition AMReX_Parser_Exe_Body.H:33
Abstract base class for memory arenas.
Definition AMReX_Arena.H:154
Multi-dimensional array class.
Definition AMReX_TableData.H:431
TableData() noexcept=default
const_table_type const_table() const noexcept
Definition AMReX_TableData.H:643
void resize(Array< int, N > const &lo, Array< int, N > const &hi, Arena *ar=nullptr)
Definition AMReX_TableData.H:544
Array< int, N > const & hi() const noexcept
Definition AMReX_TableData.H:467
friend class TableData
Definition AMReX_TableData.H:434
void copy(TableData< T, N, ORDER > const &rhs) noexcept
Definition AMReX_TableData.H:650
constexpr int dim() const noexcept
Definition AMReX_TableData.H:459
Long size() const noexcept
Definition AMReX_TableData.H:568
Array< int, N > const & lo() const noexcept
Definition AMReX_TableData.H:465
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:436
T value_type
Definition AMReX_TableData.H:435
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:440
TableData< T, N, ORDER > & operator=(TableData< T, N, ORDER > const &)=delete
~TableData() noexcept
Definition AMReX_TableData.H:533
void clear() noexcept
Definition AMReX_TableData.H:579
table_type table() noexcept
Definition AMReX_TableData.H:629
amrex_long Long
Definition AMReX_INT.H:30
std::array< T, N > Array
Definition AMReX_Array.H:31
void dtod_memcpy_async(void *p_d_dst, const void *p_d_src, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:468
void dtoh_memcpy_async(void *p_h, const void *p_d, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:454
void htod_memcpy_async(void *p_d, const void *p_h, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:440
Definition AMReX_Amr.cpp:50
Order
Definition AMReX_Order.H:7
@ F
Definition AMReX_Order.H:7
__host__ __device__ Dim3 begin(BoxND< dim > const &box) noexcept
Return the iterator begin coordinate of box as Dim3.
Definition AMReX_Box.H:2241
void Abort(const std::string &msg)
Print a fatal-error message to stderr and abort execution.
Definition AMReX.cpp:244
__host__ __device__ Dim3 end(BoxND< dim > const &box) noexcept
Return the iterator end coordinate of box as Dim3.
Definition AMReX_Box.H:2259
DataAllocator() noexcept=default
Default-construct; alloc/free will use The_Arena().
Arena * arena() const noexcept
Return the stored arena, or The_Arena() if none was set.
Definition AMReX_DataAllocator.H:41
void free(void *pt) const noexcept
Return the allocation at pt to the arena.
Definition AMReX_DataAllocator.H:36
Arena * m_arena
Stored arena pointer (may be null).
Definition AMReX_DataAllocator.H:22
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:52
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
__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
T *__restrict__ p
Definition AMReX_TableData.H:182
Long stride1
Definition AMReX_TableData.H:183
__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
Long stride2
Definition AMReX_TableData.H:184
constexpr Table3D() noexcept=default
GpuArray< int, 3 > end
Definition AMReX_TableData.H:186
GpuArray< int, 3 > begin
Definition AMReX_TableData.H:185
__host__ __device__ T & operator()(int i, int j, int k) const noexcept
Definition AMReX_TableData.H:218
Definition AMReX_TableData.H:277
Long stride3
Definition AMReX_TableData.H:281
constexpr Table4D() noexcept=default
__host__ __device__ T & operator()(int i, int j, int k, int n) const noexcept
Definition AMReX_TableData.H:317
Long stride2
Definition AMReX_TableData.H:280
T *__restrict__ p
Definition AMReX_TableData.H:278
Long stride1
Definition AMReX_TableData.H:279
__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
GpuArray< int, 4 > begin
Definition AMReX_TableData.H:282
GpuArray< int, 4 > end
Definition AMReX_TableData.H:283