CUDA Shared-Memory 1D Stencil

hardPROShared Memory & TilingCUDA Basics with Numba

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

Input

n = 512

Output

True

Reasoning

The input n = 512 is used to allocate an array of length 512 and initialize it with some values.
The stencil_kernel function is launched with a single block, where each thread computes a 3-point stencil along the 1D array: for interior points out[i] = x[i-1] + x[i] + x[i+1], with the two endpoints set to 0, 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) returns True, 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]

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 primerPREVIEW

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.

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