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
Input
n = 1000
Output
True
Reasoning
- The input
n = 1000is the array length. The array holds random float32 values, and the expected result is for every from to . - With threads per block, covering elements needs blocks, which is threads in total.
- Thread of block owns global index . Blocks to cover indices to ; block covers to , and only to (232 threads) land inside the array.
- The remaining threads have and must write nothing; otherwise they would touch memory past the end of the output.
- When every valid element equals within a tolerance of , the result is
True. - Notice how the size of the idle tail changes with : for it is threads, for it is just .
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