各位大佬们，本人智商捉鸡，一个CUDA并行规约的问题想了好几天了绕不过来，求指导！！！

这是我在网上找的一个CUDA求最值的核函数，把它调通了之后可以正常获取到最大值。
现在就是一个5000行1024列的一个矩阵，通过cuda并行规约可以计算输出每一行的最值，最后文件保存输出的是5000个最值，每一行数据的最值。
但是我想要的是每一个最值所对应的索引值，即矩阵下标值，我尝试了很多方法都不行，因为并行规约的问题，下标值一直变，没办法得到最初的行序列中的最大值对应的索引值。
但现在我想获取到每一个最值对应的索引值，但就是怎么该都获取不到，各位大佬可不可以给点想法，该怎么改啊，
对于我来说，只有最大值对应的真实索引下标值才是真正有用的数据。
下面的代码可以直接复制运行。

#include<math.h>
#include<cuda_runtime_api.h> 
#include "device_launch_parameters.h"
#include<stdio.h>
#include<stdlib.h>
#include<time.h>
#include <cuComplex.h>
#include <device_functions.h>
#include <cuda_runtime.h>
//#include <cufft.h>

#define Height  5000
#define Width   1024
const int N = Height * Width;
const int threadsPerBlock = 1024;
const int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
#define SAVE_FILE_PATH "F:\\GPU-Reduce-partialMax_device_Result.txt"

__global__ void findMaxOfMatrix(double Mat_device[Height][Width], double partialMax_device[Height])
{
    //申请共享内存，存在于每个block中（GTX1060的共享内存为48KB）
    __shared__ double partialMax[threadsPerBlock];
    //确定索引 
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    int tid = threadIdx.x;
    //double MaxValue[threadsPerBlock];
    //int index;
    //传global memory数据到shared memory 
    partialMax[tid] = Mat_device[id / Width][id % Width];
    //传输同步 
    __syncthreads();
    //在共享存储器中进行规约
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)//右移一位相当于除了个2
    {
        __syncthreads();
        if (tid < stride)
        {
            if (partialMax[tid] < partialMax[tid + stride])
            {
                int temp = partialMax[tid];
                partialMax[tid] = partialMax[tid + stride];
                partialMax[tid + stride] = temp;
            }
            __syncthreads();
        }
    }
    //将当前block的计算结果写回输出数组 
    if (tid == 0)
    {

        //partialMax_device[blockIdx.x] = index;
        partialMax_device[blockIdx.x] = partialMax[0];
    }
}

int main()
{
    float time1;
    cudaEvent_t start1, stop1;
    cudaEventCreate(&start1);
    cudaEventCreate(&stop1);
    //申请host端内存及初始化,获得随机值 
    double(*Mat_host)[Width] = new double[Height][Width];
    double *partialMax_host = new double[blocksPerGrid];

    //构建一个 Height 行 Width 列的数组矩阵
    for (int i = 0; i < Height; ++i)
        for (int j = 0; j < Width; j++)
            Mat_host[i][j] = (rand() % 2020) + 1;

    //分配Device空间矩阵
    double(*Mat_device)[Width];
    double *partialMax_device;
    cudaMalloc((void**)&Mat_device, N * sizeof(double));
    cudaMalloc((void**)&partialMax_device, blocksPerGrid*sizeof(double));

    //int *index;
    //cudaMalloc((void**)&index, blocksPerGrid * sizeof(double));
    //int window_slide_num = blocksPerGrid;

    cudaEventRecord(start1, 0);
    //把数据从Host传到Device 
    cudaMemcpy(Mat_device, Mat_host, N * sizeof(double), cudaMemcpyHostToDevice);
    //调用内核函数 
    findMaxOfMatrix << <blocksPerGrid, threadsPerBlock >> >(Mat_device, partialMax_device);
    //将结果传回到主机端 
    cudaMemcpy(partialMax_host, partialMax_device, sizeof(double)*blocksPerGrid, cudaMemcpyDeviceToHost);
    //将部分和继续作对比，求最大值
    /*double max = 0.0;
    for (int i = 0; i < blocksPerGrid; i++)
    {
        if (max < partialMax_host[i])
        {
            max = partialMax_host[i];
        }
    }*/
    cudaEventRecord(stop1, 0);
    cudaEventSynchronize(stop1);
    //printf("\n\n maxNum:%5.5lf ", max);
    cudaEventElapsedTime(&time1, start1, stop1);
    printf("\n The time of  calculating is : %f ms\n", time1);

    //创建一个用于读写的空文件，如果文件名称与已存在的文件相同，则会删除已有文件的内容，文件被视为一个新的空文件
    FILE* fp;
    char file[128] = { 0 };
    sprintf(file, SAVE_FILE_PATH);
    fp = fopen(file, "w+");
    if (fp == NULL){ printf("保存数据的文件创建失败!\n"); return 0; }
    else
    {
        printf("\n正在保存数据，请稍等.....\n");
    }
    for (int i = 0; i < blocksPerGrid; i++)
    {
        fprintf(fp, "%5.5lf\n", partialMax_host[i]);//将数据格式化写入stream文件流中
    }
    fclose(fp);
    printf("\n文件保存完成！\n\n");
    double max = 0.0;
    for (int i = 0; i < blocksPerGrid; i++)
    {
        if (max < partialMax_host[i])
        {
            max = partialMax_host[i];
        }
    }
    printf("\n\n maxNum:%5.5lf ", max);
    //释放显空间 
    system("pause");
    cudaFree(Mat_device);
    cudaFree(partialMax_device);
    free(Mat_host);
    free(partialMax_host);
    return 0;
}