CUDA Shared-Memory Block Reduction

hardPROShared Memory & TilingCUDA Basics with Numba

Problem Statement

Sum each block's chunk of a 1D array into a per-block partial sum using a shared-memory tree reduction. The kernel outputs one value per block; the host adds the partials. Verify the total equals x.sum().

Background

Each thread loads one element into shared memory. Then halve the active range each step (stride = TPB//2, TPB//4, ...), adding s[t] += s[t + stride], with a cuda.syncthreads() between steps so every add sees the previous round. Thread 0 writes s[0] to out[blockIdx.x].

Your Task

Implement block_sum_kernel and run(n=4096) returning whether the summed partials match 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 = 4096

Output

True

Reasoning

The input n = 4096 is used to create a 1D array x of length $n$ and initialize it with some values.
The block_sum_kernel function is launched, which performs a shared-memory tree reduction on the array x in parallel across multiple blocks, with each block producing a partial sum.
The partial sums from each block are then summed on the host to produce the final gpu_result.
The gpu_result is compared to the reference value calculated using x.sum() and the result of this comparison, np.allclose(gpu_result, reference), is returned as a boolean value, which in this case is True.

Constraints

Load into shared memory (use 0.0 for out-of-range threads)
Tree reduction: halve stride each step with cuda.syncthreads() between
Thread 0 writes s[0] to out[cuda.blockIdx.x]

Learn this in

Moving Data: GPU Memory with NumbaCUDA with Numba Parallel Patterns: Reductions, Atomics & TilingCUDA with Numba GPU Architecture & the Memory HierarchyGPU Basics with Triton

More from CUDA Basics with Numba

CUDA Dot Producthard CUDA Atomic Max Reductionmedium CUDA Two-Stage Reduction with Atomicshard CUDA Atomic Histogrammedium CUDA Atomic Sum Reductionmedium

Concept primerPREVIEW

Background Knowledge

The problem involves using shared-memory tree reduction to sum each block's chunk of a 1D array in parallel using CUDA. To understand this, it's essential to know how shared memory works in CUDA. Shared memory is a small amount of memory that is shared among threads within a block. It's much faster than global memory, making it ideal for temporary storage and communication between threads. In the context of this problem, each thread loads one element into shared memory, and then a reduction operation is performed to calculate the sum.

The reduction operation is a common pattern in parallel computing, where an array of elements is reduced to a single value by applying a binary operation (in this case, addition) to all elements. A tree reduction is a specific type of reduction that uses a hierarchical approach to reduce the array. This approach is particularly useful when dealing with large arrays, as it reduces the number of operations required to calculate the sum. The tree reduction works by dividing the array into smaller chunks, reducing each chunk, and then combining the results.

To implement this in CUDA, we'll use a kernel function, which is a function that runs on the GPU. The kernel function will be launched with multiple threads, each of which will perform a portion of the reduction operation. The @cuda.jit decorator is used to compile the kernel function for the GPU. We'll also use CUDA blocks, which are groups of threads that can cooperate with each other using shared memory. Each block will calculate the sum of its chunk of the array, and the results will be combined to produce the final sum.

Algorithm/Approach

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

Divide the input array into chunks, one for each block
Launch a kernel function with multiple threads, each of which loads one element into shared memory
Perform a tree reduction operation within each block to calculate the sum of its chunk
Combine the results from each block to produce the final sum This approach takes advantage of the parallel processing capabilities of the GPU to calculate the sum of the array in a efficient and scalable way.

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