Implement a CUDA kernel that returns the index of the maximum element of a 1-D float vector input of length N:
result=argi∈[0,N)maxinput[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 maximum value, the lowest index wins (this matches the default PyTorch torch.argmax 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: 5
The maximum value is 9 at index 5.
Input: input = [-3, -1, -4, -5, -2]
Output: 1
The largest value is −1 at index 1.
Carry a (value,index) pair through the reduction rather than just the value. The combine rule is: keep the pair with the larger 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].
Accepts: array
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that returns the index of the maximum element of a 1-D float vector input of length N:
result=argi∈[0,N)maxinput[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 maximum value, the lowest index wins (this matches the default PyTorch torch.argmax 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: 5
The maximum value is 9 at index 5.
Input: input = [-3, -1, -4, -5, -2]
Output: 1
The largest value is −1 at index 1.
Carry a (value,index) pair through the reduction rather than just the value. The combine rule is: keep the pair with the larger 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].
Accepts: array
Submit a correct solution to see your kernel performance.
