CUDA Two-Stage Reduction with Atomics
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:
n = 8192
True
- The input
n = 8192represents the size of the 1D array to be reduced to a single sum. - The
reduce_sum_kernelfunction 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 ins[0]. - Only thread 0 of each block then performs a
cuda.atomic.addoperation to add its block's partial sum to the global resultout[0], minimizing contention with only gridDim.x atomic operations. - The final output
Trueindicates that the global sumout[0]matches the reference sum calculated usingx.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
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.
Continue the full explanation
You're reading the free preview. Unlock the complete walkthrough, the code editor, test runner and reference solution with Premium.
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.