CUDA Synchronization: Threads, Barriers & Race Conditions

CUDA Thread Synchronization: `__syncthreads()`, Race Conditions, Atomics, and Shared Memory

Introduction

CUDA allows thousands of GPU threads to execute work in parallel. This massive parallelism is one of the main reasons GPUs are so effective for scientific computing, deep learning, image processing, simulations, and other data-intensive workloads.

For many CUDA kernels, threads can operate independently:

cpp
1output[i] = input[i] * 2;

Thread 0 can process input[0], Thread 1 can process input[1], and so on.

text
1Thread 0 → input[0] → output[0]
2Thread 1 → input[1] → output[1]
3Thread 2 → input[2] → output[2]
4Thread 3 → input[3] → output[3]

In this situation, synchronization is unnecessary because one thread does not depend on another.

However, many real CUDA algorithms require threads to cooperate.

For example:

text
1Thread 0 ──┐
2Thread 1 ──┤
3Thread 2 ──┼──→ Shared data
4Thread 3 ──┤
5Thread 4 ──┘

One thread may write data that another thread needs to read.

This creates an important question:

How can CUDA threads safely coordinate their work?

CUDA provides several synchronization mechanisms for this purpose.

In this tutorial, you will learn:

What CUDA thread synchronization means
How __syncthreads() works
Why synchronization is important with shared memory
What race conditions are
Why __syncthreads() does not make operations atomic
How atomicAdd() works
Why excessive atomic operations can hurt performance
How __syncwarp() differs from __syncthreads()
How synchronization is used in tiled matrix multiplication
How synchronization relates to AI and Transformer workloads
Common CUDA synchronization mistakes
How to design safer cooperative kernels

CUDA Thread Hierarchy and Synchronization

Before learning synchronization, it is important to understand where CUDA threads live.

The basic CUDA execution hierarchy is:

text
1Grid
2 │
3 ├── Block 0
4 │    ├── Thread 0
5 │    ├── Thread 1
6 │    ├── Thread 2
7 │    └── ...
8 │
9 ├── Block 1
10 │    ├── Thread 0
11 │    ├── Thread 1
12 │    └── ...
13 │
14 └── Block 2
15      └── ...

Threads are grouped into blocks, and blocks form a grid.

Inside a block, threads are further organized into warps.

text
1Grid
2 ↓
3Blocks
4 ↓
5Warps
6 ↓
7Threads

A key CUDA design principle is that blocks are independently schedulable.

This has an important consequence:

Normal CUDA block synchronization cannot be used as a general synchronization mechanism between different blocks.

Why Threads Need Synchronization

Consider a kernel where each thread performs an independent operation:

cpp
1__global__ void scaleArray(float* output, const float* input)
2{
3    int i = blockIdx.x * blockDim.x + threadIdx.x;
4
5    output[i] = input[i] * 2.0f;
6}

Each thread reads one element and writes one independent element.

There is no dependency between threads.

Now consider shared memory:

cpp
1__shared__ float data[256];

Suppose every thread writes one value:

cpp
1data[threadIdx.x] = threadIdx.x;

Then another part of the kernel wants to use the entire array.

The execution needs to follow this logical sequence:

text
1All threads
2    ↓
3Write shared memory
4    ↓
5Synchronization
6    ↓
7Read shared memory
8    ↓
9Continue computation

Without synchronization, one thread could begin reading the shared array while another thread is still writing it.

That is where synchronization becomes important.

The `__syncthreads()` Function

The most commonly used block-level synchronization primitive is:

cpp
1__syncthreads();

It creates a barrier for threads in the same block.

Conceptually:

text
1Thread 0 ─────────────┐
2Thread 1 ────────┐    │
3Thread 2 ────────┤    │
4Thread 3 ─────────────┤
5Thread 4 ────────┘    │
6                      ↓
7                __syncthreads()
8                      ↓
9                 Continue

Threads reaching the barrier wait until the required participating threads in that block reach the same synchronization point.

A common pattern is:

cpp
1// Phase 1
2perform_work();
3
4__syncthreads();
5
6// Phase 2
7continue_work();

The important idea is that synchronization separates two phases of cooperative work.

Important Limitation of `__syncthreads()`

__syncthreads() synchronizes threads within a block.

It does not provide a normal grid-wide barrier across arbitrary blocks in a conventional kernel launch.

For example:

text
1Grid
2│
3├── Block 0
4│    ├── Thread 0
5│    ├── Thread 1
6│    └── Thread 2
7│
8└── Block 1
9     ├── Thread 0
10     ├── Thread 1
11     └── Thread 2

If Block 0 executes:

cpp
1__syncthreads();

it does not mean that Block 1 has reached the same point.

This is fundamental to CUDA's execution model.

Why CUDA Blocks Are Independent

CUDA schedules blocks onto Streaming Multiprocessors (SMs).

A simplified example might look like:

text
1Block 0 → SM 0
2Block 1 → SM 1
3Block 2 → SM 0
4Block 3 → SM 1

The exact scheduling depends on GPU hardware, resource availability, occupancy, and runtime behavior.

Therefore, a kernel should not normally assume that:

text
1Block 0
2   ↓
3Block 1
4   ↓
5Block 2

will execute in a predictable order.

This is why many GPU algorithms are divided into multiple kernel launches:

text
1Kernel 1
2   ↓
3Kernel launch boundary
4   ↓
5Kernel 2

The separation between kernel launches provides a much stronger global execution boundary than an ordinary __syncthreads().

Shared Memory and Synchronization

Shared memory is one of the most important places where block-level synchronization becomes useful.

Shared memory belongs to a thread block and can be accessed by threads within that block.

For example:

cpp
1__shared__ float tile[256];

Threads can cooperatively populate the array:

text
1Thread 0 → tile[0]
2Thread 1 → tile[1]
3Thread 2 → tile[2]
4...
5Thread 255 → tile[255]

Then all threads can use the data.

The typical pattern is:

cpp
1tile[threadIdx.x] = value;
2
3__syncthreads();
4
5use(tile);

The barrier creates a synchronization point between the write phase and the read phase.

Complete Shared Memory Example

The following example demonstrates a simple producer-consumer pattern inside one block.

cpp
1#include <cstdio>
2#include <cuda_runtime.h>
3
4__global__ void sharedMemoryExample()
5{
6    __shared__ int data[4];
7
8    int tid = threadIdx.x;
9
10    // Phase 1: each thread writes one value.
11    if (tid < 4) {
12        data[tid] = (tid + 1) * 10;
13    }
14
15    // Wait until all participating threads have completed
16    // their shared-memory writes.
17    __syncthreads();
18
19    // Phase 2: threads read the shared data.
20    if (tid < 4) {
21        printf(
22            "Thread %d reads data[%d] = %d\n",
23            tid,
24            tid,
25            data[tid]
26        );
27    }
28}
29
30int main()
31{
32    sharedMemoryExample<<<1, 4>>>();
33
34    cudaError_t err = cudaGetLastError();
35
36    if (err != cudaSuccess) {
37        fprintf(
38            stderr,
39            "Kernel launch failed: %s\n",
40            cudaGetErrorString(err)
41        );
42        return 1;
43    }
44
45    err = cudaDeviceSynchronize();
46
47    if (err != cudaSuccess) {
48        fprintf(
49            stderr,
50            "Kernel execution failed: %s\n",
51            cudaGetErrorString(err)
52        );
53        return 1;
54    }
55
56    return 0;
57}

The important section is:

cpp
1data[tid] = (tid + 1) * 10;
2
3__syncthreads();
4
5printf("%d", data[tid]);

The synchronization divides the kernel into two logical phases:

text
1PHASE 1
2Write shared memory
3       ↓
4__syncthreads()
5       ↓
6PHASE 2
7Read shared memory

What Is a Race Condition?

A race condition occurs when multiple threads access shared data concurrently and the result depends on the timing or ordering of those accesses.

A simple example is:

cpp
1counter++;

At first glance, this appears to be one operation.

Conceptually, however, it involves multiple steps:

text
1Read counter
2     ↓
3Add 1
4     ↓
5Write counter

Suppose the initial value is:

text
1counter = 0

Two threads execute the increment at approximately the same time.

text
1Thread 0              Thread 1
2
3Read 0                Read 0
4  ↓                      ↓
5Add 1                  Add 1
6  ↓                      ↓
7Write 1                Write 1

The final result can be:

text
1counter = 1

even though two increments were requested.

The expected mathematical result would be:

text
1counter = 2

This is a classic data race.

Unsafe CUDA Counter Example

Consider:

cpp
1__global__ void badCounter(int* counter)
2{
3    (*counter)++;
4}

You might launch:

cpp
1badCounter<<<1, 100>>>(counter);

It is tempting to expect:

text
1counter = 100

but the update is not safely coordinated between the threads.

Multiple threads can read and write the same memory location concurrently.

Why `__syncthreads()` Does Not Fix a Race Condition

A common beginner mistake is to write:

cpp
1(*counter)++;
2
3__syncthreads();

and assume the race has been fixed.

It has not.

The race occurs during:

cpp
1(*counter)++;

The barrier occurs afterward.

The barrier only coordinates when threads reach a synchronization point. It does not transform an ordinary read-modify-write operation into an atomic operation.

Therefore:

text
1__syncthreads()
2       ≠
3atomic operation

This distinction is extremely important in CUDA programming.

Atomic Operations

CUDA provides atomic operations for safely updating shared memory locations when multiple threads compete for the same value.

For example:

cpp
1atomicAdd(counter, 1);

Instead of performing an ordinary increment:

cpp
1(*counter)++;

threads can perform:

cpp
1atomicAdd(counter, 1);

Example:

cpp
1__global__ void safeCounter(int* counter)
2{
3    atomicAdd(counter, 1);
4}

If the counter starts at zero and 256 threads each execute one atomic increment:

text
1Initial value
2     ↓
3     0
4     ↓
5256 atomic increments
6     ↓
7   256

The updates are serialized as necessary at the atomic operation, preventing the lost-update problem of the ordinary increment.

Complete Atomic Counter Example

A more complete CUDA program can include explicit error handling and device memory management:

cpp
1#include <cstdio>
2#include <cuda_runtime.h>
3
4__global__ void incrementCounter(int* counter)
5{
6    atomicAdd(counter, 1);
7}
8
9int main()
10{
11    constexpr int threads = 256;
12
13    int hostCounter = 0;
14    int* deviceCounter = nullptr;
15
16    cudaError_t err;
17
18    err = cudaMalloc(
19        &deviceCounter,
20        sizeof(int)
21    );
22
23    if (err != cudaSuccess) {
24        fprintf(
25            stderr,
26            "cudaMalloc failed: %s\n",
27            cudaGetErrorString(err)
28        );
29        return 1;
30    }
31
32    err = cudaMemcpy(
33        deviceCounter,
34        &hostCounter,
35        sizeof(int),
36        cudaMemcpyHostToDevice
37    );
38
39    if (err != cudaSuccess) {
40        fprintf(
41            stderr,
42            "cudaMemcpy H2D failed: %s\n",
43            cudaGetErrorString(err)
44        );
45
46        cudaFree(deviceCounter);
47        return 1;
48    }
49
50    incrementCounter<<<1, threads>>>(deviceCounter);
51
52    err = cudaGetLastError();
53
54    if (err != cudaSuccess) {
55        fprintf(
56            stderr,
57            "Kernel launch failed: %s\n",
58            cudaGetErrorString(err)
59        );
60
61        cudaFree(deviceCounter);
62        return 1;
63    }
64
65    err = cudaDeviceSynchronize();
66
67    if (err != cudaSuccess) {
68        fprintf(
69            stderr,
70            "Kernel execution failed: %s\n",
71            cudaGetErrorString(err)
72        );
73
74        cudaFree(deviceCounter);
75        return 1;
76    }
77
78    err = cudaMemcpy(
79        &hostCounter,
80        deviceCounter,
81        sizeof(int),
82        cudaMemcpyDeviceToHost
83    );
84
85    if (err != cudaSuccess) {
86        fprintf(
87            stderr,
88            "cudaMemcpy D2H failed: %s\n",
89            cudaGetErrorString(err)
90        );
91
92        cudaFree(deviceCounter);
93        return 1;
94    }
95
96    printf("Counter = %d\n", hostCounter);
97
98    cudaFree(deviceCounter);
99
100    return 0;
101}

Expected output:

text
1Counter = 256

assuming no other code modifies the counter.

Why Atomic Operations Can Be Expensive

Atomic operations solve a correctness problem, but they are not automatically the fastest solution.

Imagine:

text
11,000,000 threads
2       ↓
3atomicAdd()
4       ↓
5ONE memory location

A huge number of threads are competing for the same location.

This can create contention and reduce performance.

For high-performance CUDA kernels, a better strategy is often to reduce the number of global updates.

Instead of:

text
1Every thread
2    ↓
3Global atomic

a reduction-oriented design can look like:

text
1Thread-local computation
2        ↓
3Warp-level reduction
4        ↓
5Block-level reduction
6        ↓
7Small number of global updates

This is one of the major ideas behind optimized CUDA reductions.

Synchronization and Parallel Reduction

Suppose 256 threads need to calculate the sum of their values.

A naive approach might make every thread update one global counter:

cpp
1atomicAdd(globalSum, value);

That can create heavy contention.

A more efficient approach is to first combine values locally.

Conceptually:

text
1256 threads
2    ↓
3Local values
4    ↓
5Warp reduction
6    ↓
7Block reduction
8    ↓
9One global update

This dramatically reduces the number of global atomic operations.

Understanding synchronization is therefore an important prerequisite for learning CUDA reduction algorithms.

Shared Memory Sum Example

Consider four threads:

text
1Thread 0 → 1
2Thread 1 → 2
3Thread 2 → 3
4Thread 3 → 4

The target sum is:

text
11 + 2 + 3 + 4 = 10

A simple shared-memory example is:

cpp
1#include <cstdio>
2#include <cuda_runtime.h>
3
4__global__ void sharedSum()
5{
6    __shared__ int data[4];
7
8    int tid = threadIdx.x;
9
10    data[tid] = tid + 1;
11
12    // Make sure every thread has completed its write.
13    __syncthreads();
14
15    int sum =
16        data[0] +
17        data[1] +
18        data[2] +
19        data[3];
20
21    printf(
22        "Thread %d sees sum = %d\n",
23        tid,
24        sum
25    );
26}
27
28int main()
29{
30    sharedSum<<<1, 4>>>();
31
32    cudaError_t err = cudaGetLastError();
33
34    if (err != cudaSuccess) {
35        fprintf(
36            stderr,
37            "Launch failed: %s\n",
38            cudaGetErrorString(err)
39        );
40        return 1;
41    }
42
43    err = cudaDeviceSynchronize();
44
45    if (err != cudaSuccess) {
46        fprintf(
47            stderr,
48            "Execution failed: %s\n",
49            cudaGetErrorString(err)
50        );
51        return 1;
52    }
53
54    return 0;
55}

The synchronization is important:

cpp
1data[tid] = tid + 1;
2
3__syncthreads();
4
5int sum = data[0] + data[1] + data[2] + data[3];

Without the barrier, a thread could attempt to read the shared array before another thread has finished writing its element.

What Can Happen Without Synchronization?

Consider:

text
1Thread 0
2   ↓
3write data[0]
4   ↓
5read data[1]

while another thread is still doing:

text
1Thread 1
2   ↓
3has not written data[1] yet

The reader cannot safely assume that the other thread has completed its write.

The general pattern is:

text
1Producer threads
2      ↓
3Write shared data
4      ↓
5Synchronization
6      ↓
7Consumer threads
8      ↓
9Read shared data

A Critical `__syncthreads()` Rule

Be careful when placing __syncthreads() inside conditional control flow.

This pattern is dangerous:

cpp
1if (threadIdx.x < 16) {
2    __syncthreads();
3}

Some threads enter the barrier while other threads skip it.

A much safer beginner pattern is:

cpp
1if (threadIdx.x < 16) {
2    // Perform conditional work.
3}
4
5__syncthreads();

Now the barrier is outside the conditional block.

The general rule is:

Design block-wide barriers so that participating threads reach them consistently.

Incorrect barrier placement can cause synchronization bugs, undefined behavior, or a kernel that fails to make progress.

`__syncwarp()`

CUDA also provides warp-level synchronization:

cpp
1__syncwarp();

A warp consists of a group of threads executed together according to the GPU's SIMT execution model.

Conceptually:

text
1Warp
2 ├── Lane 0
3 ├── Lane 1
4 ├── Lane 2
5 ├── ...
6 └── Lane 31
7        ↓
8   __syncwarp()

__syncwarp() is useful when coordinating threads within a warp.

A simple example:

cpp
1__global__ void warpExample()
2{
3    int lane = threadIdx.x % 32;
4
5    // Warp-level work.
6
7    __syncwarp();
8
9    // Continue warp-level work.
10}

Warp-level synchronization becomes particularly useful when learning:

Warp shuffle operations
Parallel reductions
Prefix scans
Cooperative warp algorithms
Tensor Core programming
Highly optimized CUDA kernels

`syncthreads()` vs `syncwarp()`

Feature	`__syncthreads()`	`__syncwarp()`
Scope	Block	Warp
Main purpose	Block-wide coordination	Warp-level coordination
Typical use	Shared-memory cooperation	Warp primitives
Number of threads	Threads in the block	Participating warp lanes
Typical abstraction	Block	Warp

Think of them as:

text
1__syncwarp()
2     ↓
3  Warp scope

and:

text
1__syncthreads()
2     ↓
3  Block scope

Choosing the correct synchronization scope is important for both correctness and performance.

Synchronization Scope

CUDA programming involves multiple levels of execution:

text
1Thread
2   ↓
3Warp
4   ↓
5Block
6   ↓
7Grid
8   ↓
9Device

Different CUDA primitives operate at different scopes.

For example:

cpp
1__syncwarp();

is a warp-level primitive.

cpp
1__syncthreads();

is a block-level primitive.

Atomic operations can also have different memory-scope and ordering semantics depending on the operation and CUDA memory model being used.

As you move into advanced CUDA programming, understanding scope becomes increasingly important.

Synchronization in Tiled Matrix Multiplication

Synchronization becomes especially important in matrix multiplication.

Matrix multiplication is fundamental to:

Neural networks
Transformers
Large language models
Convolutional workloads
Scientific computing
Computer vision

For:

text
1C = A × B

a high-performance CUDA implementation often uses tiling.

Instead of repeatedly loading data from global memory, threads cooperatively load tiles into shared memory.

Conceptually:

text
1Global Memory
2      ↓
3Load tile
4      ↓
5Shared Memory
6      ↓
7__syncthreads()
8      ↓
9Compute
10      ↓
11__syncthreads()
12      ↓
13Load next tile

The synchronization points are essential because the threads are cooperating on the same shared-memory tile.

Why Two Synchronization Points Are Common

Consider this simplified tiled algorithm:

text
1Load tile A
2Load tile B
3     ↓
4Synchronize
5     ↓
6Compute using tiles
7     ↓
8Synchronize
9     ↓
10Load next tiles

The first barrier ensures that the required tile data has been loaded before computation begins.

The second barrier ensures that threads have finished using the current tile before the shared-memory storage is reused for the next tile.

Without these synchronization points, one thread could overwrite shared memory while another thread is still reading it.

Simplified Tiled Matrix Multiplication Kernel

The following example demonstrates the synchronization pattern:

cpp
1#include <cuda_runtime.h>
2
3constexpr int TILE_SIZE = 16;
4
5__global__ void matrixMultiplyTiled(
6    const float* A,
7    const float* B,
8    float* C,
9    int M,
10    int N,
11    int K
12)
13{
14    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
15    __shared__ float tileB[TILE_SIZE][TILE_SIZE];
16
17    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
18    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
19
20    float sum = 0.0f;
21
22    int tiles = (K + TILE_SIZE - 1) / TILE_SIZE;
23
24    for (int t = 0; t < tiles; ++t) {
25
26        int aCol = t * TILE_SIZE + threadIdx.x;
27        int bRow = t * TILE_SIZE + threadIdx.y;
28
29        // Load A tile.
30        if (row < M && aCol < K) {
31            tileA[threadIdx.y][threadIdx.x] =
32                A[row * K + aCol];
33        } else {
34            tileA[threadIdx.y][threadIdx.x] = 0.0f;
35        }
36
37        // Load B tile.
38        if (bRow < K && col < N) {
39            tileB[threadIdx.y][threadIdx.x] =
40                B[bRow * N + col];
41        } else {
42            tileB[threadIdx.y][threadIdx.x] = 0.0f;
43        }
44
45        // Make sure the complete tile is available.
46        __syncthreads();
47
48        // Compute using the shared-memory tiles.
49        for (int k = 0; k < TILE_SIZE; ++k) {
50            sum +=
51                tileA[threadIdx.y][k] *
52                tileB[k][threadIdx.x];
53        }
54
55        // Make sure no thread is still using the
56        // current tiles before overwriting them.
57        __syncthreads();
58    }
59
60    if (row < M && col < N) {
61        C[row * N + col] = sum;
62    }
63}

The key synchronization pattern is:

cpp
1__syncthreads();
2
3for (int k = 0; k < TILE_SIZE; ++k) {
4    sum += tileA[threadIdx.y][k] *
5           tileB[k][threadIdx.x];
6}
7
8__syncthreads();

This is a fundamental pattern in shared-memory CUDA kernels.

Why Synchronization Matters for AI Workloads

Synchronization is not just a theoretical CUDA concept.

Modern AI workloads perform enormous numbers of tensor operations.

For example, Transformer models contain operations such as:

text
1Q = XWq
2K = XWk
3V = XWv

followed by operations involving:

text
1QKᵀ
2Softmax
3Attention × V
4Output projection
5MLP

Many of these operations rely on highly optimized GPU kernels.

A simplified view of a high-performance GPU computation is:

text
1Threads
2   ↓
3Warps
4   ↓
5Shared Memory
6   ↓
7Synchronization
8   ↓
9Tensor/Matrix Computation
10   ↓
11Global Memory

Understanding CUDA synchronization therefore provides useful background for understanding how optimized deep-learning kernels operate.

Synchronization vs Atomicity

These concepts should not be confused.

Synchronization

Synchronization coordinates thread execution.

Example:

cpp
1__syncthreads();

Conceptually:

text
1Reach barrier
2     ↓
3Wait
4     ↓
5Continue

Atomicity

Atomicity applies to a specific operation on shared data.

Example:

cpp
1atomicAdd(counter, 1);

It prevents competing updates from producing the same lost-update behavior as an ordinary read-modify-write operation.

Memory Ordering and Visibility

Memory ordering concerns how memory operations are ordered and observed between threads and scopes.

CUDA provides additional primitives and memory-model semantics for advanced use cases.

The key beginner takeaway is:

text
1Synchronization
2      ≠
3Atomicity
4      ≠
5General memory ordering

They are related concepts, but they solve different problems.

Three Common CUDA Synchronization Problems

Race Condition

Multiple threads access the same data without safe coordination.

text
1Multiple threads
2      ↓
3Same memory location
4      ↓
5Concurrent update
6      ↓
7Race condition

Possible solutions include:

text
1Atomic operations
2        or
3Algorithm redesign
4        or
5Parallel reduction

Missing Synchronization

One thread writes data while another thread reads it before the cooperative write phase is complete.

text
1Thread A
2   ↓
3Write
4
5Thread B
6   ↓
7Read too early

For block-level shared-memory cooperation, the appropriate solution may be:

cpp
1__syncthreads();

Invalid Barrier Control Flow

A block-wide barrier is placed in a control-flow path that does not have valid participation.

For example:

cpp
1if (condition) {
2    __syncthreads();
3}

The correct design depends on the algorithm and control-flow structure, but beginners should generally keep block-wide barriers outside divergent conditional paths.

Production CUDA Error Checking

When learning CUDA synchronization, it is useful to distinguish kernel launch errors from execution errors.

A useful pattern is:

cpp
1kernel<<<grid, block>>>();
2
3cudaError_t err = cudaGetLastError();
4
5if (err != cudaSuccess) {
6    fprintf(
7        stderr,
8        "Kernel launch failed: %s\n",
9        cudaGetErrorString(err)
10    );
11}
12
13err = cudaDeviceSynchronize();
14
15if (err != cudaSuccess) {
16    fprintf(
17        stderr,
18        "Kernel execution failed: %s\n",
19        cudaGetErrorString(err)
20    );
21}

This is better than launching a kernel and assuming it succeeded.

For larger projects, a reusable CUDA error-checking macro or helper function can make this pattern less repetitive.

A Practical Mental Model

When designing a CUDA kernel, ask these questions in order:

Question 1: Are threads independent?

If every thread works on independent data:

text
1Thread → independent element

you may not need synchronization.

Question 2: Do threads share data?

If threads cooperate through shared memory:

text
1Thread A ──┐
2Thread B ──┼──→ Shared Memory
3Thread C ──┘

you must reason about when writes and reads occur.

Question 3: Can multiple threads update the same location?

If yes:

text
1Thread A ──┐
2Thread B ──┼──→ Same variable
3Thread C ──┘

you need to consider atomics or a different algorithm.

Question 4: What is the synchronization scope?

Ask whether coordination is required at:

text
1Warp level
2Block level
3Grid level
4Device/system level

Choosing the correct scope is essential.

CUDA Synchronization Mental Model

A useful mental model is:

text
1                    GRID
2                      │
3          ┌───────────┴───────────┐
4          ↓                       ↓
5       BLOCK 0                 BLOCK 1
6          │                       │
7      ┌───┴───┐               ┌───┴───┐
8      ↓       ↓               ↓       ↓
9    Warp    Warp            Warp    Warp
10      │       │               │       │
11   Threads Threads         Threads Threads
12      │
13      ↓
14Shared Memory
15      │
16      ↓
17__syncthreads()
18      │
19      ↓
20Cooperative Computation

The important boundaries are:

text
1Warp
2 ↓
3__syncwarp()
4
5Block
6 ↓
7__syncthreads()
8
9Multiple kernel stages
10 ↓
11Kernel launch boundary

Practice Challenge

Build a kernel with:

text
1Block size = 256

Each thread should:

Calculate its thread ID.
Store the thread ID in shared memory.
Synchronize with the block.
Read the previous thread's value.
Print the result.

For example:

text
1Thread 1 → reads shared[0]
2
3Thread 2 → reads shared[1]
4
5Thread 3 → reads shared[2]

A good starting implementation is:

cpp
1#include <cstdio>
2#include <cuda_runtime.h>
3
4__global__ void previousThreadValue()
5{
6    __shared__ int data[256];
7
8    int tid = threadIdx.x;
9
10    data[tid] = tid;
11
12    __syncthreads();
13
14    if (tid > 0) {
15        int previous = data[tid - 1];
16
17        printf(
18            "Thread %d reads %d\n",
19            tid,
20            previous
21        );
22    }
23}
24
25int main()
26{
27    previousThreadValue<<<1, 256>>>();
28
29    cudaError_t err = cudaGetLastError();
30
31    if (err != cudaSuccess) {
32        fprintf(
33            stderr,
34            "Launch failed: %s\n",
35            cudaGetErrorString(err)
36        );
37        return 1;
38    }
39
40    err = cudaDeviceSynchronize();
41
42    if (err != cudaSuccess) {
43        fprintf(
44            stderr,
45            "Execution failed: %s\n",
46            cudaGetErrorString(err)
47        );
48        return 1;
49    }
50
51    return 0;
52}

The important sequence is:

text
1Thread indexing
2      ↓
3Shared memory
4      ↓
5__syncthreads()
6      ↓
7Read another thread's data
8      ↓
9Thread cooperation

Practice Questions

Question 1

What is the primary purpose of __syncthreads()?

Question 2

Does __syncthreads() synchronize threads from different blocks?

Question 3

Why doesn't this automatically fix a race condition?

cpp
1counter++;
2
3__syncthreads();

Question 4

What CUDA primitive can be used for an atomic integer addition?

Question 5

Why can excessive atomic operations reduce performance?

Question 6

What is the primary difference between:

cpp
1__syncwarp();

and:

cpp
1__syncthreads();

Question 7

Why do tiled matrix multiplication kernels commonly use synchronization before and after computation on a shared-memory tile?

Summary

CUDA synchronization allows threads to safely cooperate when their work is no longer independent.

The most important concepts are:

text
1Independent threads
2      ↓
3No synchronization may be required

When threads cooperate:

text
1Shared memory
2      ↓
3Synchronization
4      ↓
5__syncthreads()

When multiple threads update the same memory location:

text
1Concurrent updates
2      ↓
3Potential race condition
4      ↓
5Atomic operation or algorithm redesign

At the warp level:

text
1Warp cooperation
2      ↓
3__syncwarp()

For high-performance algorithms:

text
1Threads
2   ↓
3Warps
4   ↓
5Shared Memory
6   ↓
7Synchronization
8   ↓
9Parallel Reduction / Matrix Computation

The most important distinction to remember is:

text
1__syncthreads()
2    ↓
3Synchronizes participating threads in a block
4
5atomicAdd()
6    ↓
7Provides an atomic update for the targeted memory location
8
9__syncwarp()
10    ↓
11Synchronizes participating threads at warp scope

What Comes Next?

A natural next step is:

text
1CUDA Fundamentals
2      ↓
3GPU Architecture
4      ↓
5CUDA Kernels
6      ↓
7Thread Indexing
8      ↓
9Threads / Blocks / Grids
10      ↓
11Warps / SIMT
12      ↓
13Synchronization
14      ↓
15Shared Memory
16      ↓
17CUDA Memory Hierarchy
18      ↓
19Global Memory Access
20      ↓
21Memory Coalescing
22      ↓
23Shared-Memory Bank Conflicts
24      ↓
25Warp-Level Programming
26      ↓
27Parallel Reduction
28      ↓
29Occupancy
30      ↓
31Kernel Optimization

The next lesson should focus on CUDA Memory Hierarchy and Shared Memory.

Once you understand synchronization, shared memory becomes much easier to reason about because you can see not only where data is stored, but also why threads must coordinate when loading, consuming, and reusing that data.

CUDA Thread Synchronization: __syncthreads(), Race Conditions, Atomics, and Shared Memory

Introduction

CUDA Thread Hierarchy and Synchronization

Why Threads Need Synchronization

The __syncthreads() Function

Important Limitation of __syncthreads()

Why CUDA Blocks Are Independent

Shared Memory and Synchronization

Complete Shared Memory Example

What Is a Race Condition?

Unsafe CUDA Counter Example

Why __syncthreads() Does Not Fix a Race Condition

Atomic Operations

Complete Atomic Counter Example

Why Atomic Operations Can Be Expensive

Synchronization and Parallel Reduction

Shared Memory Sum Example

What Can Happen Without Synchronization?

A Critical __syncthreads() Rule

__syncwarp()

__syncthreads() vs __syncwarp()

Synchronization Scope

Synchronization in Tiled Matrix Multiplication

Why Two Synchronization Points Are Common

Simplified Tiled Matrix Multiplication Kernel

Why Synchronization Matters for AI Workloads

Synchronization vs Atomicity

Synchronization

Atomicity

Memory Ordering and Visibility

Three Common CUDA Synchronization Problems

Race Condition

Missing Synchronization

Invalid Barrier Control Flow

Production CUDA Error Checking

A Practical Mental Model

Question 1: Are threads independent?

Question 2: Do threads share data?

Question 3: Can multiple threads update the same location?

Question 4: What is the synchronization scope?

CUDA Synchronization Mental Model

Practice Challenge

Practice Questions

Question 1

Question 2

Question 3

Question 4

Question 5

Question 6

Question 7

Summary

What Comes Next?

CUDA Thread Synchronization: `__syncthreads()`, Race Conditions, Atomics, and Shared Memory

The `__syncthreads()` Function

Important Limitation of `__syncthreads()`

Why `__syncthreads()` Does Not Fix a Race Condition

A Critical `__syncthreads()` Rule

`__syncwarp()`

`syncthreads()` vs `syncwarp()`