CUDA Atomic Histogram

mediumPROReductions, Sync & AtomicsCUDA Basics with Numba

Problem Statement

Build a histogram of values in [0, 1) into nbins bins using atomic adds. Each thread computes its bin index and does cuda.atomic.add(hist, b, 1).

Background

Many threads land in the same bin, so the increment must be atomic. Compute the bin as int(x[i] * nbins), clamped into [0, nbins-1], then atomically bump that counter.

Your Task

Implement hist_kernel and run(n=4096, nbins=10) returning whether the GPU histogram matches np.bincount.

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, nbins = 10

Output

True

Reasoning

The input values n = 4096 and nbins = 10 are used to allocate an array of n random values between 0 and 1, and a histogram array of size nbins.
Each thread in the hist_kernel computes its bin index as $b = \min(\max(int(x[i] \cdot nbins), 0), nbins-1)$ , where $x[i]$ is the $i^{th}$ value in the input array.
The threads then perform an atomic add operation on the corresponding bin in the histogram array using cuda.atomic.add(hist, b, 1), effectively counting the number of values that fall into each bin.
The resulting histogram is compared to a reference histogram computed using np.bincount, and the function returns True if the two histograms are close (i.e., np.allclose(gpu_result, reference)), indicating that the GPU histogram matches the expected result.

Constraints

b = int(x[i] * nbins), clamped to [0, nbins-1]
cuda.atomic.add(hist, b, 1) on an int32 histogram
Inputs are in [0, 1)

Learn this in

Parallel Patterns: Reductions, Atomics & TilingCUDA with Numba

More from CUDA Basics with Numba

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

Concept primerPREVIEW

Background Knowledge

The problem involves building a histogram of values using CUDA, a parallel computing platform developed by NVIDIA. To understand this problem, it's essential to have a basic knowledge of parallel programming and GPU architecture. In parallel programming, multiple threads or processes execute simultaneously, improving the overall performance of the system. In the context of CUDA, the GPU (Graphics Processing Unit) is used to perform computations, and the CPU (Central Processing Unit) is used to manage the execution of the program.

The problem also involves atomic operations, which are essential in parallel programming to ensure that multiple threads access shared resources safely. An atomic operation is a sequence of operations that is executed as a single, indivisible unit. In this problem, atomic adds are used to update the histogram bins. The cuda.atomic.add function is used to perform an atomic addition operation, which ensures that the update operation is thread-safe.

The problem requires a good understanding of histogram construction and binning. A histogram is a graphical representation of the distribution of a set of values. In this problem, the values are in the range [0, 1), and the goal is to build a histogram with nbins bins. The bin index is computed as int(x[i] * nbins), clamped into [0, nbins-1]. This means that each value is mapped to a bin based on its value, and the bin index is used to update the corresponding counter in the histogram.

Algorithm/Approach

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

Allocate memory for the input array and the histogram on the GPU
Launch a CUDA kernel that computes the bin index for each value in the input array
Use atomic adds to update the corresponding bin in the histogram
Copy the histogram from the GPU to the CPU and compare it with the reference histogram computed using np.bincount

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