CUDA 2D Grid Matrix Scale

mediumPROMemory & Data MovementCUDA Basics with Numba

Problem Statement

Scale a 2D matrix by a runtime scalar using a 2D launch grid: out[i, j] = a[i, j] * scale.

Background

cuda.grid(2) returns a pair (i, j) — the row and column this thread owns. Launch with 2D threadsperblock (e.g. (16, 16)) and a 2D blockspergrid covering both dimensions.

Your Task

Implement scale2d_kernel and run(M=64, N=48, scale=3.0) returning whether the result matches a * scale.

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

M = 64, N = 48, scale = 3.0

Output

True

Reasoning

The input values are $M = 64$ , $N = 48$ , and $scale = 3.0$ , representing the dimensions of the 2D matrix and the scaling factor.
A 2D launch grid is created with a suitable number of blocks and threads per block to cover the entire matrix, allowing each thread to calculate $out[i, j] = a[i, j] \cdot scale$ .
The scale2d_kernel function is launched on the GPU, performing the scaling operation on each element of the input matrix $a$ using the given $scale$ factor: $out[i, j] = a[i, j] \cdot 3.0$ .
The resulting scaled matrix is compared to the reference solution using np.allclose, checking if the GPU result matches the expected result within a tolerance, yielding the output True if the results match.

Constraints

i, j = cuda.grid(2)
Guard with if i < a.shape[0] and j < a.shape[1]
Launch with a 2D block and 2D grid

Learn this in

Thinking in Threads: The CUDA Model with NumbaCUDA with Numba Moving Data: GPU Memory with NumbaCUDA with Numba

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 scale a 2D matrix by a runtime scalar. To understand this problem, it's essential to have a basic knowledge of parallel computing, GPU architecture, and CUDA programming model. In CUDA, a kernel is a function that runs on the GPU, and it's launched with a specified number of threads. These threads are organized into a grid, which is a two-dimensional array of blocks. Each block contains a specified number of threads.

In this problem, we're using a 2D launch grid, where each thread owns a specific row and column of the matrix. The cuda.grid(2) function returns a pair (i, j) representing the row and column of the thread. This allows us to access the corresponding element in the matrix. To launch a kernel with a 2D grid, we need to specify the number of threads per block and the number of blocks per grid. The number of blocks per grid should cover both dimensions of the matrix.

The problem also involves memory management, as we need to allocate memory on the GPU, copy data from the host (CPU) to the device (GPU), and copy the result back to the host. We'll use Numba, a library that provides a high-level interface to CUDA, to simplify the process of launching kernels and managing memory. The @cuda.jit decorator is used to compile a Python function into a CUDA kernel.

Algorithm/Approach

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

Allocate memory on the GPU for the input matrix, output matrix, and any temporary data structures.
Copy the input matrix from the host to the device.
Launch the scale2d_kernel with a 2D grid, where each thread scales the corresponding element in the matrix.
Copy the result from the device back to the host.
Compare the result with the reference solution using np.allclose.

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