CUDA Global Thread Index

easyFREECUDA Execution ModelCUDA Basics with Numba

Problem Statement

Fill an output array with each element's own global thread index: out[i] = i. This is the CUDA equivalent of np.arange.

Background

Compute the index the long way to prove you understand the mapping: i = blockIdx.x * blockDim.x + threadIdx.x. In Numba that's cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x.

Your Task

Implement iota_kernel and run(n=1024) returning whether out equals np.arange(n).

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

Output

True

Reasoning

The input value n = 1024 determines the size of the output array, which will have 1024 elements.
The iota_kernel function is launched with a suitable block and thread configuration to cover all n elements, using the formula i = blockIdx.x * blockDim.x + threadIdx.x to calculate the global thread index i.
Each thread then assigns its global thread index i to the corresponding element in the output array out, effectively creating an array of sequential integers from 0 to n-1.
The resulting output array out is compared to the reference array np.arange(n) using np.allclose, which returns True if the two arrays are identical, thus producing the output True.

Constraints

Build i from cuda.blockIdx.x, cuda.blockDim.x, cuda.threadIdx.x
out[i] = i (stored as float32)
Bounds-check with if i < out.size

Learn this in

Thinking in Threads: The CUDA Model with NumbaCUDA with Numba Why GPUs? The Parallel Computing MindsetGPU Basics with Triton GPU Architecture & the Memory HierarchyGPU Basics with Triton

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 CUDA Execution Model is a crucial concept in parallel computing using NVIDIA GPUs. It defines how threads are organized and executed on the GPU. In this model, threads are grouped into blocks, and each block is assigned a unique block index (blockIdx.x). Within a block, each thread has a unique thread index (threadIdx.x). The global thread index is calculated by combining the block index and thread index.

To understand this problem, it's essential to grasp the concept of thread indexing in CUDA. The global thread index i can be calculated using the formula: i = blockIdx.x * blockDim.x + threadIdx.x. This formula allows each thread to determine its unique global index, which can be used to access and manipulate data in a parallel manner. In the context of this problem, we need to use this formula to fill an output array with each element's own global thread index.

The Numba library provides a Python interface to CUDA, allowing developers to write CUDA kernels using Python syntax. The @cuda.jit decorator is used to define a CUDA kernel, which can be launched on the GPU. In this problem, we need to implement a CUDA kernel (iota_kernel) that fills an output array with each element's own global thread index. We will also need to use the cuda.blockIdx.x, cuda.blockDim.x, and cuda.threadIdx.x variables to calculate the global thread index.

Algorithm/Approach

The general approach to solving this problem involves:

Defining a CUDA kernel (iota_kernel) that calculates the global thread index using the formula i = blockIdx.x * blockDim.x + threadIdx.x.
Launching the kernel on the GPU with the desired number of threads and blocks.
Allocating and initializing the output array on the GPU.
Copying the output array from the GPU to the host (CPU) after the kernel has finished executing.
Comparing the output array with the expected result (np.arange(n)) to verify correctness.

Step-by-Step Strategy

To implement the solution, follow these steps:

Define the iota_kernel function using the @cuda.jit decorator.
Calculate the global thread index i using the formula i = blockIdx.x * blockDim.x + threadIdx.x.
Use the global thread index i to fill the output array.
Define the run function, which allocates the inputs, copies them to the GPU, launches the kernel, and returns a Python bool indicating whether the output array matches the expected result.
Use the np.allclose function to compare the output array with the expected result (np.arange(n)).

Common Pitfalls

When implementing the solution, watch out for the following common pitfalls:

Incorrectly calculating the global thread index i.
Failing to allocate and initialize the output array on the GPU.
Launching the kernel with incorrect thread and block configurations.
Not copying the output array from the GPU to the host (CPU) after kernel execution.

Time & Space Complexity

The time complexity of the solution is $O(n)$ , where $n$ is the number of elements in the output array. This is because each thread calculates its global thread index and fills the corresponding element in the output array.

The space complexity is also $O(n)$ , as we need to allocate an output array of size $n$ on the GPU. Additionally, we need to allocate memory for the input array and other variables used in the kernel, but these are typically negligible compared to the output array.

solution.py

Python 3.13

GPU · T4

Test results