运算符#

运算符用于描述要解决的 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 元素组成的 forward 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() );

为了使 FFT 描述符完整，需要以下内容

一个且仅一个大小运算符。
一个且仅一个方向运算符，除非添加了 cufftdx::Type<fft_type::r2c>() 或 cufftdx::Type<fft_type::c2r>()。
一个且仅一个SM 运算符，除非添加了线程运算符。

大小运算符#

cufftdx::Size<unsigned int S>()

设置要计算的 FFT 运算的大小 S。

没有默认大小。

限制

S 必须大于 1。

方向运算符#

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

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

没有默认方向。

如果使用 Type<R2C> 运算符构造 FFT，则方向假定为 forward，并且不需要方向运算符。

如果使用 Type<C2R> 运算符构造 FFT，则方向假定为 inverse，并且不需要方向运算符。

限制

fft_direction::forward 需要 Type<C2C> 或 Type<R2C>。
fft_direction::inverse 需要 Type<C2C> 或 Type<C2R>。

cuFFTDx 执行未归一化的 FFT；也就是说，对输入数据集执行 forward FFT，然后对结果集执行 inverse FFT，产生的数据等于输入，并按 FFT 的大小缩放。是否按数据集大小的倒数缩放任一变换由用户自行决定。

类型运算符#

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

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

默认类型为 fft_type::c2c。

限制

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)。
Turing: 750 (sm_75)。
Ampere: 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_layout 值

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

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

可能的 real_mode 值

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

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

执行运算符#

运算符	默认值	描述
`线程`	未设置。	创建 FFT 线程执行对象。
`块`	未设置。	创建 FFT 块执行对象。请参阅块配置运算符。

执行运算符配置 FFT 运算在 GPU 上的运行方式。与描述运算符结合使用，它们形成完整的 FFT 描述符，可以在 GPU 上执行。

添加运算符以构造 FFT 描述符类型。例如，对于由两个每个包含 128 个 float 元素的 FFT 组成的 forward FFT 运算，它们在一个 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（使用描述运算符描述）。

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

限制

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

块运算符#

cufftdx::Block()

生成要在单个 CUDA 块中运行的集体 FFT 运算。一个或多个线程将协同计算集体 FFT 运算。

要计算的 FFT 数量以及用于计算每个 FFT 的线程数可以使用块配置运算符进行配置。

限制

与 Thread 运算符互斥
除非使用BlockDim 运算符，否则集体 FFT 运算只能在大小为 2D 的块内执行
- blockDim.x = size_of<Description>::value/Description::elements_per_thread.
- blockDim.y = Description::ffts_per_block.
- blockDim.z = 1.
BlockDim 运算符尚未实现。
运算符 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>()

注意

BlockDim 运算符尚未实现。以下详细信息显示了用于配置 FFT 运算符的块维度的规则，借助这些见解，即使 BlockDim 运算符不可用，cuFFTDx 用户也可以调整执行配置。

设置 CUDA 块大小为 (X, Y, Z) 以配置执行。

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

FFT 的默认块维度为

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

请参阅FFT::block_dim。