Implement a CUDA kernel that computes the scalar sum of a 1-D float vector input of length N:
result=i=0∑N−1input[i]The result is a single float32 scalar written to the first slot of the device buffer result (a one-element buffer). Because every element contributes to the same output, the kernel must accumulate across all N elements using a cross-thread reduction. Do not assume a single block covers N: the grid may span many blocks, and partial sums from different blocks must combine into result[0].
Input: input = [1, 2, 3, 4, 5]
Output: 15.0
Computed as 1+2+3+4+5=15.
Input: input = [-1, 2, -3, 4]
Output: 2.0
Computed as −1+2+(−3)+4=2.
A clean pattern is a per-block tree reduction in __shared__ memory: each thread loads its element into a register (or 0 when i≥N), writes the partial into shared memory, then participates in log2(blockDim) pairwise merges under __syncthreads.
To merge across blocks, have thread 0 of each block call atomicAdd(result, block_sum). The launcher in the starter pre-zeros result, so accumulating into it directly is safe.
A pure-atomics alternative is to skip shared memory and have every thread atomicAdd(result, input[i]), but contention is much higher; prefer one atomic per block.
Accepts: array
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that computes the scalar sum of a 1-D float vector input of length N:
result=i=0∑N−1input[i]The result is a single float32 scalar written to the first slot of the device buffer result (a one-element buffer). Because every element contributes to the same output, the kernel must accumulate across all N elements using a cross-thread reduction. Do not assume a single block covers N: the grid may span many blocks, and partial sums from different blocks must combine into result[0].
Input: input = [1, 2, 3, 4, 5]
Output: 15.0
Computed as 1+2+3+4+5=15.
Input: input = [-1, 2, -3, 4]
Output: 2.0
Computed as −1+2+(−3)+4=2.
A clean pattern is a per-block tree reduction in __shared__ memory: each thread loads its element into a register (or 0 when i≥N), writes the partial into shared memory, then participates in log2(blockDim) pairwise merges under __syncthreads.
To merge across blocks, have thread 0 of each block call atomicAdd(result, block_sum). The launcher in the starter pre-zeros result, so accumulating into it directly is safe.
A pure-atomics alternative is to skip shared memory and have every thread atomicAdd(result, input[i]), but contention is much higher; prefer one atomic per block.
Accepts: array
Submit a correct solution to see your kernel performance.