使用“cuFFT 设备回调”答案

【问题标题】：Using "cuFFT Device Callbacks"使用“cuFFT 设备回调”
【发布时间】：2014-11-07 10:47:41
【问题描述】：

这是我的第一个问题，所以我会尽量详细。我正在努力在 CUDA 6.5 中实现降噪算法。我的代码基于这个 Matlab 实现：http://pastebin.com/HLVq48C1.
我很想使用新的 cuFFT 设备回调功能，但我坚持使用 cufftXtSetCallback。每次我的 cufftResult 都是 CUFFT_NOT_IMPLEMENTED (14)。即使 nVidia 提供的示例也以同样的方式失败...... 我的设备回调测试代码：

__device__ void noiseStampCallback(void *dataOut,
                                size_t offset,
                                cufftComplex element,
                                void *callerInfo,
                                void *sharedPointer) {
    element.x = offset;
    element.y = 2;
    ((cufftComplex*)dataOut)[offset] = element;
}
__device__ cufftCallbackStoreC noiseStampCallbackPtr = noiseStampCallback;

CUDA 部分代码：

cufftHandle forwardFFTPlan;//RtC
//find how many windows there are
int batch = targetFile->getNbrOfNoiseWindows();
size_t worksize;

cufftCreate(&forwardFFTPlan);
cufftMakePlan1d(forwardFFTPlan, WINDOW, CUFFT_R2C, batch, &worksize); //WINDOW = 2048 

//host memory, allocate
float *h_wave;
cufftComplex *h_complex_waveSpec;
unsigned int m_num_real_elems = batch*WINDOW*2;
h_wave = (float*)malloc(m_num_real_elems * sizeof(float));
h_complex_waveSpec = (cufftComplex*)malloc((m_num_real_elems/2+1)*sizeof(cufftComplex));

//init
memset(h_wave, 0, sizeof(float) * m_num_real_elems); //last window won't probably be full of file data, so fill memory with 0
memset(h_complex_waveSpec, 0, sizeof(cufftComplex) * (m_num_real_elems/2+1));
targetFile->getNoiseFile(h_wave); //fill h_wave with samples from sound file

//device memory, allocate, copy from host
float *d_wave;
cufftComplex *d_complex_waveSpec;

cudaMalloc((void**)&d_wave, m_num_real_elems * sizeof(float));
cudaMalloc((void**)&d_complex_waveSpec, (m_num_real_elems/2+1) * sizeof(cufftComplex));

cudaMemcpy(d_wave, h_wave, m_num_real_elems * sizeof(float), cudaMemcpyHostToDevice);

//prepare callback
cufftCallbackStoreC hostNoiseStampCallbackPtr;

cudaMemcpyFromSymbol(&hostNoiseStampCallbackPtr,
                          noiseStampCallbackPtr,
                          sizeof(hostNoiseStampCallbackPtr));

cufftResult status = cufftXtSetCallback(forwardFFTPlan,
                                        (void **)&hostNoiseStampCallbackPtr,
                                        CUFFT_CB_ST_COMPLEX,
                                        NULL);
//always return status 14 - CUFFT_NOT_IMPLEMENTED

//run forward plan
cufftResult result = cufftExecR2C(forwardFFTPlan, d_wave, d_complex_waveSpec);
//result seems to be okay without cufftXtSetCallback

我知道我只是 CUDA 的初学者。我的问题是：
如何正确调用 cufftXtSetCallback 或导致此错误的原因是什么？

【问题讨论】：

标签： c++ cuda callback nvidia cufft

【解决方案1】：

参考documentation：

回调 API 仅在静态链接的 cuFFT 库中可用，并且仅在 64 位 LINUX 操作系统上可用。使用此 API 需要当前的许可证。注册开发人员可在 2015 年 6 月 30 日之前免费获得评估许可证。要了解更多信息，请访问cuFFT developer page。

我认为您遇到了未实现的错误，因为您不是在 Linux 64 位平台上，或者您没有明确链接到 CUFFT 静态库。 cufft callback sample 中的 Makefile 将给出正确的链接方法。

即使您解决了该问题，您也可能会遇到CUFFT_LICENSE_ERROR，除非您已获得评估许可证之一。

请注意，有各种device limitations as well 用于链接到 cufft 静态库。应该可以构建一个静态链接的 CUFFT 应用程序，该应用程序将在 cc 2.0 和更高版本的设备上运行。

【讨论】：

你是对的。我错过的这篇笔记的唯一部分是 64 位 LINUX。好的，谢谢你的帮助！
我们知道这些限制很烦人，我们正在努力在未来的版本中删除其中的一些。它们的存在是出于技术原因，因此需要一些工程来解决它们。敬请期待……
我发现这个网页更明确地有助于设置静态链接代码以及获得许可证。该网页还表示，未来将取消许可证。 devblogs.nvidia.com/parallelforall/…
我刚碰到这个。在 Cuda 7.0 中无需许可即可使用回调（我是世界上最后一个了解这一点的人）吗？ devtalk.nvidia.com/default/topic/833368/…
@JonathanCohen：您是否知道 64 位 Linux 的静态链接限制是否计划在不久的将来随时更改？ 64 位和静态链接并不是什么大问题，但如果能在 OS X/Windows 上获得支持就好了。

【解决方案2】：

一种新的（2019 年）可能性是 cuFFT 设备扩展 (cuFFTDX)。作为 Math Library Early Access 的一部分，它们是设备 FFT 函数，可以内联到用户内核中。

cuFFTDX 公告：

https://developer.download.nvidia.com/video/gputechconf/gtc/2019/presentation/s9240-cuda-new-features-and-beyond.pdf

数学库抢先体验：

https://developer.nvidia.com/cuda-math-library-early-access-program-page

示例代码：

https://github.com/mnicely/cufft_examples

【讨论】：