atomic.cuh File Reference

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes
union	uint32_short2
union	uint32_char2

Functions
static __device__ double2	atomicAdd (double2 *addr, double2 val)
	Implementation of double2 atomic addition using two double-precision additions. More...
static __device__ float2	atomicAdd (float2 *addr, float2 val)
	Implementation of float2 atomic addition using two single-precision additions. More...
static __device__ int2	atomicAdd (int2 *addr, int2 val)
	Implementation of int2 atomic addition using two int additions. More...
static __device__ short2	atomicAdd (short2 *addr, short2 val)
	Implementation of short2 atomic addition using compare and swap. More...
static __device__ char2	atomicAdd (char2 *addr, char2 val)
	Implementation of char2 atomic addition using compare and swap. More...
static __device__ float	atomicMax (float *addr, float val)
	Implementation of single-precision atomic max using compare and swap. May not support NaNs properly... More...

Detailed Description

Description

Provides definitions of atomic functions that are not native to CUDA. These are intentionally not declared in the namespace to avoid confusion when resolving the native atomicAdd functions.

Definition in file atomic.cuh.

Function Documentation

atomicAdd() [1/5]

static __device__ double2 atomicAdd ( double2 *  addr, double2  val  )

inlinestatic

Implementation of double2 atomic addition using two double-precision additions.

Parameters

addr	Address that stores the atomic variable to be updated
val	Value to be added to the atomic

Definition at line 51 of file atomic.cuh.

Referenced by quda::gauge::Accessor< Float, nColor, QUDA_QDP_GAUGE_ORDER, storeFloat, use_tex >::atomic_add(), quda::gauge::Accessor< Float, nColor, QUDA_MILC_GAUGE_ORDER, storeFloat, use_tex >::atomic_add(), quda::gauge::Accessor< Float, nColor, QUDA_FLOAT2_GAUGE_ORDER, storeFloat, use_tex >::atomic_add(), atomicAdd(), quda::computeVUV(), quda::GaugeFixHit_AtomicAdd(), quda::isUnitary(), and quda::ProjectSU3kernel().

Here is the caller graph for this function:

atomicAdd() [2/5]

static __device__ float2 atomicAdd ( float2 *  addr, float2  val  )

inlinestatic

Implementation of float2 atomic addition using two single-precision additions.

Parameters

addr	Address that stores the atomic variable to be updated
val	Value to be added to the atomic

Definition at line 72 of file atomic.cuh.

References atomicAdd().

Here is the call graph for this function:

atomicAdd() [3/5]

static __device__ int2 atomicAdd ( int2 *  addr, int2  val  )

inlinestatic

Implementation of int2 atomic addition using two int additions.

Parameters

addr	Address that stores the atomic variable to be updated
val	Value to be added to the atomic

Definition at line 86 of file atomic.cuh.

References atomicAdd().

Here is the call graph for this function:

atomicAdd() [4/5]

static __device__ short2 atomicAdd ( short2 *  addr, short2  val  )

inlinestatic

Implementation of short2 atomic addition using compare and swap.

Parameters

addr	Address that stores the atomic variable to be updated
val	Value to be added to the atomic

Definition at line 102 of file atomic.cuh.

References uint32_short2::i, and uint32_short2::s.

atomicAdd() [5/5]

static __device__ char2 atomicAdd ( char2 *  addr, char2  val  )

inlinestatic

Implementation of char2 atomic addition using compare and swap.

Parameters

addr	Address that stores the atomic variable to be updated
val	Value to be added to the atomic

Definition at line 123 of file atomic.cuh.

References uint32_char2::i, and uint32_char2::s.

atomicMax()

static __device__ float atomicMax ( float *  addr, float  val  )

inlinestatic

Implementation of single-precision atomic max using compare and swap. May not support NaNs properly...

Parameters

addr	Address that stores the atomic variable to be updated
val	Value to be added to the atomic

Definition at line 142 of file atomic.cuh.

Referenced by quda::CalculateYhatGPU().

Here is the caller graph for this function: