     __device__ MinAndArgMin Identity() { return {CudaType{}, -1}; }
     __device__ MinAndArgMin MapIn(CudaType in, int64_t index) { return {in, index}; }
     __device__ void Reduce(MinAndArgMin next, MinAndArgMin& accum) {
+        // Note that `next` can be the return value of `Identity()` in which case `accum` should not be updated.
         if (next.argmin != -1 && (accum.argmin == -1 || accum.min > next.min)) {
             accum = next;
         }

    }
};

CHAINERX_REGISTER_OP_CUDA(ArgMinOp, CudaArgMinOp);

}  // namespace


    __device__ int64_t MapOut(MinAndArgMin accum) { return accum.argmin; }
};

class CudaArgMinKernel : public ArgMinKernel {
protected:
    void Call(const Array& a, const Axes& axis, const Array& out) override {
        Device& device = a.device();
        device.CheckDevicesCompatible(a, out);
        CudaSetDeviceScope scope{device.index()};
