Implement a CUDA kernel that returns the index of the minimum element of a 1-D float vector input of length N:
result=argi∈[0,N)mininput[i]The result is a single int32 value written to slot 0 of the one-element device buffer result. When two or more elements share the minimum value, the lowest index wins (this matches the default PyTorch torch.argmin tie-break).
Because the answer depends on every element, the kernel must reduce across all N entries. The grid can span many blocks, so per-block partial answers must combine into the single integer at result[0].
Input: input = [3, 1, 4, 1, 5, 9, 2, 6]
Output: 1
The minimum value is 1 and first occurs at index 1.
Input: input = [-3, -1, -4, -5, -2]
Output: 3
The smallest value is −5 at index 3.
Carry a (value,index) pair through the reduction rather than just the value. The combine rule is: keep the pair with the smaller value; if values tie, keep the pair with the smaller index.
Use two __shared__ arrays per block, one for values and one for indices. After the per-block tree reduce, write the block's best pair into two global scratch arrays sized by the number of blocks.
Launch a second kernel with one block to reduce the per-block scratch arrays into the final pair, then write only the index to result[0]. Seed inactive lanes with +FLT_MAX so they never win.
Accepts: array
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that returns the index of the minimum element of a 1-D float vector input of length N:
result=argi∈[0,N)mininput[i]The result is a single int32 value written to slot 0 of the one-element device buffer result. When two or more elements share the minimum value, the lowest index wins (this matches the default PyTorch torch.argmin tie-break).
Because the answer depends on every element, the kernel must reduce across all N entries. The grid can span many blocks, so per-block partial answers must combine into the single integer at result[0].
Input: input = [3, 1, 4, 1, 5, 9, 2, 6]
Output: 1
The minimum value is 1 and first occurs at index 1.
Input: input = [-3, -1, -4, -5, -2]
Output: 3
The smallest value is −5 at index 3.
Carry a (value,index) pair through the reduction rather than just the value. The combine rule is: keep the pair with the smaller value; if values tie, keep the pair with the smaller index.
Use two __shared__ arrays per block, one for values and one for indices. After the per-block tree reduce, write the block's best pair into two global scratch arrays sized by the number of blocks.
Launch a second kernel with one block to reduce the per-block scratch arrays into the final pair, then write only the index to result[0]. Seed inactive lanes with +FLT_MAX so they never win.
Accepts: array
Submit a correct solution to see your kernel performance.