Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
Macros
AMReX_GpuLaunchMacrosG.nolint.H File Reference

Go to the source code of this file.

Macros

#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE(TN, TI, block)
 
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_2(TN1, TI1, block1, TN2, TI2, block2)
 
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_3(TN1, TI1, block1, TN2, TI2, block2, TN3, TI3, block3)
 
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE(TN, TI, block)
 
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_2(TN1, TI1, block1, TN2, TI2, block2)
 
#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_3(TN1, TI1, block1, TN2, TI2, block2, TN3, TI3, block3)
 
#define AMREX_HOST_DEVICE_FOR_1D(n, i, block)
 
#define AMREX_HOST_DEVICE_PARALLEL_FOR_1D(n, i, block)
 
#define AMREX_FOR_1D(n, i, block)
 
#define AMREX_PARALLEL_FOR_1D(n, i, block)
 
#define AMREX_HOST_DEVICE_FOR_3D(box, i, j, k, block)
 
#define AMREX_HOST_DEVICE_PARALLEL_FOR_3D(box, i, j, k, block)
 
#define AMREX_FOR_3D(box, i, j, k, block)
 
#define AMREX_PARALLEL_FOR_3D(box, i, j, k, block)
 
#define AMREX_HOST_DEVICE_FOR_4D(box, ncomp, i, j, k, n, block)
 
#define AMREX_HOST_DEVICE_PARALLEL_FOR_4D(box, ncomp, i, j, k, n, block)
 
#define AMREX_FOR_4D(box, ncomp, i, j, k, n, block)
 
#define AMREX_PARALLEL_FOR_4D(box, ncomp, i, j, k, n, block)
 

Macro Definition Documentation

◆ AMREX_FOR_1D

#define AMREX_FOR_1D (   n,
  i,
  block 
)
Value:
{ \
using amrex_i_inttype = std::remove_const_t<decltype(n)>; \
amrex::ParallelFor(n,[=] AMREX_GPU_DEVICE (amrex_i_inttype i) noexcept block); \
}
AMREX_GPU_DEVICE
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18

◆ AMREX_FOR_3D

#define AMREX_FOR_3D (   box,
  i,
  j,
  k,
  block 
)
Value:
{ \
amrex::ParallelFor(box,[=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept block); \
}

◆ AMREX_FOR_4D

#define AMREX_FOR_4D (   box,
  ncomp,
  i,
  j,
  k,
  n,
  block 
)
Value:
{ \
amrex::ParallelFor(box,ncomp,[=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept block); \
}

◆ AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE

#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE (   TN,
  TI,
  block 
)
Value:
{ auto const& amrex_i_tn = TN; \
if (!amrex::isEmpty(amrex_i_tn)) { \
if (amrex::Gpu::inLaunchRegion()) \
{ \
auto amrex_i_ec = amrex::Gpu::ExecutionConfig(amrex_i_tn); \
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_ec.numBlocks, amrex_i_ec.numThreads, amrex_i_ec.sharedMem, amrex::Gpu::gpuStream(), \
[=] AMREX_GPU_DEVICE () noexcept { \
for (auto const TI : amrex::Gpu::Range(amrex_i_tn)) { \
block \
} \
}); \
AMREX_GPU_ERROR_CHECK(); \
} \
else { \
amrex::Abort("AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE: cannot call device function from host"); \
}}}
amrex::Gpu::inLaunchRegion
bool inLaunchRegion() noexcept
Definition AMReX_GpuControl.H:92
amrex::Gpu::gpuStream
gpuStream_t gpuStream() noexcept
Definition AMReX_GpuDevice.H:241
amrex
Definition AMReX_Amr.cpp:49
amrex::isEmpty
bool isEmpty(T n) noexcept
Definition AMReX_GpuRange.H:14
amrex::Gpu::ExecutionConfig
Definition AMReX_GpuLaunch.H:132

◆ AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_2

#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_2 (   TN1,
  TI1,
  block1,
  TN2,
  TI2,
  block2 
)
Value:
{ auto const& amrex_i_tn1 = TN1; auto const& amrex_i_tn2 = TN2; \
if (!amrex::isEmpty(amrex_i_tn1) || !amrex::isEmpty(amrex_i_tn2)) { \
if (amrex::Gpu::inLaunchRegion()) \
{ \
const auto amrex_i_ec1 = amrex::Gpu::ExecutionConfig(amrex_i_tn1); \
const auto amrex_i_ec2 = amrex::Gpu::ExecutionConfig(amrex_i_tn2); \
dim3 amrex_i_nblocks = amrex::max(amrex_i_ec1.numBlocks.x, \
amrex_i_ec2.numBlocks.x); \
amrex_i_nblocks.y = 2; \
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_nblocks, amrex_i_ec1.numThreads, 0, amrex::Gpu::gpuStream(), \
[=] AMREX_GPU_DEVICE () noexcept { \
switch (blockIdx.y) { \
case 0: for (auto const TI1 : amrex::Gpu::Range(amrex_i_tn1)) { \
block1 \
} \
break; \
case 1: for (auto const TI2 : amrex::Gpu::Range(amrex_i_tn2)) { \
block2 \
} \
} \
}); \
AMREX_GPU_ERROR_CHECK(); \
} \
else { \
amrex::Abort("AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_2: cannot call device function from host"); \
}}}
amrex::max
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:35

◆ AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_3

#define AMREX_GPU_LAUNCH_DEVICE_LAMBDA_RANGE_3 (   TN1,
  TI1,
  block1,
  TN2,
  TI2,
  block2,
  TN3,
  TI3,
  block3 
)

◆ AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE

#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE (   TN,
  TI,
  block 
)
Value:
{ auto const& amrex_i_tn = TN; \
if (!amrex::isEmpty(amrex_i_tn)) { \
if (amrex::Gpu::inLaunchRegion()) \
{ \
const auto amrex_i_ec = amrex::Gpu::ExecutionConfig(amrex_i_tn); \
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, amrex_i_ec.numBlocks, amrex_i_ec.numThreads, amrex_i_ec.sharedMem, amrex::Gpu::gpuStream(), \
[=] AMREX_GPU_DEVICE () noexcept { \
for (auto const TI : amrex::Gpu::Range(amrex_i_tn)) { \
block \
} \
}); \
AMREX_GPU_ERROR_CHECK(); \
} \
else { \
for (auto const TI : amrex::Gpu::Range(amrex_i_tn)) { \
block \
} \
}}}
amrex::Gpu::Range
__host__ __device__ range_detail::range_impl< T > Range(T const &b) noexcept
Definition AMReX_GpuRange.H:125

◆ AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_2

#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_2 (   TN1,
  TI1,
  block1,
  TN2,
  TI2,
  block2 
)

◆ AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_3

#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE_3 (   TN1,
  TI1,
  block1,
  TN2,
  TI2,
  block2,
  TN3,
  TI3,
  block3 
)

◆ AMREX_HOST_DEVICE_FOR_1D

#define AMREX_HOST_DEVICE_FOR_1D (   n,
  i,
  block 
)
Value:
{ \
auto const& amrex_i_n = n; \
using amrex_i_inttype = std::remove_const_t<decltype(n)>; \
if (amrex::Gpu::inLaunchRegion()) { \
amrex::ParallelFor(amrex_i_n,[=] AMREX_GPU_DEVICE (amrex_i_inttype i) noexcept block); \
} else { \
auto amrex_i_lambda = [=] (amrex_i_inttype i) noexcept block; \
AMREX_PRAGMA_SIMD \
for (amrex_i_inttype i = 0; i < amrex_i_n; ++i) amrex_i_lambda(i); \
} \
}

◆ AMREX_HOST_DEVICE_FOR_3D

#define AMREX_HOST_DEVICE_FOR_3D (   box,
  i,
  j,
  k,
  block 
)
Value:
{ \
auto const& amrex_i_box = box; \
if (amrex::Gpu::inLaunchRegion()) { \
amrex::ParallelFor(amrex_i_box,[=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept block); \
} else { \
amrex::LoopConcurrentOnCpu(amrex_i_box,[=] (int i, int j, int k) noexcept block); \
} \
}

◆ AMREX_HOST_DEVICE_FOR_4D

#define AMREX_HOST_DEVICE_FOR_4D (   box,
  ncomp,
  i,
  j,
  k,
  n,
  block 
)
Value:
{ \
auto const& amrex_i_box = box; \
auto const& amrex_i_ncomp = ncomp; \
if (amrex::Gpu::inLaunchRegion()) { \
amrex::ParallelFor(amrex_i_box,amrex_i_ncomp,[=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept block); \
} else { \
amrex::LoopConcurrentOnCpu(amrex_i_box,amrex_i_ncomp,[=] (int i, int j, int k, int n) noexcept block); \
} \
}

◆ AMREX_HOST_DEVICE_PARALLEL_FOR_1D

#define AMREX_HOST_DEVICE_PARALLEL_FOR_1D (   n,
  i,
  block 
)
Value:
{ \
auto const& amrex_i_n = n; \
using amrex_i_inttype = std::remove_const_t<decltype(n)>; \
if (amrex::Gpu::inLaunchRegion()) { \
amrex::ParallelFor(amrex_i_n,[=] AMREX_GPU_DEVICE (amrex_i_inttype i) noexcept block); \
} else { \
auto amrex_i_lambda = [=] (amrex_i_inttype i) noexcept block; \
AMREX_PRAGMA_SIMD \
for (amrex_i_inttype i = 0; i < amrex_i_n; ++i) amrex_i_lambda(i); \
} \
}

◆ AMREX_HOST_DEVICE_PARALLEL_FOR_3D

#define AMREX_HOST_DEVICE_PARALLEL_FOR_3D (   box,
  i,
  j,
  k,
  block 
)
Value:
{ \
auto const& amrex_i_box = box; \
if (amrex::Gpu::inLaunchRegion()) { \
amrex::ParallelFor(amrex_i_box,[=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept block); \
} else { \
amrex::LoopConcurrentOnCpu(amrex_i_box,[=] (int i, int j, int k) noexcept block); \
} \
}

◆ AMREX_HOST_DEVICE_PARALLEL_FOR_4D

#define AMREX_HOST_DEVICE_PARALLEL_FOR_4D (   box,
  ncomp,
  i,
  j,
  k,
  n,
  block 
)
Value:
{ \
auto const& amrex_i_box = box; \
auto const& amrex_i_ncomp = ncomp; \
if (amrex::Gpu::inLaunchRegion()) { \
amrex::ParallelFor(amrex_i_box,amrex_i_ncomp,[=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept block); \
} else { \
amrex::LoopConcurrentOnCpu(amrex_i_box,amrex_i_ncomp,[=] (int i, int j, int k, int n) noexcept block); \
} \
}

◆ AMREX_PARALLEL_FOR_1D

#define AMREX_PARALLEL_FOR_1D (   n,
  i,
  block 
)
Value:
{ \
using amrex_i_inttype = std::remove_const_t<decltype(n)>; \
amrex::ParallelFor(n,[=] AMREX_GPU_DEVICE (amrex_i_inttype i) noexcept block); \
}

◆ AMREX_PARALLEL_FOR_3D

#define AMREX_PARALLEL_FOR_3D (   box,
  i,
  j,
  k,
  block 
)
Value:
{ \
amrex::ParallelFor(box,[=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept block); \
}

◆ AMREX_PARALLEL_FOR_4D

#define AMREX_PARALLEL_FOR_4D (   box,
  ncomp,
  i,
  j,
  k,
  n,
  block 
)
Value:
{ \
amrex::ParallelFor(box,ncomp,[=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept block); \
}