CUDA Grid-Stride Loop

mediumPROCUDA Execution ModelCUDA Basics with Numba

Problem Statement

Copy a 1D array using a grid-stride loop so the kernel is correct even when there are fewer total threads than elements: out = x.

Background

A grid-stride loop lets each thread process multiple elements, stepping by the total number of threads (gridDim.x * blockDim.x). This decouples the launch configuration from the data size — the same kernel works for any n.

Your Task

Implement copy_stride_kernel and run(n=4096) (launched with deliberately few blocks) that returns whether the copy is exact.

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 = 4096 (launched with 8 blocks x 128 threads)

Output

True

Reasoning

Input: n = 4096 is the length of a random float32 array $x$ , and the launch has 8 blocks of 128 threads, so only $8 \times 128 = 1024$ threads exist for $4096$ elements.
Thread $t$ of block $b$ has global index $g = 128b + t$ , from $0$ to $1023$ . The stride is the total thread count, $S = 1024$ .
Thread $g$ handles the indices $g, g + S, g + 2S, \dots$ while they stay below $n$ . For example thread $5$ copies elements $5, 1029, 2053, 3077$ ; the next candidate, $4101$ , is past the end, so it stops.
Every index $i$ belongs to exactly one thread, $g = i \bmod 1024$ , so all $4096$ elements are written once, $4$ per thread.
Output: True, because $\text{out}_i = x_i$ for every $i$ . Copying does no arithmetic, so values match exactly.
Notice the uneven cases in the tests: for $n = 100000 = 97 \cdot 1024 + 672$ , threads $0$ to $671$ copy $98$ elements and the rest copy $97$ ; for $n = 1000$ , threads $1000$ to $1023$ copy nothing.

Constraints

start = cuda.grid(1); stride = cuda.gridDim.x * cuda.blockDim.x
Loop i from start to n stepping by stride
Must work when total threads < n

Learn this in

Thinking in Threads: The CUDA Model 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 CUDA Execution Model is a key concept in parallel computing using NVIDIA GPUs. It defines how threads are organized and executed on the GPU. In CUDA, threads are grouped into blocks, and blocks are grouped into a grid. Each thread has a unique identifier, threadIdx.x, which can be used to identify the thread within a block. The number of threads in a block is defined by blockDim.x, and the number of blocks in a grid is defined by gridDim.x. Understanding these concepts is crucial for writing efficient CUDA kernels.

A grid-stride loop is a technique used in CUDA programming to decouple the launch configuration from the data size. This means that the same kernel can be used to process arrays of different sizes, without having to adjust the launch configuration. In a grid-stride loop, each thread processes multiple elements, stepping by the total number of threads (gridDim.x * blockDim.x). This allows the kernel to handle arrays of any size, even if there are fewer threads than elements.

The CUDA Basics collection provides a foundation for understanding CUDA programming. It covers topics such as memory management, thread synchronization, and kernel optimization. In the context of this problem, understanding how to launch a kernel, manage memory, and use a grid-stride loop are essential skills. The convolution and gradient descent concepts are not directly applicable to this problem, but understanding parallel computing and thread management are crucial.

Algorithm/Approach

The general approach to solving this problem involves using a grid-stride loop to copy a 1D array. The kernel will be launched with a deliberate few blocks, and each thread will process multiple elements using the grid-stride loop technique. The algorithm can be broken down into the following steps:

Allocate memory for the input and output arrays on the GPU
Copy the input array to the GPU
Launch the kernel with the desired launch configuration
Use a grid-stride loop to copy the input array to the output array
Copy the output array back to the host and compare it with the reference array

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