CUDA Two-Stage Reduction with Atomics

hardPROReductions, Sync & AtomicsCUDA Basics with Numba

Problem Statement

Reduce a 1D array to a single sum by combining a shared-memory tree reduction within each block with a single cuda.atomic.add per block into the global result. Verify out[0] equals x.sum().

Background

Per-block: load into shared memory, tree-reduce to s[0]. Then only thread 0 of each block does cuda.atomic.add(out, 0, s[0]) — so there are just gridDim.x atomics total instead of n, which is far less contention than one atomic per element.

Your Task

Implement reduce_sum_kernel and run(n=8192) returning whether the global sum matches x.sum().

How it is tested

Your solution must define a top-level function run(...) that allocates the inputs, copies them to the GPU, launches your @cuda.jit kernel, and returns a Python bool from np.allclose(gpu_result, reference). The grader prints run(...); the expected output is True.

Example

Input

n = 8192

Output

True

Reasoning

The input n = 8192 represents the size of the 1D array to be reduced to a single sum.
The reduce_sum_kernel function is launched with this input, where each block performs a shared-memory tree reduction to calculate the sum of its assigned elements, resulting in a partial sum stored in s[0].
Only thread 0 of each block then performs a cuda.atomic.add operation to add its block's partial sum to the global result out[0], minimizing contention with only $gridDim.x$ atomic operations.
The final output True indicates that the global sum out[0] matches the reference sum calculated using x.sum(), verifying the correctness of the CUDA two-stage reduction implementation.

Constraints

Shared-memory tree reduction within the block (cuda.syncthreads() between steps)
Only thread 0 calls cuda.atomic.add(out, 0, s[0])
Initialize the device accumulator to 0

Learn this in

Parallel Patterns: Reductions, Atomics & TilingCUDA with Numba ML InfrastructureSystems & Networking

More from CUDA Basics with Numba

CUDA Dot Producthard CUDA Atomic Max Reductionmedium CUDA Atomic Histogrammedium CUDA Atomic Sum Reductionmedium CUDA Shared-Memory Tiled Transposehard

Concept primerPREVIEW

Background Knowledge

The problem involves using CUDA, a parallel computing platform and programming model developed by NVIDIA, to perform a reduction operation on a large dataset. Reduction operations involve combining elements of an array using a binary operator, such as sum or product, to produce a single output value. In this case, we're tasked with implementing a two-stage reduction using shared-memory tree reduction and atomic operations.

A key concept in this problem is the use of shared memory, a small amount of memory that is shared among threads within a block. Shared memory is much faster than global memory, making it ideal for temporary storage and intermediate calculations. Tree reduction is a technique used to reduce the number of operations required to compute the sum of an array by dividing the array into smaller segments and combining them in a hierarchical manner. Atomic operations, such as cuda.atomic.add, are used to update a variable in a thread-safe manner, ensuring that only one thread can modify the variable at a time.

The problem also involves understanding CUDA's execution model, which consists of a grid of blocks, each containing a thread block of threads. Threads within a block can cooperate using shared memory and synchronization primitives, while blocks are executed independently. The gridDim.x variable represents the number of blocks in the grid, and threadIdx.x represents the index of a thread within a block. Understanding these concepts is crucial to implementing an efficient and correct solution.

Algorithm/Approach

The general approach to solving this problem involves the following steps:

Divide the input array into smaller segments, each processed by a separate block
Use shared-memory tree reduction to compute the sum of each segment
Use atomic operations to combine the partial sums from each block into a single global sum This approach takes advantage of the hierarchical nature of the CUDA execution model to minimize contention and maximize parallelism.

PRO

Continue the full explanation

You're reading the free preview. Unlock the complete walkthrough, the code editor, test runner and reference solution with Premium.

PRO

Editor locked

The code editor is locked for Pro problems. It is only available for free problems. Please upgrade to gain access to the code editor for all problems.

solution.py

Python 3.13

GPU · T4

Test results