cugar::cuda::warp_static_atomic Struct Reference

Detailed description

An efficient warp-synchronous atomic adder, to add or subtract compile-time constants to a shared integer.

Given a pointer to an integer (e.g. representing a "pool"), this class allows the threads in a warp to add (allocate) or subtract (deallocate) a compile-time constant to that integer in a predicated fashion.

#include <warp_atomics.h>

Public Methods
__device__ __forceinline__	warp_static_atomic (uint32 *pool)
template<uint32 N>
__device__ __forceinline__ void	add (bool p)
template<uint32 N>
__device__ __forceinline__ void	sub (bool p)
template<uint32 N>
__device__ __forceinline__ void	add (bool p, uint32 *result)
template<uint32 N>
__device__ __forceinline__ void	sub (bool p, uint32 *result)

Static Public Methods
template<uint32 N>
__device__ static __forceinline__ void	static_add (bool p, uint32 *dest)
template<uint32 N>
__device__ static __forceinline__ void	static_sub (bool p, uint32 *dest)
template<uint32 N>
__device__ static __forceinline__ void	static_add (bool p, uint32 *dest, uint32 *result)
template<uint32 N>
__device__ static __forceinline__ void	static_sub (bool p, uint32 *dest, uint32 *result)

Constructor & Destructor Documentation

warp_static_atomic()

__device__ __forceinline__ cugar::cuda::warp_static_atomic::warp_static_atomic ( uint32 *  pool )

inline

stateful constructor

Member Function Documentation

add() [1/2]

template<uint32 N>

__device__ __forceinline__ void cugar::cuda::warp_static_atomic::add ( bool  p )

inline

add zero or exactly N per thread to a shared value without waiting for the result: useful to alloc N entries from a common pool

Parameters

n	number of elements to alloc

add() [2/2]

template<uint32 N>

__device__ __forceinline__ void cugar::cuda::warp_static_atomic::add ( bool  p, uint32 *  result  )

inline

add zero or exactly N per thread to a shared value: useful to alloc N entries from a common pool

Parameters

n	number of elements to alloc

static_add() [1/2]

template<uint32 N>

__device__ static __forceinline__ void cugar::cuda::warp_static_atomic::static_add ( bool  p, uint32 *  dest  )

inlinestatic

add zero or exactly N per thread to a shared value without waiting for the result: useful to alloc N entries from a common pool

Parameters

p	allocation predicate
dest	the destination of the atomic

static_add() [2/2]

template<uint32 N>

__device__ static __forceinline__ void cugar::cuda::warp_static_atomic::static_add ( bool  p, uint32 *  dest, uint32 *  result  )

inlinestatic

add zero or exactly N per thread to a shared value: useful to alloc N entries from a common pool

Parameters

p	allocation predicate
dest	the destination of the atomic
result	output result

static_sub() [1/2]

template<uint32 N>

__device__ static __forceinline__ void cugar::cuda::warp_static_atomic::static_sub ( bool  p, uint32 *  dest  )

inlinestatic

subtract zero or exactly N per thread to a shared value without waiting for the result: useful to dealloc N entries from a common pool

Parameters

p	allocation predicate
dest	the destination of the atomic

static_sub() [2/2]

template<uint32 N>

__device__ static __forceinline__ void cugar::cuda::warp_static_atomic::static_sub ( bool  p, uint32 *  dest, uint32 *  result  )

inlinestatic

subtract zero or exactly N per thread to a shared value: useful to dealloc N entries from a common pool

Parameters

p	allocation predicate
dest	the destination of the atomic
result	output result

sub() [1/2]

template<uint32 N>

__device__ __forceinline__ void cugar::cuda::warp_static_atomic::sub ( bool  p )

inline

subtract zero or exactly N per thread to a shared value without waiting for the result: useful to alloc N entries from a common pool

Parameters

p	allocation predicate

sub() [2/2]

template<uint32 N>

__device__ __forceinline__ void cugar::cuda::warp_static_atomic::sub ( bool  p, uint32 *  result  )

inline

subtract zero or exactly N per thread to a shared value: useful to alloc N entries from a common pool

Parameters

p	allocation predicate

---

The documentation for this struct was generated from the following file:

• C:/p4research/research/jpantaleoni/Fermat-Public/contrib/cugar/basic/cuda/warp_atomics.h