特征

特征为用户提供关于使用 算子 构建的 FFT 描述的信息。它们分为三类

• 描述特征,

• 执行特征，以及

• 其他特征.

---

描述特征

特征	默认值	描述
size_of<Description>::value	无。	要计算的 FFT 的大小。
type_of<Description>::value	fft_type::c2c	FFT 操作的类型，可以是 fft_type::c2c、 fft_type::r2c 或 fft_type::c2r。
direction_of<Description>::value	请参阅 方向特征。	FFT 操作的方向，可以是 fft_direction::inverse 或 fft_direction::forward。
precision_of<Description>::type	float	用于计算 FFT 的底层浮点值的类型： double、 float 或 __half。
is_fft<Description>::value	无。	如果 Description 是一个 FFT 描述，使用 描述算子 构成，则为 true。
is_fft_execution<Description>::value	无。	如果 Description 是一个 FFT 描述，使用 执行算子 配置，则为 true。
is_complete_fft<Description>::value	无。	如果 Description 是一个有效的 FFT 描述，使用 描述算子 构成，则为 true。
is_complete_fft_execution<Description>::value	无。	如果 is_complete_fft<Description>::value 为 true 且 is_fft_execution<Description>::value 为 true，则为 true。
real_fft_layout_of<Description>::value	complex_layout::natural	FFT 的输入和输出数据布局。请参阅 复杂元素布局。
real_fft_mode_of<Description>::value	real_mode::normal	real_mode::normal 用于非优化，real_mode::folded 用于优化内核执行。请参阅 值格式。

描述特征可以使用提供的辅助函数从 FFT 描述符中检索。例如

#include <iostream>
#include <cufftdx.hpp>

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

if(cufftdx::is_complete<FFT>::value)
  std::cout << "Size of the FFT operation: " << cufftdx::size_of<FFT>::value << std::endl;

大小特征

cufftdx::size_of<FFT>::value
cufftdx::size_of_v<FFT>

要计算的 FFT 的大小，由 大小算子 设置。

没有默认大小。如果描述符不是使用 大小算子 创建的，编译将失败并显示错误消息。

类型特征

cufftdx::type_of<FFT>::value
cufftdx::type_of_v<FFT>

FFT 操作的类型，由 类型算子 设置。

默认类型是复数到复数， fft_type::c2c。

方向特征

cufftdx::direction_of<FFT>::value
cufftdx::direction_of_v<FFT>

FFT 操作的方向，由 方向算子 设置。

默认方向

• 如果 FFT 类型是 fft_type::r2c，默认方向是 fft_direction::forward。

• 如果 FFT 类型是 fft_type::c2r，默认方向是 fft_direction::inverse。

• 对于任何其他类型，没有默认方向。如果描述符不是使用 方向算子 创建的，编译将失败并显示错误消息。

精度特征

cufftdx::precision_of<FFT>::type
cufftdx::precision_of_t<FFT>

FFT 操作的浮点精度，由 精度算子 设置。

默认精度是 float。

是 FFT 吗？特征

cufftdx::is_fft<FFT>::value
cufftdx::is_fft_v<FFT>

如果描述符是一个 FFT 描述，使用 描述算子 构成，则该特征为 true。

没有默认值。描述符要么是 FFT 描述，要么不是。

是 FFT 执行吗？特征

cufftdx::is_fft_execution<FFT>::value
cufftdx::is_fft_execution_v<FFT>

如果描述符是一个 FFT 描述，使用 描述算子 和一个 执行算子 构成，则该特征为 true。

没有默认值。描述符要么是 FFT 描述，包括执行算子，要么不是。

是 FFT 完成？特征

cufftdx::is_complete_fft<FFT>::value
cufftdx::is_complete_fft_v<FFT>

如果描述符是一个完整的 FFT 描述，使用 描述算子 构成，则该特征为 true。

注意

在此上下文中，“完整”意味着描述符已使用所有必要的 描述算子 构成，并且只缺少一个 执行算子 才能运行。

为了使 FFT 描述符完整，需要满足以下条件：

• 一个且仅一个 大小算子。

• 一个且仅一个 方向算子，除非添加了 cufftdx::Type<fft_type::r2c>() 或 cufftdx::Type<fft_type::c2r>()。

• 一个且仅一个 SM 算子，除非添加了 线程算子。

没有默认值。描述符要么是 FFT 完成描述，要么不是。

是 FFT 完成执行？特征

cufftdx::is_complete_fft_execution<FFT>::value
cufftdx::is_complete_fft_execution_v<FFT>

如果 cufftdx::is_fft_execution 和 cufftdx::is_complete_fft 都为 true，则该特征为 true。

注意

如果描述符 FFT 的 cufftdx::is_complete_fft_execution 特征为 true，那么我们可以使用 执行方法 来计算 FFT。

没有默认值。

实数 FFT 布局特征

用于实数到复数 (fft_type::r2c) 和复数到实数 (fft_type::c2r) FFT 类型的输入和输出数据布局，使用 RealFFTOptions 算子 设置。

默认值是 complex_layout::natural。

实数 FFT 模式特征

该值表示 FFT 内核将要执行的优化模式，使用 RealFFTOptions 算子 设置。

默认值是 real_mode::normal。

---

执行特征

执行特征可以直接从已使用 执行算子 配置的 FFT 描述符中检索。可用的执行特征取决于用于构建描述符的算子；可以是 线程算子 或 块算子。

线程特征

特征	默认值	描述
Description::value_type	detail::complex<float>	用于计算 FFT 的底层数据的复数类型。
Description::input_ept	与 每线程元素特征 相同。	要为此 FFT 加载的 输入类型特征 类型的元素计数
Description::output_ept	与 每线程元素特征 相同。	在此 FFT 之后要存储的 输出类型特征 类型的元素计数
Description::input_type	无。	要加载的元素的类型。提供最佳的向量化和正确性。
Description::output_type	无。	要存储的元素的类型。提供最佳的向量化和正确性。
Description::input_length	无。	要作为一个批次加载的 输入类型特征 类型的输入的长度。
Description::output_length	无。	要作为一个批次存储的 输出类型特征 类型的输入的长度。
Description::implicit_type_batching	如果 cufftdx::precision_of<FFT>::type 是 __half，则为 2，否则为 1	来自不同 FFT 的值批处理到一个 Description::value_type 类型的元素中的数量。
Description::elements_per_thread	size_of<Description>::value	每个线程要计算的 FFT 元素数量。
Description::storage_size	Description::elements_per_thread	每个线程必须分配以计算 FFT 的 Description::value_type 元素数量。
Description::stride	始终为 1	每个线程在其 input 中持有的线程 FFT 元素之间的步幅

线程特征可以从使用 线程算子 构建的描述符中检索。

例如

#include <cufftdx.hpp>

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


// Retrieve the FFT data type
using complex_type = typename FFT::value_type;

// Retrieve the number of elements per thread
auto elements_per_thread = FFT::elements_per_thread;

值类型特征

FFT::value_type

用于 FFT 计算的底层数据的复数类型。

默认类型是 cufftdx::detail::complex<float>，如 types.hpp 头文件中所定义。

输入 EPT 特征

FFT::input_ept

单个线程为 FFT 执行提供的 输入类型特征 类型的最大元素计数。

默认值与 每线程元素特征 相同。

输出 EPT 特征

FFT::output_ept

FFT 执行后，单个线程将返回的 输出类型特征 类型的最大元素计数。

默认值与 每线程元素特征 相同。

输入类型特征

FFT::input_type

为 FFT 执行提供的元素的类型。对于 C2C 和 C2R 配置，这与 值类型特征 相同，但在 R2C 中，它直接取决于使用的 RealFFTOptions 算子 值。

默认类型与 值类型特征 相同。

输出类型特征

FFT::output_type

FFT 执行后将返回的元素的类型。对于 C2C 和 R2C 配置，这与 值类型特征 相同，但在 C2R 中，它直接取决于使用的 RealFFTOptions 算子 值。

默认类型与 值类型特征 相同。

输入长度特征

FFT::input_length

此 FFT 的单个批次输入的完整长度。对于 C2C，此值等效于 大小特征，但对于 R2C，它取决于 RealFFTOptions 算子 real_mode 值，对于 C2R，它取决于 RealFFTOptions 算子 complex_layout 值。

默认值与 大小特征 相同。

输出长度特征

FFT::output_length

此 FFT 的单个批次输出的完整长度。对于 C2C，此值等效于 大小特征，但对于 C2R，它取决于 RealFFTOptions 算子 real_mode 值，对于 R2C，它取决于 RealFFTOptions 算子 complex_layout 值。

默认值与 大小特征 相同。

隐式类型批处理特征

FFT::implicit_type_batching

来自不同 FFT 的值批处理到一个 Description::value_type 类型的元素中的数量，用于 FFT 计算。如果它高于 1，则意味着线程 FFT 对象一次计算多个 FFT。

如果 cufftdx::precision_of<FFT>::type 是 __half，则值为 2，否则为 1。

注意

请注意，在 cuFFTDx 的未来版本中，FFT::implicit_type_batching 可能会被替换和/或扩展。

每线程元素特征

FFT::elements_per_thread

每个线程将计算的 FFT 元素的逻辑数量。这可能与实际元素数量不同，因为 RealFFTOptions 算子 可能会更改输入或输出元素的数量及其逻辑布局。请参考 输入 EPT 特征 和 输出 EPT 特征 以了解输入和输出内存操作，并参考 存储大小特征 以分配寄存器空间。

默认值与 大小特征 相同。

存储大小特征

FFT::storage_size

每个线程必须分配以计算 FFT 的 Description::value_type 元素数量。

默认值与 输入 EPT 特征 相同。

步幅大小特征

FFT::stride

每个线程在其输入中持有的 FFT 元素之间的步幅。

对于线程 FFT，FFT::stride 始终为 1。

块特征

特征	默认值	描述
Description::value_type	detail::complex<float>	用于计算 FFT 的底层数据的复数类型。
Description::input_ept	无。	每个线程要加载的 输入类型特征 类型的元素计数。
Description::output_ept	无。	每个线程要存储的 输出类型特征 类型的元素计数。
Description::input_type	无。	每个线程要加载的元素的类型。提供最佳的向量化和正确性。
Description::output_type	无。	每个线程要存储的元素的类型。提供最佳的向量化和正确性。
Description::input_length	无。	要作为一个批次加载的 输入类型特征 类型的输入的长度。
Description::output_length	无。	要作为一个批次存储的 输出类型特征 类型的输入的长度。
Description::workspace_type	Description::workspace_type	FFT 计算所需的设备端工作区类型。
Description::implicit_type_batching	如果 cufftdx::precision_of<FFT>::type 是 __half，则为 2，否则为 1	来自不同 FFT 的值批处理到一个 Description::value_type 类型的元素中的数量。
Description::elements_per_thread	启发式。	每个线程要计算的 FFT 元素数量。
Description::storage_size	由 Description::elements_per_thread 确定	每个线程必须分配以计算 FFT 的 Description::value_type 元素数量。
Description::stride	由 Description::elements_per_thread 和 FFT 的大小确定	每个线程在其输入中持有的块 FFT 元素之间的步幅
Description::ffts_per_block	1	在此 FFT 操作中，一个 CUDA 块要计算的 FFT 数量。
Description::suggested_ffts_per_block	启发式。	为了达到最佳性能，建议一个 CUDA 块计算的 FFT 数量。
Description::shared_memory_size	由 Description::ffts_per_block 和 Description::elements_per_thread 确定	共享内存的大小（以字节为单位）。
Description::block_dim	请参阅 块维度特征。	用于计算 FFT 操作的 CUDA 块的 dim3。
Description::max_threads_per_block	由 Description::block_dim 确定	CUDA 块中的线程总数。
Description::requires_workspace	如果 FFT 实现需要额外的工作区，则为 true；否则为 false。	确定是否需要在全局内存中使用 cufftdx::make_workspace(cudaError_t&, cudaStream_t) 分配额外的工作区。
Description::workspace_size	如果 Description::requires_workspace 为 false，则为 0，否则 > 0。	工作区所需的全局内存大小（以字节为单位）。

块特征可以从使用 块算子 构建的描述符中检索。

例如

#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>() );

// Retrieve the FFT data type
using complex_type = typename FFT::value_type;

// Allocate managed memory for input/output
complex_type* data;
auto          size       = FFT::ffts_per_block * cufftdx::size_of<FFT>::value;
auto          size_bytes = size * sizeof(complex_type);

cudaMallocManaged(&data, size_bytes);

值类型特征

FFT::value_type

用于 FFT 计算的底层数据的复数类型。

默认类型是 cufftdx::detail::complex<float>，如 types.hpp 头文件中所定义。

输入 EPT 特征

FFT::input_ept

单个线程为寄存器 API FFT 执行提供的 输入类型特征 类型的最大元素计数。

输出 EPT 特征

FFT::output_ept

寄存器 API FFT 执行后，单个线程将返回的 输出类型特征 类型的最大元素计数。

输入类型特征

FFT::input_type

为 FFT 执行提供的元素的类型。对于 C2C 和 C2R 配置，这与 值类型特征 相同，但在 R2C 中，它直接取决于使用的 RealFFTOptions 算子 值。

输出类型特征

FFT::output_type

FFT 执行后将返回的元素的类型。对于 C2C 和 R2C 配置，这与 值类型特征 相同，但在 C2R 中，它直接取决于使用的 RealFFTOptions 算子 值。

输入长度特征

FFT::input_length

此 FFT 的单个批次输入的完整长度。对于 C2C，此值等效于 大小特征，但对于 R2C，它取决于 RealFFTOptions 算子 real_mode 值，对于 C2R，它取决于 RealFFTOptions 算子 complex_layout 值。

输出长度特征

FFT::output_length

此 FFT 的单个批次输出的完整长度。对于 C2C，此值等效于 大小特征，但对于 C2R，它取决于 RealFFTOptions 算子 real_mode 值，对于 R2C，它取决于 RealFFTOptions 算子 complex_layout 值。

工作区类型特征

FFT::workspace_type

由 FFT 的 execute(...) 函数所需的工作空间类型。用户应检查 FFT 是否需要使用 Description::requires_workspace 特征的工作空间，并使用 cufftdx::make_workspace<FFT>(cudaError_t&, cudaStream_t) 创建一个。

有关工作空间的更多详细信息，请参阅 创建工作空间函数。

警告

FFT::workspace_type 对象不跟踪底层内存的生命周期，并且仅在其转换自的工作空间对象的生命周期内有效。

警告

由 cufftdx::make_workspace<FFT>(cudaError_t&, cudaStream_t) 返回的类型可能因不同的 FFT 描述而异，并且与 FFT::workspace_type 不同。用户在创建工作空间对象时应使用 auto。

隐式类型批处理特征

FFT::implicit_type_batching

从不同 FFT 批量处理到 Description::value_type 类型的一个元素中的值的数量，用于 FFT 计算。如果它高于 1，则意味着 Block FFT 对象一次计算多个 FFT。

如果 cufftdx::precision_of<FFT>::type 是 __half，则值为 2，否则为 1。

注意

请注意，在 cuFFTDx 的未来版本中，FFT::implicit_type_batching 可能会被替换和/或扩展。

每线程元素特征

FFT::elements_per_thread

每个线程将计算的 FFT 元素的逻辑数量。这可能与物理元素的数量不同，因为 RealFFTOptions 运算符 可能会更改输入或输出元素的数量及其在线程之间的逻辑布局。有关输入和输出内存操作，请参阅 输入 EPT 特征 和 输出 EPT 特征，有关分配寄存器空间，请参阅 存储大小特征。

默认值与 大小特征 相同。

存储大小特征

FFT::storage_size

每个线程必须分配以计算 FFT 的 Description::value_type 元素数量。

默认值与 输入 EPT 特征 相同。

步幅大小特征

FFT::stride

每个线程在其 input 中持有的块 FFT 元素之间的步幅。

另请参阅 预期的输入数据格式。

示例

步幅 FFT::stride 等于 2 的 8 点 FFT 的第 0 个线程应在其 input 中具有值 0、2、4 和 6。

每块 FFT 特征

FFT::ffts_per_block

在 CUDA 块内并行计算的 FFT 数量，作为集体 FFT 操作的一部分。

默认值为 1。

建议的每块 FFT 特征

FFT::suggested_ffts_per_block

建议在 CUDA 块内并行计算的 FFT 数量，作为集体 FFT 操作的一部分，以最大化性能。

默认值是启发式的，并且取决于 FFT 的大小、每个线程的元素数量和其他参数。

共享内存大小特征

FFT::shared_memory_size

FFT 操作执行所需的共享内存大小，以字节为单位。

默认值由 每块 FFT 特征 和 每线程元素特征 确定。

块维度特征

FFT::block_dim

所需的 CUDA 块维度，其中

• x = (size_of<FFT>::value / FFT::elements_per_thread),

• y = (FFT::ffts_per_block / FFT::implicit_type_batching)，以及

• z = 1.

每块最大线程数特征

FFT::max_threads_per_block

CUDA 块中 FFT 的最大线程数。

默认值由 每块 FFT 特征 和 每线程元素特征 确定。

需要工作空间特征

FFT::requires_workspace

布尔值。如果为 true，则必须创建工作空间并将其传递给 FFT::execute(...) 方法（请参阅 块执行方法）。否则，无需创建和传递工作空间。可以使用 cufftdx::make_workspace<FFT>(cudaError_t&, cudaStream_t) 函数创建工作空间。为不需要工作空间的 FFT 创建的工作空间将为空，并且不会分配任何全局内存。

工作空间大小特征

FFT::workspace_size

告知所需工作空间将分配多少全局内存。如果 Description::requires_workspace 为 false，则为 0；否则大于零。

---

其他特征

特征	默认值	描述
is_supported<FFT, Architecture>::value	false	如果 FFT 在提供的 CUDA 架构 (Architecture) 上受支持，则为 true。
sm_of<Description>::value	无。	底层 FFT 函数使用的目标架构。请参阅 SM 运算符。

cufftdx::is_supported

// FFT - FFT description without CUDA architecture defined using SM operator
// Architecture - unsigned integer representing CUDA architecture (SM)
template<class FFT, unsigned int Architecture>
struct is_supported : std::bool_constant<...> { };

// Helper variable template
template<class FFT, unsigned int Architecture>
inline constexpr bool is_supported_v<FFT, Architecture> = is_supported<FFT, Architecture>::value;

// true if FFT is supported on the provided CUDA architecture
cufftdx::is_supported<FFT, Architecture>::value;

cufftdx::is_supported 检查 FFT 是否在 Architecture CUDA 架构上受支持。

// true if FFT is supported on the provided CUDA architecture
cufftdx::is_supported<FFT, Architecture>::value;

系统要求

• FFT 必须定义大小、方向（如果无法从类型推断）。请参阅 描述运算符 部分。

• FFT 必须包含 Block 运算符的 Thread。

• FFT 不能通过 SM 运算符定义目标 CUDA 架构。

• 如果 FFT 描述包含 ElementsPerThread 运算符，cufftdx::is_supported 在验证支持时会考虑这一点。

示例

using FFT = decltype(Size<32768>() + Type<fft_type::c2c>() + Direction<fft_direction::inverse>() + Block() + Precision<float>());
cufftdx::is_supported<FFT, 800>::value; // true
cufftdx::is_supported<FFT, 700>::value; // false

using FFT = decltype(Size<8192>() + Type<fft_type::c2c>() + Direction<fft_direction::forward>() + Block() + Precision<double>());
cufftdx::is_supported<FFT, 800>::value; // true
cufftdx::is_supported<FFT, 750>::value; // false
cufftdx::is_supported<FFT, 700>::value; // true

using FFT = decltype(Size<4095>() + Type<fft_type::c2c>() + Direction<fft_direction::inverse>() + Block() + Precision<float>());
cufftdx::is_supported_v<FFT, 800>; // true
cufftdx::is_supported_v<FFT, 750>; // false
cufftdx::is_supported_v<FFT, 700>; // true

cufftdx::sm_of 特征

使用 SM 运算符 定义的底层 FFT 函数使用的架构

没有默认值。