Implement a CUDA kernel that computes the scalar max of a 1-D float vector input of length N:
result=i=0maxN−1input[i]The result is a single float32 scalar written to the first slot of the device buffer result (a one-element buffer). The kernel must perform a cross-thread reduction with the max combiner: a single thread cannot scan all of input. Because the grid may span many blocks, partial maxima from different blocks must combine into result[0].
Input: input = [3, 1, 4, 1, 5, 9, 2, 6]
Output: 9.0
The maximum element is 9 at index 5.
Input: input = [-3, -1, -4, -5]
Output: -1.0
All values are negative; the largest (least-negative) is −1.
A clean pattern is a per-block tree reduction in __shared__ memory: each thread loads its element into a register (or -FLT_MAX when i≥N), writes it into shared memory, then participates in log2(blockDim) pairwise merges with fmaxf under __syncthreads.
For the cross-block combine, write each block's partial maximum to a scratch buffer of size numBlocks, then launch a second kernel that reduces that scratch buffer to result[0].
Pre-fill result[0] (and any scratch slots that may not be written) with -FLT_MAX, since it is the identity for fmaxf. Include <float.h> for the constant.
Accepts: array
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that computes the scalar max of a 1-D float vector input of length N:
result=i=0maxN−1input[i]The result is a single float32 scalar written to the first slot of the device buffer result (a one-element buffer). The kernel must perform a cross-thread reduction with the max combiner: a single thread cannot scan all of input. Because the grid may span many blocks, partial maxima from different blocks must combine into result[0].
Input: input = [3, 1, 4, 1, 5, 9, 2, 6]
Output: 9.0
The maximum element is 9 at index 5.
Input: input = [-3, -1, -4, -5]
Output: -1.0
All values are negative; the largest (least-negative) is −1.
A clean pattern is a per-block tree reduction in __shared__ memory: each thread loads its element into a register (or -FLT_MAX when i≥N), writes it into shared memory, then participates in log2(blockDim) pairwise merges with fmaxf under __syncthreads.
For the cross-block combine, write each block's partial maximum to a scratch buffer of size numBlocks, then launch a second kernel that reduces that scratch buffer to result[0].
Pre-fill result[0] (and any scratch slots that may not be written) with -FLT_MAX, since it is the identity for fmaxf. Include <float.h> for the constant.
Accepts: array
Submit a correct solution to see your kernel performance.