cuFFTDx

https://docs.nvidia.com/cuda/cufftdx/
Admin
cuFFTDx is a header-only C++ library that enables developers to implement Fast Fourier Transform...
Tokens:35,771
Snippets:409
Trust Score:10
Update:2 months ago
Show doc for...
Tokens:
Raw
### Build examples with optimization

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

Build command for provided examples using the separate twiddles optimization.

```bash
make all_twiddles
```

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

### Run cufftdx_cufft_lto_helper

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

Execution syntax and example usage for the helper tool.

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

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

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

### Compile and link with LTO support

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

Example command for compiling a cuFFTDX project with the necessary dependencies and macros.

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

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

### Example usage of is_supported

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

Demonstrates checking support for various FFT configurations on specific 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
```

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

### FFT workspace creation and kernel execution

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

Example demonstrating the creation of a workspace object and its usage in a CUDA kernel.

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

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

### Implementing a Custom FFT Kernel

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

A complete example demonstrating the definition of an FFT, workspace allocation, and kernel invocation using cuFFTDx.

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

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

### Define FFT configurations in CSV

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

Example CSV file structure defining multiple FFT configurations for the LTO helper.

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

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

### Dynamic Batching Shared Memory Calculation Example

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

Demonstrates calculating shared memory requirements using both the FFT description and raw trait values.

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

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

### Retrieve FFT Description Traits

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

Example of using helper traits to inspect an FFT descriptor constructed with operators.

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

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

### Launch FFT Kernel with cuFFTDx

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

Implements a CUDA kernel for block-level FFT execution and demonstrates the host-side setup, including workspace allocation and kernel invocation.

```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;
  auto workspace = make_workspace<FFT>(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);
  cudaDeviceSynchronize();

  cudaFree(data);
}
```

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

### Retrieve LTO Database and LTOIRs

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

Example usage of cufftdx::utils::get_database_and_ltoir to obtain the database string and LTOIRs for a specific FFT configuration.

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

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

### Constructing an FFT Descriptor

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

Example of combining multiple operators to define a forward FFT operation with 8 double-precision 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() );
```

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

### cufftDescriptionGetTraitInt64

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

Gets a trait value from an FFT description.

```APIDOC
## cufftDescriptionGetTraitInt64

### Description
Gets a trait value from an FFT description.

### Parameters
- **desc_handle** (cufftDescriptionHandle) - In - cufftDescriptionHandle returned by cufftDescriptionCreate.
- **trait** (cufftDescriptionTrait) - In - The trait identifier to query.
- **value** (long long int*) - Out - Pointer to store the trait value.

### Return Values
- **CUFFT_SUCCESS** - Trait value was successfully retrieved.
- **CUFFT_INVALID_VALUE** - Invalid handle or trait.
```

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

### Shared Memory Alignment Examples

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

Demonstrates various ways to ensure shared memory pointers are aligned to 128 bits (16 bytes) as required by the library.

```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
extern __shared__ __align__(16)
  value_type shared_mem[];                                 // 2
extern __shared__ __align__(alignof(float4))
  value_type shared_mem[];                                 // 3
extern __shared__ float4 shared_mem[];                     // 4

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

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

### Constructing an FFT Descriptor

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

Example of defining an FFT descriptor type using various operators for a forward FFT operation.

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

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

### Block-level FFT Execution

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

Example of defining an FFT descriptor and executing it within a CUDA kernel using shared memory.

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

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

### Get real FFT layout

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

Retrieves the input and output data layout for real-to-complex or complex-to-real FFT types.

```cpp
cufftdx::real_fft_layout_of<FFT>::value
cufftdx::real_fft_layout_of_v<FFT>
```

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

### Display ptxas warning message

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

Example of the ptxas warning message that may appear during compilation.

```text
ptxas warning : Program uses 32-bit address on line 'XXX'
which is conflicting with .address_size 64
```

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

### Create Device Handle

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

Initializes a device handle from the previously created description handle.

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

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

### Get code type

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

Retrieves the type of code used for FFT compute device functions, defaulting to PTX.

```cpp
cufftdx::experimental::code_type_of<Description>::value
cufftdx::experimental::code_type_of_v<Description>
```

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

### Configuring FFT Kernel with Dynamic Shared Memory

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

Demonstrates defining an FFT configuration using expression templates and setting the required dynamic shared memory size for the kernel launch.

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

  (...)

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

  (...)
}

void introduction_example() {
  // 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>());

  (...)

  // Increases the max dynamic shared memory size to match FFT
  // requirements
  cudaFuncSetAttribute(block_fft_kernel<FFT>,
      cudaFuncAttributeMaxDynamicSharedMemorySize,
      FFT::shared_memory_size);

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

  (...)
}
```

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

### Get real FFT mode

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

Retrieves the optimization mode for the FFT kernel execution as defined by the RealFFTOptions operator.

```cpp
cufftdx::real_fft_mode_of<FFT>::value
cufftdx::real_fft_mode_of_v<FFT>
```

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

### Build with online kernel generation

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

Compile applications requiring cuFFT and nvJitLink dependencies with necessary macros defined.

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

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

### Build cufftdx_cufft_lto_helper

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

Commands to build the helper tool using direct g++ compilation or CMake.

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

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

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

### Frontend to Backend Traits Conversion

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

Demonstrates the structure and usage of the frontend_to_backend function for converting FFT descriptions into backend implementation traits.

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

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

### Host Link Executable

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

Link the host and device object files to create the final executable.

```bash
g++ -L<cuda_lib64_path> -lcudart \
    introduction_lto_dlink.o \
    introduction_lto_example.o \
    -o introduction_lto_example
```

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

### Compile Database Creation Executable

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

Command to compile the helper source file into an 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
```

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

### Get Database and LTOIR Function Signature

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

The function signature for retrieving the LTO database and LTOIRs for a specified FFT operation.

```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 */);
```

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

### Execute the LTO Helper

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

Command to run the LTO Helper to generate the database files for specific CUDA architectures.

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

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

### cufftDescriptionCreate

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

Creates a new FFT description handle in memory owned and provided by the user.

```APIDOC
## cufftDescriptionCreate

### Description
Creates a new FFT description handle in memory owned and provided by the user.

### Parameters
- **desc_handle** (cufftDescriptionHandle*) - In/Out - Pointer to a cufftDescriptionHandle object. Contains a cufftDescriptionHandle handle value upon success.

### Return Values
- **CUFFT_SUCCESS** - Description handle was successfully created.
- **CUFFT_INVALID_VALUE** - The handle pointer is NULL.
- **CUFFT_INTERNAL_ERROR** - The internal handle factory reaches its maximum capacity.
```

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

### Retrieve FFT Traits

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

Demonstrates how to define an FFT descriptor and retrieve traits like value type and memory requirements.

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

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

### Get Shared Memory Size for Dynamic Batching (FFT Description)

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

Computes shared memory requirements using an FFT description template parameter.

```cpp
template<class FFT>
constexpr unsigned int
get_shared_memory_size_for_dynamic_batching(
    const unsigned int ffts_per_block);
```

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

### cufftDeviceCreate

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

Creates a new device handle with the specified FFT descriptions. It searches for matching device functions and building blocks for each provided description.

```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 as well as the building blocks needed to generate the device functions.

### Parameters
- **handle** (cufftDeviceHandle*) - In/Out - Pointer to a cufftDeviceHandle object.
- **desc_count** (size_t) - In - Number of descriptions.
- **desc_handles** (cufftDescriptionHandle*) - In - Pointer to an array of description handles with length desc_count.

### Return Values
- **CUFFT_SUCCESS** - Device handle was successfully created.
- **CUFFT_INVALID_VALUE** - The handle pointer or the pointer to the description handles pointer is NULL, or any of the description handles used in the array is invalid.
- **CUFFT_INTERNAL_ERROR** - The internal handle factory reaches its maximum capacity.
```

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

### Write Database Header to Disk

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

Retrieves the database string from the device handle and prepares it for disk storage.

```cpp
size_t database_str_size = 0;
cufftDeviceGetDatabaseStrSize(device_handle, &database_str_size);
char* database_str = (char*)malloc(database_str_size * sizeof(char));
cufftDeviceGetDatabaseStr(device_handle, database_str_size, database_str);

// Write database_str to disk as "lto_database.hpp.inc" and free all memory
// associated with database_str
```

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

### Get Shared Memory Size for Dynamic Batching (Runtime Values)

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

Computes shared memory requirements using explicit trait values instead of an FFT description.

```cpp
constexpr unsigned int
get_shared_memory_size_for_dynamic_batching(
    const unsigned int shared_memory_size_per_fft,
    const unsigned int ffts_per_block,
    const unsigned int implicit_type_batching);
```

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

### run_cufft_lto_helper()

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

Runs the LTO helper to generate LTO artifacts and creates an object library for linking.

```APIDOC
## run_cufft_lto_helper()

### Description
Runs the helper to generate LTO artifacts and creates an object library. This function also automatically triggers the build of the LTO Helper tool.

### Parameters
- **SRC_DIR** (string) - Required - Directory containing the helper source files.
- **OUTPUT_DIR** (string) - Required - Directory where the LTO database will be generated.
- **OUTPUT_NAME** (string) - Required - Base name for the generated library (creates ${OUTPUT_NAME}_lto_lib).
- **DESCS** (string) - Required - Path to the input CSV file with FFT configurations.
- **ARCHITECTURES** (string) - Required - Target CUDA architectures (comma-separated, e.g., 75,80).
```

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

### Device link with LTOIRs

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

Perform device linking using the -dlto flag to combine LTOIR and fatbin files.

```bash
nvcc -dlink \
        -dlto \
        --generate-code arch=compute_75,code=sm_75 \
        database_X.fatbin ... \
        database_X.ltoir ... \
        <your_source_file>.o -o <your_source_file>_dlink.o
```

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

### Configure Dynamic Shared Memory for Large FFTs

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

Demonstrates how to set the maximum dynamic shared memory size and launch a kernel for large FFTs that exceed 48 KB.

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

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

__global__ void block_fft_kernel(FFT::value_type* data) {
  // dynamic shared memory
  extern __shared__ __align__(alignof(float4)) FFT::value_type shared_mem[];

  (...)
}

void example() {
  (...)

  // Increases the max dynamic shared memory size to match FFT
  // requirements
  cudaFuncSetAttribute(block_fft_kernel,
      cudaFuncAttributeMaxDynamicSharedMemorySize,
      FFT::shared_memory_size)
  // Invokes kernel with FFT::block_dim threads in CUDA block
  block_fft_kernel<FFT><<<1, FFT::block_dim,
                          FFT::shared_memory_size>>>(data, workspace);

  (...)
}
```

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

### Online cuFFT LTO Database Integration

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

Configure CMake for online kernel generation by defining necessary macros and linking required dependencies.

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

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

### Check Device Handle Support

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

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

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

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

### Define a Basic FFT Description

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

Initializes a 128-point single precision complex-to-complex forward FFT using cuFFTDx operators.

```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 implementation of a given FFT configuration.

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

### Description
Returns the optimal implementation for a specified FFT configuration as a `std::optional<cufftdx::utils::frontend_impl_traits>`. Returns `std::nullopt` if no implementation is found.

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

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

### Compile with binary size optimization

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

Compilation command using separate twiddles to optimize binary size for projects with multiple translation units.

```bash
nvcc -DCUFFTDX_USE_SEPARATE_TWIDDLES -std=c++17 -arch sm_XY \
     (...) -rdc=true -I<mathdx_include_dir> \
     <cufftdx_include_dir>/liblut/lut.cu \
     <your_source_file>.cu <mathdx_include_dir> \
     -o <your_binary>
```

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

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

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

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

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

### Description
Returns a `std::vector<cufftdx::utils::frontend_impl_traits>` containing all matching implementations for the provided configuration. Returns an empty vector if no implementations are found.

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

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

### Convert Frontend to Backend Traits

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

Maps FFT operation traits to effective backend traits for LTO database inclusion.

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

auto backend_traits = cufftdx::utils::frontend_to_backend(
      cufftdx::utils::algorithm::ct,
      cufftdx::utils::execution_type::block,
      128 /* fft_size */,
      cufftdx::fft_type::c2c,
      cufftdx::fft_direction::forward,
      750 /* sm */,
      cufftdx::real_mode::normal,
      8 /* elements_per_thread */,
      0 /* block_dim_x */,
      cufftdx::experimental::code_type::ltoir
);
```

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

### cufftdx::utils::frontend_to_backend

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

Converts a frontend FFT description to effective backend traits used to query cuFFTDx databases.

```APIDOC
## cufftdx::utils::frontend_to_backend

### Description
Converts a frontend FFT description to effective backend traits. This is used to query the cuFFTDx databases and can be used to generate LTO databases containing device function code and metadata.

### Signature
`backend_impl_traits frontend_to_backend(algorithm algo, execution_type exec_type, unsigned int fft_size, fft_type type, fft_direction direction, unsigned int sm, real_mode real_mode, unsigned int elements_per_thread, unsigned int block_dim_x, experimental::code_type code_type)`

### Parameters
- **algo** (algorithm) - The FFT algorithm (ct or bluestein).
- **exec_type** (execution_type) - The execution type (thread or block).
- **fft_size** (unsigned int) - The size of the FFT.
- **type** (fft_type) - The type of the FFT.
- **direction** (fft_direction) - The direction of the FFT.
- **sm** (unsigned int) - The SM version.
- **real_mode** (real_mode) - The real mode configuration.
- **elements_per_thread** (unsigned int) - Elements per thread or 0 if not set.
- **block_dim_x** (unsigned int) - Block dimension X or 0 if not set.
- **code_type** (experimental::code_type) - The experimental code type.

### Returns
- **backend_impl_traits** (struct) - A structure containing size, type, direction, sm, elements_per_thread, and min_elements_per_thread.
```

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

### Launch FFT Kernel with cuFFTDx

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

Defines a device-side FFT kernel and a host-side function to configure and launch it using specified FFT parameters.

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

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

### Generate LTO Database and Link Library in CMake

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

Use run_cufft_lto_helper to generate artifacts and link the resulting library to an executable. Ensure the include directory is set to the generated artifacts path.

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

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

### Query all matching FFT implementations

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

Returns a vector of all matching implementation traits, or an empty vector if none 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);
```

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

### Set Traits for Description Handle

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

Configures a description handle with specific backend traits using cufftDescriptionSetTraitInt64.

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

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

### Execute FFT in a CUDA Kernel

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

Demonstrates the structure of a kernel using an FFT descriptor handle to perform the operation.

```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? */);
}
```

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

### Host link object files

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

Finalize the executable by linking the object files with the CUDA runtime library.

```bash
g++ -L${CUDA_PATH}/lib64 -lcudart \
       <your_source_file>.o <your_source_file>_dlink.o -o <your_binary>
```

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

### Query optimal FFT implementation

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

Returns the optimal implementation as an optional trait object, or std::nullopt if no match 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);
```

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

### build_cufft_lto_helper()

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

Builds the LTO Helper tool. This is typically called automatically by run_cufft_lto_helper().

```APIDOC
## build_cufft_lto_helper()

### Description
Builds the LTO Helper tool required for generating LTO databases.

### Parameters
- **SRC_DIR** (string) - Required - Directory containing the helper source files.
- **OUTPUT_DIR** (string) - Required - Directory where the helper executable will be built.
```

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

### Retrieve LTO Database and LTOIRs

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

Use the get_database_and_ltoir function to obtain pointers to the C++ header and LTOIRs required for online kernel generation.

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

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

### Compile the LTO Helper

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

Command to compile the LTO Helper utility using g++.

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

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

### Compile cuFFTDx Program

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

Command-line instruction for compiling a cuFFTDx application using nvcc.

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