CUDA Shared-Memory Block Reverse

mediumFREEShared Memory & TilingCUDA Basics with Numba

Problem Statement

Reverse a 1D array (length ≤ block size) using shared memory: out[i] = x[n - 1 - i]. Stage the whole array in shared memory, synchronize, then read it back reversed.

Background

cuda.shared.array(SIZE, dtype) allocates fast on-chip memory shared by all threads in a block. The pattern is: every thread copies one element into shared memory, cuda.syncthreads() so all writes are visible, then each thread reads the mirrored position.

Your Task

Implement reverse_kernel (single block of TPB threads) and run(n=256) comparing to x[::-1]. Assume n ≤ TPB.

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 = 256

Output

True

Reasoning

The input array x of length n = 256 is allocated and copied to the GPU.
The reverse_kernel function is launched with a single block of TPB threads, where each thread copies one element of x into shared memory, and then cuda.syncthreads() is called to ensure all writes are visible.
After synchronization, each thread reads the mirrored position from shared memory, effectively reversing the array, and stores it in the output array out.
The run function compares the reversed array out from the GPU with the reference array x[::-1] using np.allclose, and returns True if they are equal within a tolerance, which is the case for the given input.

Constraints

s = cuda.shared.array(TPB, dtype=float32); TPB is a module-level constant
Load s[t] = x[t], then cuda.syncthreads(), then out[t] = s[n - 1 - t]
Single block; guard threads with if t < n

Learn this in

Moving Data: GPU Memory with NumbaCUDA 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 primer

Background Knowledge

The problem involves using shared memory in CUDA, a type of memory that is shared among all threads in a block. This allows for faster data access and manipulation compared to global memory. In CUDA, cuda.shared.array(SIZE, dtype) is used to allocate shared memory, and cuda.syncthreads() is used to synchronize threads, ensuring that all writes to shared memory are visible to all threads.

The concept of tiling is also relevant here, where a larger array is divided into smaller blocks that can fit into shared memory. This allows for efficient processing of large datasets by breaking them down into smaller, more manageable chunks. In this problem, we're working with a 1D array that fits entirely within a block, so we can stage the whole array in shared memory and then reverse it.

The use of CUDA kernels is also crucial in this problem. A kernel is a function that runs on the GPU, and it's where the actual computation takes place. In this case, we need to implement a kernel that reverses the array using shared memory. The kernel will be launched with a single block of threads, and each thread will be responsible for copying one element into shared memory, synchronizing with other threads, and then reading the reversed element.

Algorithm/Approach

The general approach to solving this type of problem involves the following pattern:

Allocate shared memory to store the array
Copy the array from global memory to shared memory, with each thread responsible for one element
Synchronize threads using cuda.syncthreads() to ensure all writes are visible
Read the reversed array from shared memory, with each thread reading the mirrored position
Store the reversed array in global memory

This approach takes advantage of the fast access times of shared memory and the parallel processing capabilities of the GPU.

Step-by-Step Strategy

To implement the solution, follow these steps:

Define the kernel function reverse_kernel that takes the input array and output array as arguments
Allocate shared memory within the kernel using cuda.shared.array(SIZE, dtype)
Copy the input array from global memory to shared memory, with each thread responsible for one element
Synchronize threads using cuda.syncthreads() to ensure all writes are visible
Read the reversed array from shared memory, with each thread reading the mirrored position
Store the reversed array in global memory
Define the run function that allocates the input array, copies it to the GPU, launches the kernel, and compares the result to the reference solution

Common Pitfalls

Some common pitfalls to watch out for when implementing this solution include:

Forgetting to synchronize threads using cuda.syncthreads() before reading from shared memory
Incorrectly indexing the shared memory array when reading the reversed elements
Failing to allocate sufficient shared memory to store the entire array
Not checking the bounds of the array when accessing elements

Time & Space Complexity

The time complexity of this solution is $O(n)$ , where $n$ is the length of the array, since each thread is responsible for copying one element into shared memory and reading one element from shared memory. The space complexity is also $O(n)$ , since we need to allocate shared memory to store the entire array. Note that $n$ is assumed to be less than or equal to the block size TPB.

solution.py

Python 3.13

GPU · T4

Test results