CUDA Shared-Memory 1D Stencil
Problem Statement
Compute a 3-point stencil along a 1D array (length ≤ block size): for interior points out[i] = x[i-1] + x[i] + x[i+1], with the two endpoints set to 0. Use shared memory so each value is read from DRAM only once.
Background
Neighboring threads need overlapping inputs. Stage the array in shared memory and cuda.syncthreads(); then each thread reads its neighbors s[t-1] and s[t+1] straight from shared memory instead of going back to global memory.
Your Task
Implement stencil_kernel (single block) and run(n=512) comparing to a NumPy reference. 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 = 512
True
- The input
n = 512is used to allocate an array of length 512 and initialize it with some values. - The
stencil_kernelfunction is launched with a single block, where each thread computes a 3-point stencil along the 1D array: for interior pointsout[i] = x[i-1] + x[i] + x[i+1], with the two endpoints set to0, using shared memory to minimize global memory access. - The result from the GPU is compared to a NumPy reference implementation using
np.allclose(gpu_result, reference), which checks if the two arrays are element-wise equal within a tolerance. - Since the GPU result matches the reference implementation, the function
run(n=512)returnsTrue, indicating that the GPU result is correct.
Constraints:
- Stage x into shared memory, then cuda.syncthreads()
- Endpoints (t == 0 or t == n-1) -> 0.0
- Interior -> s[t-1] + s[t] + s[t+1]
Background Knowledge
The problem involves using CUDA, a parallel computing platform developed by NVIDIA, to compute a 3-point stencil along a 1D array. This requires understanding of shared memory, a type of memory in CUDA that is shared among threads within a block. Shared memory is much faster than global memory, but it has limited capacity. The problem also involves tiling, a technique used to divide a large dataset into smaller blocks that can fit into shared memory.
To solve this problem, you need to understand how threads and blocks work in CUDA. A block is a group of threads that can cooperate with each other, and each block can have its own shared memory. The thread block size (TPB) is the maximum number of threads that can be executed concurrently within a block. In this problem, the length of the 1D array is less than or equal to the block size, which means that each thread can process one element of the array.
The problem also involves understanding of synchronization in CUDA, specifically the cuda.syncthreads() function, which is used to synchronize threads within a block. This function ensures that all threads have completed their previous tasks before proceeding to the next task. In this problem, synchronization is necessary to ensure that each thread has loaded its data into shared memory before computing the stencil.
Algorithm/Approach
The general approach to solving this type of problem involves the following steps:
- Divide the 1D array into smaller blocks that can fit into shared memory
- Load each block into shared memory
- Compute the 3-point stencil for each element in the block using the data in shared memory
- Store the results in a separate array The key insight is to use shared memory to reduce the number of global memory accesses, which can significantly improve performance.
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.