Atomics#

Single OS thread + cooperative fibers: fibers only ever switch at an explicit __syncthreads()/warp sync, never mid-expression, so a plain read-modify-write IS atomic with respect to every other thread in the block. Each returns the OLD value, like the real CUDA intrinsics. The value argument is a non-deduced context (type_identity_t) so atomicAdd(&ull, 1) etc. convert like CUDA’s per-type overloads instead of failing template deduction.

Defined in code/simulator.h.

This page documents 11 symbols: 11 functions.

Functions#

Signature

Description

Availability

Location

Example

inline T atomicAdd (T* a, std::type_identity_t<T> v)

Single OS thread + cooperative fibers: fibers only ever switch at an explicit __syncthreads()/warp sync, never mid-expression, so a plain read-modify-write IS atomic with respect to every other …

sim only

simulator.h:963

View

inline T atomicAnd (T* a, std::type_identity_t<T> v)

sim only

simulator.h:975

View

inline T atomicCAS (T* a, std::type_identity_t<T> compare, std::type_identity_t<T> v)

A FAILED atomicCAS is the signature of a spin-wait, so it yields: that is what lets a device spin-lock (while (atomicCAS(mutex, 0, 1) != 0);) make progress.

sim only

simulator.h:981

View

inline unsigned atomicDec(unsigned* a, unsigned v)

sim only

simulator.h:996

View

inline T atomicExch(T* a, std::type_identity_t<T> v)

atomicExch yields when the exchange changed nothing, i.e.

sim only

simulator.h:967

View

inline unsigned atomicInc(unsigned* a, unsigned v)

atomicInc/Dec are unsigned-only wrap counters, matching CUDA’s semantics.

sim only

simulator.h:993

View

inline T atomicMax (T* a, std::type_identity_t<T> v)

sim only

simulator.h:974

View

inline T atomicMin (T* a, std::type_identity_t<T> v)

sim only

simulator.h:973

View

inline T atomicOr (T* a, std::type_identity_t<T> v)

sim only

simulator.h:976

View

inline T atomicSub (T* a, std::type_identity_t<T> v)

sim only

simulator.h:964

View

inline T atomicXor (T* a, std::type_identity_t<T> v)

sim only

simulator.h:977

View

Examples#

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

atomicAdd (test_atomics.cu:38)
28    } while (0)
29
30constexpr int N      = 4096;   // number of contributing threads
31constexpr int BLOCK  = 128;
32constexpr int BLOCKS = N / BLOCK;
33constexpr int BINS   = 16;
34
35// ---- Kernels ---------------------------------------------------------------
36
37__global__ void kCounterInt(int* c) {
38    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicAdd(c, 1);
39}
40__global__ void kCounterUll(unsigned long long* c) {
41    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicAdd(c, 1ULL);
42}
43__global__ void kCounterFloat(float* c) {
44    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicAdd(c, 1.0f);
45}
46__global__ void kCounterDouble(double* c) {
47    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicAdd(c, 1.0);
48}
atomicAnd (test_atomics.cu:78)
68}
69
70// 32 threads each set their own bit via atomicOr -> full mask.
71__global__ void kBitmaskOr(unsigned* mask) {
72    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
73    if (i < 32u) atomicOr(mask, 1u << i);
74}
75// Start all-ones; each of 32 threads clears its bit via atomicAnd -> 0.
76__global__ void kBitmaskAnd(unsigned* mask) {
77    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
78    if (i < 32u) atomicAnd(mask, ~(1u << i));
79}
80// XOR each value in; final = running parity of all values.
81__global__ void kXor(const unsigned* data, unsigned* acc) {
82    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
83    if (i < static_cast<unsigned>(N)) atomicXor(acc, data[i]);
84}
85
86__global__ void kSub(int* c) {
87    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicSub(c, 1);
88}
atomicCAS (test_atomics.cu:67)
57    if (i < static_cast<unsigned>(N)) atomicMax(out, data[i]);
58}
59__global__ void kMin(const int* data, int* out) {
60    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
61    if (i < static_cast<unsigned>(N)) atomicMin(out, data[i]);
62}
63
64// Exactly one thread should win the compare-and-swap for the sentinel.
65__global__ void kClaimOnce(int* owner, int* winners) {
66    const int i = static_cast<int>(blockIdx.x * blockDim.x + threadIdx.x);
67    if (i < N && atomicCAS(owner, -1, i) == -1) atomicAdd(winners, 1);
68}
69
70// 32 threads each set their own bit via atomicOr -> full mask.
71__global__ void kBitmaskOr(unsigned* mask) {
72    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
73    if (i < 32u) atomicOr(mask, 1u << i);
74}
75// Start all-ones; each of 32 threads clears its bit via atomicAnd -> 0.
76__global__ void kBitmaskAnd(unsigned* mask) {
77    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
78    if (i < 32u) atomicAnd(mask, ~(1u << i));
atomicDec (test_web_cs_atomics.cu:44)
34__global__ void testKernel(int *g_odata)
35{
36    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
37
38    atomicAdd(&g_odata[0], 10);
39    atomicSub(&g_odata[1], 10);
40    atomicExch(&g_odata[2], tid);
41    atomicMax(&g_odata[3], tid);
42    atomicMin(&g_odata[4], tid);
43    atomicInc((unsigned int *)&g_odata[5], 17);
44    atomicDec((unsigned int *)&g_odata[6], 137);
45    atomicCAS(&g_odata[7], tid - 1, tid);
46    atomicAnd(&g_odata[8], 2 * tid + 7);
47    atomicOr(&g_odata[9], 1 << tid);
48    atomicXor(&g_odata[10], tid);
49}
50
51// verbatim CPU reference from simpleAtomicIntrinsics_cpu.cpp (computeGold), with
52// its printf diagnostics kept. Returns true iff every atomic matched.
53// [SIM] max/min are provided by simulator.h (exact-type overloads; builtins under nvcc).
54static bool computeGold(int *gpuData, const int len)
55{
atomicExch (test_web_cbe_hashtable.cu:55)
45    Lock() {
46        HANDLE_ERROR(cudaMalloc((void**)&mutex, sizeof(int)));
47        HANDLE_ERROR(cudaMemset(mutex, 0, sizeof(int)));
48    }
49    __device__ void lock() {
50        while (atomicCAS(mutex, 0, 1) != 0);
51        __threadfence();
52    }
53    __device__ void unlock() {
54        __threadfence();
55        atomicExch(mutex, 0);
56    }
57};
58
59#define SIZE          (64 * 1024)                 // [SIM] was 100*1024*1024
60#define ELEMENTS      (SIZE / sizeof(unsigned int))
61#define HASH_ENTRIES  1024
62
63struct Entry {
64    unsigned int key;
65    void*        value;
66    Entry*       next;
atomicInc (test_web_cs_atomics.cu:43)
33// verbatim device kernel from simpleAtomicIntrinsics_kernel.cuh
34__global__ void testKernel(int *g_odata)
35{
36    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
37
38    atomicAdd(&g_odata[0], 10);
39    atomicSub(&g_odata[1], 10);
40    atomicExch(&g_odata[2], tid);
41    atomicMax(&g_odata[3], tid);
42    atomicMin(&g_odata[4], tid);
43    atomicInc((unsigned int *)&g_odata[5], 17);
44    atomicDec((unsigned int *)&g_odata[6], 137);
45    atomicCAS(&g_odata[7], tid - 1, tid);
46    atomicAnd(&g_odata[8], 2 * tid + 7);
47    atomicOr(&g_odata[9], 1 << tid);
48    atomicXor(&g_odata[10], tid);
49}
50
51// verbatim CPU reference from simpleAtomicIntrinsics_cpu.cpp (computeGold), with
52// its printf diagnostics kept. Returns true iff every atomic matched.
53// [SIM] max/min are provided by simulator.h (exact-type overloads; builtins under nvcc).
54static bool computeGold(int *gpuData, const int len)
atomicMax (test_atomics.cu:57)
47    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicAdd(c, 1.0);
48}
49
50__global__ void kHistogram(const int* data, unsigned* hist) {
51    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
52    if (i < static_cast<unsigned>(N)) atomicAdd(&hist[data[i] % BINS], 1u);
53}
54
55__global__ void kMax(const int* data, int* out) {
56    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
57    if (i < static_cast<unsigned>(N)) atomicMax(out, data[i]);
58}
59__global__ void kMin(const int* data, int* out) {
60    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
61    if (i < static_cast<unsigned>(N)) atomicMin(out, data[i]);
62}
63
64// Exactly one thread should win the compare-and-swap for the sentinel.
65__global__ void kClaimOnce(int* owner, int* winners) {
66    const int i = static_cast<int>(blockIdx.x * blockDim.x + threadIdx.x);
67    if (i < N && atomicCAS(owner, -1, i) == -1) atomicAdd(winners, 1);
68}
atomicMin (test_atomics.cu:61)
51    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
52    if (i < static_cast<unsigned>(N)) atomicAdd(&hist[data[i] % BINS], 1u);
53}
54
55__global__ void kMax(const int* data, int* out) {
56    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
57    if (i < static_cast<unsigned>(N)) atomicMax(out, data[i]);
58}
59__global__ void kMin(const int* data, int* out) {
60    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
61    if (i < static_cast<unsigned>(N)) atomicMin(out, data[i]);
62}
63
64// Exactly one thread should win the compare-and-swap for the sentinel.
65__global__ void kClaimOnce(int* owner, int* winners) {
66    const int i = static_cast<int>(blockIdx.x * blockDim.x + threadIdx.x);
67    if (i < N && atomicCAS(owner, -1, i) == -1) atomicAdd(winners, 1);
68}
69
70// 32 threads each set their own bit via atomicOr -> full mask.
71__global__ void kBitmaskOr(unsigned* mask) {
72    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
atomicOr (test_atomics.cu:73)
63// Exactly one thread should win the compare-and-swap for the sentinel.
64__global__ void kClaimOnce(int* owner, int* winners) {
65    const int i = static_cast<int>(blockIdx.x * blockDim.x + threadIdx.x);
66    if (i < N && atomicCAS(owner, -1, i) == -1) atomicAdd(winners, 1);
67}
68
69// 32 threads each set their own bit via atomicOr -> full mask.
70__global__ void kBitmaskOr(unsigned* mask) {
71    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
72    if (i < 32u) atomicOr(mask, 1u << i);
73}
74// Start all-ones; each of 32 threads clears its bit via atomicAnd -> 0.
75__global__ void kBitmaskAnd(unsigned* mask) {
76    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
77    if (i < 32u) atomicAnd(mask, ~(1u << i));
78}
79// XOR each value in; final = running parity of all values.
80__global__ void kXor(const unsigned* data, unsigned* acc) {
81    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
82    if (i < static_cast<unsigned>(N)) atomicXor(acc, data[i]);
83}
atomicSub (test_atomics.cu:87)
77    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
78    if (i < 32u) atomicAnd(mask, ~(1u << i));
79}
80// XOR each value in; final = running parity of all values.
81__global__ void kXor(const unsigned* data, unsigned* acc) {
82    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
83    if (i < static_cast<unsigned>(N)) atomicXor(acc, data[i]);
84}
85
86__global__ void kSub(int* c) {
87    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicSub(c, 1);
88}
89
90// ---- Host driver -----------------------------------------------------------
91
92template <class T>
93static T* devFrom(const T& init) {
94    T* d = nullptr;
95    CUDA_CHECK(cudaMalloc(&d, sizeof(T)));
96    CUDA_CHECK(cudaMemcpy(d, &init, sizeof(T), cudaMemcpyHostToDevice));
97    return d;
98}
atomicXor (test_atomics.cu:83)
73    if (i < 32u) atomicOr(mask, 1u << i);
74}
75// Start all-ones; each of 32 threads clears its bit via atomicAnd -> 0.
76__global__ void kBitmaskAnd(unsigned* mask) {
77    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
78    if (i < 32u) atomicAnd(mask, ~(1u << i));
79}
80// XOR each value in; final = running parity of all values.
81__global__ void kXor(const unsigned* data, unsigned* acc) {
82    const unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
83    if (i < static_cast<unsigned>(N)) atomicXor(acc, data[i]);
84}
85
86__global__ void kSub(int* c) {
87    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicSub(c, 1);
88}
89
90// ---- Host driver -----------------------------------------------------------
91
92template <class T>
93static T* devFrom(const T& init) {
94    T* d = nullptr;
On this page