CUDA In-Place Device Update

easyFREEMemory & Data MovementCUDA Basics with Numba

Problem Statement

Scale a device array in place: x[i] = x[i] * scale, writing back into the same buffer rather than a second output array.

Background

A kernel can read and write the same device array. Allocate it once with cuda.to_device, mutate it, then copy it back — no separate output buffer.

Your Task

Implement scale_inplace_kernel and run(n=1024, scale=2.0) returning whether the mutated array equals the original times 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

n = 1024, scale = 2.0

Output

True

Reasoning

The input values n = 1024 and scale = 2.0 are used to create an array of length n and scale it in place using the scale_inplace_kernel function.
The scale_inplace_kernel function applies the transformation $x[i] = x[i] \cdot scale$ to each element in the array, effectively doubling each value since $scale = 2.0$ .
The mutated array is then compared to the original array scaled by scale using np.allclose, checking if the two arrays are element-wise equal within a small tolerance.
The comparison yields True because the in-place scaling operation correctly doubles each element in the array, resulting in an array that matches the original array scaled by $2.0$ .

Constraints

Mutate the input array directly: x[i] = x[i] * scale
No separate output device array
Capture the reference before the kernel runs

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 problem involves using CUDA, a parallel computing platform developed by NVIDIA, to scale a device array in place. This means that the kernel will read and write to the same device array, modifying its contents directly. To achieve this, we need to understand how CUDA kernels work and how to use Numba's CUDA interface to launch these kernels. A CUDA kernel is a function that runs on the GPU, and it can be launched with multiple threads, allowing for parallel execution.

In the context of this problem, we will use Numba's @cuda.jit decorator to define our kernel function. This decorator allows us to write Python functions that can be compiled to CUDA kernels. We will also use cuda.to_device to allocate memory on the GPU and copy data to it. Understanding how to manage memory and launch kernels is crucial for solving this problem.

The problem also involves comparing the mutated array with the original array scaled by a factor. This comparison will be done using np.allclose, which checks if two arrays are element-wise equal within a tolerance. This is important because floating-point operations can introduce small errors due to rounding.

Algorithm/Approach

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

Allocate memory on the GPU for the input array.
Copy the input array to the GPU.
Launch a CUDA kernel that scales the array in place.
Copy the scaled array back to the host.
Compare the scaled array with the original array scaled by the same factor.

This approach requires understanding how to manage memory, launch kernels, and perform comparisons between arrays.

Step-by-Step Strategy

To implement the solution, follow these steps:

Import necessary libraries, including numba and numpy.
Define the scale_inplace_kernel function using the @cuda.jit decorator.
Allocate memory on the GPU for the input array using cuda.to_device.
Copy the input array to the GPU.
Launch the scale_inplace_kernel function with the appropriate number of threads and blocks.
Copy the scaled array back to the host.
Compare the scaled array with the original array scaled by the same factor using np.allclose.
Return the result of the comparison.

Common Pitfalls

When implementing the solution, watch out for the following:

Incorrect memory allocation or deallocation, which can lead to memory leaks or crashes.
Insufficient synchronization between kernel launches, which can cause incorrect results.
Incorrect use of thread indices or block dimensions, which can lead to incorrect scaling or out-of-bounds access.

Time & Space Complexity

The time complexity of the solution will depend on the number of elements in the input array and the number of threads used to launch the kernel. Assuming a linear scaling of the array, the time complexity will be $O(n)$ , where $n$ is the number of elements in the array. The space complexity will be $O(n)$ as well, since we need to allocate memory for the input array on the GPU.

solution.py

Python 3.13

GPU · T4

Test results