示例#

cuFFTDx 库提供了多个线程和块级 FFT 示例，涵盖所有支持的精度和类型，以及一些突出 cuFFTDx 性能优势的特殊示例。

示例

组

示例

描述

子组

介绍示例

introduction_example

cuFFTDx API 介绍

简单 FFT 示例

线程 FFT 示例

simple_fft_thread

复数到复数线程 FFT

simple_fft_thread_fp16

复数到复数线程 FFT 半精度

块 FFT 示例

simple_fft_block

复数到复数块 FFT

simple_fft_block_r2c

实数到复数块 FFT

simple_fft_block_c2r

复数到实数块 FFT

simple_fft_block_half2

使用 __half2 作为数据类型的复数到复数块 FFT

simple_fft_block_fp16

复数到复数块 FFT 半精度

simple_fft_block_r2c_fp16

实数到复数块 FFT 半精度

simple_fft_block_c2r_fp16

复数到实数块 FFT 半精度

额外的块 FFT 示例

simple_fft_block_shared

复数到复数块 FFT 共享内存 API

simple_fft_block_std_complex

使用 cuda::std::complex 作为数据类型的复数到复数块 FFT

simple_fft_block_cub_io

使用 CUB 加载/存储数据的复数到复数块 FFT

NVRTC 示例

nvrtc_fft_thread

复数到复数线程 FFT

nvrtc_fft_block

复数到复数块 FFT

FFT 性能

block_fft_performance

C2C 块 FFT 的基准测试

block_fft_performance_many

C2C/R2C/C2R 块 FFT 的基准测试

卷积示例

convolution

简化 FFT 卷积

convolution_r2c_c2r

简化 R2C-C2R FFT 卷积

convolution_padded

带有优化和零填充的 R2C-C2R FFT 卷积

convolution_performance

使用 cuFFTDx 和 cuFFT 进行 FFT 卷积的基准测试

conv_3d/convolution_3d

cuFFTDx 融合 3D 卷积，带有预处理、滤波和后处理

conv_3d/convolution_3d_r2c

cuFFTDx 融合 3D R2C/C2R FFT 卷积

conv_3d/convolution_3d_c2r

cuFFTDx 融合 3D C2R/R2C FFT 卷积

conv_3d/convolution_3d_padded

cuFFTDx 融合 3D FFT 卷积，使用零填充

conv_3d/convolution_3d_padded_r2c

cuFFTDx 融合 3D R2C/C2R FFT 卷积，带有零填充

2D/3D FFT 高级示例

fft_2d

示例展示如何使用 cuFFTDx 执行 2D FP32 C2C FFT

fft_2d_r2c_c2r

示例展示如何使用 cuFFTDx 执行 2D FP32 R2C/C2R 卷积

fft_2d_single_kernel

在单个内核中使用 Cooperative Groups 内核启动的 2D FP32 FFT

fft_3d_box_single_block

适合单个块的小型 3D FP32 FFT，每个维度都不同

fft_3d_cube_single_block

适合单个块的小型 3D（相等维度）FP32 FFT

fft_3d

示例展示如何使用 cuFFTDx 执行 3D FP32 C2C FFT

混合精度示例

mixed_precision_fft_1d

示例展示如何使用单独的存储和计算精度

mixed_precision_fft_2d

带有基准测试和精度比较的混合精度 2D FFT

介绍示例#

  • introduction_example

文档中使用的示例，用于解释 cuFFTDx 库及其 API 的基础知识。 introduction_example 用于 cuFFTDx API 的入门指南：使用 cuFFTDx 的第一个 FFT。

简单 FFT 示例#

simple_fft_thread* 示例#

  • simple_fft_thread

  • simple_fft_thread_fp16

在上面列出的每个示例中，一维复数到复数 FFT 例程由单个 CUDA 线程执行。在这两个示例中，运行了多个线程，每个线程计算一个 FFT。输入数据在主机上生成，复制到设备缓冲区，然后将最终结果复制回主机。

simple_fft_thread_fp16 示例展示了 cuFFTDx 中对半精度 (fp16) 的支持。请注意，对于半精度，cuFFTDx 以两个 FFT 的隐式批次处理值，即每个线程处理两个 FFT。另请参阅半精度隐式批处理部分。

simple_fft_block* 示例#

  • simple_fft_block

  • simple_fft_block_r2c

  • simple_fft_block_c2r

  • simple_fft_block_half2

  • simple_fft_block_fp16

  • simple_fft_block_r2c_fp16

  • simple_fft_block_c2r_fp16

在上面列出的每个示例中，一维复数到复数、实数到复数或复数到实数 FFT 在 CUDA 块中执行。这些示例展示了如何创建完整的 FFT 描述，然后设置正确的块维度和必要的共享内存量。在内核中，分配了每个线程寄存器中所需的数组 (thread_data)，将输入数据复制到其中，执行 FFT，并将结果传输回全局内存。所有示例都使用来自 block_io.hpp 的输入/输出函数。输入数据在主机上生成，复制到设备缓冲区，然后将最终结果复制回主机。

simple_fft_block_(*)_fp16 示例展示了 cuFFTDx 中对半精度 (fp16) 的支持。请注意，在半精度中，值以两个 FFT 的隐式批次处理，即每个线程处理两个 FFT。另请参阅半精度隐式批处理部分。

simple_fft_block_half2 示例与 simple_fft_block_fp16 不同，因为它使用 __half2 类型而不是 cufftdx::complex<__half2> 用于半精度复数值，这意味着数据在类型级别上没有隐式批处理。因此，此示例使用特殊的加载函数（以及相应的存储函数），该函数将输入缓冲区中的值加载并重新排列为 cufftdx::complex<__half2> 值，从而引入隐式批处理。另请参阅半精度隐式批处理部分。

额外的 simple_fft_block(*) 示例#

  • simple_fft_block_shared

  • simple_fft_block_std_complex

  • simple_fft_block_cub_io

simple_fft_block_shared 与其他 simple_fft_block_(*) 示例不同，因为它使用了共享内存 cuFFTDx API，请参阅块执行方法部分中的方法 #3 和 #4。

simple_fft_block_std_complex 示例表明，cuda::std::complex 类型可以用作传递给 cuFFTDx 的数据的复数值类型。它可以工作，因为它与 cufftdx::complex 具有相同的大小和对齐方式。

在 simple_fft_block_cub_io 中，NVIDIA CUB 库 (NVIDIA/cub) 用于输入/输出函数，而不是来自 block_io.hpp 的函数。它需要 1.13 或更高版本的 CUB。

NVRTC 示例#

  • nvrtc_fft_thread

  • nvrtc_fft_block

NVRTC 示例展示了如何在线程和块级别上使用 cuFFTDx 以及 NVRTC 运行时编译。使用 cuFFTDx 运算符创建的 FFT 描述仅在设备代码中定义。头文件 cufftdx.hpp 也仅包含在传递给 NVRTC 的设备代码中。只要 FFT 不需要额外的工作空间，它就可以工作，请参阅创建工作空间函数部分和FFT::requires_workspace。

注意

自 0.3.0 版本起，cuFFTDx 实验性地支持使用 NVRTC 进行编译。请参阅要求和功能部分。

FFT 性能#

  • block_fft_performance

  • block_fft_performance_many

上面列出的示例报告了计算 FFT 的 cuFFTDx 设备函数的性能。用户可以轻松修改 block_fft_performance 来测试他们想要使用的特定 FFT 的性能。 block_fft_performance_many 示例运行多个不同单精度 FFT 问题的基准测试，以显示性能如何随 FFT 的大小和类型而变化。

卷积示例#

  • convolution

  • convolution_r2c_c2r

  • convolution_padded

  • convolution_performance

卷积示例执行简化的 FFT 卷积，可以使用复数到复数正向和反向 FFT (convolution)，也可以使用实数到复数和复数到实数 FFT (convolution_r2c_c2r)。最详细的示例 (convolution_padded) 以 3 种方式执行实数卷积

  • 通过将输入填充 0 到最接近的 2 的幂，并执行优化的 cuFFTDx R2C / C2R 卷积

  • 通过保持输入不变并执行非优化的 cuFFTDx R2C / C2R 卷积

  • 通过使用 3 核 cuFFT 卷积方法

并比较它们在 8 种不同 FFT 大小上的精度和性能，以指出这种优化可能最有用的地方。

FFT padded convolution performance with cuFFT and cuFFTDx on H100 80GB with maximum clocks set.

图 3 比较了在 H100 80GB 上使用最大时钟设置的 cuFFT、具有默认设置和不变输入的 cuFFTDx 以及具有零填充输入到最接近的 2 的幂且启用 real_mode::folded 优化的 cuFFTDx 执行的批处理实数到实数卷积与逐点缩放（正向 FFT、缩放、反向 FFT）。图表显示了相对于 cuFFT 的相对性能（浅蓝色）。#

convolution_performance 示例报告了 3 种选项之间的性能差异：使用 cuFFTDx 的单核路径（正向 FFT、逐点运算、单核中的反向 FFT）、使用 cuFFT 调用和用于逐点运算的自定义内核的 3 核路径、使用 cuFFT 回调 API 的 2 核路径（需要将 CUFFTDX_EXAMPLES_CUFFT_CALLBACK cmake 选项设置为 ON：-DCUFFTDX_EXAMPLES_CUFFT_CALLBACK=ON）。根据设备、精度和给定 FFT 的大小，使用 cuFFTDx 的改进范围为 45% 到高达 3 倍的加速。 cuFFTDx 和 cuFFT convolution_performance NVIDIA H100 80GB HBM3 GPU 结果之间的性能比较如图 图 4 所示。

FFT convolution performance with cuFFT and cuFFTDx on H100 80GB HBM3 with maximum clocks set.

图 4 比较了在 H100 80GB HBM3 上使用最大时钟设置的 cuFFT 和 cuFFTDx 执行的批处理复数到复数卷积与逐点缩放（正向 FFT、缩放、反向 FFT）。图表显示了相对于 cuFFT 的相对性能（浅蓝色）。#

3D 卷积示例#

  • conv_3d/convolution_3d

  • conv_3d/convolution_3d_r2c

  • conv_3d/convolution_3d_c2r

  • conv_3d/convolution_3d_padded

  • conv_3d/convolution_3d_padded_r2c

在 conv_3d 文件夹中，有几个 3D FFT 卷积示例以及所需的 I/O 函数。

在 conv_3d/convolution_3d 中，我们介绍了一种惯用的方法来执行不适合共享内存的大型 3D FFT 卷积，带有：* 逐元素预处理函数 * 逐元素滤波函数 * 逐元素后处理函数。

预处理与最外层维度的第一次执行合并，后处理与其第二次执行合并。最内层维度的 FFT 与滤波函数合并在一起。这可以将内核数量从 9 个减少到仅 5 个。由于步幅内核对于外部维度是必需的，因此需要添加输入和输出处理，以允许通过填充的共享内存暂存这些传输，从而最大程度地提高全局内存合并并最大程度地减少共享内存库冲突。

conv_3d/convolution_3d_padded 示例展示了与 3D FFT 卷积类似的方法，但将维度填充为不同的大小。这可以提高需要回退到较慢路径的非常规维度的性能，例如素数。通常，填充到最接近的 2 的幂可以提供最佳的加速，但是填充到最接近的优化大小可以提供更高的结果精度。所有支持的优化（不需要工作空间）大小都可以在此处查看。

当将使用 cuFFTDx 的 3D FFT 与使用 cuFFT+Thrust（用于预处理/后处理和滤波函数）实现的 3D FFT 进行比较时，cuFFTDx 根据大小获得了 3.9x 到 1.3x 之间的加速。 cuFFTDx 和 cuFFT 之间的比较如图 图 5 在 H100 80GB HBM3 上所示。

FFT 3D convolution performance with cuFFT+Thrust and cuFFTDx on H100 80GB HBM3 with maximum clocks set.

图 5 比较了在 H100 80GB HBM3 上使用最大时钟设置的 cuFFT+Thrust 和 cuFFTDx 以及默认 EPT 和建议的每个块 FFT 执行的 3D 复数到复数卷积与预处理/后处理和滤波。对于非优化大小，还显示了使用 cuFFTDx 和零填充输入到最接近的 2 的幂的 3D FFT 的结果。图表显示了相对于 cuFFT 的相对加速（浅蓝色）。#

其他示例展示了如何执行正向-反向 3D R2C/C2R FFT 卷积 (conv_3d/convolution_3d_r2c) 或反向-正向 3D C2R/R2C FFT 卷积 (conv_3d/convolution_3d_c2r)，其中卷积与 X 或 Z 维度 FFT 内核融合。与之前的复数到复数卷积示例的主要区别在于执行 C2R/R2C 更改的 Z 维度内核的配置。如果大小是 2 的幂，则此内核可以从 real_mode::folded 优化中受益。

注意

3D 卷积示例对其形状和选择的参数有限制。假设卷积大小为 X/Y/Z，则每个维度的子批次数（因此是剩余维度的乘积）必须能被所讨论维度的 ffts-per-block 参数整除。

对于 C2R/R2C，子批次数取决于输入/输出长度，这与 FFT 大小不同。有关更多详细信息，请参阅示例中的注释。

2D/3D FFT 高级示例#

  • fft_2d

  • fft_2d_r2c_c2r

  • fft_2d_single_kernel

  • fft_3d

  • fft_3d_box_single_block

  • fft_3d_cube_single_block

在上面列出的每个示例中，cuFFTDx 用于执行多维 FFT。此外，其中一些示例包括与 cuFFT 的性能比较。使用 cuFFTDx 进行 2D 或 3D FFT 的最终性能将取决于输入/输出函数、FFT 的确切定义（精度、大小等）以及可以融合到内核中的自定义预处理和后处理函数。

fft_2d、fft_2d_r2c_c2r 和 fft_2d_single_kernel 示例展示了如何使用 cuFFTDx 块级执行 (cufftdx::Block) 计算 2D FFT。维度足够大，数据无法放入共享内存，因此同步和数据交换必须通过全局内存完成。 fft_2d_r2c_c2r 示例类似于 convolution_r2c_c2r，因为它使用实数到复数 FFT 转换输入，然后使用复数到实数 FFT 转换回来。 fft_2d_single_kernel 尝试在单个内核中使用 Cooperative Groups 网格启动和全网格同步来执行 2D FFT。

fft_3d 示例展示了如何使用 cuFFTDx 块级执行 (cufftdx::Block) 计算 3D FFT。维度足够大，数据无法放入共享内存，因此同步和数据交换必须通过全局内存完成。 fft_3d 示例与 fft_2d 类似，因为它是未融合内核一次处理 1D FFT 的想法的演变。它显示了一般概念，但为了获得最佳性能，应考虑下面提到的其余 3D FFT 示例。

fft_3d 示例可以用作数据和计算的不同精度实验的基础，或者在单个内核中融合预处理和后处理。

在 fft_3d_box_single_block 和 fft_3d_cube_single_block 示例中，cuFFTDx 在线程级别 (cufftdx::Thread) 上用于在单个块中执行小型 3D FFT。

混合精度 FFT 示例#

  • mixed_precision_fft_1d

  • mixed_precision_fft_2d

混合示例演示了如何将 cuFFTDx 计算精度与输入和输出全局内存缓冲区的类型分离。上述两个文件中介绍的技术使用户可以最大程度地减少 I/O，从而实现显着的速度提升。结果精度的降低可能是最小的，但可能取决于输入数据和精确算法。这两个示例都包括性能基准测试和与等效全精度计算的精度比较。重要的是要注意，这并不意味着利用任何 CUDA 混合精度功能。

mixed_precision_fft_1d 示例展示了如何使用 cuFFTDx 在双精度下计算 1D FFT（正向 FFT、复数到复数），其中数据以单精度存储在全局内存中。它显示了使用寄存器和共享内存 API 的这种方法的性能。该示例还将结果的精度与 I/O 和计算都使用相同精度的典型方法进行了比较。通过使用混合精度获得的加速取决于多个因素，对于 512 个元素的 FFT，加速达到 1.7 倍。性能结果如图 图 6 所示。

FFT performance for mixed precision with cuFFT and cuFFTDx on H100 80GB HBM3 with maximum clocks set.

图 6 比较了在 H100 80GB HBM3 上使用最大时钟设置的 cuFFT 和使用寄存器 API 的 cuFFTDx 执行的批处理复数到复数 FFT 与混合精度（存储：float，计算：double）。图表显示了相对于 cuFFT 的相对性能（浅蓝色）。#

在 mixed_precision_fft_2d 示例中，采用类似的方法来使用 cuFFTDx 在双精度下计算 2D FFT（正向 FFT、复数到复数），其中数据以单精度存储在全局内存中。 2D FFT 在单个内核中实现，就像在 fft_2d_single_kernel 示例中一样。

这两个示例都使用特定的 I/O 函数，这些函数在从/向全局内存加载和存储时，将原始精度转换为目标精度（请参阅输入/输出辅助函数）。

输入/输出辅助函数#

注意

包含的 I/O 函数不能保证为每种 FFT 配置提供最佳性能。用户可能必须编写自己的函数以满足他们的需求。

  • block_io.hpp

block_io.hpp 包含示例内核中使用的所有辅助输入/输出函数。它们是根据输入/输出数据格式和值格式部分中描述的数据布局要求实现的。

  • mixed_io.hpp

此外，mixed_io.hpp 包含混合精度 cuFFTDx 使用所需的辅助输入/输出函数和结构，并在需要时执行数据类型转换。目前仅处理几种存储/计算精度组合：fp16/fp32、bf16/fp32、fp32/fp64。

  • padded_io.hpp

padded_io.hpp 包含示例 (convolution_padded) 中提出的零填充卷积所需的辅助输入/输出函数和结构。那里的实用程序根据信号长度有条件地加载数据，并相应地偏移内存。