CUDA NPP中值滤波函数调用内存越界问题排查求助

调用CUDA NPP的nppiFilterMedian_32f_C1R函数实现中值滤波时，compute-sanitizer检测到大量无效全局内存读取的越界错误。本人怀疑问题出在cudaMallocPitch与cudaMemcpy2D的内存尺寸设置上，现求助熟悉NPP函数调用的专业人士帮忙定位错误原因。

实现代码 #

#define gpuErrchk(ans) gpuAssert((ans), __FILE__, __LINE__)
inline int gpuAssert(cudaError_t code, const char *file, int line)
{
    if (code != cudaSuccess) 
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        return 1;
    }
    return 0;
}

int gpuMedfilt2(const float* pSrc, float* pDst, int M, int N, int winSize)
{

    NppStatus status;
    Npp32f* d_in, *d_out;
    Npp32s nSrcStep = N * sizeof(float), nDstStep = N * sizeof(float);
    NppiSize oSizeROI = {N, M};
    NppiSize oMaskSize = {winSize, winSize};
    NppiPoint oAnchor = {oMaskSize.width / 2, oMaskSize.height / 2};
    Npp8u* pBuffer;
    Npp32u pBufferSize;
    size_t d_in_pitch, d_out_pitch;


    if (gpuErrchk(cudaMallocPitch((void**)&d_in, &d_in_pitch, N * sizeof(float), M)))
        return 0;

    if (gpuErrchk(cudaMallocPitch((void**)&d_out, &d_out_pitch, N * sizeof(float), M)))
    {
        cudaFree((void*)d_in);
        return 0;
    }
    if (gpuErrchk(cudaMemcpy2D((void*)d_in, d_in_pitch, (const void*)pSrc, nSrcStep, N * sizeof(float), M, cudaMemcpyHostToDevice)))
    {
        cudaFree((void*)d_in);
        cudaFree((void*)d_out);
        return 0;
    }

    if ((status = nppiFilterMedianGetBufferSize_32f_C1R(oSizeROI, oMaskSize, &pBufferSize)) != NPP_SUCCESS)
    {
        fprintf(stderr, "NPP Error: Failed to calculate buffer space for median filter operation\n");
        cudaFree((void*)d_in);
        cudaFree((void*)d_out);
        return 0;
    }

    if (gpuErrchk(cudaMalloc((void**)&pBuffer, pBufferSize)))
    {
        fprintf(stderr, "NPP Error: Failed to allocate buffer space for median filter operation\n");
        cudaFree((void*)d_in);
        cudaFree((void*)d_out);
        return 0;
    }

    if ((status = nppiFilterMedian_32f_C1R(d_in, d_in_pitch, d_out, d_out_pitch, oSizeROI, oMaskSize, oAnchor, pBuffer)) != NPP_SUCCESS)
    {
        fprintf(stderr, "NPP Error: Failed to call nppiFilterMedian_32f_C1R function\n");
        cudaFree((void*)pBuffer);
        cudaFree((void*)d_in);
        cudaFree((void*)d_out);
        return 0;
    }

    if (gpuErrchk(cudaMemcpy2D((void*)pDst, nDstStep, (const void*)d_out, d_out_pitch, sizeof(float) * N, M, cudaMemcpyDeviceToHost)))
    {
        cudaFree((void*)pBuffer);
        cudaFree((void*)d_in);
        cudaFree((void*)d_out);
        return 0;
    }

    return 1;
}

主函数调用代码 #

float* in = malloc(sizeof(float) * M * N);
float* out = malloc(sizeof(float) * M * N);
gpuMedfilt2(in, out, M, N, 5);

compute-sanitizer报错信息 #

========= 无效的__global__读取，大小为4字节
=========     位于0x2b8，函数为void FilterMedianKernelSortingNetworkShared::RunKernel5x5<float, (int)1, (int)1, (int)25>(Pixel<T1, T2> *, int, NppiSize, NppiSize, const Pixel<T1, T2> *, int, int)
=========     由线程(1,0,0)在块(0,2,0)中触发
=========     地址0x701d1fffc超出边界
=========     距离最近的内存分配地址0x701d20000（大小25,600字节）仅差4字节
=========     保存了内核启动时驱动入口点之前的主机回溯信息
=========     主机帧: [0x30b442]
=========                在/lib/x86_64-linux-gnu/libcuda.so.1中
=========     主机帧: [0x393adb]
=========                在/home/rctodd/cuda11.7/lib/libnppif.so.11中
=========     主机帧: [0x3ef278]
=========                在/home/rctodd/cuda11.7/lib/libnppif.so.11中
=========     主机帧: [0x140046]
=========                在/home/rctodd/cuda11.7/lib/libnppif.so.11中
=========     主机帧: [0x1401cb]
=========                在/home/rctodd/cuda11.7/lib/libnppif.so.11中
=========     主机帧:nppiFilterMedian_32f_C1R [0x12266f]
=========                在/home/rctodd/cuda11.7/lib/libnppif.so.11中
=========     主机帧:gpuMedfilt2 [0x1601]
=========                在/home/rctodd/code/cuda/cuMedfilt2/libgpuMedfilt2.so中
=========     主机帧:main [0x127e]
=========                在/home/rctodd/code/cuda/cuMedfilt2/./app中
=========     主机帧:../csu/libc-start.c:342:__libc_start_main [0x24083]
=========                在/lib/x86_64-linux-gnu/libc.so.6中
=========     主机帧:_start [0x110e]
=========                在/home/rctodd/code/cuda/cuMedfilt2/./app中
=========

共享库依赖（ldd输出） #

ldd libgpuMedfilt2.so 
    linux-vdso.so.1 (0x00007fff6533d000)
    libcudart.so.11.0 => /home/rctodd/cuda11.7/lib/libcudart.so.11.0 (0x00007f1bb2911000)
    libnppif.so.11 => /home/rctodd/cuda11.7/lib/libnppif.so.11 (0x00007f1badfbc000)
    libstdc++.so.6 => /lib/x86_64-linux-gnu/libstdc++.so.6 (0x00007f1baddc8000)
    libc.so.6 => /lib/x86_64-linux-gnu/libc.so.6 (0x00007f1badbd6000)
    /lib64/ld-linux-x86-64.so.2 (0x00007f1bb2bbd000)
    libdl.so.2 => /lib/x86_64-linux-gnu/libdl.so.2 (0x00007f1badbd0000)
    libpthread.so.0 => /lib/x86_64-linux-gnu/libpthread.so.0 (0x00007f1badbad000)
    librt.so.1 => /lib/x86_64-linux-gnu/librt.so.1 (0x00007f1badba1000)
    libnppc.so.11 => /home/rctodd/cuda11.7/lib/libnppc.so.11 (0x00007f1bad813000)
    libm.so.6 => /lib/x86_64-linux-gnu/libm.so.6 (0x00007f1bad6c4000)
    libgcc_s.so.1 => /home/rctodd/cuda11.7/lib/libgcc_s.so.1 (0x00007f1bad6ab000)

编译命令 #

# 编译共享库
/home/rctodd/cuda11.7/bin/nvcc -o libgpuMedfilt2.so -shared gpuMedfilt2.cu --compiler-options '-fPIC' -Xlinker -L/home/rctodd/cuda11.7/lib -Xlinker -rpath=/home/rctodd/cuda11.7/lib -lcudart -lnppif -arch=sm_50
# 编译应用程序
gcc -o app main.c -L$(pwd) -Wl,-rpath=$(pwd) -lgpuMedfilt2

内容的提问来源于stack exchange，提问作者Batmanslow23