CUDA Vector Addition Kernel

easyFREECUDA Execution ModelCUDA Basics with Numba

Problem Statement

Write a CUDA kernel that adds two 1D float arrays element-wise: c = a + b.

Background

Every CUDA thread computes its own global index. With Numba you get it from cuda.grid(1) (equivalent to blockIdx.x * blockDim.x + threadIdx.x). One thread handles one element, so the launch needs ceil(n / threads) blocks.

Your Task

Implement add_kernel and a run(n=1024) that launches it over a 1D grid and returns whether the result matches a + b.

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

Output

True

Reasoning

The input n = 1024 is the length of both vectors $a$ and $b$ (1024 random 32-bit floats each). The target is $c_i = a_i + b_i$ for every $i$ from $0$ to $1023$ .
With 128 threads per block, the grid needs $\lceil 1024 / 128 \rceil = 8$ blocks, giving exactly $8 \times 128 = 1024$ threads, one per element.
Each thread's global index is block index $\times$ block size $+$ thread index. Thread 5 of block 3 gets $i = 3 \cdot 128 + 5 = 389$ , so it alone computes $c_{389} = a_{389} + b_{389}$ .
Notice the test with $n = 1000$ : it still needs $\lceil 1000/128 \rceil = 8$ blocks, so 1024 threads run and indices $1000$ to $1023$ fall past the end of the arrays. Those 24 threads must write nothing.
The output True means all $n$ GPU sums agree with the same sums computed on the CPU, within an absolute tolerance of $10^{-5}$ . A skipped or misplaced element would mismatch and the check would fail.

Constraints

Use @cuda.jit and i = cuda.grid(1)
Guard with if i < a.size before writing
blocks = (n + threads - 1) // threads

Learn this in

Thinking in Threads: The CUDA Model with NumbaCUDA with Numba Why GPUs? The Parallel Computing MindsetGPU Basics with Triton Introduction to TritonGPU 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 CUDA Execution Model is a parallel computing platform and programming model developed by NVIDIA. It allows developers to harness the power of Graphics Processing Units (GPUs) to perform general-purpose computing tasks. In the context of this problem, we're dealing with a CUDA kernel, which is a small program that runs on the GPU. The kernel is responsible for performing a specific task, in this case, element-wise addition of two 1D float arrays.

The CUDA Execution Model is based on a hierarchical structure, consisting of blocks, threads, and grids. A grid is a collection of blocks, and each block is a collection of threads. In this problem, we're working with a 1D grid, where each thread handles one element of the input arrays. The global index of a thread is calculated using the formula blockIdx.x * blockDim.x + threadIdx.x, which is equivalent to cuda.grid(1) in Numba. This index is used to access the corresponding elements of the input arrays.

To launch a CUDA kernel, we need to specify the number of blocks and threads per block. The number of blocks is typically calculated as ceil(n / threads), where n is the length of the input arrays and threads is the number of threads per block. This ensures that each element of the input arrays is processed by one thread. The CUDA kernel is launched using the @cuda.jit decorator in Numba, which compiles the kernel function to CUDA code.

Algorithm/Approach

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

Define a CUDA kernel function that performs the element-wise addition of two 1D float arrays.
Launch the kernel over a 1D grid, specifying the number of blocks and threads per block.
Allocate memory on the GPU for the input and output arrays.
Copy the input data from the host (CPU) to the device (GPU).
Launch the kernel and perform the computation on the GPU.
Copy the result from the device back to the host.
Verify the result by comparing it with the reference solution.

Step-by-Step Strategy

To implement the solution, follow these steps:

Define the add_kernel function using the @cuda.jit decorator.
Calculate the global index of each thread using cuda.grid(1).
Access the corresponding elements of the input arrays using the global index.
Perform the element-wise addition and store the result in the output array.
Define the run function that allocates memory on the GPU, copies the input data, launches the kernel, and verifies the result.
Use np.allclose to compare the result with the reference solution and return a Python bool.

Common Pitfalls

When implementing the solution, watch out for the following common pitfalls:

Incorrect calculation of the global index or block dimensions.
Insufficient memory allocation on the GPU.
Failure to copy data between the host and device.
Incorrect kernel launch configuration.
Not verifying the result against the reference solution.

Time & Space Complexity

The expected time complexity of the solution is $O(n)$ , where $n$ is the length of the input arrays. This is because each element of the input arrays is processed by one thread, and the number of threads is proportional to the length of the input arrays. The space complexity is also $O(n)$ , as we need to allocate memory on the GPU for the input and output arrays.

solution.py

Python 3.13

GPU · T4

Test results