CUDA Atomic Max Reduction

mediumPROReductions, Sync & AtomicsCUDA Basics with Numba

Problem Statement

Find the maximum of a 1D array using cuda.atomic.max: every thread atomically maxes its element into out[0].

Background

cuda.atomic.max(out, 0, value) keeps the larger of the current value and value, atomically. Initialize the accumulator to -inf so the first real value always wins.

Your Task

Implement atomic_max_kernel and run(n=4096) returning whether out[0] matches x.max().

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

Output

True

Reasoning

The input n = 4096 represents the size of the 1D array to find the maximum value from.
The atomic_max_kernel function is launched, where every thread atomically updates the maximum value in out[0] using cuda.atomic.max.
The initial value of out[0] is set to $-\infty$ , ensuring the first real value from the array wins and becomes the new maximum.
After all threads have executed, out[0] holds the maximum value of the array, which is then compared to the maximum value calculated using x.max() to produce the output True if they match.

Constraints

cuda.atomic.max(out, 0, x[i]) for each in-range thread
Initialize the accumulator to -inf
A max reduction is exact (no float-order error)

Learn this in

Parallel Patterns: Reductions, Atomics & TilingCUDA with Numba

More from CUDA Basics with Numba

CUDA Dot Producthard 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 using CUDA, a parallel computing platform developed by NVIDIA, to find the maximum value in a 1D array. CUDA allows for parallel processing on the GPU, which can significantly speed up certain types of computations. In this case, we're using CUDA atomic operations to update a shared variable (out) in a thread-safe manner. The cuda.atomic.max function is used to atomically update the maximum value, ensuring that only one thread can modify the value at a time.

To understand this problem, it's essential to have a basic understanding of parallel computing and GPU architecture. In parallel computing, multiple threads or processes work together to complete a task, often by dividing the work into smaller, independent chunks. In the context of CUDA, threads are organized into blocks, which are executed on the GPU. Each block can contain multiple threads, and threads within a block can share memory and synchronize with each other using barriers or atomic operations.

The concept of atomic operations is crucial in parallel computing, as it ensures that shared variables are updated consistently and correctly. In this problem, we're using cuda.atomic.max to update the maximum value in a thread-safe manner. This operation is atomic, meaning that it's executed as a single, uninterruptible unit, ensuring that only one thread can modify the value at a time.

Algorithm/Approach

The general approach to solving this problem involves using a parallel reduction algorithm, where each thread contributes to finding the maximum value in the array. The key idea is to have each thread update the shared maximum value using cuda.atomic.max, ensuring that the update is thread-safe. The algorithm can be broken down into the following high-level steps:

Initialize the shared maximum value to negative infinity (-inf)
Launch a CUDA kernel with multiple threads, where each thread updates the shared maximum value using cuda.atomic.max
After all threads have finished executing, retrieve the final maximum value from the shared variable

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