示例

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性能基准测试

卷积示例

卷积

简化FFT卷积

convolution_r2c_c2r

简化的R2C-C2R FFT卷积

convolution_padded

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

convolution_performance

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

2D/3D FFT高级示例

fft_2d

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

fft_2d_r2c_c2r

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

fft_2d_single_kernel

使用协作组内核启动在单个内核中实现的2D FP32 FFT

fft_3d_box_single_block

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

fft_3d_cube_single_block

适用于单个块的小型3D(等维度)FP32 FFT

混合精度示例

mixed_precision_fft_1d

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

mixed_precision_fft_2d

混合精度2D FFT，包含基准测试和精度比较

入门示例

  • introduction_example

文档中用于解释cuFFTDx库及其API基础知识的示例。introduction_example被用在cuFFTDx API入门指南中：First FFT Using cuFFTDx。

简单FFT示例

simple_fft_thread* 示例

  • simple_fft_thread

  • simple_fft_thread_fp16

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

simple_fft_thread_fp16示例展示了cuFFTDx对半精度(fp16)的支持。请注意，对于半精度运算，cuFFTDx以两个FFT的隐式批次处理数值，即每个线程处理两个FFT。另请参阅Half-Precision Implicit Batching章节。

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

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

simple_fft_block_(*)_fp16示例展示了cuFFTDx对半精度(fp16)的支持。请注意，在半精度处理过程中，每个线程会处理两个FFT的隐式批次值，即每个线程处理两个FFT。另请参阅Half-Precision Implicit Batching章节。

simple_fft_block_half2示例与simple_fft_block_fp16的不同之处在于，它使用__half2类型而非cufftdx::complex<__half2> 来表示半精度复数，这意味着数据不会在类型层面隐式批处理。因此，该示例使用了一个特殊的加载函数（以及相应的存储函数），将输入缓冲区的值加载并重新排列为cufftdx::complex<__half2>值， 从而引入隐式批处理。另请参阅 Half-Precision Implicit Batching章节。

额外简单的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，具体请参阅Block Execute Method章节中的方法#3和#4。

simple_fft_block_std_complex示例展示了cuda::std::complex类型可以作为传递给cuFFTDx数据的复数值类型使用。它能正常工作是因为其大小和对齐方式与cufftdx::complex相同。

在simple_fft_block_cub_io中，使用NVIDIA CUB库(https://github.com/NVIDIA/cub)替代了block_io.hpp中的输入/输出函数。这需要CUB 1.13或更高版本。

NVRTC 示例

  • nvrtc_fft_thread

  • nvrtc_fft_block

NVRTC示例展示了如何在使用NVRTC运行时编译时，在线程和块级别上使用cuFFTDx。通过cuFFTDx运算符创建的FFT描述仅在设备代码中定义。头文件cufftdx.hpp也仅包含在传递给NVRTC的设备代码中。只要FFT不需要额外的工作空间，这种方法就有效，具体请参阅Make Workspace Function章节和FFT::requires_workspace。

注意

自0.3.0版本起，cuFFTDx实验性支持通过NVRTC进行编译。详情请参阅需求与功能章节。

FFT性能

  • block_fft_performance

  • block_fft_performance_many

上述示例展示了cuFFTDx设备函数计算FFT的性能表现。用户可以轻松修改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.

图1 展示了在H100 80GB显卡上（时钟频率设为最大值）使用cuFFT、默认设置且输入未改变的cuFFTDx，以及将输入零填充至最接近2的幂次方并启用real_mode::folded优化的cuFFTDx进行批量实数-实数卷积（包含正向FFT、缩放、逆向FFT操作）的性能对比。图表以浅蓝色柱状图表示相对于cuFFT的相对性能。

convolution_performance示例报告了三种方案的性能差异：使用cuFFTDx的单内核路径（单个内核中执行前向FFT、逐点操作、反向FFT）、使用cuFFT调用和自定义逐点操作内核的3内核路径、使用cuFFT回调API的2内核路径（需要将CUFFTDX_EXAMPLES_CUFFT_CALLBACK cmake选项设置为ON：-DCUFFTDX_EXAMPLES_CUFFT_CALLBACK=ON）。根据设备、精度和给定FFT大小的不同，使用cuFFTDx带来的性能提升从45%到最高3倍不等。cuFFTDx与cuFFTconvolution_performance在NVIDIA H100 80GB HBM3 GPU上的性能对比结果如图2所示。

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

图2 在H100 80GB HBM3显卡（设置为最高时钟频率）上，使用cuFFT和cuFFTDx进行批处理复数到复数卷积（包含前向FFT、缩放、逆向FFT操作）的性能对比。图表显示的是相对于cuFFT（浅蓝色）的相对性能。

2D/3D FFT 高级示例

  • fft_2d

  • fft_2d_r2c_c2r

  • fft_2d_single_kernel

  • 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)计算二维FFT。由于数据维度较大，无法完全放入共享内存，因此需要通过全局内存进行同步和数据交换。 fft_2d_r2c_c2r示例与convolution_r2c_c2r类似，它先使用实数到复数FFT转换输入，然后再用复数到实数FFT转换回来。 fft_2d_single_kernel尝试在单个内核中使用协作组网格启动和网格级同步来完成二维FFT计算。

在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混合精度的任何功能。

输入/输出辅助函数

注意

内置的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）所必需的。其中的工具会根据信号长度有条件地加载数据，并相应地偏移内存。