sycl_kernel_utils.h File Reference

#include <config.h>
#include <type_traits>
#include "gmxsycl.h"
#include "gputraits_sycl.h"

Include dependency graph for sycl_kernel_utils.h:

Description

SYCL kernel helper functions.

Author

Andrey Alekseenko al42a.nosp@m.nd@g.nosp@m.mail..nosp@m.com

Classes
struct	NullptrWrapper< T >
	Shim to smooth over implementation differences. More...
struct	detail::UnderlyingType< T >
	Type trait for implementing StaticLocalStorage. More...
class	StaticLocalStorage< T, size, condition >
	Shim for the two implementations of local storage. More...

Macros
#define	SYCL_ASSERT(condition)

Functions
constexpr bool	compilingForHost ()
template<int expectedSubGroupSize>
constexpr bool	compilingForSubGroupSize ()
template<int expectedSubGroupSize>
constexpr bool	skipKernelCompilation ()
template<typename T , sycl::memory_scope MemoryScope, sycl::access::address_space AddressSpace>
static void	atomicAddDefault (T &val, const T delta)
template<typename T , sycl::memory_scope MemoryScope = sycl::memory_scope::device, sycl::access::address_space AddressSpace = sycl::access::address_space::global_space>
static void	atomicFetchAdd (T &val, const T delta)
	Convenience wrapper to do atomic addition to a global buffer.
template<typename T >
static void	atomicFetchAddLocal (T &val, const T delta)
	Portability wrapper for performing an atomic value addition. More...
template<typename T >
static void	atomicFetchAddLocal (T *values, const int index, const T delta)
	Portability wrapper for performing atomic add on a value in an array. More...
template<typename T , typename W >
static void	atomicFetchAddLocal (T &val, const int index, const W delta)
	Portability wrapper for performing atomic add on a single value of a composite type like RVec. More...
static void	staggeredAtomicAddForce (sycl::global_ptr< Float3 > gm_f, Float3 f, const int index, const int localId)
	Staggered atomic force component accumulation into global memory to reduce clashes. More...
template<typename T , sycl::memory_scope MemoryScope = sycl::memory_scope::device>
static T	atomicLoad (T &val)
	Convenience wrapper to do atomic loads from a global buffer.

Variables
static constexpr unsigned int	c_cudaFullWarpMask = 0xffffffff
	Full warp active thread mask used in CUDA warp-level primitives.
static constexpr bool	sc_useCommandGroupHandler = true

Function Documentation

template<typename T >

static void atomicFetchAddLocal ( T &  val, const T  delta  )

inlinestatic

Portability wrapper for performing an atomic value addition.

Template Parameters

T	Type to perform atomic add for.

Parameters

[in,out]	val	Destination memory for atomic add.
[in]	delta	Source value to add.

template<typename T >

static void atomicFetchAddLocal ( T *  values, const int  index, const T  delta  )

inlinestatic

Portability wrapper for performing atomic add on a value in an array.

Template Parameters

T	Type to perform atomic add for.

Parameters

[in,out]	values	Array of values to perform addition in.
[in]	index	Which address in the array should be used for add.
[in]	delta	Source value to add.

template<typename T , typename W >

static void atomicFetchAddLocal ( T &  val, const int  index, const W  delta  )

inlinestatic

Portability wrapper for performing atomic add on a single value of a composite type like RVec.

Template Parameters

T	Type for composite type to use.
W	Type for actual additional operation.

Parameters

[in,out]	val	Composite data to perform atomic add on.
[in]	index	Which field in the composite data should be used for the add.
[in]	delta	Source value to add.

static void staggeredAtomicAddForce ( sycl::global_ptr< Float3 >  gm_f, Float3  f, const int  index, const int  localId  )

inlinestatic

Staggered atomic force component accumulation into global memory to reduce clashes.

Reduce the number of atomic clashes by a theoretical max 3x by having consecutive threads accumulate different force components at the same time.

Parameters

[in,out]	gm_f	Global array of forces.
[in]	f	Force to add.
[in]	index	Which value in the global force array we want to add to.
[in]	localId	Which position in the thread we are in.