2

所以我将 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]);

    }

然后我更改了我的代码，以便我完成所有内存副本（同步）并将所有内核一次发送到流，我得到了以下分析结果：

在此处输入图像描述

然后我被确认内核没有以并发方式运行。

我查看了一个链接，该链接详细解释了如何通过在#include 或代码中传递“–default-stream per-thread”命令行参数或#define CUDA_API_PER_THREAD_DEFAULT_STREAM 来设置利用完全并发。这是 CUDA 7 中引入的一个功能。我在我的 MacBook Pro Retina 15' 上使用 GeForce GT750M（与上面链接中使用的机器相同）运行了上面链接中的示例代码，并且我能够获得并发内核运行。但是我无法让我的 cuFFT 内核并行运行。

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

4

1 回答 1

2

我假设您cufftSetStream()在显示的代码之前调用了适用于 each 的代码planr2c[j]，以便每个计划都与单独的流相关联。我在您发布的代码中没有看到它。如果您确实希望 cufft 内核与其他 cufft 内核重叠，则有必要启动这些内核以分离流。因此，例如，图像 0 的 cufft exec 调用必须启动到与图像 1 的 cufft exec 调用不同的流中。

为了使任何两个 CUDA 操作有可能重叠，它们必须被启动到不同的流中。

话虽如此，具有内核执行但不是并发内核的并发内存副本是我对合理大小的 FFT 的期望。

一阶近似值的 128x128 FFT 将启动约 15,000 个线程，因此如果我的线程块每个有约 500 个线程，那将是 30 个线程块，这将使 GPU 保持相当的占用，为额外的内核留下太多“空间”。（您实际上可以在分析器本身中发现内核的总块和线程。）您的 GT750m可能有 2 个 Kepler SM，每个 SM 最多 16 个块，因此最大瞬时容量为 32 个块。由于共享内存使用、寄存器使用或其他因素，特定内核的此容量数可能会减少。

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

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

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

于 2016-04-16T01:59:02.650 回答