cuFFT 流的并发
Concurrency of cuFFT streams
所以我正在结合使用 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，它详细解释了如何通过在你之前传递“–default-stream per-thread”命令行参数或#define CUDA_API_PER_THREAD_DEFAULT_STREAM 来设置利用完全并发#include 或在您的代码中。这是 CUDA 7 中引入的功能。我 运行 在我的 MacBook Pro Retina 15' 和 GeForce GT750M 上 运行 上面的示例代码 link （与上面 link 使用的机器相同） ), 而且我能够获得并发内核 运行s。但是我无法并行获得我的 cuFFT 内核 运行ning。
然后我发现这个 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 保持相当的占用，不会留下太多 "room" 用于额外的内核。 （您实际上可以在分析器本身中发现内核的总块数和线程数。）您的 GT750m probably has 2 Kepler SMs with a maximum of 16 blocks per SM 因此最大瞬时容量为 32 个块。由于共享内存使用、寄存器使用或其他因素，对于特定内核，此容量数可能会减少。
无论您 运行 正在使用什么 GPU 的瞬时容量（每个 SM 的最大块数 * SM 的数量）将决定内核重叠（并发）的可能性。如果您在启动单个内核时超出了该容量，那么 "fill" GPU 会在一段时间内阻止内核并发。
理论上CUFFT内核应该可以同时运行。但就像任何内核并发场景（CUFFT 或其他）一样，这些内核的资源使用率必须非常低才能实际见证并发。通常，当您的资源使用率较低时，这意味着内核的 threads/threadblocks 数量相对较少。这些内核通常不会花费很长时间来执行，这使得实际见证并发变得更加困难（因为启动延迟和其他延迟因素可能会阻碍）。见证并发内核的最简单方法是让内核具有异常低的资源需求和异常长的 运行 时间。对于 CUFFT 内核或任何其他内核，这通常不是典型场景。
复制和计算的重叠仍然是 CUFFT 流的一个有用特性。而并发的想法，如果没有对机器容量和资源约束的理解基础，本身就有些不合理。例如，如果内核并发是任意可实现的 ("I should be able to make any 2 kernels run concurrently")，而不考虑容量或资源的具体情况，那么在您同时获得两个内核 运行ning 之后，下一个合乎逻辑的步骤是转到 4 、8、16核并发。但现实是机器无法同时处理那么多的工作。一旦您在单个内核启动中暴露了足够的并行性（松散地翻译为 "enough threads"），通过额外的内核启动来暴露额外的工作并行性通常不能使机器 运行 更快，或更快地处理工作。
所以我正在结合使用 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，它详细解释了如何通过在你之前传递“–default-stream per-thread”命令行参数或#define CUDA_API_PER_THREAD_DEFAULT_STREAM 来设置利用完全并发#include 或在您的代码中。这是 CUDA 7 中引入的功能。我 运行 在我的 MacBook Pro Retina 15' 和 GeForce GT750M 上 运行 上面的示例代码 link （与上面 link 使用的机器相同） ), 而且我能够获得并发内核 运行s。但是我无法并行获得我的 cuFFT 内核 运行ning。
然后我发现这个 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 保持相当的占用，不会留下太多 "room" 用于额外的内核。 （您实际上可以在分析器本身中发现内核的总块数和线程数。）您的 GT750m probably has 2 Kepler SMs with a maximum of 16 blocks per SM 因此最大瞬时容量为 32 个块。由于共享内存使用、寄存器使用或其他因素，对于特定内核，此容量数可能会减少。
无论您 运行 正在使用什么 GPU 的瞬时容量（每个 SM 的最大块数 * SM 的数量）将决定内核重叠（并发）的可能性。如果您在启动单个内核时超出了该容量，那么 "fill" GPU 会在一段时间内阻止内核并发。
理论上CUFFT内核应该可以同时运行。但就像任何内核并发场景（CUFFT 或其他）一样，这些内核的资源使用率必须非常低才能实际见证并发。通常，当您的资源使用率较低时，这意味着内核的 threads/threadblocks 数量相对较少。这些内核通常不会花费很长时间来执行，这使得实际见证并发变得更加困难（因为启动延迟和其他延迟因素可能会阻碍）。见证并发内核的最简单方法是让内核具有异常低的资源需求和异常长的 运行 时间。对于 CUFFT 内核或任何其他内核，这通常不是典型场景。
复制和计算的重叠仍然是 CUFFT 流的一个有用特性。而并发的想法，如果没有对机器容量和资源约束的理解基础，本身就有些不合理。例如，如果内核并发是任意可实现的 ("I should be able to make any 2 kernels run concurrently")，而不考虑容量或资源的具体情况，那么在您同时获得两个内核 运行ning 之后，下一个合乎逻辑的步骤是转到 4 、8、16核并发。但现实是机器无法同时处理那么多的工作。一旦您在单个内核启动中暴露了足够的并行性（松散地翻译为 "enough threads"），通过额外的内核启动来暴露额外的工作并行性通常不能使机器 运行 更快，或更快地处理工作。