Implement a CUDA kernel that computes the scalar minimum of a 1-D float vector input of length N:
result=i=0minN−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 must be considered, the kernel must reduce across all N elements using a cross-thread reduction. Do not assume one block covers N: the grid may span many blocks, and partial minima from different blocks must combine into result[0]. The launcher pre-fills result[0] with +FLT_MAX, the identity for min.
Input: input = [3, 1, 4, 1, 5, 9, 2, 6]
Output: 1.0
The minimum of the eight elements is 1.
Input: input = [-3, -1, -4, -5]
Output: -5.0
The smallest element is −5.
A clean pattern is a per-block tree reduction in __shared__ memory using fminf: each thread loads its element (or +∞ when i≥N) into a register, writes the partial into shared memory, then participates in log2(blockDim) pairwise fminf merges under __syncthreads.
To merge across blocks, have thread 0 of each block atomically fold its block-local min into result[0]. Use the atomicCAS trick on a unsigned int view of the float to implement atomic min.
The identity for min is +∞, available as FLT_MAX from <float.h>. The launcher pre-fills result[0] with FLT_MAX, so accumulating into it via atomic-min is safe.
Accepts: array
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that computes the scalar minimum of a 1-D float vector input of length N:
result=i=0minN−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 must be considered, the kernel must reduce across all N elements using a cross-thread reduction. Do not assume one block covers N: the grid may span many blocks, and partial minima from different blocks must combine into result[0]. The launcher pre-fills result[0] with +FLT_MAX, the identity for min.
Input: input = [3, 1, 4, 1, 5, 9, 2, 6]
Output: 1.0
The minimum of the eight elements is 1.
Input: input = [-3, -1, -4, -5]
Output: -5.0
The smallest element is −5.
A clean pattern is a per-block tree reduction in __shared__ memory using fminf: each thread loads its element (or +∞ when i≥N) into a register, writes the partial into shared memory, then participates in log2(blockDim) pairwise fminf merges under __syncthreads.
To merge across blocks, have thread 0 of each block atomically fold its block-local min into result[0]. Use the atomicCAS trick on a unsigned int view of the float to implement atomic min.
The identity for min is +∞, available as FLT_MAX from <float.h>. The launcher pre-fills result[0] with FLT_MAX, so accumulating into it via atomic-min is safe.
Accepts: array
Submit a correct solution to see your kernel performance.