amrex::Gpu Namespace Reference

Namespaces
namespace	Atomic

Classes
class	AsyncArray
struct	AtomicAdd
struct	AtomicMin
struct	AtomicMax
struct	AtomicLogicalAnd
struct	AtomicLogicalOr
class	Buffer
struct	HostToDevice
struct	DeviceToHost
struct	DeviceToDevice
struct	LaunchSafeGuard
struct	GraphSafeGuard
struct	SingleStreamRegion
struct	NoSyncRegion
class	StreamManager
class	Device
struct	ExternalGpuStreamRegion
struct	SyncAtExitOnly
class	Elixir
class	KernelInfo
struct	ExecutionConfig
struct	ExecConfig
struct	Managed
struct	Pinned
struct	Deleter
struct	DeviceScalar
struct	SharedMemory
struct	SharedMemory< double >
struct	SharedMemory< float >
struct	SharedMemory< long >
struct	SharedMemory< long long >
struct	SharedMemory< int >
struct	SharedMemory< short >
struct	SharedMemory< char >
struct	SharedMemory< unsigned long >
struct	SharedMemory< unsigned long long >
struct	SharedMemory< unsigned int >
struct	SharedMemory< unsigned short >
struct	SharedMemory< unsigned char >
struct	SharedMemory< bool >
struct	warpReduce
struct	Handler
struct	StreamItInfo
class	StreamIter
struct	TrackedVector
struct	SharedMemory< openbc::Moments::array_type >

Typedefs
template<class T>
using	DeviceVector = PODVector<T, ArenaAllocator<T> >
	A PODVector that uses the standard memory Arena. Note that the memory might or might not be managed depending on the amrex.the_arena_is_managed ParmParse parameter.
template<class T>
using	NonManagedDeviceVector = PODVector<T, DeviceArenaAllocator<T> >
	A PODVector that uses the non-managed device memory arena.
template<class T>
using	ManagedVector = PODVector<T, ManagedArenaAllocator<T> >
	A PODVector that uses the managed memory arena.
template<class T>
using	PinnedVector = PODVector<T, PinnedArenaAllocator<T> >
	A PODVector that uses the pinned memory arena.
template<class T>
using	AsyncVector = PODVector<T, AsyncArenaAllocator<T> >
	A PODVector that uses the async memory arena. Maybe useful for temporary vectors inside MFIters that are accessed on the device.
template<class T>
using	HostVector = PinnedVector<T>
	A PODVector that uses pinned host memory. Same as PinnedVector. For a vector class that uses std::allocator by default, see amrex::Vector.
template<class T>
using	ManagedDeviceVector = PODVector<T, ManagedArenaAllocator<T> >
	This is identical to ManagedVector<T>. The ManagedDeviceVector form is deprecated and will be removed in a future release.

Enumerations
enum class	ExternalStreamSync { Yes , No }

Functions
template<class InIter, class OutIter>
void	copy (HostToDevice, InIter begin, InIter end, OutIter result) noexcept
	A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.
template<class InIter, class OutIter>
void	copy (DeviceToHost, InIter begin, InIter end, OutIter result) noexcept
	A device-to-host copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.
template<class InIter, class OutIter>
void	copy (DeviceToDevice, InIter begin, InIter end, OutIter result) noexcept
	A device-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.
template<class InIter, class OutIter>
void	copyAsync (HostToDevice, InIter begin, InIter end, OutIter result) noexcept
	A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.
template<class InIter, class OutIter>
void	copyAsync (DeviceToHost, InIter begin, InIter end, OutIter result) noexcept
	A device-to-host copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.
template<class InIter, class OutIter>
void	copyAsync (DeviceToDevice, InIter begin, InIter end, OutIter result) noexcept
	A device-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.
template<class Iter>
void	prefetchToHost (Iter begin, Iter end) noexcept
	Migrate elements of a container from device to host. This is a no-op for host-only code.
template<class Iter>
void	prefetchToDevice (Iter begin, Iter end) noexcept
	Migrate elements of a container from host to device. This is a no-op for host-only code.
template<typename IT, typename F, typename T = typename std::iterator_traits<IT>::value_type> requires ((sizeof(T) <= 36*8) && // so there is enough shared memory std::is_trivially_copyable_v<T> && amrex::IsCallable<F, T&, Long>::value)
void	fillAsync (IT first, IT last, F const &f) noexcept
	Fill the elements in the given range using the given calllable.
bool	inLaunchRegion () noexcept
bool	notInLaunchRegion () noexcept
bool	setLaunchRegion (bool launch) noexcept
bool	inGraphRegion ()
bool	notInGraphRegion ()
bool	setGraphRegion (bool graph)
bool	inSingleStreamRegion () noexcept
bool	inNoSyncRegion () noexcept
bool	setSingleStreamRegion (bool b) noexcept
bool	setNoSyncRegion (bool b) noexcept
int	SelectDeviceBeforeMPIInit ()
	Select this process's GPU before MPI is initialized.
bool	isStreamActive (gpuStream_t s)
	Check if the given stream is active.
gpuStream_t	gpuStream () noexcept
int	numGpuStreams () noexcept
void	synchronize () noexcept
void	streamSynchronize () noexcept
void	streamSynchronize (gpuStream_t s) noexcept
void	streamSynchronizeActive () noexcept
void	streamSynchronizeAll () noexcept
void	freeAsync (Arena *arena, void *mem) noexcept
bool	clearFreeAsyncBuffer () noexcept
void	setExternalGpuStream (gpuStream_t stream)
	Provide a user-supplied GPU stream.
void	resetExternalGpuStream (ExternalStreamSync sync_stream=ExternalStreamSync::Yes) noexcept
	Pop the current user-supplied GPU stream.
void	htod_memcpy_async (void *p_d, const void *p_h, const std::size_t sz) noexcept
void	dtoh_memcpy_async (void *p_h, const void *p_d, const std::size_t sz) noexcept
void	dtod_memcpy_async (void *p_d_dst, const void *p_d_src, const std::size_t sz) noexcept
void	htod_memcpy (void *p_d, const void *p_h, const std::size_t sz) noexcept
void	dtoh_memcpy (void *p_h, const void *p_d, const std::size_t sz) noexcept
void	dtod_memcpy (void *p_d_dst, const void *p_d_src, const std::size_t sz) noexcept
template<typename T>
void	memcpy_from_host_to_device_global_async (T &dg, const void *src, std::size_t nbytes, std::size_t offset=0)
template<typename T>
void	memcpy_from_device_global_to_host_async (void *dst, T const &dg, std::size_t nbytes, std::size_t offset=0)
void	ErrorCheck (const char *file, int line)
constexpr std::size_t	numThreadsPerBlockParallelFor ()
__host__ __device__ Box	getThreadBox (const Box &bx, Long offset) noexcept
template<int MT>
ExecutionConfig	makeExecutionConfig (Long N) noexcept
template<int MT>
ExecutionConfig	makeExecutionConfig (const Box &box) noexcept
template<int MT>
Vector< ExecConfig >	makeNExecutionConfigs (Long N) noexcept
template<int MT, int dim>
Vector< ExecConfig >	makeNExecutionConfigs (BoxND< dim > const &box) noexcept
template<typename T>
__host__ __device__ range_detail::range_impl< T >	Range (T const &b) noexcept
template<typename T>
__device__ void	deviceReduceSum (T *dest, T source, Gpu::Handler const &h) noexcept
template<typename T>
__device__ void	deviceReduceMin (T *dest, T source, Gpu::Handler const &h) noexcept
template<typename T>
__device__ void	deviceReduceMax (T *dest, T source, Gpu::Handler const &h) noexcept
__device__ void	deviceReduceLogicalAnd (int *dest, int source, Gpu::Handler const &h) noexcept
__device__ void	deviceReduceLogicalOr (int *dest, int source, Gpu::Handler const &h) noexcept
template<int warpSize, typename T, typename WARPREDUCE>
__device__ T	blockReduce (T x, WARPREDUCE &&warp_reduce, T x0)
template<int warpSize, typename T, typename WARPREDUCE, typename ATOMICOP>
__device__ void	blockReduce_partial (T *dest, T x, WARPREDUCE &&warp_reduce, ATOMICOP &&atomic_op, Gpu::Handler const &handler)
template<typename T>
__device__ T	blockReduceSum (T source) noexcept
template<typename T>
__device__ void	deviceReduceSum_full (T *dest, T source) noexcept
template<int BLOCKDIMX, typename T>
__device__ T	blockReduceSum (T source) noexcept
template<int BLOCKDIMX, typename T>
__device__ void	deviceReduceSum_full (T *dest, T source) noexcept
template<typename T>
__device__ T	blockReduceMin (T source) noexcept
template<typename T>
__device__ void	deviceReduceMin_full (T *dest, T source) noexcept
template<int BLOCKDIMX, typename T>
__device__ T	blockReduceMin (T source) noexcept
template<int BLOCKDIMX, typename T>
__device__ void	deviceReduceMin_full (T *dest, T source) noexcept
template<typename T>
__device__ T	blockReduceMax (T source) noexcept
template<typename T>
__device__ void	deviceReduceMax_full (T *dest, T source) noexcept
template<int BLOCKDIMX, typename T>
__device__ T	blockReduceMax (T source) noexcept
template<int BLOCKDIMX, typename T>
__device__ void	deviceReduceMax_full (T *dest, T source) noexcept
__device__ int	blockReduceLogicalAnd (int source) noexcept
__device__ void	deviceReduceLogicalAnd_full (int *dest, int source) noexcept
template<int BLOCKDIMX>
__device__ int	blockReduceLogicalAnd (int source) noexcept
template<int BLOCKDIMX>
__device__ void	deviceReduceLogicalAnd_full (int *dest, int source) noexcept
__device__ int	blockReduceLogicalOr (int source) noexcept
__device__ void	deviceReduceLogicalOr_full (int *dest, int source) noexcept
template<int BLOCKDIMX>
__device__ int	blockReduceLogicalOr (int source) noexcept
template<int BLOCKDIMX>
__device__ void	deviceReduceLogicalOr_full (int *dest, int source) noexcept
template<typename T>
__host__ __device__ T	LDG (Array4< T > const &a, int i, int j, int k) noexcept
template<typename T>
__host__ __device__ T	LDG (Array4< T > const &a, int i, int j, int k, int n) noexcept
bool	isManaged (void const *p) noexcept
bool	isDevicePtr (void const *p) noexcept
bool	isPinnedPtr (void const *p) noexcept
bool	isGpuPtr (void const *p) noexcept
__host__ __device__ void *	memcpy (void *dest, const void *src, std::size_t count)
template<class InIter, class OutIter>
OutIter	inclusive_scan (InIter begin, InIter end, OutIter result)
template<class InIter, class OutIter>
OutIter	exclusive_scan (InIter begin, InIter end, OutIter result)
template<typename T>
__host__ __device__ bool	isnan (T x) noexcept
	Return true if x is NaN. Works under fast math.
template<typename T>
__host__ __device__ bool	isinf (T x) noexcept
	Return true if x is +/-infinity. Works under fast math.
template<typename T>
__host__ __device__ bool	isfinite (T x) noexcept
	Return true if x is neither NaN nor infinity. Works under fast math.

Variables
static constexpr HostToDevice	hostToDevice {}
static constexpr DeviceToHost	deviceToHost {}
static constexpr DeviceToDevice	deviceToDevice {}
bool	in_launch_region = true
bool	in_graph_region = false
bool	in_single_stream_region = false
bool	in_nosync_region = false

Enumeration Type Documentation

ExternalStreamSync

enum class amrex::Gpu::ExternalStreamSync

strong

Enumerator
Yes
No

Function Documentation

blockReduce()

template<int warpSize, typename T, typename WARPREDUCE>

__device__ T amrex::Gpu::blockReduce ( T x, WARPREDUCE && warp_reduce, T x0 )

inline

blockReduce_partial()

template<int warpSize, typename T, typename WARPREDUCE, typename ATOMICOP>

__device__ void amrex::Gpu::blockReduce_partial ( T * dest, T x, WARPREDUCE && warp_reduce, ATOMICOP && atomic_op, Gpu::Handler const & handler )

inline

blockReduceLogicalAnd() [1/2]

template<int BLOCKDIMX>

__device__ int amrex::Gpu::blockReduceLogicalAnd ( int source )

inlinenoexcept

blockReduceLogicalAnd() [2/2]

__device__ int amrex::Gpu::blockReduceLogicalAnd ( int source )

inlinenoexcept

blockReduceLogicalOr() [1/2]

template<int BLOCKDIMX>

__device__ int amrex::Gpu::blockReduceLogicalOr ( int source )

inlinenoexcept

blockReduceLogicalOr() [2/2]

__device__ int amrex::Gpu::blockReduceLogicalOr ( int source )

inlinenoexcept

blockReduceMax() [1/2]

template<int BLOCKDIMX, typename T>

__device__ T amrex::Gpu::blockReduceMax ( T source )

inlinenoexcept

blockReduceMax() [2/2]

template<typename T>

__device__ T amrex::Gpu::blockReduceMax ( T source )

inlinenoexcept

blockReduceMin() [1/2]

template<int BLOCKDIMX, typename T>

__device__ T amrex::Gpu::blockReduceMin ( T source )

inlinenoexcept

blockReduceMin() [2/2]

template<typename T>

__device__ T amrex::Gpu::blockReduceMin ( T source )

inlinenoexcept

blockReduceSum() [1/2]

template<int BLOCKDIMX, typename T>

__device__ T amrex::Gpu::blockReduceSum ( T source )

inlinenoexcept

blockReduceSum() [2/2]

template<typename T>

__device__ T amrex::Gpu::blockReduceSum ( T source )

inlinenoexcept

clearFreeAsyncBuffer()

bool amrex::Gpu::clearFreeAsyncBuffer ( )

inlinenoexcept

Clear the temporary buffer used by freeAsync. This is done by synchronizing the stream in case there is memory in the buffer. Returns true if some memory could be freed.

copy() [1/3]

template<class InIter, class OutIter>

void amrex::Gpu::copy ( DeviceToDevice , InIter begin, InIter end, OutIter result )

noexcept

A device-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.

This version is blocking - CPU execution will halt until the copy is finished.

Template Parameters

InIter	The input iterator type
OutIter	The output iterator type

Parameters

begin	Where in the input to start reading
end	Where in the input to stop reading
result	Where in the output to start writing

Example usage:

Gpu::copy(Gpu::deviceToDevice, a.begin(), a.end(), b.begin());

copy() [2/3]

template<class InIter, class OutIter>

void amrex::Gpu::copy ( DeviceToHost , InIter begin, InIter end, OutIter result )

noexcept

A device-to-host copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.

This version is blocking - CPU execution will halt until the copy is finished.

Template Parameters

InIter	The input iterator type
OutIter	The output iterator type

Parameters

begin	Where in the input to start reading
end	Where in the input to stop reading
result	Where in the output to start writing

Example usage:

Gpu::copy(Gpu::deviceToHost, a.begin(), a.end(), b.begin());

copy() [3/3]

template<class InIter, class OutIter>

void amrex::Gpu::copy ( HostToDevice , InIter begin, InIter end, OutIter result )

noexcept

A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.

This version is blocking - CPU execution will halt until the copy is finished.

Template Parameters

InIter	The input iterator type
OutIter	The output iterator type

Parameters

begin	Where in the input to start reading
end	Where in the input to stop reading
result	Where in the output to start writing

Example usage:

Gpu::copy(Gpu::hostToDevice, a.begin(), a.end(), b.begin());

copyAsync() [1/3]

template<class InIter, class OutIter>

void amrex::Gpu::copyAsync ( DeviceToDevice , InIter begin, InIter end, OutIter result )

noexcept

A device-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.

This version is asynchronous - CPU execution will continue, whether or not the copy is finished.

Template Parameters

InIter	The input iterator type
OutIter	The output iterator type

Parameters

begin	Where in the input to start reading
end	Where in the input to stop reading
result	Where in the output to start writing

Example usage:

Gpu::copyAsync(Gpu::deviceToDevice, a.begin(), a.end(), b.begin());

copyAsync() [2/3]

template<class InIter, class OutIter>

void amrex::Gpu::copyAsync ( DeviceToHost , InIter begin, InIter end, OutIter result )

noexcept

A device-to-host copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.

This version is asynchronous - CPU execution will continue, whether or not the copy is finished.

Template Parameters

InIter	The input iterator type
OutIter	The output iterator type

Parameters

begin	Where in the input to start reading
end	Where in the input to stop reading
result	Where in the output to start writing

Example usage:

Gpu::copyAsync(Gpu::deviceToHost, a.begin(), a.end(), b.begin());

copyAsync() [3/3]

template<class InIter, class OutIter>

void amrex::Gpu::copyAsync ( HostToDevice , InIter begin, InIter end, OutIter result )

noexcept

A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous storage. The amrex-provided containers like Gpu::HostVector, Gpu::DeviceVector, etc. meet this requirement.

This version is asynchronous - CPU execution will continue, whether or not the copy is finished.

Template Parameters

InIter	The input iterator type
OutIter	The output iterator type

Parameters

begin	Where in the input to start reading
end	Where in the input to stop reading
result	Where in the output to start writing

Example usage:

Gpu::copyAsync(Gpu::hostToDevice, a.begin(), a.end(), b.begin());

deviceReduceLogicalAnd()

__device__ void amrex::Gpu::deviceReduceLogicalAnd ( int * dest, int source, Gpu::Handler const & h )

inlinenoexcept

deviceReduceLogicalAnd_full() [1/2]

template<int BLOCKDIMX>

__device__ void amrex::Gpu::deviceReduceLogicalAnd_full ( int * dest, int source )

inlinenoexcept

deviceReduceLogicalAnd_full() [2/2]

__device__ void amrex::Gpu::deviceReduceLogicalAnd_full ( int * dest, int source )

inlinenoexcept

deviceReduceLogicalOr()

__device__ void amrex::Gpu::deviceReduceLogicalOr ( int * dest, int source, Gpu::Handler const & h )

inlinenoexcept

deviceReduceLogicalOr_full() [1/2]

template<int BLOCKDIMX>

__device__ void amrex::Gpu::deviceReduceLogicalOr_full ( int * dest, int source )

inlinenoexcept

deviceReduceLogicalOr_full() [2/2]

__device__ void amrex::Gpu::deviceReduceLogicalOr_full ( int * dest, int source )

inlinenoexcept

deviceReduceMax()

template<typename T>

__device__ void amrex::Gpu::deviceReduceMax ( T * dest, T source, Gpu::Handler const & h )

inlinenoexcept

deviceReduceMax_full() [1/2]

template<int BLOCKDIMX, typename T>

__device__ void amrex::Gpu::deviceReduceMax_full ( T * dest, T source )

inlinenoexcept

deviceReduceMax_full() [2/2]

template<typename T>

__device__ void amrex::Gpu::deviceReduceMax_full ( T * dest, T source )

inlinenoexcept

deviceReduceMin()

template<typename T>

__device__ void amrex::Gpu::deviceReduceMin ( T * dest, T source, Gpu::Handler const & h )

inlinenoexcept

deviceReduceMin_full() [1/2]

template<int BLOCKDIMX, typename T>

__device__ void amrex::Gpu::deviceReduceMin_full ( T * dest, T source )

inlinenoexcept

deviceReduceMin_full() [2/2]

template<typename T>

__device__ void amrex::Gpu::deviceReduceMin_full ( T * dest, T source )

inlinenoexcept

deviceReduceSum()

template<typename T>

__device__ void amrex::Gpu::deviceReduceSum ( T * dest, T source, Gpu::Handler const & h )

inlinenoexcept

deviceReduceSum_full() [1/2]

template<int BLOCKDIMX, typename T>

__device__ void amrex::Gpu::deviceReduceSum_full ( T * dest, T source )

inlinenoexcept

deviceReduceSum_full() [2/2]

template<typename T>

__device__ void amrex::Gpu::deviceReduceSum_full ( T * dest, T source )

inlinenoexcept

dtod_memcpy()

void amrex::Gpu::dtod_memcpy ( void * p_d_dst, const void * p_d_src, const std::size_t sz )

inlinenoexcept

dtod_memcpy_async()

void amrex::Gpu::dtod_memcpy_async ( void * p_d_dst, const void * p_d_src, const std::size_t sz )

inlinenoexcept

dtoh_memcpy()

void amrex::Gpu::dtoh_memcpy ( void * p_h, const void * p_d, const std::size_t sz )

inlinenoexcept

dtoh_memcpy_async()

void amrex::Gpu::dtoh_memcpy_async ( void * p_h, const void * p_d, const std::size_t sz )

inlinenoexcept

ErrorCheck()

void amrex::Gpu::ErrorCheck ( const char * file, int line )

inline

fillAsync()

template<typename IT, typename F, typename T = typename std::iterator_traits<IT>::value_type>
requires ((sizeof(T) <= 36*8) && // so there is enough shared memory std::is_trivially_copyable_v<T> && amrex::IsCallable<F, T&, Long>::value)

void amrex::Gpu::fillAsync ( IT first, IT last, F const & f )

noexcept

Fill the elements in the given range using the given calllable.

This function is asynchronous for GPU builds.

Template Parameters

IT	the iterator type
F	the callable type

Parameters

first	the inclusive first in the range [first, last)
last	the exclusive last in the range [first, last)
f	the callable with the function signature of void(T&, Long), where T is the element type and the Long parameter is the index for the element to be filled.

freeAsync()

void amrex::Gpu::freeAsync ( Arena * arena, void * mem )

inlinenoexcept

Deallocate memory belonging to an arena asynchronously. Memory deallocated in this way is held in a pool and will not be reused until the next amrex::Gpu::streamSynchronize(). GPU kernels that were already launched on the currently active stream can still continue to use the memory after this function is called. There is no need to use this function for CPU-only memory or with The_Async_Arena.

Parameters

[in]	arena	the arena the memory belongs to
[in]	mem	pointer to the memory to be freed

getThreadBox()

__host__ __device__ Box amrex::Gpu::getThreadBox ( const Box & bx, Long offset )

inlinenoexcept

gpuStream()

gpuStream_t amrex::Gpu::gpuStream ( )

inlinenoexcept

htod_memcpy()

void amrex::Gpu::htod_memcpy ( void * p_d, const void * p_h, const std::size_t sz )

inlinenoexcept

htod_memcpy_async()

void amrex::Gpu::htod_memcpy_async ( void * p_d, const void * p_h, const std::size_t sz )

inlinenoexcept

inGraphRegion()

bool amrex::Gpu::inGraphRegion ( )

inlinenodiscard

inLaunchRegion()

bool amrex::Gpu::inLaunchRegion ( )

inlinenodiscardnoexcept

inNoSyncRegion()

bool amrex::Gpu::inNoSyncRegion ( )

inlinenodiscardnoexcept

inSingleStreamRegion()

bool amrex::Gpu::inSingleStreamRegion ( )

inlinenodiscardnoexcept

isDevicePtr()

bool amrex::Gpu::isDevicePtr ( void const * p )

inlinenoexcept

isfinite()

template<typename T>

__host__ __device__ bool amrex::Math::isfinite ( T x )

inlinenodiscardnoexcept

Return true if x is neither NaN nor infinity. Works under fast math.

isGpuPtr()

bool amrex::Gpu::isGpuPtr ( void const * p )

inlinenoexcept

isinf()

template<typename T>

__host__ __device__ bool amrex::Math::isinf ( T x )

inlinenodiscardnoexcept

Return true if x is +/-infinity. Works under fast math.

isManaged()

bool amrex::Gpu::isManaged ( void const * p )

inlinenoexcept

isnan()

template<typename T>

__host__ __device__ bool amrex::Math::isnan ( T x )

inlinenodiscardnoexcept

Return true if x is NaN. Works under fast math.

isPinnedPtr()

bool amrex::Gpu::isPinnedPtr ( void const * p )

inlinenoexcept

isStreamActive()

bool amrex::Gpu::isStreamActive

(

gpuStream_t

s

)

Check if the given stream is active.

LDG() [1/2]

template<typename T>

__host__ __device__ T amrex::Gpu::LDG ( Array4< T > const & a, int i, int j, int k )

inlinenoexcept

LDG() [2/2]

template<typename T>

__host__ __device__ T amrex::Gpu::LDG ( Array4< T > const & a, int i, int j, int k, int n )

inlinenoexcept

makeExecutionConfig() [1/2]

template<int MT>

ExecutionConfig amrex::Gpu::makeExecutionConfig ( const Box & box )

noexcept

makeExecutionConfig() [2/2]

template<int MT>

ExecutionConfig amrex::Gpu::makeExecutionConfig ( Long N )

noexcept

makeNExecutionConfigs() [1/2]

template<int MT, int dim>

Vector< ExecConfig > amrex::Gpu::makeNExecutionConfigs ( BoxND< dim > const & box )

noexcept

makeNExecutionConfigs() [2/2]

template<int MT>

Vector< ExecConfig > amrex::Gpu::makeNExecutionConfigs ( Long N )

noexcept

memcpy()

__host__ __device__ void * amrex::Gpu::memcpy ( void * dest, const void * src, std::size_t count )

inline

memcpy_from_device_global_to_host_async()

template<typename T>

void amrex::Gpu::memcpy_from_device_global_to_host_async	(	void *	dst,
		T const &	dg,
		std::size_t	nbytes,
		std::size_t	offset = 0 )

Copy nbytes bytes from device global variable to host. offset is the offset in bytes from the start of the device global variable.

memcpy_from_host_to_device_global_async()

template<typename T>

void amrex::Gpu::memcpy_from_host_to_device_global_async	(	T &	dg,
		const void *	src,
		std::size_t	nbytes,
		std::size_t	offset = 0 )

Copy nbytes bytes from host to device global variable. offset is the offset in bytes from the start of the device global variable.

notInGraphRegion()

bool amrex::Gpu::notInGraphRegion ( )

inlinenodiscard

notInLaunchRegion()

bool amrex::Gpu::notInLaunchRegion ( )

inlinenodiscardnoexcept

numGpuStreams()

int amrex::Gpu::numGpuStreams ( )

inlinenoexcept

numThreadsPerBlockParallelFor()

std::size_t amrex::Gpu::numThreadsPerBlockParallelFor ( )

constexpr

prefetchToDevice()

template<class Iter>

void amrex::Gpu::prefetchToDevice ( Iter begin, Iter end )

noexcept

Migrate elements of a container from host to device. This is a no-op for host-only code.

This version is blocking - CPU execution will halt until the migration is finished.

Warning

For GPU builds this helper only supports unified allocations (for example Gpu::ManagedVector). Passing iterators to non-managed memory, like the default DeviceVector or PinnedVector, will trigger a runtime error (cudaErrorInvalidValue) from cudaMemPrefetchAsync.

prefetchToHost()

template<class Iter>

void amrex::Gpu::prefetchToHost ( Iter begin, Iter end )

noexcept

Migrate elements of a container from device to host. This is a no-op for host-only code.

This version is blocking - CPU execution will halt until the migration is finished.

Warning

For GPU builds this helper only supports unified allocations (for example Gpu::ManagedVector). Passing iterators to non-managed memory, like the default DeviceVector or PinnedVector, will trigger a runtime error (cudaErrorInvalidValue) from cudaMemPrefetchAsync.

Range()

template<typename T>

__host__ __device__ range_detail::range_impl< T > amrex::Gpu::Range ( T const & b )

noexcept

resetExternalGpuStream()

void amrex::Gpu::resetExternalGpuStream ( ExternalStreamSync sync_stream = ExternalStreamSync::Yes )

inlinenoexcept

Pop the current user-supplied GPU stream.

Must be called once for every preceding successful setExternalGpuStream invocation (RAII helpers do this automatically). When sync_stream is :cpp:ExternalStreamSync::No AMReX skips the stream-wide synchronize unless deferred frees recorded by The_Async_Arena are pending, in which case a synchronization is still forced for correctness. If an older external stream was active before the current one, it becomes current again after the pop. Not supported inside OpenMP parallel regions.

SelectDeviceBeforeMPIInit()

int amrex::Gpu::SelectDeviceBeforeMPIInit

(

)

Select this process's GPU before MPI is initialized.

GPU-aware MPI may bind to the current GPU in MPI_Init. amrex::Initialize calls this function when it initializes MPI. An application that initializes MPI itself should call it before MPI_Init; amrex::Initialize then uses the selected GPU.

The GPU is chosen from the node-local rank set by the job launcher (e.g., SLURM_LOCALID). This function does nothing in CPU and SYCL builds, after MPI is initialized, or if the local rank is unknown.

Returns

The selected GPU, or -1 if none was selected.

setExternalGpuStream()

void amrex::Gpu::setExternalGpuStream ( gpuStream_t stream )

inline

Provide a user-supplied GPU stream.

Nested calls form a stack. Each new call makes the supplied stream current for AMReX GPU launches until resetExternalGpuStream pops it and restores the previous external stream, if any. Calling this function inside an OpenMP parallel region will trigger an assertion.

setGraphRegion()

bool amrex::Gpu::setGraphRegion ( bool graph )

inline

setLaunchRegion()

bool amrex::Gpu::setLaunchRegion ( bool launch )

inlinenoexcept

Enable/disable GPU kernel launches.

Note

This will only switch from GPU to CPU for kernels launched with macros. Functions like amrex::ParallelFor will be unaffected. Therefore it should not be used for comparing GPU to non-GPU performance or behavior.

Gpu::setLaunchRegion(0);
 
//...
 
Gpu::setLaunchRegion(1);

Will disable the launching of GPU kernels between the calls.

setNoSyncRegion()

bool amrex::Gpu::setNoSyncRegion ( bool b )

inlinenoexcept

setSingleStreamRegion()

bool amrex::Gpu::setSingleStreamRegion ( bool b )

inlinenoexcept

streamSynchronize() [1/2]

void amrex::Gpu::streamSynchronize ( )

inlinenoexcept

streamSynchronize() [2/2]

void amrex::Gpu::streamSynchronize ( gpuStream_t s )

inlinenoexcept

streamSynchronizeActive()

void amrex::Gpu::streamSynchronizeActive ( )

inlinenoexcept

streamSynchronizeAll()

void amrex::Gpu::streamSynchronizeAll ( )

inlinenoexcept

synchronize()

void amrex::Gpu::synchronize ( )

inlinenoexcept

Variable Documentation

deviceToDevice

DeviceToDevice amrex::Gpu::deviceToDevice {}

staticconstexpr

deviceToHost

DeviceToHost amrex::Gpu::deviceToHost {}

staticconstexpr

hostToDevice

HostToDevice amrex::Gpu::hostToDevice {}

staticconstexpr

in_graph_region

bool amrex::Gpu::in_graph_region = false

in_launch_region

bool amrex::Gpu::in_launch_region = true

in_nosync_region

bool amrex::Gpu::in_nosync_region = false

in_single_stream_region

bool amrex::Gpu::in_single_stream_region = false