Implement a CUDA kernel that computes the scalar dot product of two equal-length 1-D float vectors A and B:
dot(A,B)=i=0∑N−1A[i]⋅B[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: A = [1, 2, 3, 4], B = [5, 6, 7, 8]
Output: 70.0
Computed as 1⋅5+2⋅6+3⋅7+4⋅8=70.
Input: A = [0.5, -1, 2], B = [4, 3, -1]
Output: -3.0
Computed as 0.5⋅4+(−1)⋅3+2⋅(−1)=2−3−2=−3.
A clean pattern is a per-block tree reduction in __shared__ memory: each thread folds a strided slice of A[i]*B[i] into a register, 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, A[i]*B[i]), but contention is much higher; prefer one atomic per block.
Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that computes the scalar dot product of two equal-length 1-D float vectors A and B:
dot(A,B)=i=0∑N−1A[i]⋅B[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: A = [1, 2, 3, 4], B = [5, 6, 7, 8]
Output: 70.0
Computed as 1⋅5+2⋅6+3⋅7+4⋅8=70.
Input: A = [0.5, -1, 2], B = [4, 3, -1]
Output: -3.0
Computed as 0.5⋅4+(−1)⋅3+2⋅(−1)=2−3−2=−3.
A clean pattern is a per-block tree reduction in __shared__ memory: each thread folds a strided slice of A[i]*B[i] into a register, 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, A[i]*B[i]), but contention is much higher; prefer one atomic per block.
Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.