CUDA Bounds-Checked Increment
Problem Statement
Add one to every element of a 1D array, out = x + 1, with a length that is not a multiple of the block size.
Background
When n isn't a multiple of blockDim.x, the last block has threads whose global index runs past the end of the array. The if i < n guard is what stops those threads from writing out of bounds.
Your Task
Implement inc_kernel and run(n=1000) (deliberately non-power-of-two) returning whether the result is exact.
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:
n = 1000
True
- The input
n = 1000is used to create a 1D array of length 1000, with each element initialized to a value. - The
inc_kernelfunction is launched with this array, adding 1 to each element using the formulaout = x + 1, while using a bounds checkif i < nto prevent out-of-bounds writes. - The resulting array is compared to a reference array, where each element is the original value plus 1, using
np.allcloseto check for exactness, considering floating point precision issues with a tolerance. - The comparison yields
True, indicating that the result of the CUDA kernel is exact, which is the final output.
Constraints:
- out[i] = x[i] + 1.0
- The if i < x.size guard must protect the tail block
- n is not a multiple of the block size
Background Knowledge
The problem revolves around the CUDA Execution Model, which is a key concept in parallel computing using NVIDIA's CUDA architecture. In CUDA, the execution model refers to how threads are organized and executed on the GPU. Threads are grouped into blocks, and each block is executed on a multiprocessor. The number of threads in a block is defined by the blockDim.x variable. When the length of the array is not a multiple of the block size, the last block will have threads that exceed the array bounds, leading to potential out-of-bounds access.
In CUDA, each thread has a unique global index, which is used to access the corresponding element in the array. The global index is calculated using the threadIdx.x and blockIdx.x variables, which represent the thread's position within the block and the block's position within the grid, respectively. To prevent out-of-bounds access, a bounds check is necessary to ensure that the global index is within the valid range of the array.
The CUDA kernel is a small program that runs on the GPU, and it is responsible for performing the desired operation on the array. In this case, the kernel needs to increment each element of the array by 1. The kernel is launched using the @cuda.jit decorator, which compiles the kernel function for execution on the GPU. The run function will allocate the input array, copy it to the GPU, launch the kernel, and verify the result using np.allclose.
Algorithm/Approach
The general approach to solving this problem involves:
Continue the full explanation
You're reading the free preview. Unlock the complete walkthrough, the code editor, test runner and reference solution with Premium.
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.