为啥 cuda-memcheck racecheck 报告 cufft 错误？

Posted 2023-02-16

【中文标题】为啥 cuda-memcheck racecheck 报告 cufft 错误？

【英文标题】：Why does cuda-memcheck racecheck report errors with cufft?为什么 cuda-memcheck racecheck 报告 cufft 错误？

【发布时间】：2016-10-05 11:46:17

【问题描述】：

racecheck 工具报告了我的应用程序的内存争用情况。我已将其与 CUFFT 执行功能隔离开来。

我做错了吗？如果没有，我怎样才能让racecheck 忽略这个？

这是一个最小的示例，当在 cuda-memcheck --tool racecheck 中运行时会产生一堆“危险”，例如

========= Race reported between Write access at 0x00000a30 in void spVector0128C::kernelTex<unsigned int, float, fftDirection_t=-1, unsigned int=8, unsigned int=8, LUT, ALL, WRITEBACK>(kernel_parameters_t<fft_tex_t, unsigned int, float>)
=========     and Read access at 0x00000a70 in void spVector0128C::kernelTex<unsigned int, float, fftDirection_t=-1, unsigned int=8, unsigned int=8, LUT, ALL, WRITEBACK>(kernel_parameters_t<fft_tex_t, unsigned int, float>) [4 hazards]

例子

#include <cufft.h>
#include <iostream>

#define ck(cmd) if ( cmd)  std::cerr << "error at line " << __LINE__ << std::endl;exit(1);

int main(int argc,char ** argv)

    int nfft=128;
    cufftComplex * ibuf;
    cufftComplex * obuf;
    ck( cudaMalloc((void**)&ibuf, sizeof(cufftComplex)*nfft) );
    ck( cudaMalloc((void**)&obuf, sizeof(cufftComplex)*nfft) );
    ck( cudaMemset( ibuf,0,sizeof(cufftComplex)*nfft) );

    cufftHandle fft;
    ck( cufftPlanMany(&fft,1,&nfft,
                NULL,1,nfft,
                NULL,1,nfft,
                CUFFT_C2C,1) );

    ck( cufftExecC2C(fft,ibuf,obuf,CUFFT_FORWARD) );

    ck( cudaDeviceSynchronize() );
    cufftDestroy( fft );
    ck(cudaFree(ibuf));
    ck(cudaFree(obuf));
    return 0;

【问题讨论】：

FWIW，我提交了一个针对 cuFFT 的 nVidia 错误 #1823484。也许它会被重新分配给 cuda-memcheck。

【参考方案1】：

你没有做错任何事。我不认为它可以被禁用类似于 nvprof - cudaProfilerStart/cudaProfilerStop

请注意 __syncthreads 和 BAR.SYNC 指令描述之间的细微差别：

__syncthreads - http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#synchronization-functions

等到线程块中的所有线程都到达这个点

BAR.SYNC - http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#texture-instructions

Barriers 是在 per-warp 的基础上执行的，就好像一个线程中的所有线程一样 经线处于活动状态。”

这不是完全相同的行为。 cuda-memcheck 竞赛检查可能遵循 __syncthreads 定义和 cuFFT 内核 BAR.SYNC 之一

这很可能会在下一个版本中得到修复。

【讨论】：

__syncthreads() 编译为bar.sync，所以它们的效果是一样的。区别仅在于文档中，__syncthreads() 的描述被简化以省略条件代码中的行为。

"这很可能会在下一个版本中得到修复。" - 你有任何内幕消息吗？这种文档细节级别的差异在 7 个左右的主要 CUDA 版本中一直存在（基本上自从 PTX 正式记录以来），因此仅观察过去的行为，在我看来，Nvidia 并不打算很快改变这种情况。跨度>

让我向您指出未来可能会发生变化的公开信息。请看演示文稿的最后一部分是关于“合作团体”的：on-demand.gputechconf.com/gtc/2016/presentation/…