c++ - 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 中引入的一项功能。我在配备 GeForce GT750M 的 MacBook Pro Retina 15'(与上述链接使用的同一台机器)上运行上述链接中的示例代码，并且我能够获得并发内核运行。但是我无法让我的 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 个线程，因此如果我的线程 block 每个约有 500 个线程，那将是 30 个线程 block ，这将保持 GPU 相当占用，没有太多“空间”用于额外的内核。 (您实际上可以在分析器本身中发现内核的总 block 数和线程数。)您的 GT750m probably has 2 Kepler SMs与 a maximum of 16 blocks per SM所以最大瞬时容量为 32 个 block 。由于共享内存使用、寄存器使用或其他因素，对于特定内核，此容量数可能会减少。

您运行的任何 GPU 的瞬时容量(每个 SM 的最大块数 * SM 的数量)将决定内核重叠(并发)的可能性。如果您在启动单个内核时超出了该容量，那么这将“填满”GPU，从而在一段时间内阻止内核并发。

理论上CUFFT内核应该可以并发运行。但就像任何内核并发场景(CUFFT 或其他)一样，这些内核的资源使用率必须非常低才能实际见证并发。通常，当您的资源使用率较低时，这意味着内核的线程/线程 block 数量相对较少。这些内核通常不会花费很长时间来执行，这使得实际见证并发变得更加困难(因为启动延迟和其他延迟因素可能会阻碍)。见证并发内核的最简单方法是让内核具有异常低的资源需求和异常长的运行时间。这通常不是 CUFFT 内核或任何其他内核的典型场景。

复制和计算的重叠仍然是 CUFFT 流的一个有用特性。而并发的想法，如果没有对机器容量和资源约束的理解基础，本身就有些不合理。例如，如果内核并发是任意可实现的(“我应该能够让任何 2 个内核同时运行”)，而不考虑容量或资源细节，那么在你同时运行两个内核之后，下一个合乎逻辑的步骤是同时转到4、8、16个内核。但现实是机器无法同时处理那么多的工作。一旦您在单个内核启动中暴露了足够的并行性(松散地翻译为“足够的线程”)，通过额外的内核启动来暴露额外的工作并行性通常不能使机器运行得更快，或更快地处理工作。