CUDA Shared-Memory Block Reverse
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:
n = 256
True
- The input array
xof lengthn = 256is allocated and copied to the GPU. - The
reverse_kernelfunction is launched with a single block ofTPBthreads, where each thread copies one element ofxinto shared memory, and thencuda.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
runfunction compares the reversed arrayoutfrom the GPU with the reference arrayx[::-1]usingnp.allclose, and returnsTrueif 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
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.