cuFFT 流的并发性

Question

所以我将 cuFFT 与 CUDA 流功能结合使用。我遇到的问题是我似乎无法让 cuFFT 内核完全并发运行。以下是我从 nvvp 得到的结果。每个流都在 128 个大小为 128x128 的图像上运行 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 中引入的一个功能。我在我的 MacBook Pro Retina 15' 上使用 GeForce GT750M（与上面链接中使用的机器相同）运行了上面链接中的示例代码，并且我能够获得并发内核运行。但我无法让我的 cuFFT 内核并行运行。

然后我发现这个link 有人说 cuFFT 内核将占据整个 GPU，因此没有两个 cuFFT 内核并行运行。然后我被卡住了。因为我没有找到任何正式的文档来说明 CUFFT 是否启用并发内核。这是真的吗？有没有办法解决这个问题？

Answer 1

我假设您在显示代码之前调用了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 的流的一个有用功能。而并发的想法，在没有了解机器容量和资源约束的基础上，本身就有些不合理。例如，如果内核并发是任意可实现的（“我应该能够让任意 2 个内核同时运行”），而不考虑容量或资源细节，那么在你让两个内核同时运行之后，下一个合乎逻辑的步骤是同时转到 4、8、16 个内核。但现实情况是，这台机器无法同时处理那么多工作。一旦您在单个内核启动中暴露了足够的并行性（大致翻译为“足够多的线程”），通过额外的内核启动暴露额外的工作并行性通常不会使机器运行得更快，或者处理工作更快。