### Introduction to cuFFTDx API

Source: https://docs.nvidia.com/cuda/cufftdx/examples.html

This example is used in the introductory guide to explain the basics of the cuFFTDx library and its API.

```cpp
introduction_example
```

--------------------------------

### Example LTO Helper Execution

Source: https://docs.nvidia.com/cuda/cufftdx/lto_helper.html

An example command demonstrating how to run the LTO Helper tool with a specific output folder, input CSV file, and CUDA architectures.

```bash
./cufftdx_cufft_lto_helper output_folder input.csv --CUDA_ARCHITECTURES=75

```

--------------------------------

### Build Examples with Optimized Binary Size

Source: https://docs.nvidia.com/cuda/cufftdx/installation.html

This command builds the cuFFTDx examples using the optimization for binary size, which involves linking to a shared object with pre-calculated constants.

```makefile
make all_twiddles

```

--------------------------------

### cuFFTDx Execution Example

Source: https://docs.nvidia.com/cuda/cufftdx/api/methods.html

An example demonstrating the setup and execution of an FFT operation using cuFFTDx within a CUDA kernel.

```APIDOC
## cuFFTDx Execution Setup

### Description
This section shows how to define an FFT operation using cuFFTDx descriptors and how to set up a CUDA kernel for execution.

### Code Example
```cpp
#include <cufftdx.hpp>

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

using complex_type = typename FFT::value_type;

__global__ kernel(... /* arguments */) {

  // Shared memory pointer
  extern __shared__ __align__(alignof(float4)) complex_type shared_mem[];

  // Register data
  complex_type thread_data[FFT::storage_size];

  // Load data into registers (thread_data)
  // ...

  FFT().execute(thread_data, shared_mem);

  // Store results (thread_data) into global memory
}
```
```

--------------------------------

### Example FFT Descriptor Configuration and Trait Retrieval

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

This example demonstrates how to configure an FFT descriptor using various operators and then retrieve traits like value_type and elements_per_thread.

```c++
#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;
```

--------------------------------

### Example Usage of Traits

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Example demonstrating how to use cuFFTDx traits to query FFT properties.

```APIDOC
```cpp
#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 Operator Combination Example

Source: https://docs.nvidia.com/cuda/cufftdx/api/operators.html

Example demonstrating how to combine Description Operators to construct an FFT descriptor.

```APIDOC
## Operator Combination Example

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:

```cpp
#include <cufftdx.hpp>

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

### Completeness Requirements for an FFT Descriptor

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.
```

--------------------------------

### Dynamic Batching Shared Memory Usage Example

Source: https://docs.nvidia.com/cuda/cufftdx/api/other_methods.html

Demonstrates how to query traits and compute total shared memory size for a dynamic batching configuration.

```cpp
using FFT = decltype(Size<128>() + Precision<float>() +
               Type<fft_type::c2c>() +
               Direction<fft_direction::forward>() +
               ElementsPerThread<8>() + DynamicBatching());

// Query compile-time traits
constexpr auto shared_memory_size_per_fft =
    FFT::shared_memory_size;
constexpr auto implicit_type_batching =
    FFT::implicit_type_batching;

// Runtime value chosen by user
unsigned int ffts_per_block = 8;

// Compute total shared memory requirement
auto total_shared_memory_size_bytes =
    cufftdx::experimental::utils::
        get_shared_memory_size_for_dynamic_batching(
            shared_memory_size_per_fft,
            ffts_per_block,
            implicit_type_batching);

auto total_shared_memory_size_bytes_with_desc =
    cufftdx::experimental::utils::
        get_shared_memory_size_for_dynamic_batching<FFT>(
            ffts_per_block);
```

--------------------------------

### Example Usage of is_supported

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Demonstrates checking support for various FFT configurations across different CUDA architectures.

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

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

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
```

--------------------------------

### Execute FFT with Workspace

Source: https://docs.nvidia.com/cuda/cufftdx/api/methods.html

This example demonstrates how to create a workspace using cufftdx::make_workspace and then launch a kernel that utilizes this workspace for executing the FFT. Ensure the workspace is created for the correct FFT descriptor.

```cpp
// Kernel
template<class FFT>
__launch_bounds__(FFT::max_threads_per_block)
__global__ void block_fft_kernel(
    typename FFT::value_type* data,
    typename FFT::workspace_type workspace) {
    // ...

    // Execute FFT
    FFT().execute(thread_data, shared_mem, workspace);
}

// Create workspace
cudaError_t error = 0;
auto workspace = cufftdx::make_workspace<FFT>(error, stream);

// ...

// Run kernel with FFT
block_fft_kernel<FFT><<<1, FFT::block_dim,
    FFT::shared_memory_size, stream>>>(data, workspace);


```

--------------------------------

### Retrieve FFT Description Traits

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Example demonstrating how to check if an FFT descriptor is complete and retrieve its size using traits.

```cpp
#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;
```

--------------------------------

### CMake Example for cuFFTDx LTO Database Generation

Source: https://docs.nvidia.com/cuda/cufftdx/lto_helper.html

Demonstrates how to use `run_cufft_lto_helper` to generate LTO artifacts and link the resulting library to an executable. The `build_cufft_lto_helper` is called automatically.

```cmake
include(/path/to/lto_helper/lto_helper.cmake)

# Generate LTO database
run_cufft_lto_helper(
   SRC_DIR ${CMAKE_SOURCE_DIR}/lto_helper
   OUTPUT_DIR ${CMAKE_BINARY_DIR}/lto_helper
   OUTPUT_NAME my_fft
   DESCS ${CMAKE_SOURCE_DIR}/input.csv
   ARCHITECTURES 75;80
)

# Link the generated LTO library to your executable
add_executable(my_executable main.cpp)
target_link_libraries(my_executable PRIVATE my_fft_lto_lib)

# Include the LTO database header
target_include_directories(my_executable
   PRIVATE
      ${CMAKE_BINARY_DIR}/lto_helper/my_fft_artifacts
)

```

--------------------------------

### cuFFTDx FFT Initialization and Kernel Launch Example

Source: https://docs.nvidia.com/cuda/cufftdx/introduction2.html

This C++ function demonstrates how to define an FFT execution descriptor using cuFFTDx operators, allocate managed memory, create a workspace, and launch the custom FFT kernel. It shows how to obtain the suggested FFTs per block for optimal performance.

```C++
void introduction_example(cudaStream_t stream = 0) {
  // Base of the FFT description
  using FFT_base = decltype(Size<128>() + Precision<float>() +
      Type<fft_type::c2c>() +
      Direction<fft_direction::forward>()
      /* Notice lack of ElementsPerThread and FFTsPerBlock
       * operators */
      + SM<750>() + Block());
  // FFT description with suggested FFTs per CUDA block for the
  // default (optimal) elements per thread
  using FFT = decltype(FFT_base() +
      FFTsPerBlock<FFT_base::suggested_ffts_per_block>());

  // 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);
  // Generate data
  for (size_t i = 0; i < size; i++) {
      data[i] = complex_type {float(i), -float(i)};
  }

  cudaError_t error_code = cudaSuccess;


  auto workspace = make_workspace<FFT>(error_code, stream);

  // Invokes kernel with FFT::block_dim threads in CUDA block
  block_fft_kernel<FFT><<<1, FFT::block_dim,
                          FFT::shared_memory_size,
                          stream>>>(data, workspace);
  cudaDeviceSynchronize();

  cudaFree(data);
}

```

--------------------------------

### FFT Descriptor Construction

Source: https://docs.nvidia.com/cuda/cufftdx/api/operators.html

Example of constructing a complete FFT descriptor type using various configuration operators.

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

--------------------------------

### Example CSV for FFT Configurations

Source: https://docs.nvidia.com/cuda/cufftdx/lto_helper.html

Defines three distinct FFT configurations for inclusion in the LTO database. Ensure all required fields are present for each configuration.

```csv
size,direction,precision,type,real_mode,execution_type,elements_per_thread
4,fft_direction::forward,float,fft_type::c2c,real_mode::normal,execution_type::block,
1024,,,fft_type::c2c,real_mode::normal,execution_type::block,32
2048,,double,fft_type::r2c,real_mode::folded,execution_type::thread,16

```

--------------------------------

### Define FFT configuration in CSV

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_lto/augment_with_lto.html

Example CSV file format for specifying FFT parameters to the LTO Helper.

```csv
size,direction,execution_type
128,fft_direction::forward,execution_type::block
```

--------------------------------

### Shared Memory Alignment Examples

Source: https://docs.nvidia.com/cuda/cufftdx/api/methods.html

Examples demonstrating how to ensure shared memory pointers are aligned to 128 bits (16 bytes) for optimal performance. Proper alignment can be achieved using __align__ or alignas directives.

```cpp
// Examples of how to make sure shared memory pointer is aligned
// to 128 bits (16 bytes):
using value_type = typename FFT::value_type;
extern __shared__ alignas(float4)
  unsigned char shared_mem[];                              // 1
```

```cpp
extern __shared__ __align__(16)
  value_type shared_mem[];                                 // 2
```

```cpp
extern __shared__ __align__(alignof(float4))
  value_type shared_mem[];                                 // 3
```

```cpp
extern __shared__ float4 shared_mem[];                     // 4
```

```cpp
// Warning: std::aligned_storage became deprecated in C++23
extern __shared__ std::aligned_storage_t<sizeof(float4),
                                         alignof(float4)> 
  shared_mem[];                                            // 5
```

--------------------------------

### Define FFT with Block Traits

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

This example demonstrates how to define an FFT configuration using various block traits such as size, type, direction, precision, and elements per thread. It also shows how to allocate managed memory for input/output data based on the FFT configuration.

```cpp
#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);

```

--------------------------------

### cufftDeviceGetNumDeviceFunctions

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_api/device_api.html

Gets the number of device functions that match a given description.

```APIDOC
## cufftDeviceGetNumDeviceFunctions

### Description
Gets the number of device functions that match a given description from the associated device handle.

### Parameters
#### Path Parameters
- **handle** (cufftDeviceHandle) - Required - Device handle.
- **desc_handle** (cufftDescriptionHandle) - Required - Description handle.
- **func_count** (size_t*) - Required - Pointer to store the number of functions (Out).

### Response
#### Success Response (CUFFT_SUCCESS)
- **func_count** (size_t) - Number of matching functions.

#### Error Handling
- **CUFFT_INVALID_VALUE**: Returned if handle/desc_handle is invalid or func_count pointer is NULL.
```

--------------------------------

### Get cuFFTDx LTO Database and LTOIRs

Source: https://docs.nvidia.com/cuda/cufftdx/api/other_methods.html

Use this function to obtain the database string and LTOIRs for a specified FFT operation. Ensure `CUFFTDX_ENABLE_CUFFT_DEPENDENCY` is defined and link against the cuFFT library.

```cpp
std::tuple<std::string, std::vector<std::vector<char>>, Dim3,
           unsigned int>
cufftdx::utils::get_database_and_ltoir(
    unsigned int                    fft_size,
    cufftdx::fft_direction          dir,
    cufftdx::fft_type               type,
    unsigned int                    sm,
    cufftdx::detail::execution_type execution,
    cufftdx::precision              prec = 
        cufftdx::precision::f32,
    cufftdx::complex_layout         layout = 
        cufftdx::complex_layout::natural,
    cufftdx::real_mode              rmode = 
        cufftdx::real_mode::normal,
    unsigned int                    fft_ept = 0
        /* use heuristic */,
    unsigned int                    ffts_per_block = 1
        /* 0: use suggested ffts_per_block */);
```

```cpp
// Assuming the following FFT operator is defined in the NVRTC-compiled code:
// using FFT = decltype(cufftdx::Block() +
//                      cufftdx::Size<128>() +
//                      cufftdx::Type<cufftdx::fft_type::c2c>() +
//                      cufftdx::Direction<cufftdx::fft_direction::forward>() +
//                      cufftdx::Precision<float>() +
//                      cufftdx::ElementsPerThread<8>() +
//                      cufftdx::FFTsPerBlock<2>() +
//                      cufftdx::SM<750>));

// You can get the database string and LTOIRs for the FFT operation by calling:
auto [lto_db, ltoirs, block_dim, sm_size] =
   cufftdx::utils::get_database_and_ltoir(128,
                                          cufftdx::fft_direction::forward,
                                          cufftdx::fft_type::c2c,
                                          750,
                                          cufftdx::detail::execution_type::block,
                                          cufftdx::precision::f32,
                                          cufftdx::complex_layout::natural,
                                          cufftdx::real_mode::normal,
                                          8,
                                          2);
```

--------------------------------

### Get All Matching cuFFTDx FFT Implementations

Source: https://docs.nvidia.com/cuda/cufftdx/api/other_methods.html

Use `get_all_implementations` to retrieve a vector of all FFT implementations matching the specified configuration. Returns an empty vector if no implementations are found. Requires `CUFFTDX_ENABLE_RUNTIME_DATABASE` to be defined.

```cpp
std::vector<cufftdx::utils::frontend_impl_traits>
cufftdx::experimental::utils::get_all_implementations(
    unsigned int                    fft_size,
    cufftdx::fft_direction          dir,
    cufftdx::fft_type               type,
    unsigned int                    sm,
    cufftdx::utils::execution_type  execution,
    cufftdx::precision              prec = 
        cufftdx::precision::f32,
    unsigned int                    fft_ept = 0,
    unsigned int                    ffts_per_block = 0,
    std::tuple<unsigned int, unsigned int, unsigned int>
        block_dim = {0, 0, 0},
    cufftdx::complex_layout         layout = 
        cufftdx::complex_layout::natural,
    cufftdx::real_mode              rmode = 
        cufftdx::real_mode::normal,
    cufftdx::experimental::code_type code_type = 
        cufftdx::experimental::code_type::ptx);

```

--------------------------------

### Get Database String Size (cufftDeviceGetDatabaseStrSize)

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_api/device_api.html

Obtains the required size for the database string. This function is useful for allocating the correct buffer size before retrieving the string itself. An optional function handle can be provided.

```c
cufftResult cufftDeviceGetDatabaseStrSize( 
    _cufftDeviceHandle handle_,     _size_t *db_str_size_,     _cufftDeviceFunctionHandle func_handle_, );
```

--------------------------------

### Get LTO Database and LTOIRs using NVRTC

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_lto/augment_with_lto.html

Use cufftdx::utils::get_database_and_ltoir to obtain pointers to the C++ header and LTOIRs for online kernel generation. This function wraps the cuFFT Device API and requires CUFFTDX_ENABLE_CUFFT_DEPENDENCY to be defined and cuFFT library to be linked.

```cpp
auto [lto_db, ltoirs, block_dim, shared_memory_size] =
   cufftdx::utils::get_database_and_ltoir(
      fft_size,
      cufftdx::fft_direction::inverse,
      cufftdx::fft_type::c2c,
      example::nvrtc::get_device_architecture(current_device) * 10,
      cufftdx::detail::execution_type::block,
      cufftdx::precision::f64,
      cufftdx::complex_layout::natural,
      cufftdx::real_mode::normal,
      fft_ept
   );

```

--------------------------------

### Create and configure description handles

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_lto/custom_lto_helper.html

Initializes a description handle and sets specific traits using the mapped backend traits.

```cpp
cufftDescriptionHandle desc_handle;
cufftDescriptionCreate(&desc_handle);
cufftDescriptionSetTraitInt64(desc_handle, CUFFT_DESC_TRAIT_SIZE,
   static_cast<long long int>(backend_traits.size));
cufftDescriptionSetTraitInt64(desc_handle, CUFFT_DESC_TRAIT_TYPE,
   static_cast<long long int>(backend_traits.type));
cufftDescriptionSetTraitInt64(desc_handle, CUFFT_DESC_TRAIT_DIRECTION,
   static_cast<long long int>(backend_traits.direction));
cufftDescriptionSetTraitInt64(desc_handle, CUFFT_DESC_TRAIT_SM,
   static_cast<long long int>(backend_traits.sm));
cufftDescriptionSetTraitInt64(desc_handle, CUFFT_DESC_TRAIT_PRECISION,
   static_cast<long long int>(backend_traits.precision));
cufftDescriptionSetTraitInt64(desc_handle, CUFFT_DESC_TRAIT_ELEMENTS_PER_THREAD,
   static_cast<long long int>(backend_traits.elements_per_thread));
```

--------------------------------

### Run LTO Helper

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_lto/augment_with_lto.html

Command to execute the LTO helper to generate the database files.

```bash
./cufftdx_cufft_lto_helper <output_dir_path>
                           introduction_lto_example.csv \
                           --CUDA_ARCHITECTURES=XX;YY;...
```

--------------------------------

### Define and Launch FFT Kernel (cuFFTDx)

Source: https://docs.nvidia.com/cuda/cufftdx/introduction1.html

Defines a generic FFT kernel and demonstrates its invocation from the host. Requires cuFFTDx library and CUDA. The kernel requires block size and shared memory, which are determined by the FFT description.

```cpp
#include <cufftdx.hpp>
using namespace cufftdx;

template<class FFT>
__global__ void block_fft_kernel(FFT::value_type* data,
                                  typename FFT::workspace_type workspace) {
  using complex_type = typename FFT::value_type;
  // Registers
  complex_type thread_data[FFT::storage_size];

  // FFT::shared_memory_size bytes of shared memory
  extern __shared__ __align__(alignof(float4))
    complex_type shared_mem[];

  // Execute FFT
  FFT().execute(thread_data, shared_mem, workspace);
}

// CUDA_CHECK_AND_EXIT - marco checks if function returns
// cudaSuccess; if not it prints the error code and exits the
// program
void introduction_example(FFT::value_type* data
    /* data is a manage memory pointer*/) {
  // FFT description
  using FFT = decltype(Size<128>() + Precision<float>() + Type<fft_type::c2c>()
                       + Direction<fft_direction::forward>() + FFTsPerBlock<1>()
                       + ElementsPerThread<8>() + SM<750>() + Block());

  cudaError_t error_code = cudaSuccess;
  auto workspace = make_workspace<FFT>(error_code);
  CUDA_CHECK_AND_EXIT(error_code);

  // Invokes kernel with FFT::block_dim threads in CUDA block
  block_fft_kernel<FFT><<<1, FFT::block_dim,
                          FFT::shared_memory_size>>>(data, workspace);
  CUDA_CHECK_AND_EXIT(cudaPeekAtLastError());
  CUDA_CHECK_AND_EXIT(cudaDeviceSynchronize());
}

```

--------------------------------

### Get Real FFT Mode Trait

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Retrieves the optimization mode for FFT kernel execution, set by the RealFFTOptions Operator. The default is real_mode::normal.

```c++
cufftdx::real_fft_mode_of<FFT>::value
cufftdx::real_fft_mode_of_v<FFT>
```

--------------------------------

### Create device handle

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_lto/custom_lto_helper.html

Creates a device handle from the configured description handle.

```cpp
cufftDeviceHandle device_handle;
cufftDeviceCreate(&device_handle, 1, &desc_handle);
```

--------------------------------

### Get Real FFT Layout Trait

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Retrieves the input and output data layout for real-to-complex and complex-to-real FFT types. The default is complex_layout::natural.

```c++
cufftdx::real_fft_layout_of<FFT>::value
cufftdx::real_fft_layout_of_v<FFT>
```

--------------------------------

### Get Number of LTOIRs (cufftDeviceGetNumLTOIRs)

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_api/device_api.html

Retrieves the count of LTOIRs associated with a device handle. An optional function handle can be provided to query for a specific function.

```c
cufftResult cufftDeviceGetNumLTOIRs( 
    _cufftDeviceHandle handle_,     _size_t *code_count_,     _cufftDeviceFunctionHandle func_handle = -1_, );
```

--------------------------------

### Configure cuFFT LTO Database for Online Generation

Source: https://docs.nvidia.com/cuda/cufftdx/installation.html

Required macros, module paths, and linking instructions for online kernel generation using cuFFT LTO.

```cmake
target_compile_definitions(YourProgram PRIVATE CUFFTDX_ENABLE_CUFFT_DEPENDENCY)
target_compile_definitions(YourProgram PRIVATE CUFFT_ENABLE_EXPERIMENTAL_API)
```

```cmake
list(APPEND CMAKE_MODULE_PATH "${cufftdx_cufft_MODULE_PATH}") # Path to Findcufft.cmake
find_package(cufft REQUIRED)
```

```cmake
cmake -Dcufft_ROOT="<your_ctk_directory>/" (...)
```

```cmake
target_link_libraries(YourProgram
  PRIVATE
    cufft::cufft_static # or, cufft::cufft
    CUDA::nvJitLink
)
```

--------------------------------

### Compile database creation tool

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_lto/custom_lto_helper.html

Command to compile the helper executable for generating the LTO database.

```bash
g++ -I<cufft_include_path> \
    -L<cufft_lib_path> \
    -I<cuda_include_path> \
    cufft_device_api_lto_helper.cpp -o cufft_device_api_lto_helper
```

--------------------------------

### Retrieve FFT Output Type

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Access the `output_type` trait from an FFT configuration to get the type of elements returned from FFT execution. This can differ from `value_type` for C2R configurations.

```cpp
FFT::output_type

```

--------------------------------

### Retrieve FFT Input Type

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Access the `input_type` trait from an FFT configuration to get the type of elements to be provided for FFT execution. This can differ from `value_type` for R2C configurations.

```cpp
FFT::input_type

```

--------------------------------

### Query cuFFTDx Database for Optimal FFT Implementation

Source: https://docs.nvidia.com/cuda/cufftdx/api/other_methods.html

Use `query_database` to find the best FFT implementation for a given configuration without requiring a workspace. Returns an optional containing implementation traits or nullopt if none is found. Requires `CUFFTDX_ENABLE_RUNTIME_DATABASE` to be defined.

```cpp
std::optional<cufftdx::utils::frontend_impl_traits>
cufftdx::experimental::utils::query_database(
    unsigned int                    fft_size,
    cufftdx::fft_direction          dir,
    cufftdx::fft_type               type,
    unsigned int                    sm,
    cufftdx::utils::execution_type  execution,
    cufftdx::precision              prec = 
        cufftdx::precision::f32,
    unsigned int                    fft_ept = 0,
    unsigned int                    ffts_per_block = 0,
    std::tuple<unsigned int, unsigned int, unsigned int>
        block_dim = {0, 0, 0},
    cufftdx::complex_layout         layout = 
        cufftdx::complex_layout::natural,
    cufftdx::real_mode              rmode = 
        cufftdx::real_mode::normal,
    cufftdx::experimental::code_type code_type = 
        cufftdx::experimental::code_type::ptx);

```

--------------------------------

### Get Trait Value from Device Function (cufftDeviceGetDeviceFunctionTraitInt64)

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_api/device_api.html

Retrieves a specific trait value from a device function. Ensure valid handles and a non-NULL value pointer are provided.

```c
cufftResult cufftDeviceGetDeviceFunctionTraitInt64( 
    _cufftDeviceHandle handle_,     _cufftDeviceFunctionHandle func_handle_,     _cufftDeviceFunctionTrait trait_,     _long long int *value_, );
```

--------------------------------

### Get FFT Code Type Trait

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Retrieves the type of code used for FFT compute device functions, as set by the CodeType Operator. The default is experimental::code_type::ptx.

```c++
cufftdx::experimental::code_type_of<Description>::value
cufftdx::experimental::code_type_of_v<Description>
```

--------------------------------

### Check device support

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_lto/custom_lto_helper.html

Verifies if the provided description handle is supported by the device.

```cpp
cufftDeviceIsSupported(device_handle, desc_handle);
```

--------------------------------

### Run LTO Helper Tool

Source: https://docs.nvidia.com/cuda/cufftdx/lto_helper.html

Execute the compiled LTO Helper tool. Provide the output directory and the input CSV file path. Optionally specify target CUDA architectures.

```bash
./cufftdx_cufft_lto_helper <output_dir> <csv_file> [--CUDA_ARCHITECTURES=XX;YY;...]

```

--------------------------------

### Configure cuFFT LTO Database for Offline Generation

Source: https://docs.nvidia.com/cuda/cufftdx/installation.html

Steps to generate an LTO database and configure target properties for link-time optimization.

```cmake
include(${cufftdx_ROOT}/cufftdx/example/lto_helper/lto_helper.cmake)

run_cufft_lto_helper(
  SRC_DIR ${cufftdx_ROOT}/cufftdx/example/lto_helper/
  BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}
  OUTPUT_NAME my_fft
  DESCS ${CMAKE_SOURCE_DIR}/input.csv
  ARCHITECTURES 70,80
)

target_link_libraries(YourProgram PRIVATE my_fft_lto_lib)

set_target_properties(YourProgram
  PROPERTIES
    CUDA_SEPARABLE_COMPILATION ON
    INTERPROCEDURAL_OPTIMIZATION ON
)
```

--------------------------------

### Retrieve FFT Input Length

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Access the `input_length` trait from an FFT configuration to get the complete length of a single batch input. This value can depend on `real_mode` or `complex_layout` for R2C/C2R configurations.

```cpp
FFT::input_length

```

--------------------------------

### Retrieve FFT Value Type

Source: https://docs.nvidia.com/cuda/cufftdx/api/traits.html

Access the `value_type` trait from an FFT configuration to get the complex data type used for FFT computation. The default is `cufftdx::detail::complex<float>`.

```cpp
FFT::value_type

```

--------------------------------

### cufftDeviceCreate

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_api/device_api.html

Creates a new device handle with the specified FFT descriptions.

```APIDOC
## cufftDeviceCreate

### Description
Creates a new device handle with the specified FFT descriptions. For each FFT description provided, cuFFT will search for device functions with matching traits.

### Parameters
#### Path Parameters
- **handle** (cufftDeviceHandle*) - Required - Pointer to a cufftDeviceHandle object (Out: contains handle value).
- **desc_count** (size_t) - Required - Number of descriptions.
- **desc_handles** (cufftDescriptionHandle*) - Required - Pointer to an array of description handles.

### Response
#### Success Response (CUFFT_SUCCESS)
- **handle** (cufftDeviceHandle) - The created device handle.

#### Error Handling
- **CUFFT_INVALID_VALUE**: Returned if handle pointer or desc_handles pointer is NULL, or if any description handle is invalid.
- **CUFFT_INTERNAL_ERROR**: Returned if the internal handle factory reaches maximum capacity.
```

--------------------------------

### Instantiate and Execute FFT (Register Only)

Source: https://docs.nvidia.com/cuda/cufftdx/introduction1.html

Instantiates an FFT descriptor and calls its execute method. The input data is expected in thread_data registers, and results are stored there. This is a basic execution without explicit shared memory or workspace arguments.

```cpp
#include <cufftdx.hpp>
using namespace cufftdx;

// FFT description
using FFT = decltype(Size<128>() + Precision<float>() + Type<fft_type::c2c>()
                     + Direction<fft_direction::forward>() + FFTsPerBlock<1>()
                     + ElementsPerThread<8>() + SM<750>() + Block());
using complex_type = typename FFT::value_type;

__global__ void block_fft_kernel(complex_type* data) {
  // Execute FFT
  FFT().execute(/* What are the arguments? */);
}
```

--------------------------------

### Full FFT Kernel with Global Memory I/O (cuFFTDx)

Source: https://docs.nvidia.com/cuda/cufftdx/introduction1.html

Implements a complete FFT kernel that includes loading data from global memory, executing the FFT, and saving results back to global memory. This version is functionally equivalent to cuFFT's complex-to-complex kernel for size 128 and single precision.

```cpp
#include <cufftdx.hpp>
using namespace cufftdx;

template<class FFT>
__global__ void block_fft_kernel(FFT::value_type* data,
                                  typename FFT::workspace_type workspace) {
  using complex_type = typename FFT::value_type;

  // Registers
  complex_type thread_data[FFT::storage_size];

  // Local batch id of this FFT in CUDA block, in range
  // [0; FFT::ffts_per_block)
  const unsigned int local_fft_id = threadIdx.y;
  // Global batch id of this FFT in CUDA grid is equal to number
  // of batches per CUDA block (ffts_per_block) times CUDA block
  // id, plus local batch id.
  const unsigned int global_fft_id =
      (blockIdx.x * FFT::ffts_per_block) + local_fft_id;

  // Load data from global memory to registers
  const unsigned int offset = cufftdx::size_of<FFT>::value * global_fft_id;
  const unsigned int stride = FFT::stride;
  unsigned int       index  = offset + threadIdx.x;
  for (unsigned int i = 0; i < FFT::elements_per_thread; i++) {
      // Make sure not to go out-of-bounds
      if ((i * stride + threadIdx.x) < cufftdx::size_of<FFT>::value) {
          thread_data[i] = data[index];
          index += stride;
      }
  }

  // FFT::shared_memory_size bytes of shared memory
  extern __shared__ __align__(alignof(float4)) complex_type shared_mem[];

  // Execute FFT
  FFT().execute(thread_data, shared_mem, workspace);

  // Save results
  index = offset + threadIdx.x;
  for (unsigned int i = 0; i < FFT::elements_per_thread; i++) {
      if ((i * stride + threadIdx.x) < cufftdx::size_of<FFT>::value) {
          data[index] = thread_data[i];
          index += stride;
      }
  }
}

void introduction_example() {
  // FFT description
  using FFT = decltype(Size<128>() + Precision<float>() +
                       Type<fft_type::c2c>() +
                       Direction<fft_direction::forward>() +
                       FFTsPerBlock<1>() + ElementsPerThread<8>() +
                       SM<750>() + Block());

  // 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);
  // Generate data
  for (size_t i = 0; i < size; i++) {
      data[i] = complex_type {float(i), -float(i)};
  }

  cudaError_t error_code = cudaSuccess;

```

--------------------------------

### Real Input/Output Layouts for Thread Execution

Source: https://docs.nvidia.com/cuda/cufftdx/api/methods.html

Explains the 'normal' and 'folded' real input layouts for thread execution, detailing how data is partitioned into thread-local arrays. This applies to both R2C input and C2R output. Half-precision FFTs account for implicit batching.

```c++
thread / element | r0 | r1 | r2 | r3 | r4 | r5 | r6 | r7  
--|---|---|---|---|---|---|---|---
0 |  **X** |  |  |  |  **X** |  |  |   
1 |  |  **X** |  |  |  |  **X** |  |   
2 |  |  |  **X** |  |  |  |  **X** |   
3 |  |  |  |  **X** |  |  |  |  **X**
```

```c++
thread / element | (r0, r1) | (r2, r3) | (r4, r5) | (r6, r7)  
--|---|---|---|---
0 |  **X** |  |  |   
1 |  |  **X** |  |   
2 |  |  |  **X** |   
3 |  |  |  |  **X**
```

--------------------------------

### Convert Frontend FFT Description to Backend Traits

Source: https://docs.nvidia.com/cuda/cufftdx/api/other_methods.html

Use this function to map frontend FFT parameters to backend implementation traits for database queries.

```cpp
#include "cufftdx/utils.hpp"

namespace cufftdx {
   namespace utils {
      struct backend_impl_traits {
         unsigned int  size;
         fft_type      type;
         fft_direction direction;
         unsigned int  sm;
         unsigned int  elements_per_thread;
         unsigned int  min_elements_per_thread;
      };

      enum class algorithm {
         ct,
         bluestein,
      };

      enum class execution_type {
         thread,
         block,
      };

      backend_impl_traits backend_traits =
         frontend_to_backend(
            algorithm algo,
            execution_type exec_type,
            unsigned int fft_size
                /* size_of<FFT>::value */,
            fft_type type
                /* type_of<FFT>::value */,
            fft_direction direction
                /* direction_of<FFT>::value */,
            unsigned int sm
                /* sm_of<FFT>::value */,
            real_mode real_mode
                /* real_mode_of<FFT>::value */,
            unsigned int elements_per_thread
                /* elements_per_thread_of<FFT>::value or
                 * 0 if not set */,
            unsigned int block_dim_x
                /* block_dim_of<FFT>::x or 0 if not set */,
            experimental::code_type code_type
                /* experimental::code_type_of<FFT>::value */
         );
   } // namespace utils
} // namespace cufftdx
```

--------------------------------

### Build LTO Helper with CMake

Source: https://docs.nvidia.com/cuda/cufftdx/lto_helper.html

Build the LTO Helper using CMake. This method utilizes the provided CMakeLists.txt file. Specify CUDA architectures and build type as needed. Requires CMake 3.24+.

```bash
cd cufftdx/example/lto_helper
mkdir build && cd build
cmake -DCMAKE_CUDA_ARCHITECTURES=75 -DCMAKE_BUILD_TYPE=Release ..
cmake --build .

```

--------------------------------

### Compiling cuFFTDx Programs with nvcc

Source: https://docs.nvidia.com/cuda/cufftdx/introduction1.html

This command demonstrates how to compile a C++ program that includes the cuFFTDx library using nvcc. Ensure the include directory for cuFFTDx is correctly specified.

```bash
nvcc -std=c++17 -arch sm_75 -O3 -I<mathdx_include_dir> \
     introduction_example.cu -o introduction_example
```

--------------------------------

### Define Basic FFT Properties

Source: https://docs.nvidia.com/cuda/cufftdx/introduction1.html

Initializes a 128-point single precision complex-to-complex forward FFT description.

```cpp
// cuFFTDx header
#include <cufftdx.hpp>
using namespace cufftdx;

// FFT description:
// A 128-point single precision complex-to-complex forward FFT description
using FFT = decltype(Size<128>() + Precision<float>() +
                     Type<fft_type::c2c>() +
                     Direction<fft_direction::forward>());
```

--------------------------------

### cufftdx::experimental::utils::query_database

Source: https://docs.nvidia.com/cuda/cufftdx/api/other_methods.html

Queries the internal database for the optimal FFT implementation matching the provided configuration.

```APIDOC
## FUNCTION cufftdx::experimental::utils::query_database

### Description
Returns the optimal implementation for a given FFT configuration as a std::optional. If no implementation is found, it returns std::nullopt.

### Parameters
- **fft_size** (unsigned int) - Required - Size of the FFT.
- **dir** (cufftdx::fft_direction) - Required - Direction of the FFT.
- **type** (cufftdx::fft_type) - Required - Type of the FFT.
- **sm** (unsigned int) - Required - Streaming Multiprocessor version.
- **execution** (cufftdx::utils::execution_type) - Required - Execution type.
- **prec** (cufftdx::precision) - Optional - Precision (default: f32).
- **fft_ept** (unsigned int) - Optional - FFTs per thread (default: 0).
- **ffts_per_block** (unsigned int) - Optional - FFTs per block (default: 0).
- **block_dim** (std::tuple<unsigned int, unsigned int, unsigned int>) - Optional - Block dimensions (default: {0, 0, 0}).
- **layout** (cufftdx::complex_layout) - Optional - Complex layout (default: natural).
- **rmode** (cufftdx::real_mode) - Optional - Real mode (default: normal).
- **code_type** (cufftdx::experimental::code_type) - Optional - Code type (default: ptx).

### Response
- **Return Type** (std::optional<cufftdx::utils::frontend_impl_traits>) - The optimal implementation or std::nullopt if not found.
```

--------------------------------

### Implement convolution with nvmath-python

Source: https://docs.nvidia.com/cuda/cufftdx/python_bindings.html

Demonstrates a CUDA kernel using nvmath-python's FFT base classes to perform forward and inverse transforms.

```python
FFT = FFT_base(direction="forward")
IFFT = FFT_base(direction="inverse")

@cuda.jit(link=FFT.files + IFFT.files)
def f(data):
    thread_data = cuda.local.array(shape=(storage_size,), dtype=value_type)

    local_fft_id = cuda.threadIdx.y
    fft_id = cuda.blockIdx.x * ffts_per_block + local_fft_id

    index = cuda.threadIdx.x
    for i in range(elements_per_thread):
        thread_data[i] = data[fft_id, index]
        index += stride

    shared_mem = cuda.shared.array(shape=(0,), dtype=value_type)
    FFT(thread_data, shared_mem)

    for i in range(elements_per_thread):
        thread_data[i] = thread_data[i] / size

    IFFT(thread_data, shared_mem)

    index = cuda.threadIdx.x
    for i in range(elements_per_thread):
        data[fft_id, index] = thread_data[i]
        index += stride
```

--------------------------------

### Compile cuFFTDx with NVRTC via Command Line

Source: https://docs.nvidia.com/cuda/cufftdx/installation.html

Use this command to compile source files requiring cuFFT and nvJitLink dependencies with the necessary experimental API and LTO macros.

```bash
g++ -DCUFFTDX_ENABLE_CUFFT_DEPENDENCY \
    -DCUFFT_ENABLE_EXPERIMENTAL_API \
    -L<cufft_lib_path> -I<cufft_include_path> \
    -I<cufftdx_include_path> \
    -lcufft \
    -lnvJitLink \
    -lnvrtc \
    -L<cuda_lib_path> -lcuda <your_source_file>.cu -o <your_binary>
```

--------------------------------

### Compile LTO Helper with g++

Source: https://docs.nvidia.com/cuda/cufftdx/lto_helper.html

Compile the LTO Helper directly using g++. Ensure include and library paths for CUDA and cuFFT are correctly specified. Requires cuFFT 12.1.0+ and CUDA Toolkit 13.0+.

```bash
g++ -I<cuda_include_path> \
    -I<cufft_include_path> \
    -L<cufft_lib_path> \
    -lcufft cufftdx_cufft_lto_helper.cpp -o cufftdx_cufft_lto_helper

```

--------------------------------

### Compile CUDA with cuFFTDX and cuFFT Dependencies

Source: https://docs.nvidia.com/cuda/cufftdx/cufft_lto/augment_with_lto.html

Use this command to compile CUDA code that utilizes cuFFTDX with cuFFT experimental APIs and the get_database_and_ltoir function. Ensure all include and library paths are correctly set.

```bash
nvcc -DCUDA_INCLUDE_DIR="<cuda_include_path>" \
     -DCUFFTDX_INCLUDE_DIRS="<cufftdx_include_path>" \
     -DCOMMONDX_INCLUDE_DIR="<cufftdx_include_path>" \
     -DCUFFTDX_ENABLE_CUFFT_DEPENDENCY \
     -DCUFFT_ENABLE_EXPERIMENTAL_API \
     -L<cufft_lib_path> -I<cufft_include_path> \
     -I<cufftdx_include_path> \
     -lcufft \
     -lnvJitLink \
     -lnvrtc \
     -lcuda \
     nvrtc_fft_block_lto.cu -o nvrtc_fft_block_lto

```

--------------------------------

### RealFFTOptions Configuration

Source: https://docs.nvidia.com/cuda/cufftdx/api/methods.html

Defines the layout options for Real-to-Complex and Complex-to-Real FFT operations.

```APIDOC
## RealFFTOptions Configuration

### Description
Configures the data layout for R2C and C2R FFTs using the RealFFTOptions operator.

### Parameters
- **complex_layout** (enum) - Optional - Options: `complex_layout::natural` (default), `complex_layout::packed`, `complex_layout::full`.
- **real_mode** (enum) - Optional - Options: `real_mode::normal` (default), `real_mode::folded`.
```

--------------------------------

### cufftdx::experimental::utils::get_all_implementations

Source: https://docs.nvidia.com/cuda/cufftdx/api/other_methods.html

Retrieves all matching FFT implementations from the internal database for a given configuration.

```APIDOC
## FUNCTION cufftdx::experimental::utils::get_all_implementations

### Description
Returns a vector of all matching implementations for the provided configuration. Returns an empty vector if no matches are found.

### Parameters
- **fft_size** (unsigned int) - Required - Size of the FFT.
- **dir** (cufftdx::fft_direction) - Required - Direction of the FFT.
- **type** (cufftdx::fft_type) - Required - Type of the FFT.
- **sm** (unsigned int) - Required - Streaming Multiprocessor version.
- **execution** (cufftdx::utils::execution_type) - Required - Execution type.
- **prec** (cufftdx::precision) - Optional - Precision (default: f32).
- **fft_ept** (unsigned int) - Optional - FFTs per thread (default: 0).
- **ffts_per_block** (unsigned int) - Optional - FFTs per block (default: 0).
- **block_dim** (std::tuple<unsigned int, unsigned int, unsigned int>) - Optional - Block dimensions (default: {0, 0, 0}).
- **layout** (cufftdx::complex_layout) - Optional - Complex layout (default: natural).
- **rmode** (cufftdx::real_mode) - Optional - Real mode (default: normal).
- **code_type** (cufftdx::experimental::code_type) - Optional - Code type (default: ptx).

### Response
- **Return Type** (std::vector<cufftdx::utils::frontend_impl_traits>) - A list of all matching implementations.
```