Operators#

Operators are used to describe the FFT operation to solve, and to configure the execution. They are divided into Description Operators and Execution Operators.

Description Operators#

Operator	Default value	Description
`Size<unsigned int S>`	Not set.	Size `S` of the FFT to calculate.
`Direction<fft_direction>`	Not set.	Direction of the FFT, either `fft_direction::inverse` or `fft_direction::forward`.
`Type<fft_type>`	`fft_type::c2c`	Types of input and output data (C2C, R2C, C2R).
`Precision<P>`	`float`	Precision `P` of the floating-point values used to compute the FFT: `double`, `float` or `__half`.
`SM<unsigned int CC>`	Not set.	Target CUDA architecure for which the FFT function should be generated.
`RealFFTOptions<complex_layout, real_mode>`	`complex_layout::natural` and `real_mode::normal`.	Complex data layout choice and execution mode for R2C/C2R FFT types

Description operators define the FFT operation to be solved. Combined with Execution Operators, they form a complete FFT descriptor that can be executed on a GPU.

Operators are added to construct the FFT descriptor type. For example, for a forward FFT operation consisting of a FFT with 8 double elements per thread:

#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:

One, and only one, Size Operator.
One, and only one, Direction Operator unless either cufftdx::Type<fft_type::r2c>() or cufftdx::Type<fft_type::c2r>() are added.
One, and only one, SM Operator unless a Thread Operator is added.

Size Operator#

cufftdx::Size<unsigned int S>()

Sets the size S of the FFT operation to compute.

There is no default size.

Restrictions:

S must be greater than 1.

Direction Operator#

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

Sets the direction of the FFT, either fft_direction::inverse or fft_direction::forward.

There is no default direction.

If the FFT is constructed with the Type<R2C> operator, direction is assumed to be forward and a direction operator is not necessary.

If the FFT is constructed with the Type<C2R> operator, direction is assumed to be inverse and a direction operator is not necessary.

Restrictions:

fft_direction::forward requires Type<C2C> or Type<R2C>.
fft_direction::inverse requires Type<C2C> or Type<C2R>.

cuFFTDx performs un-normalized FFTs; that is, performing a forward FFT on an input data set followed by an inverse FFT on the resulting set yields data that is equal to the input, scaled by the size of the FFT. Scaling either transform by the reciprocal of the size of the data set is left for the user to perform as seen fit.

Type Operator#

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

Sets the type of the FFT to compute, either fft_type::c2c for complex-to-complex; fft_type::r2c for real-to-complex; or fft_type::c2r for complex-to-real.

The default type is fft_type::c2c.

Restrictions:

fft_type::r2c requires fft_direction::forward. If no direction is specified, it is assumed to be fft_direction::forward.
fft_type::c2r requires fft_direction::inverse. If no direction is specified, it is assumed to be fft_direction::inverse.
cuFFDx performs unnormalized Fast Fourier Transform calculations.

Precision Operator#

cufftdx::Precision<__half>()

cufftdx::Precision<float>()

cufftdx::Precision<double>()

Sets the floating-point precision used to compute the FFT. This is the type of the values used for input and output, as well as the underlying type of the values used to compute the FFT.

The default precision is float.

SM Operator#

cufftdx::SM<unsigned int CC>()

Sets the target architecture CC for the underlying FFT function to use. Supported architectures are:

Volta: 700 and 720 (sm_70, sm_72).
Turing: 750 (sm_75).
Ampere: 800, 860 and 870 (sm_80, sm_86, sm_87).
Ada: 890 (sm_89).
Hopper: 900 (sm_90, sm_90a).

Note

When compiling cuFFTDx for 9.0a compute capability use 900 in the SM operator (see also CUDA C++ Programming Guide: Feature Availability).

Warning

It is not guaranteed that executions of exactly the same FFTs on GPUs of different CUDA architectures will produce bit-identical results.

RealFFTOptions Operator#

cufftdx::RealFFTOptions<complex_layout, real_mode>()

The first argument defines the input and output data layouts for real-to-complex (fft_type::r2c) and complex-to-real (fft_type::c2r) FFT types, which is explained in detail in Complex Element Layouts for complex inputs and outputs and Real Element Layouts for real inputs and outputs.

Possible complex_layout values:

complex_layout::natural (default)
complex_layout::packed
complex_layout::full

The second argument, a real_mode enum value, allows to choose between optimized and non optimized execution kernel. There are two modes: normal (default), and optimized, which is called folded. The optimized execution is an opt-in, as it changes how the data is partitioned between threads in block execution.

Possible real_mode values:

real_mode::normal (default)
real_mode::folded

All traits necessary for correct I/O are shared by the topmost FFT type. Please refer to examples (e.g. simple_fft_block_r2c and simple_fft_block_c2r) for a direct showcase of proper idioms.

Execution Operators#

Operator	Default value	Description
`Thread`	Not set.	Creates FFT thread execution object.
`Block`	Not set.	Creates FFT block execution object. See Block Configuration Operators.

Execution operators configure how the FFT operation will run on the GPU. Combined with Description Operators, they form a complete FFT descriptor that can be executed on a GPU.

Operators are added to construct the FFT descriptor type. For example, for a forward FFT operation consisting of two FFTs with 128 float elements each, running simultaneously in one CUDA block:

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

Thread Operator#

cufftdx::Thread()

Sets the FFT operation to run in a thread context. The FFT operation will simultaneously run a single, independent FFT (described using Description Operators) per thread.

Each thread will compute one FFT of the size defined by the Size Operator.

Restrictions:

Is mutually exclusive with Block operator
Compilation will fail when used with block-only operators: FFTsPerBlock, ElementsPerThread, BlockDim.
With Precision<__half> restricts Size to range \([2, 32]\).
With Precision<float> restricts Size to range \([2, 32]\).
With Precision<double> restricts Size to range \([2, 16]\).

Block Operator#

cufftdx::Block()

Generates a collective FFT operation to run in a single CUDA block. One or more threads will cooperate to compute the collective FFT operation.

The number of FFTs to compute, as well as the number of threads used to calculate each FFT, can be configured using Block Configuration Operators.

Restrictions:

Is mutually exclusive with Thread operator
Unless a BlockDim Operator is used, the collective FFT operation can only be executed inside a 2D block of sizes:
- blockDim.x = size_of<Description>::value/Description::elements_per_thread.
- blockDim.y = Description::ffts_per_block.
- blockDim.z = 1.
BlockDim Operator is not implemented yet.
Operator cufftdx::Precision<__half>() restricts cufftdx::Size<U>() to range \([2, 32768]\).
Operator cufftdx::Precision<float>() restricts cufftdx::Size<U>() to range \([2, 32768]\).
Operator cufftdx::Precision<double>() restricts cufftdx::Size<U>() to range \([2, 16384]\).

Block Configuration Operators#

Operators	Default value	Description
`FFTsPerBlock<unsigned int F>`	1	Number `F` of FFTs calculated per CUDA block.
`ElementsPerThread<unsigned int E>`	Heuristic.	Number `E` of FFT values per CUDA thread.
`BlockDim<unsigned int X, Y, Z>`	Not set.	Required for executing block FFT within block with custom dimensions.

Block-configuration operators allow the user to tune how the collective FFT operation will run on a single CUDA block.

Note

Block configuration operators can only be used with Block Operator.

Warning

It is not guaranteed that executions of the same FFTs (size, direction, type, precision) but with different

number of elements per thread (ElementsPerThread),
number of FFTs calculated per CUDA block (FFTsPerBlock), or
block dimension (BlockDim),

will produce bit-identical results.

FFTs Per Block Operator#

cufftdx::FFTsPerBlock<unsigned int>()

Sets the number of FFT to compute in parallel within a single CUDA block. Each FFT is computed concurrently by a separate group of threads.

The default is one FFT per block.

Elements Per Thread Operator#

cufftdx::ElementsPerThread<unsigned int>()

Sets the number of FFT elements to be computed by each thread.

The default is determined heuristically to target performance.

Restrictions:

If FFT::requires_workspace is false, it must be a divisor of the requested FFT size.
If FFT::requires_workspace is true, it must be a power of two smaller than the size of the FFT.
Must be in range \([2; 32]\) for cufftdx::Precision<float>() and cufftdx::Precision<__half>().
Must be in range \([2; 16]\) for cufftdx::Precision<double>().

BlockDim Operator#

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

Note

BlockDim operator is not implemented yet. Details below show rules used for configuring block dimensions for FFT operators and with these insights cuFFTDx users can adjust execution configuration even though BlockDim operator is not available.

Sets the CUDA block size to (X, Y, Z), to configure the execution.

Using this operator, the user can run the collective FFT operation with 2D or 3D CUDA blocks.

The default block dimensions for an FFT are:

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

See FFT::block_dim.