12#ifndef MFEM_FORALL_HPP
13#define MFEM_FORALL_HPP
23#include <_hypre_utilities.h>
48struct DofQuadLimits_CUDA
50 static constexpr int MAX_D1D = 14;
51 static constexpr int MAX_Q1D = 14;
52 static constexpr int HCURL_MAX_D1D = 5;
53 static constexpr int HCURL_MAX_Q1D = 6;
54 static constexpr int HDIV_MAX_D1D = 5;
55 static constexpr int HDIV_MAX_Q1D = 6;
56 static constexpr int MAX_INTERP_1D = 8;
57 static constexpr int MAX_DET_1D = 6;
60struct DofQuadLimits_HIP
62 static constexpr int MAX_D1D = 10;
63 static constexpr int MAX_Q1D = 10;
64 static constexpr int HCURL_MAX_D1D = 5;
65 static constexpr int HCURL_MAX_Q1D = 5;
66 static constexpr int HDIV_MAX_D1D = 5;
67 static constexpr int HDIV_MAX_Q1D = 6;
68 static constexpr int MAX_INTERP_1D = 8;
69 static constexpr int MAX_DET_1D = 6;
72struct DofQuadLimits_CPU
75 static constexpr int MAX_D1D = 24;
76 static constexpr int MAX_Q1D = 24;
78 static constexpr int MAX_D1D = 14;
79 static constexpr int MAX_Q1D = 14;
81 static constexpr int HCURL_MAX_D1D = 10;
82 static constexpr int HCURL_MAX_Q1D = 10;
83 static constexpr int HDIV_MAX_D1D = 10;
84 static constexpr int HDIV_MAX_Q1D = 10;
85 static constexpr int MAX_INTERP_1D = MAX_D1D;
86 static constexpr int MAX_DET_1D = MAX_D1D;
99#if defined(__CUDA_ARCH__)
101#elif defined(__HIP_DEVICE_COMPILE__)
131 return dof_quad_limits;
140 else { Populate<internal::DofQuadLimits_CPU>(); }
147 template <
typename T>
void Populate()
161#define MFEM_PRAGMA(X) _Pragma(#X)
164#if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__)
165#define MFEM_UNROLL(N) MFEM_PRAGMA(unroll(N))
167#define MFEM_UNROLL(N)
173#if defined(MFEM_USE_CUDA)
174#define MFEM_GPU_FORALL(i, N,...) CuWrap1D(N, [=] MFEM_DEVICE \
175 (int i) {__VA_ARGS__})
176#elif defined(MFEM_USE_HIP)
177#define MFEM_GPU_FORALL(i, N,...) HipWrap1D(N, [=] MFEM_DEVICE \
178 (int i) {__VA_ARGS__})
180#define MFEM_GPU_FORALL(i, N,...) do { } while (false)
187#define MFEM_FORALL(i,N,...) \
188 ForallWrap<1>(true,N,[=] MFEM_HOST_DEVICE (int i) {__VA_ARGS__})
191#define MFEM_FORALL_2D(i,N,X,Y,BZ,...) \
192 ForallWrap<2>(true,N,[=] MFEM_HOST_DEVICE (int i) {__VA_ARGS__},X,Y,BZ)
195#define MFEM_FORALL_3D(i,N,X,Y,Z,...) \
196 ForallWrap<3>(true,N,[=] MFEM_HOST_DEVICE (int i) {__VA_ARGS__},X,Y,Z)
200#define MFEM_FORALL_3D_GRID(i,N,X,Y,Z,G,...) \
201 ForallWrap<3>(true,N,[=] MFEM_HOST_DEVICE (int i) {__VA_ARGS__},X,Y,Z,G)
206#define MFEM_FORALL_SWITCH(use_dev,i,N,...) \
207 ForallWrap<1>(use_dev,N,[=] MFEM_HOST_DEVICE (int i) {__VA_ARGS__})
211template <
typename HBODY>
214#ifdef MFEM_USE_OPENMP
215 #pragma omp parallel for
216 for (
int k = 0; k < N; k++)
221 MFEM_CONTRACT_VAR(N);
222 MFEM_CONTRACT_VAR(h_body);
223 MFEM_ABORT(
"OpenMP requested for MFEM but OpenMP is not enabled!");
229#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_CUDA)
231 RAJA::LaunchPolicy<RAJA::cuda_launch_t<true>>;
233 RAJA::LoopPolicy<RAJA::cuda_block_x_direct>;
235 RAJA::LoopPolicy<RAJA::cuda_thread_z_direct>;
238#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_HIP)
240 RAJA::LaunchPolicy<RAJA::hip_launch_t<true>>;
242 RAJA::LoopPolicy<RAJA::hip_block_x_direct>;
244 RAJA::LoopPolicy<RAJA::hip_thread_z_direct>;
247#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_CUDA)
248template <const
int BLOCKS = MFEM_CUDA_BLOCKS,
typename DBODY>
252 RAJA::forall<RAJA::cuda_exec<BLOCKS,true>>(RAJA::RangeSegment(0,N),d_body);
255template <
typename DBODY>
257 const int X,
const int Y,
const int BZ)
259 MFEM_VERIFY(BZ>0,
"");
260 const int G = (N+BZ-1)/BZ;
262 using namespace RAJA;
263 using RAJA::RangeSegment;
265 launch<cuda_launch_policy>
266 (LaunchParams(Teams(G), Threads(X, Y, BZ)),
267 [=] RAJA_DEVICE (LaunchContext
ctx)
270 loop<cuda_teams_x>(
ctx, RangeSegment(0, G), [&] (
const int n)
273 loop<cuda_threads_z>(
ctx, RangeSegment(0, BZ), [&] (
const int tz)
276 const int k = n*BZ + tz;
277 if (k >= N) {
return; }
286 MFEM_GPU_CHECK(cudaGetLastError());
289template <
typename DBODY>
291 const int X,
const int Y,
const int Z,
const int G)
293 const int GRID = G == 0 ? N : G;
294 using namespace RAJA;
295 using RAJA::RangeSegment;
297 launch<cuda_launch_policy>
298 (LaunchParams(Teams(GRID), Threads(X, Y, Z)),
299 [=] RAJA_DEVICE (LaunchContext
ctx)
302 loop<cuda_teams_x>(
ctx, RangeSegment(0, N), d_body);
306 MFEM_GPU_CHECK(cudaGetLastError());
315 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
316 static void run(
const int N, DBODY &&d_body,
317 const int X,
const int Y,
const int Z,
const int G)
326 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
327 static void run(
const int N, DBODY &&d_body,
328 const int X,
const int Y,
const int Z,
const int G)
337 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
338 static void run(
const int N, DBODY &&d_body,
339 const int X,
const int Y,
const int Z,
const int G)
347#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_HIP)
348template <const
int BLOCKS = MFEM_HIP_BLOCKS,
typename DBODY>
352 RAJA::forall<RAJA::hip_exec<BLOCKS,true>>(RAJA::RangeSegment(0,N),d_body);
355template <
typename DBODY>
357 const int X,
const int Y,
const int BZ)
359 MFEM_VERIFY(BZ>0,
"");
360 const int G = (N+BZ-1)/BZ;
362 using namespace RAJA;
363 using RAJA::RangeSegment;
365 launch<hip_launch_policy>
366 (LaunchParams(Teams(G), Threads(X, Y, BZ)),
367 [=] RAJA_DEVICE (LaunchContext
ctx)
370 loop<hip_teams_x>(
ctx, RangeSegment(0, G), [&] (
const int n)
373 loop<hip_threads_z>(
ctx, RangeSegment(0, BZ), [&] (
const int tz)
376 const int k = n*BZ + tz;
377 if (k >= N) {
return; }
386 MFEM_GPU_CHECK(hipGetLastError());
389template <
typename DBODY>
391 const int X,
const int Y,
const int Z,
const int G)
393 const int GRID = G == 0 ? N : G;
394 using namespace RAJA;
395 using RAJA::RangeSegment;
397 launch<hip_launch_policy>
398 (LaunchParams(Teams(GRID), Threads(X, Y, Z)),
399 [=] RAJA_DEVICE (LaunchContext
ctx)
402 loop<hip_teams_x>(
ctx, RangeSegment(0, N), d_body);
406 MFEM_GPU_CHECK(hipGetLastError());
415 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
416 static void run(
const int N, DBODY &&d_body,
417 const int X,
const int Y,
const int Z,
const int G)
426 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
427 static void run(
const int N, DBODY &&d_body,
428 const int X,
const int Y,
const int Z,
const int G)
437 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
438 static void run(
const int N, DBODY &&d_body,
439 const int X,
const int Y,
const int Z,
const int G)
448#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_OPENMP)
450template <
typename HBODY>
453 RAJA::forall<RAJA::omp_parallel_for_exec>(RAJA::RangeSegment(0,N), h_body);
460template <
typename HBODY>
465#if (RAJA_VERSION_MAJOR >= 2023)
468 using raja_forall_pol = RAJA::seq_exec;
470 using raja_forall_pol = RAJA::loop_exec;
473 RAJA::forall<raja_forall_pol>(RAJA::RangeSegment(0,N), h_body);
475 MFEM_CONTRACT_VAR(N);
476 MFEM_CONTRACT_VAR(h_body);
477 MFEM_ABORT(
"RAJA requested but RAJA is not enabled!");
485template <
typename BODY> __global__
static
486void CuKernel1D(
const int N, BODY body)
488 const int k = blockDim.x*blockIdx.x + threadIdx.x;
489 if (k >= N) {
return; }
493template <
typename BODY> __global__
static
494void CuKernel2D(
const int N, BODY body)
496 const int k = blockIdx.x*blockDim.z + threadIdx.z;
497 if (k >= N) {
return; }
501template <
typename BODY> __global__
static
502void CuKernel3D(
const int N, BODY body)
504 for (
int k = blockIdx.x; k < N; k += gridDim.x) { body(k); }
507template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
510 if (N==0) {
return; }
511 const int GRID = (N+BLCK-1)/BLCK;
512 CuKernel1D<<<GRID,BLCK>>>(N, d_body);
513 MFEM_GPU_CHECK(cudaGetLastError());
516template <
typename DBODY>
518 const int X,
const int Y,
const int BZ)
520 if (N==0) {
return; }
521 MFEM_VERIFY(BZ>0,
"");
522 const int GRID = (N+BZ-1)/BZ;
523 const dim3 BLCK(X,Y,BZ);
524 CuKernel2D<<<GRID,BLCK>>>(N,d_body);
525 MFEM_GPU_CHECK(cudaGetLastError());
528template <
typename DBODY>
530 const int X,
const int Y,
const int Z,
const int G)
532 if (N==0) {
return; }
533 const int GRID = G == 0 ? N : G;
534 const dim3 BLCK(X,Y,Z);
535 CuKernel3D<<<GRID,BLCK>>>(N,d_body);
536 MFEM_GPU_CHECK(cudaGetLastError());
545 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
546 static void run(
const int N, DBODY &&d_body,
547 const int X,
const int Y,
const int Z,
const int G)
556 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
557 static void run(
const int N, DBODY &&d_body,
558 const int X,
const int Y,
const int Z,
const int G)
567 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
568 static void run(
const int N, DBODY &&d_body,
569 const int X,
const int Y,
const int Z,
const int G)
581template <
typename BODY> __global__
static
582void HipKernel1D(
const int N, BODY body)
584 const int k = hipBlockDim_x*hipBlockIdx_x + hipThreadIdx_x;
585 if (k >= N) {
return; }
589template <
typename BODY> __global__
static
590void HipKernel2D(
const int N, BODY body)
592 const int k = hipBlockIdx_x*hipBlockDim_z + hipThreadIdx_z;
593 if (k >= N) {
return; }
597template <
typename BODY> __global__
static
598void HipKernel3D(
const int N, BODY body)
600 for (
int k = hipBlockIdx_x; k < N; k += hipGridDim_x) { body(k); }
603template <const
int BLCK = MFEM_HIP_BLOCKS,
typename DBODY>
606 if (N==0) {
return; }
607 const int GRID = (N+BLCK-1)/BLCK;
608 hipLaunchKernelGGL(HipKernel1D,GRID,BLCK,0,
nullptr,N,d_body);
609 MFEM_GPU_CHECK(hipGetLastError());
612template <
typename DBODY>
614 const int X,
const int Y,
const int BZ)
616 if (N==0) {
return; }
617 const int GRID = (N+BZ-1)/BZ;
618 const dim3 BLCK(X,Y,BZ);
619 hipLaunchKernelGGL(HipKernel2D,GRID,BLCK,0,
nullptr,N,d_body);
620 MFEM_GPU_CHECK(hipGetLastError());
623template <
typename DBODY>
625 const int X,
const int Y,
const int Z,
const int G)
627 if (N==0) {
return; }
628 const int GRID = G == 0 ? N : G;
629 const dim3 BLCK(X,Y,Z);
630 hipLaunchKernelGGL(HipKernel3D,GRID,BLCK,0,
nullptr,N,d_body);
631 MFEM_GPU_CHECK(hipGetLastError());
640 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
641 static void run(
const int N, DBODY &&d_body,
642 const int X,
const int Y,
const int Z,
const int G)
651 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
652 static void run(
const int N, DBODY &&d_body,
653 const int X,
const int Y,
const int Z,
const int G)
662 template <const
int BLCK = MFEM_CUDA_BLOCKS,
typename DBODY>
663 static void run(
const int N, DBODY &&d_body,
664 const int X,
const int Y,
const int Z,
const int G)
674template <const
int DIM,
typename d_lambda,
typename h_lambda>
676 d_lambda &&d_body, h_lambda &&h_body,
677 const int X=0,
const int Y=0,
const int Z=0,
680 MFEM_CONTRACT_VAR(X);
681 MFEM_CONTRACT_VAR(Y);
682 MFEM_CONTRACT_VAR(Z);
683 MFEM_CONTRACT_VAR(G);
684 MFEM_CONTRACT_VAR(d_body);
685 if (!use_dev) {
goto backend_cpu; }
687#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_CUDA)
695#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_HIP)
722#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_OPENMP)
727#ifdef MFEM_USE_OPENMP
741 for (
int k = 0; k < N; k++) { h_body(k); }
744template <const
int DIM,
typename lambda>
745inline void ForallWrap(
const bool use_dev,
const int N, lambda &&body,
746 const int X=0,
const int Y=0,
const int Z=0,
752template<
typename lambda>
755template<
typename lambda>
761template<
typename lambda>
762inline void forall_2D(
int N,
int X,
int Y, lambda &&body)
767template<
typename lambda>
773template<
typename lambda>
774inline void forall_3D(
int N,
int X,
int Y,
int Z, lambda &&body)
779template<
typename lambda>
790template<
typename lambda>
793#ifdef HYPRE_USING_OPENMP
794 #pragma omp parallel for HYPRE_SMP_SCHEDULE
796 for (
int i = 0; i < N; i++) { body(i); }
801#if defined(HYPRE_USING_GPU)
802template<
typename lambda>
805#if defined(HYPRE_USING_CUDA)
807#elif defined(HYPRE_USING_HIP)
810#error Unknown HYPRE GPU backend!
821template<
typename lambda>
824#if !defined(HYPRE_USING_GPU)
826#elif MFEM_HYPRE_VERSION < 23100
861template<
class B,
class R>
struct reduction_kernel
864 using value_type =
typename R::value_type;
866 mutable value_type *work;
872 int items_per_thread;
874 constexpr static MFEM_HOST_DEVICE
int max_blocksize() {
return 256; }
877 static int block_log2(
unsigned N)
879#if defined(__GNUC__) or defined(__clang__)
880 return N ? (
sizeof(unsigned) * 8 - __builtin_clz(N)) : 0;
881#elif defined(_MSC_VER)
882 return sizeof(unsigned) * 8 - __lzclz(N);
894 MFEM_HOST_DEVICE
void operator()(
int work_idx)
const
896 MFEM_SHARED value_type buffer[max_blocksize()];
897 reducer.SetInitialValue(buffer[MFEM_THREAD_ID(x)]);
899 for (
int idx = 0; idx < items_per_thread; ++idx)
901 int i = MFEM_THREAD_ID(x) +
902 (idx + work_idx * items_per_thread) * MFEM_THREAD_SIZE(x);
905 body(i, buffer[MFEM_THREAD_ID(x)]);
913 for (
int i = (MFEM_THREAD_SIZE(x) >> 1); i > 0; i >>= 1)
916 if (MFEM_THREAD_ID(x) < i)
918 reducer.Join(buffer[MFEM_THREAD_ID(x)], buffer[MFEM_THREAD_ID(x) + i]);
921 if (MFEM_THREAD_ID(x) == 0)
923 work[work_idx] = buffer[0];
941template <
class T,
class B,
class R>
942void reduce(
int N, T &res, B &&body,
const R &reducer,
bool use_dev,
950#if defined(MFEM_USE_HIP) || defined(MFEM_USE_CUDA)
955 using red_type = internal::reduction_kernel<typename std::decay<B>::type,
956 typename std::decay<R>::type>;
958 int block_size = std::min<int>(red_type::max_blocksize(),
959 1ll << red_type::block_log2(N));
962#if defined(MFEM_USE_CUDA)
964 constexpr int mp_sat = 8;
965#elif defined(MFEM_USE_HIP)
967 constexpr int mp_sat = 4;
970 constexpr int mp_sat = 1;
974 int nblocks = std::min(mp_sat * num_mp, (N + block_size - 1) / block_size);
975 int items_per_thread =
976 (N + block_size * nblocks - 1) / (block_size * nblocks);
978 red_type red{
nullptr, std::forward<B>(body), reducer, N, items_per_thread};
980 auto mt = workspace.
GetMemory().GetMemoryType();
985 workspace.
SetSize(nblocks, mt);
988 forall_2D(nblocks, block_size, 1, std::move(red));
991 for (
int i = 0; i < nblocks; ++i)
993 reducer.Join(res, work[i]);
999 for (
int i = 0; i < N; ++i)
Memory< T > & GetMemory()
Return a reference to the Memory object used by the Array.
void SetSize(int nsize)
Change the logical size of the array, keep existing entries.
T * HostWrite()
Shortcut for mfem::Write(a.GetMemory(), a.Size(), false).
static int NumMultiprocessors()
Same as NumMultiprocessors(int), for the currently active device.
static bool Allows(unsigned long b_mask)
Return true if any of the backends in the backend mask, b_mask, are allowed.
static int GetId()
Get the device id of the configured device.
RAJA::LaunchPolicy< RAJA::cuda_launch_t< true > > cuda_launch_policy
RAJA Cuda and Hip backends.
void RajaOmpWrap(const int N, HBODY &&h_body)
RAJA OpenMP backend.
RAJA::LaunchPolicy< RAJA::hip_launch_t< true > > hip_launch_policy
void RajaSeqWrap(const int N, HBODY &&h_body)
RAJA sequential loop backend.
MemoryClass GetHypreForallMemoryClass()
void RajaHipWrap2D(const int N, DBODY &&d_body, const int X, const int Y, const int BZ)
void hypre_forall_cpu(int N, lambda &&body)
void reduce(int N, T &res, B &&body, const R &reducer, bool use_dev, Array< T > &workspace)
Performs a 1D reduction on the range [0,N). res initial value and where the result will be written....
void RajaCuWrap2D(const int N, DBODY &&d_body, const int X, const int Y, const int BZ)
RAJA::LoopPolicy< RAJA::cuda_thread_z_direct > cuda_threads_z
void CuWrap1D(const int N, DBODY &&d_body)
MemoryClass
Memory classes identify sets of memory types.
void RajaCuWrap3D(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
void RajaHipWrap1D(const int N, DBODY &&d_body)
void forall_2D_batch(int N, int X, int Y, int BZ, lambda &&body)
void CuWrap2D(const int N, DBODY &&d_body, const int X, const int Y, const int BZ)
void hypre_forall_gpu(int N, lambda &&body)
internal::DofQuadLimits_CUDA DofQuadLimits
Maximum number of 1D DOFs or quadrature points for the architecture currently being compiled for (use...
void ForallWrap(const bool use_dev, const int N, d_lambda &&d_body, h_lambda &&h_body, const int X=0, const int Y=0, const int Z=0, const int G=0)
The forall kernel body wrapper.
void forall_2D(int N, int X, int Y, lambda &&body)
void HipWrap3D(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
RAJA::LoopPolicy< RAJA::hip_thread_z_direct > hip_threads_z
void CuWrap3D(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
void HipWrap2D(const int N, DBODY &&d_body, const int X, const int Y, const int BZ)
void forall_3D(int N, int X, int Y, int Z, lambda &&body)
void hypre_forall(int N, lambda &&body)
void OmpWrap(const int N, HBODY &&h_body)
OpenMP backend.
bool HypreUsingGPU()
Return true if HYPRE is configured to use GPU.
void forall_3D_grid(int N, int X, int Y, int Z, int G, lambda &&body)
RAJA::LoopPolicy< RAJA::hip_block_x_direct > hip_teams_x
void RajaCuWrap1D(const int N, DBODY &&d_body)
@ HOST_PINNED
Host memory: pinned (page-locked)
void HipWrap1D(const int N, DBODY &&d_body)
void forall(int N, lambda &&body)
void forall_switch(bool use_dev, int N, lambda &&body)
RAJA::LoopPolicy< RAJA::cuda_block_x_direct > cuda_teams_x
void RajaHipWrap3D(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
struct s_NavierContext ctx
@ RAJA_OMP
[host] RAJA OpenMP backend. Enabled when MFEM_USE_RAJA = YES and MFEM_USE_OPENMP = YES.
@ RAJA_CUDA
[device] RAJA CUDA backend. Enabled when MFEM_USE_RAJA = YES and MFEM_USE_CUDA = YES.
@ DEBUG_DEVICE
[device] Debug backend: host memory is READ/WRITE protected while a device is in use....
@ RAJA_CPU
[host] RAJA CPU backend: sequential execution on each MPI rank. Enabled when MFEM_USE_RAJA = YES.
@ OMP
[host] OpenMP backend. Enabled when MFEM_USE_OPENMP = YES.
@ HIP
[device] HIP backend. Enabled when MFEM_USE_HIP = YES.
@ RAJA_HIP
[device] RAJA HIP backend. Enabled when MFEM_USE_RAJA = YES and MFEM_USE_HIP = YES.
@ CUDA
[device] CUDA backend. Enabled when MFEM_USE_CUDA = YES.
@ HIP_MASK
Biwise-OR of all HIP backends.
@ CUDA_MASK
Biwise-OR of all CUDA backends.
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
Maximum number of 1D DOFs or quadrature points for the current runtime configuration of the Device (u...
static const DeviceDofQuadLimits & Get()
Return a const reference to the DeviceDofQuadLimits singleton.
int HCURL_MAX_D1D
Maximum number of 1D nodal points for H(curl).
int HCURL_MAX_Q1D
Maximum number of 1D quadrature points for H(curl).
int HDIV_MAX_Q1D
Maximum number of 1D quadrature points for H(div).
int MAX_INTERP_1D
Maximum number of points for use in QuadratureInterpolator.
int HDIV_MAX_D1D
Maximum number of 1D nodal points for H(div).
int MAX_DET_1D
Maximum number of points for determinant computation in QuadratureInterpolator.
int MAX_D1D
Maximum number of 1D nodal points.
int MAX_Q1D
Maximum number of 1D quadrature points.
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)
static void run(const int N, DBODY &&d_body, const int X, const int Y, const int Z, const int G)