因此，我将cuFFT与CUDA流功能结合使用。我的问题是，我似乎无法使cuFFT内核完全并发运行。以下是我从nvvp获得的结果。每个流在大小为128x128的128个图像上运行2D批处理FFT内核。我设置了3个流以运行3个独立的FFT批处理计划。
从图中可以看出，一些内存副本(黄色条)与一些内核计算(紫色，棕色和粉红色条)同时存在。但是内核运行根本不是并发的。如您所见，每个内核都严格遵循。以下是我用于将内存复制到设备和启动内核的代码。
for (unsigned int j = 0; j < NUM_IMAGES; j++ ) {
gpuErrchk( cudaMemcpyAsync( dev_pointers_in[j],
image_vector[j],
NX*NY*NZ*sizeof(SimPixelType),
cudaMemcpyHostToDevice,
streams_fft[j]) );
gpuErrchk( cudaMemcpyAsync( dev_pointers_out[j],
out,
NX*NY*NZ*sizeof(cufftDoubleComplex),
cudaMemcpyHostToDevice,
streams_fft[j] ) );
cufftExecD2Z( planr2c[j],
(SimPixelType*)dev_pointers_in[j],
(cufftDoubleComplex*)dev_pointers_out[j]);
}
然后，我更改了代码，以便完成所有内存副本(同步)并将所有内核一次发送到流，并获得以下分析结果:
然后，我确认内核不是以并发方式运行。
我查看了一个link，它详细说明了如何设置在通过#include或在代码中之前传递“–default-stream per-thread”命令行参数或#define CUDA_API_PER_THREAD_DEFAULT_STREAM来利用全部并发的设置。它是CUDA 7中引入的一项功能。我在上面的链接中使用GeForce GT750M(与上面的链接使用的机器相同)在MacBook Pro Retina 15'上运行了示例代码，并且能够并行运行内核。但是我无法使cuFFT内核并行运行。
然后我发现这个link，有人说cuFFT内核将占据整个GPU，因此没有两个cuFFT内核并行运行。然后我被困住了。由于我没有找到任何有关CUFFT是否启用并发内核的正式文档。这是真的吗？有办法解决这个问题吗？
最佳答案
我假设您在显示的代码之前调用了cufftSetStream()，它适合于每个planr2c[j]，因此每个计划都与一个单独的流关联。我在您发布的代码中看不到它。如果您实际上希望cufft内核与其他cufft内核重叠，则有必要启动这些内核以分隔流。因此，例如，镜像0的cufft exec调用必须与镜像1的cufft exec调用启动到不同的流中。
为了使任何两个CUDA操作都可以重叠，必须将它们启动到不同的流中。
话虽这么说，具有内核执行功能的并发内存副本(而不是并发内核)与我期望的合理大小的FFT有关。
一阶近似的128x128 FFT将旋转约15,000个线程，因此，如果我的线程块每个都是约500个线程，那将是30个线程块，这将使GPU保持相当的占用率，而为其他内核留出了很多“空间”。 (实际上，您可以在探查器本身中发现内核的总块和线程。)您的GT750m probably has 2 Kepler SMs与a maximum of 16 blocks per SM在一起，因此最大瞬时容量为32个块。由于共享内存的使用，寄存器的使用或其他因素，可以针对特定内核减少此容量数。
您正在运行的任何GPU的瞬时容量(每SM最大块数* SM数量)将确定内核重叠(并发)的可能性。如果通过单个内核启动就超过了该容量，则将“填满” GPU，从而在一段时间内阻止内核并发。
从理论上讲，CUFFT内核可以同时运行。但是，就像任何内核并发方案，CUFFT或其他情况一样，这些内核的资源使用量必须非常低才能实际见证并发性。通常，当您的资源使用率较低时，这意味着线程/线程块数量相对较少的内核。这些内核通常不需要很长时间就能执行，这使得实际见证并发变得更加困难(因为启动延迟和其他延迟因素可能会阻碍)。见证并发内核的最简单方法是使内核具有异常低的资源需求以及异常长的运行时间。对于CUFFT内核或任何其他内核，这通常不是典型情况。
复制和计算的重叠仍然是带有CUFFT的流的有用功能。而并发思想，如果没有对机器容量和资源限制的理解，就其本身而言是不合理的。例如，如果内核并发是可以实现的(“我应该能够使任何两个内核同时运行”)，而不考虑容量或资源的具体情况，那么在使两个内核同时运行之后，下一个逻辑步骤是同时进入4、8、16个内核。但是现实是机器无法同时处理大量工作。一旦在单个内核启动中公开了足够的并行度(松散地翻译为“足够的线程”)，通过其他内核启动来公开其他工作并行度通常就无法使机器运行得更快，也不能更快地处理工作。