特性¶
Traits（特性）为用户提供有关使用Operators构建的FFT描述信息。它们分为三大类：
描述特性¶
特性 |
默认值 |
描述 |
|---|---|---|
无。 |
要计算的FFT大小。 |
|
|
FFT操作的类型，可以是 |
|
FFT运算的方向，可以是 |
||
|
用于计算FFT的底层浮点数值类型： |
|
无。 |
|
|
无。 |
|
|
无。 |
|
|
无。 |
|
可以使用提供的辅助函数从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
要计算的FFT大小，由Size Operator设置。
没有默认大小。如果描述符不是使用Size Operator创建的，编译将失败并显示错误消息。
方向特性¶
cufftdx::direction_of<FFT>::value
FFT运算的方向，由Direction Operator设置。
默认方向：
如果FFT类型是
fft_type::r2c，默认方向为fft_direction::forward。如果FFT类型是
fft_type::c2r，默认方向为fft_direction::inverse。对于其他任何类型，没有默认方向。如果描述符不是使用方向运算符创建的，编译将失败并显示错误消息。
是否为FFT特性¶
cufftdx::is_fft<FFT>::value
如果描述符是通过Description Operators形成的FFT描述，则Trait为true。
没有默认值。描述符要么是FFT描述，要么不是。
FFT是否正在执行？特性¶
cufftdx::is_fft_execution<FFT>::value
如果描述符是通过描述运算符和执行运算符形成的FFT描述，则Trait为true。
没有默认值。该描述符要么是包含Execution Operators的FFT描述，要么不是。
FFT是否完备？特性¶
cufftdx::is_complete_fft<FFT>::value
如果描述符是通过Description Operators形成的完整FFT描述，则Trait为true。
要使FFT描述符完整，需要满足以下条件：
有且仅有一个Size Operator。
有且仅有一个方向运算符，除非添加了
cufftdx::Type<fft_type::r2c>()或cufftdx::Type<fft_type::c2r>()。一个且仅有一个SM Operator，除非添加了Thread Operator。
没有默认值。该描述符要么是FFT完整描述，要么不是。
FFT是否完全执行？特性¶
cufftdx::is_complete_fft_execution<FFT>::value
当cufftdx::is_fft_execution和cufftdx::is_complete_fft都为true时，该特性为true。
注意
如果描述符FFT的cufftdx::is_complete_fft_execution特性为true，那么我们可以使用Execution Methods来计算FFT。
没有默认值。
执行特性¶
执行特性可以直接从已配置Execution Operators的FFT描述符中获取。 可用的执行特性取决于用于构建描述符的运算符；可以是Thread Operator或Block Operator。
线程特性¶
特性 |
默认值 |
描述 |
|---|---|---|
|
用于计算FFT的基础数据的复杂类型。 |
|
|
用作FFT输入的底层数据类型。 |
|
|
用作FFT输出的基础数据类型。 |
|
为此FFT加载的Input Type Trait类型元素的数量 |
||
在此FFT之后要存储的Output Type Trait类型元素的数量 |
||
无。 |
要加载的元素类型。提供最佳的向量化和正确性。 |
|
无。 |
要存储的元素类型。提供最佳的向量化和正确性。 |
|
无。 |
作为单批次加载的输入类型特征输入长度。 |
|
无。 |
要存储为一个批次的Output Type Trait类型输入的长度。 |
|
|
将来自不同FFT的数值批量合并到 |
|
|
每个线程需要计算的FFT元素数量。 |
|
|
每个线程必须分配的 |
|
始终 |
每个线程在其 |
线程特性可以从使用Thread Operator构建的描述符中获取。
例如：
#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;
输出EPT特性¶
FFT::output_ept
FFT执行后，单个线程将返回的Output Type Trait类型元素的最大数量。
默认值与Elements Per Thread Trait相同。
输入类型特征¶
FFT::input_type
FFT执行时需提供的元素类型。对于C2C和C2R配置，这与值类型特征相同，但在R2C中会直接取决于所使用的实数FFT选项运算符值。
默认类型与Value Type Trait相同。
输出类型特征¶
FFT::output_type
FFT执行返回的元素类型。对于C2C和R2C配置，这与值类型特征相同，但在C2R中，它直接取决于所使用的实数FFT选项运算符值。
默认类型与Value Type Trait相同。
输入长度特征¶
FFT::input_length
此FFT单批次输入的完整长度。该值对于C2C等同于Size Trait，但对于R2C取决于RealFFTOptions Operator的real_mode值，对于C2R则取决于RealFFTOptions Operator的complex_layout值。
默认值与Size Trait相同。
输出长度特性¶
FFT::output_length
此FFT单批次输出的完整长度。该值对于C2C等同于Size Trait，但对于C2R则取决于RealFFTOptions Operator的real_mode值，对于R2C则取决于RealFFTOptions Operator的complex_layout值。
默认值与Size Trait相同。
隐式类型批处理特性¶
FFT::implicit_type_batching
在FFT计算中，将来自不同FFT的多个值批量处理为Description::value_type类型的一个元素的数量。如果该值大于1，则表示线程FFT对象一次性计算多个FFT。
如果cufftdx::precision_of<FFT>::type是__half，则该值为2，否则为1。
注意
请注意，在未来的cuFFTDx版本中，FFT::implicit_type_batching可能会被替换和/或扩展。
每线程元素特性¶
FFT::elements_per_thread
每个线程将计算的FFT元素的逻辑数量。这可能与实际元素数量不同，因为RealFFTOptions Operator 可能会改变输入或输出元素的数量及其逻辑布局。关于输入和输出的内存操作，请参考Input EPT Trait和Output EPT Trait； 关于寄存器空间分配，请参考Storage Size Trait。
默认值与Size Trait相同。
块特性¶
特性 |
默认值 |
描述 |
|---|---|---|
|
用于计算FFT的基础数据的复杂类型。 |
|
|
用作FFT输入的底层数据类型。 |
|
|
用作FFT输出的基础数据类型。 |
|
无。 |
每个线程要加载的Input Type Trait类型元素数量。 |
|
无。 |
每个线程要存储的Output Type Trait类型元素的数量。 |
|
无。 |
每个线程加载的元素类型。提供最佳的向量化和正确性。 |
|
无。 |
每个线程存储的元素类型。提供最佳的向量化和正确性。 |
|
无。 |
作为单批次加载的输入类型特征输入长度。 |
|
无。 |
要存储为一个批次的Output Type Trait类型输入的长度。 |
|
|
FFT计算所需的设备端工作空间类型。 |
|
|
将来自不同FFT的数值批量合并到 |
|
启发式。 |
每个线程需要计算的FFT元素数量。 |
|
由 |
每个线程必须分配的 |
|
由 |
每个线程在其 |
|
|
在此FFT运算中，由CUDA块计算的FFT数量。 |
|
启发式。 |
建议每个CUDA块计算的FFT数量，以瞄准最佳性能。 |
|
由 |
共享内存的大小（以字节为单位）。 |
|
参见 Block Dim Trait。 |
|
|
由 |
CUDA块中的线程总数。 |
|
|
确定是否需要使用 |
|
|
工作区所需的全局内存大小（以字节为单位）。 |
块特征可以从使用Block Operator构建的描述符中获取。
例如：
#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::input_length
此FFT单批次输入的完整长度。该值对于C2C等同于Size Trait，但对于R2C取决于RealFFTOptions Operator的real_mode值，对于C2R则取决于RealFFTOptions Operator的complex_layout值。
输出长度特性¶
FFT::output_length
此FFT单批次输出的完整长度。该值对于C2C等同于Size Trait，但对于C2R则取决于RealFFTOptions Operator的real_mode值，对于R2C则取决于RealFFTOptions Operator的complex_layout值。
工作区类型特征¶
FFT::workspace_type
execute(...)函数所需的FFT工作空间类型。用户应通过Description::requires_workspace特性检查FFT是否需要工作空间，并使用cufftdx::make_workspace<FFT>(cudaError_t&, cudaStream_t)创建。
有关工作区的更多详情，请参阅创建工作区函数。
警告
FFT::workspace_type 对象不会跟踪底层内存的生命周期，仅在转换来源的workspace对象生命周期内有效。
警告
由cufftdx::make_workspace<FFT>(cudaError_t&, cudaStream_t)返回的类型可能因不同的FFT描述而异，
且与FFT::workspace_type不同。用户在创建工作区对象时应使用auto。
隐式类型批处理特性¶
FFT::implicit_type_batching
在一次FFT计算中，从不同FFT批量合并到Description::value_type类型一个元素中的数值数量。如果该值大于1，表示Block FFT对象会一次性计算多个FFT。
如果cufftdx::precision_of<FFT>::type是__half，则该值为2，否则为1。
注意
请注意，在未来的cuFFTDx版本中，FFT::implicit_type_batching可能会被替换和/或扩展。
每线程元素特性¶
FFT::elements_per_thread
每个线程将计算的FFT元素的逻辑数量。这可能与实际元素数量不同，因为RealFFTOptions Operator可能会改变输入或输出元素的数量及其在线程间的逻辑布局。关于输入和输出内存操作，请参考Input EPT Trait和Output EPT Trait；关于寄存器空间分配，请参考Storage Size Trait。
默认值与Size Trait相同。
步幅大小特性¶
FFT::stride
每个线程在其input中保存的块FFT元素之间的步长。
另请参阅预期输入数据格式。
示例
0-第8点FFT的第0个线程，其FFT::stride等于2，应在input中包含值0、2、4和6。
建议的每块FFT特性¶
FFT::suggested_ffts_per_block
建议在每个CUDA块内并行计算的FFT数量，作为集体FFT操作的一部分，以最大化性能。
默认值为启发式，取决于FFT的大小、每个线程的元素数量以及其他参数。
Block Dim 特性¶
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::requires_workspace
布尔值。如果为true，则必须创建并传递工作空间给FFT::execute(...)方法（参见block execute methods）。
否则无需创建和传递工作空间。工作空间可通过
cufftdx::make_workspace
工作空间大小特性¶
FFT::workspace_size
告知所需工作区将分配的全局内存量。如果Description::requires_workspace为false，则值为0；否则该值大于零。
其他特性¶
特性 |
默认值 |
描述 |
|---|---|---|
|
|
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无法通过 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