CUDA Shared-Memory Block Reduction
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:
n = 4096
True
- The input
n = 4096is used to create a 1D arrayxof length n and initialize it with some values. - The
block_sum_kernelfunction is launched, which performs a shared-memory tree reduction on the arrayxin 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_resultis compared to thereferencevalue calculated usingx.sum()and the result of this comparison,np.allclose(gpu_result, reference), is returned as a boolean value, which in this case isTrue.
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]
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.
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.