CUDA Dot Product

hardPROReductions, Sync & AtomicsCUDA Basics with Numba

Problem Statement

Compute the dot product of two 1D arrays, sum(a * b), by multiplying elementwise and reducing. Use a shared-memory tree reduction per block plus one atomic add per block into the global result.

Background

This fuses an elementwise multiply with a reduction: each thread loads a[i] * b[i] into shared memory, the block tree-reduces, and thread 0 atomically adds the block's partial into out[0].

Your Task

Implement dot_kernel and run(n=8192) returning whether out[0] matches np.dot(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 = 8192

Output

True

Reasoning

The input n = 8192 determines the size of the 1D arrays a and b, which are used to compute the dot product.
Each element of a is multiplied by the corresponding element of b, resulting in an array of products: $a_i \cdot b_i$ .
The products are then reduced using a tree reduction per block, followed by an atomic add per block into the global result, to compute the dot product: $\sum_{i=0}^{n-1} a_i \cdot b_i$ .
The computed dot product is compared to the reference result obtained using np.dot(a, b), and the function returns True if the two results are close, indicating a successful computation.

Constraints

Load s[t] = a[i] * b[i] (0.0 for out-of-range threads)
Tree-reduce within the block, then thread 0 atomic-adds the partial
Initialize the device accumulator to 0

Learn this in

Parallel Patterns: Reductions, Atomics & TilingCUDA with Numba

More from CUDA Basics with Numba

CUDA Atomic Max Reductionmedium CUDA Two-Stage Reduction with Atomicshard CUDA Atomic Histogrammedium CUDA Atomic Sum Reductionmedium CUDA Shared-Memory Tiled Transposehard

Concept primerPREVIEW

Background Knowledge

The problem involves computing the dot product of two 1D arrays using CUDA, a parallel computing platform developed by NVIDIA. To solve this problem, it's essential to understand the basics of CUDA programming, including kernel launches, thread blocks, and shared memory. In CUDA, a kernel is a function that runs on the GPU, and it's launched with a specified number of thread blocks, each containing a certain number of threads. Shared memory is a small, on-chip memory space that's shared among threads within a block, allowing for efficient communication and data exchange.

The problem also requires knowledge of reduction operations, which involve combining multiple values into a single output. In this case, the reduction operation is a tree reduction, where each thread loads a value into shared memory, and then the block reduces these values using a tree-like pattern. This approach helps to minimize the number of operations required to compute the final result. Additionally, the problem involves using atomic operations, such as atomic adds, to update the global result variable.

To understand the problem, it's also important to be familiar with NumPy and Numba, which are used to create and manipulate arrays, as well as to launch CUDA kernels. Numba is a just-in-time compiler that allows you to write CUDA kernels in Python, making it easier to develop and test CUDA applications.

Algorithm/Approach

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

Launch a CUDA kernel with multiple thread blocks, each containing a certain number of threads.
Within each block, perform an element-wise multiplication of the input arrays and store the results in shared memory.
Perform a tree reduction within each block to compute the partial dot product.
Use atomic operations to update the global result variable with the partial dot product from each block.
Compare the final result with the reference result computed using NumPy.

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