操作符¶

运算符用于描述要解决的FFT运算，并配置执行过程。它们分为描述运算符和执行运算符。

描述运算符¶

操作符	默认值	描述
`Size<unsigned int S>`	未设置。	要计算的FFT大小`S`。
`Direction<fft_direction>`	未设置。	FFT的方向，可以是`fft_direction::inverse`或`fft_direction::forward`。
`Type<fft_type>`	`fft_type::c2c`	输入和输出数据的类型（C2C、R2C、C2R）。
`Precision<P>`	`float`	用于计算FFT的浮点数值精度`P`：`double`、`float`或`__half`。
`SM<unsigned int CC>`	未设置。	生成FFT函数的目标CUDA架构。
`RealFFTOptions<complex_layout, real_mode>`	`complex_layout::natural` 和 `real_mode::normal`。	R2C/C2R FFT类型的复杂数据布局选择与执行模式

描述运算符定义了待解决的FFT运算。与执行运算符结合使用时，它们构成一个完整的FFT描述符，可在GPU上执行。

通过添加运算符来构建FFT描述符类型。例如，对于一个由每线程8个double元素组成的正向FFT操作：

#include <cufftdx.hpp>

using FFT = decltype( cufftdx::Size<8>() + cufftdx::Type<fft_type::c2c>()
                    + cufftdx::Direction<fft_direction::forward>()
                    + cufftdx::Precision<double>() + cufftdx::Thread() );

For an FFT descriptor to be complete, the following is required:

有且仅有一个Size Operator。
有且仅有一个方向运算符，除非添加了cufftdx::Type<fft_type::r2c>()或cufftdx::Type<fft_type::c2r>()。
一个且仅有一个SM Operator，除非添加了Thread Operator。

Size 操作符¶

cufftdx::Size<unsigned int S>()

设置用于计算FFT操作的尺寸S。

没有默认大小。

Restrictions:

S 必须大于 1。

方向操作符¶

cufftdx::Direction<cufftdx::fft_direction>()

设置FFT的方向，可以是fft_direction::inverse或fft_direction::forward。

没有默认方向。

如果FFT是使用Type运算符构建的，则默认方向为正向，不需要方向运算符。

如果FFT使用Type运算符构建，则假定方向为逆方向，不需要方向运算符。

Restrictions:

fft_direction::forward 需要 Type 或 Type。
fft_direction::inverse 需要 Type 或 Type。

cuFFTDx执行非归一化的FFT变换；也就是说，对输入数据集执行正向FFT后再对结果集执行反向FFT，将得到与输入数据相等的结果（需乘以FFT的尺寸系数）。用户可根据需要自行对任一变换结果乘以数据集尺寸的倒数来进行缩放处理。

类型操作符¶

cufftdx::Type<cufftdx::fft_type>()

设置要计算的FFT类型，可以是fft_type::c2c表示复数到复数；fft_type::r2c表示实数到复数；或者fft_type::c2r表示复数到实数。

默认类型为 fft_type::c2c。

Restrictions:

fft_type::r2c 需要 fft_direction::forward。如果未指定方向，则默认为 fft_direction::forward。
fft_type::c2r 需要 fft_direction::inverse。如果未指定方向，则默认为 fft_direction::inverse。
cuFFDx执行非归一化的快速傅里叶变换计算。

精度运算符¶

cufftdx::Precision<__half>()

cufftdx::Precision<float>()

cufftdx::Precision<double>()

设置用于计算FFT的浮点精度。这是用于输入和输出的数值类型，也是用于计算FFT的基础数值类型。

默认精度为float。

SM 操作符¶

cufftdx::SM<unsigned int CC>()

设置底层FFT函数使用的目标架构CC。支持的架构包括：

Volta架构: 700 和 720 (sm_70, sm_72)。
图灵架构: 750 (sm_75).
安培架构: 800, 860 和 870 (sm_80, sm_86, sm_87).
Ada: 890 (sm_89).
Hopper: 900 (sm_90, sm_90a).

注意

在编译支持9.0a计算能力的cuFFTDx时，请在SM运算符中使用900（另请参阅CUDA C++编程指南：功能可用性）。

警告

不能保证在不同CUDA架构的GPU上执行完全相同的FFT会产生完全一致的比特结果。

RealFFTOptions 运算符¶

cufftdx::RealFFTOptions<complex_layout, real_mode>()

第一个参数定义了实数到复数(fft_type::r2c)和复数到实数(fft_type::c2r)FFT类型的输入输出数据布局，具体说明请参阅复数输入输出的Complex Element Layouts和实数输入输出的Real Element Layouts。

可能的 complex_layout 取值：

complex_layout::natural (默认值)
complex_layout::packed
complex_layout::full

第二个参数real_mode枚举值，允许在优化和非优化的执行内核之间进行选择。有两种模式：普通（默认）和优化（称为folded）。优化执行是可选的，因为它会改变块执行中线程之间的数据分区方式。

real_mode 可能的取值：

real_mode::normal (默认)
real_mode::folded

最顶层的FFT类型共享了正确I/O所需的所有特性。请参考示例（例如simple_fft_block_r2c和simple_fft_block_c2r）直接展示正确的用法。

执行操作符¶

操作符	默认值	描述
`Thread`	未设置。	创建FFT线程执行对象。
`Block`	未设置。	创建FFT块执行对象。参见块配置操作符。

执行操作符配置FFT运算如何在GPU上运行。与描述操作符结合使用时，它们形成一个可在GPU上执行的完整FFT描述符。

通过添加运算符来构建FFT描述符类型。例如，对于一个由两个FFT组成的正向FFT操作，每个FFT包含128个float元素，在一个CUDA块中同时运行：

#include <cufftdx.hpp>

using FFT = decltype( cufftdx::Size<128>() + cufftdx::Type<fft_type::c2c>()
                    + cufftdx::Direction<fft_direction::forward>()
                    + cufftdx::Precision<float>() + cufftdx::Block()
                    + cufftdx::ElementsPerThread<8>() + cufftdx::FFTsPerBlock<2>() );

线程操作符¶

cufftdx::Thread()

设置FFT操作在单线程上下文中运行。该FFT操作将在每个线程上同时运行一个独立的FFT（使用描述运算符进行描述）。

每个线程将计算一个由Size Operator定义大小的FFT。

Restrictions:

与 Block 操作符互斥
与仅支持块级操作的运算符一起使用时编译将失败：FFTsPerBlock, ElementsPerThread, BlockDim。
使用 Precision<__half> 将 Size 限制在 \([2, 32]\) 范围内。
使用 Precision 将 Size 限制在 \([2, 32]\) 范围内。
使用 Precision 将 Size 限制在 \([2, 16]\) 范围内。

块操作符¶

cufftdx::Block()

生成一个集体FFT操作，在单个CUDA块中运行。一个或多个线程将协作计算该集体FFT操作。

要计算的FFT数量，以及用于计算每个FFT的线程数，可以通过Block Configuration Operators进行配置。

Restrictions:

与Thread操作符互斥
除非使用了BlockDim Operator，否则集体FFT操作只能在以下尺寸的2D块内执行：
- blockDim.x = size_of<Description>::value/Description::elements_per_thread.
- blockDim.y = Description::ffts_per_block.
- blockDim.z = 1.
BlockDim Operator 尚未实现。
运算符 cufftdx::Precision<__half>() 将 cufftdx::Size<U>() 限制在范围 \([2, 32768]\) 内。
运算符 cufftdx::Precision<float>() 将 cufftdx::Size<U>() 限制在范围 \([2, 32768]\) 内。
运算符 cufftdx::Precision<double>() 将 cufftdx::Size<U>() 限制在范围 \([2, 16384]\) 内。

块配置操作符¶

运算符	默认值	描述
`FFTsPerBlock<unsigned int F>`	1	每个CUDA块计算的FFT数量`F`。
`ElementsPerThread<unsigned int E>`	启发式。	每个CUDA线程的FFT值数量`E`。
`BlockDim<unsigned int X, Y, Z>`	未设置。	在块内执行自定义维度的块FFT所需。

块配置运算符允许用户调整集体FFT操作如何在单个CUDA块上运行。

注意

区块配置运算符只能与区块运算符一起使用。

警告

不能保证执行相同的FFT（大小、方向、类型、精度）但具有不同

每个线程处理的元素数量 (ElementsPerThread),
每个CUDA块计算的FFT数量 (FFTsPerBlock), 或
块维度 (BlockDim),

将产生完全一致的结果。

每块FFT运算¶

cufftdx::FFTsPerBlock<unsigned int>()

设置单个CUDA块内并行计算的FFT数量。每个FFT由独立的线程组并发计算。

默认每个块执行一次FFT。

每线程元素操作符¶

cufftdx::ElementsPerThread<unsigned int>()

设置每个线程需要计算的FFT元素数量。

默认值是通过启发式方法确定的，以优化性能为目标。

限制条件：

如果FFT::requires_workspace为false，则它必须是请求的FFT大小的除数。
如果FFT::requires_workspace为true，则它必须是小于FFT大小的2的幂次方。
对于cufftdx::Precision<float>()和cufftdx::Precision<__half>()，取值范围必须在\([2; 32]\)之间。
对于cufftdx::Precision<double>()，取值范围必须为\([2; 16]\)。

BlockDim 操作符¶

struct cufftdx::BlockDim<unsigned int X, unsigned int Y, unsigned int Z>()

设置CUDA块大小为(X, Y, Z)，用于配置执行参数。

使用此运算符，用户可以通过2D或3D CUDA块运行集体FFT操作。

默认的 BlockDim 大小：

blockDim.x = size_of<Description>::value/Description::elements_per_thread.
blockDim.y = Description::ffts_per_block.
blockDim.z = 1.

参见 FFT::block_dim。

注意

BlockDim 运算符尚未实现。