cuFFTDx

API reference: cuFFTDx C++ documentation.

Operators

Operators are set with cufftdxSetOperatorInt64() or cufftdxSetOperatorInt64s():

• CUFFTDX_OPERATOR_SIZE — FFT size (int64).

• CUFFTDX_OPERATOR_API — CUFFTDX_API_LMEM (register + shared scratch) or CUFFTDX_API_SMEM (shared only).

• CUFFTDX_OPERATOR_EXECUTION — COMMONDX_EXECUTION_BLOCK or COMMONDX_EXECUTION_THREAD.

• CUFFTDX_OPERATOR_TYPE — CUFFTDX_TYPE_C2C, CUFFTDX_TYPE_R2C, CUFFTDX_TYPE_C2R.

• CUFFTDX_OPERATOR_DIRECTION — CUFFTDX_DIRECTION_FORWARD, CUFFTDX_DIRECTION_INVERSE.

• CUFFTDX_OPERATOR_PRECISION — e.g. COMMONDX_PRECISION_F32.

• CUFFTDX_OPERATOR_SM — target SM (int64).

• CUFFTDX_OPERATOR_ELEMENTS_PER_THREAD — EPT (block execution).

• CUFFTDX_OPERATOR_FFTS_PER_BLOCK — FFTs per block (block execution).

• CUFFTDX_OPERATOR_BLOCK_DIM — optional block dim (3 int64s).

• CUFFTDX_OPERATOR_REAL_FFT_OPTIONS — 2 int64s: cufftdxComplexLayout_t followed by cufftdxRealMode_t (optional, for R2C/C2R).

• CUFFTDX_OPERATOR_CODE_TYPE — cufftdxCodeType_t (optional).

Options

Use cufftdxSetOptionStr() to set COMMONDX_OPTION_SYMBOL_NAME for the device function name.

Heuristics

Use cufftdxGetKnobInt64Size() and cufftdxGetKnobInt64s() (knob type CUFFTDX_KNOB_ELEMENTS_PER_THREAD or CUFFTDX_KNOB_FFTS_PER_BLOCK) to get valid EPT/FPB values. Then set one and call cufftdxIsSupported() to check validity.

Traits

Query after the descriptor is configured:

• cufftdxGetTraitInt64() — CUFFTDX_TRAIT_SHARED_MEMORY_SIZE, CUFFTDX_TRAIT_STORAGE_SIZE, CUFFTDX_TRAIT_ELEMENTS_PER_THREAD, CUFFTDX_TRAIT_STRIDE, CUFFTDX_TRAIT_FFTS_PER_BLOCK, CUFFTDX_TRAIT_SUGGESTED_FFTS_PER_BLOCK, CUFFTDX_TRAIT_IMPLICIT_TYPE_BATCHING, CUFFTDX_TRAIT_INPUT_LENGTH, CUFFTDX_TRAIT_OUTPUT_LENGTH, CUFFTDX_TRAIT_INPUT_ELEMENTS_PER_THREAD, CUFFTDX_TRAIT_OUTPUT_ELEMENTS_PER_THREAD.

• cufftdxGetTraitInt64s() — CUFFTDX_TRAIT_BLOCK_DIM.

• cufftdxGetTraitCommondxDataType() — CUFFTDX_TRAIT_VALUE_TYPE, CUFFTDX_TRAIT_INPUT_TYPE, CUFFTDX_TRAIT_OUTPUT_TYPE.

• cufftdxGetTraitStrSize() + cufftdxGetTraitStr() — CUFFTDX_TRAIT_SYMBOL_NAME.

Device function signatures

LMEM API:

• Block execution: void symbol(value_type* rmem, value_type* smem)

• Thread execution: void symbol(value_type* rmem)

Allocate CUFFTDX_TRAIT_SHARED_MEMORY_SIZE bytes for the smem argument.

SMEM API:

• void symbol(value_type* smem)

Examples

cuFFTDx example

/*
 * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

#include <cuda.h>
#include <cuda_runtime.h>
#include <libcufftdx.h>
#include <nvJitLink.h>
#include <nvrtc.h>

#include <array>
#include <cassert>
#include <cmath>
#include <complex>
#include <cstdlib>
#include <string>
#include <vector>

#include "common_examples.hpp"
#include "macros.hpp"

using namespace examples;

int main() {

    int size = 32;
    int ept = 4;
    int bpb = 2;

    arch_t dx_sm = get_dx_sm();
    arch_t target_sm = get_target_sm();

    /**
     * Create a descriptor
     * This is equivalent to `using FFT = ...` in cuFFTDx C++
     */

    cufftdxDescriptor h { 0 };
    LIBMATHDX_CHECK(cufftdxCreateDescriptor(&h));
    // CUFFTDX_API_LMEM means the function will be of signature:
    //     void(value_type*, value_type*)
    //   with the first argument being local memory ("registers"), with each thread holding "EPT" elements
    //   and the second being a pointer to a shared memory scratch buffer
    // CUFFTDX_API_SMEM would mean that the function will be of signature:
    //     void(value_type*)
    //   and takes a shared memory pointer with all the elements laid out in natural order
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_API, CUFFTDX_API_LMEM));
    // COMMONDX_EXECUTION_BLOCK means multiple threads in a block participate in the FFT
    // COMMONDX_EXECUTION_THREAD would mean that each thread computes a single FFT
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_EXECUTION, COMMONDX_EXECUTION_BLOCK));
    // This is the FFT size
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_SIZE, size));
    // CUFFTDX_TYPE_C2C means complex-to-complex FFT type
    // CUFFTDX_TYPE_R2C means real-to-complex FFT type
    // CUFFTDX_TYPE_C2R means complex-to-real FFT type
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_TYPE, CUFFTDX_TYPE_C2C));
    // CUFFTDX_DIRECTION_FORWARD means a forward FFT
    // CUFFTDX_DIRECTION_INVERSE means an inverse FFT
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_DIRECTION, CUFFTDX_DIRECTION_FORWARD));
    // COMMONDX_PRECISION_F16 for half precision
    // COMMONDX_PRECISION_F32 for single precision
    // COMMONDX_PRECISION_F64 for double precision
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_PRECISION, COMMONDX_PRECISION_F32));
    // Compute capability to target
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_SM, dx_sm.operator_sm()));
    // Number of elements per thread (only for BLOCK execution)
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_ELEMENTS_PER_THREAD, ept));
    // Number of ffts per block (only for BLOCK execution)
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_FFTS_PER_BLOCK, bpb));

    // COMMONDX_OPTION_SYMBOL_NAME indicates the required name for the device function.
    LIBMATHDX_CHECK(cufftdxSetOptionStr(h, COMMONDX_OPTION_SYMBOL_NAME, "my_fft"));

    /**
     * Compile the device function
     */

    commondxCode code;
    LIBMATHDX_CHECK(commondxCreateCode(&code));
    // Specify arch to compile to
    LIBMATHDX_CHECK(commondxSetCodeOptionInt64(code, COMMONDX_OPTION_TARGET_SM, target_sm.operator_sm()));
    LIBMATHDX_CHECK(cufftdxFinalizeCode(code, h));
    size_t lto_size = 0;
    LIBMATHDX_CHECK(commondxGetCodeLTOIRSize(code, &lto_size));
    std::vector<char> lto(lto_size);
    LIBMATHDX_CHECK(commondxGetCodeLTOIR(code, lto.size(), lto.data()));
    LIBMATHDX_CHECK(commondxDestroyCode(code));

    printf("Successfully generated LTOIR, %zu bytes for FFT of size %d (%d elements per thread, %d FFTs per block)\n",
           lto_size,
           size,
           ept,
           bpb);

    /**
     * Query traits
     */

    // How much shared memory to allocate?
    long long int shared_memory_size = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_SHARED_MEMORY_SIZE, &shared_memory_size));

    // What size of blocks to launch?
    std::array<long long int, 3> block_dim = { 0, 0, 0 };
    LIBMATHDX_CHECK(cufftdxGetTraitInt64s(h, CUFFTDX_TRAIT_BLOCK_DIM, block_dim.size(), block_dim.data()));

    // How much local memory to allocate?
    long long int storage_size = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_STORAGE_SIZE, &storage_size));

    // Stride between elements?
    long long int stride = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_STRIDE, &stride));

    // What is the name of the function?
    size_t str_size = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitStrSize(h, CUFFTDX_TRAIT_SYMBOL_NAME, &str_size));
    std::vector<char> symbol_name(str_size);
    LIBMATHDX_CHECK(cufftdxGetTraitStr(h, CUFFTDX_TRAIT_SYMBOL_NAME, str_size, symbol_name.data()));

    printf(
        "Function %s requires %lld B of shared memory %lld # of local memory elements and a block_dim of %lld %lld %lld\n",
        symbol_name.data(),
        shared_memory_size,
        storage_size,
        block_dim[0],
        block_dim[1],
        block_dim[2]);

    LIBMATHDX_CHECK(cufftdxDestroyDescriptor(h));

    int total_size = size * bpb;

    const char kernel_template[] = R"(
    // traits from cufftdx descriptor
    #define fft_function %s
    constexpr unsigned size = %d;
    constexpr unsigned storage_size = %d;
    constexpr unsigned ept = %d;
    constexpr unsigned stride = %d;

    // function from cufftdx descriptor
    extern "C" __device__ void fft_function(float2* rmem, float2* smem);

    extern "C" __global__ void simple_fft_kernel(float2* in, float2* out) {
        extern __shared__ float2 smem[];
        float2 rmem[storage_size]= {};

        // indexing for per thread FFT data copy
        int fft_id = threadIdx.y;
        int thread_id = threadIdx.x;
        int base_index = fft_id * size + thread_id;

        // copy input to rmem
        for (int i = 0; i < ept; i++) {
            rmem[i] = in[base_index + i * stride];
        }
    
        // call JIT compiled FFT device function
        fft_function(rmem, smem);

        // copy output to out
        for (int i = 0; i < ept; i++) {
            out[base_index + i * stride] = rmem[i];
        }
    }
    )";

    std::string final_kernel =
        strprintf(kernel_template, symbol_name.data(), size, (int)storage_size, (int)ept, (int)stride);
    std::vector<char> cubin = compile_and_link(final_kernel, lto, target_sm);

    CUmodule module {};
    CUfunction kernel {};
    CUDA_CHECK(cudaSetDevice(0));
    CU_CHECK(cuModuleLoadDataEx(&module, cubin.data(), 0, 0, 0));
    CU_CHECK(cuModuleGetFunction(&kernel, module, "simple_fft_kernel"));

    std::complex<float>*in = nullptr, *out = nullptr;
    CUDA_CHECK(cudaMallocManaged(&in, total_size * sizeof(std::complex<float>)));
    CUDA_CHECK(cudaMallocManaged(&out, total_size * sizeof(std::complex<float>)));
    std::fill(in, in + total_size, std::complex<float>(1.0f, 0.0f));
    std::fill(out, out + total_size, std::complex<float>(1.0f, 0.0f));

    std::vector<void*> kernel_args = { reinterpret_cast<void*>(&in), reinterpret_cast<void*>(&out) };
    CU_CHECK(cuLaunchKernel(kernel,
                            1,
                            1,
                            1,
                            static_cast<unsigned int>(block_dim[0]),
                            static_cast<unsigned int>(block_dim[1]),
                            static_cast<unsigned int>(block_dim[2]),
                            static_cast<unsigned int>(shared_memory_size),
                            0,
                            kernel_args.data(),
                            nullptr));
    CUDA_CHECK(cudaDeviceSynchronize());

    // with an input of 1.0f, the output of the FFT should be the sum of inputs and equaling the FFT size
    // we are performing a batch of size 2 so we should have 4 values (real, imag) equalling 32 with remaining values
    // being 0
    for (int i = 0; i < total_size; i++) {
        std::complex<float> expected =
            ((i % size == 0) ? std::complex<float>(static_cast<float>(size), 0.0f) : std::complex<float>(0.0f, 0.0f));

        if (std::abs(out[i] - expected) > 1e-7) {
            printf("Error: out[%d] = %f, %f, expected %f, %f\n",
                   i,
                   out[i].real(),
                   out[i].imag(),
                   expected.real(),
                   expected.imag());
            abort();
        }
    }
    printf("Successfully ran the kernel\n");

    CU_CHECK(cuModuleUnload(module));
    CUDA_CHECK(cudaFree(in));
    CUDA_CHECK(cudaFree(out));

    return 0;
}

cuFFTDx heuristics example

/*
 * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

#include <libcufftdx.h>

#include <array>
#include <vector>

#include "common_examples.hpp"
#include "macros.hpp"

using namespace examples;

int main() {

    long long int size = 2048;
    long long int arch = 890;

    /**
     * Create a descriptor
     * This is equivalent to `using FFT = ...` in cuFFTDx C++
     */

    cufftdxDescriptor h { 0 };
    LIBMATHDX_CHECK(cufftdxCreateDescriptor(&h));
    // CUFFTDX_API_LMEM means the function will be of signature:
    //     void(value_type*, value_type*)
    //   with the first argument being local memory ("registers"), with each thread holding "EPT" elements
    //   and the second being a pointer to a shared memory scratch buffer
    // CUFFTDX_API_SMEM would mean that the function will be of signature:
    //     void(value_type*)
    //   and takes a shared memory pointer with all the elements laid out in natural order
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_API, cufftdxApi::CUFFTDX_API_LMEM));
    // COMMONDX_EXECUTION_BLOCK means multiple threads in a block participate in the FFT
    // COMMONDX_EXECUTION_THREAD would mean that each thread computes a single FFT
    LIBMATHDX_CHECK(
        cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_EXECUTION, commondxExecution::COMMONDX_EXECUTION_BLOCK));
    // This is the FFT size
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_SIZE, size));
    // CUFFTDX_TYPE_C2C means complex-to-complex FFT type
    // CUFFTDX_TYPE_R2C means real-to-complex FFT type
    // CUFFTDX_TYPE_C2R means complex-to-real FFT type
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_TYPE, cufftdxType::CUFFTDX_TYPE_C2C));
    // CUFFTDX_DIRECTION_FORWARD means a forward FFT
    // CUFFTDX_DIRECTION_INVERSE means an inverse FFT
    LIBMATHDX_CHECK(
        cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_DIRECTION, cufftdxDirection::CUFFTDX_DIRECTION_FORWARD));
    // COMMONDX_PRECISION_F16 for half precision
    // COMMONDX_PRECISION_F32 for single precision
    // COMMONDX_PRECISION_F64 for double precision
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_PRECISION, commondxPrecision::COMMONDX_PRECISION_F32));
    // Compute capability to target
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_SM, arch));

    // COMMONDX_OPTION_SYMBOL_NAME indicates the required name for the device function.
    LIBMATHDX_CHECK(cufftdxSetOptionStr(h, commondxOption::COMMONDX_OPTION_SYMBOL_NAME, "my_fft"));

    /**
     * Find valid EPTs
     */

    cufftdxKnobType_t knobs = CUFFTDX_KNOB_ELEMENTS_PER_THREAD;
    size_t num_epts = 0;
    LIBMATHDX_CHECK(cufftdxGetKnobInt64Size(h, 1, &knobs, &num_epts));
    std::vector<long long int> epts(num_epts, 0);
    LIBMATHDX_CHECK(cufftdxGetKnobInt64s(h, 1, &knobs, epts.size(), epts.data()));

    for (auto ept : epts) {
        printf("Valid ept: %lld\n", ept);
    }

    /**
     * Pick of one of them directly
     */

    if (epts.empty()) {
        printf("No valid EPTs!");
        return 1;
    }
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_ELEMENTS_PER_THREAD, epts.front()));

    /**
     * Confirms that this descriptor is valid and will compile
     */

    int valid = -1;
    LIBMATHDX_CHECK(cufftdxIsSupported(h, &valid));
    printf("Descriptor is %s\n", valid ? "valid" : "invalid");
    if (!valid) {
        return 1;
    }

    /**
     * Compile the device function
     */

    commondxCode code;
    LIBMATHDX_CHECK(commondxCreateCode(&code));
    // Specify arch to compile to
    LIBMATHDX_CHECK(commondxSetCodeOptionInt64(code, COMMONDX_OPTION_TARGET_SM, 800ll));
    LIBMATHDX_CHECK(cufftdxFinalizeCode(code, h));
    size_t lto_size = 0;
    LIBMATHDX_CHECK(commondxGetCodeLTOIRSize(code, &lto_size));
    std::vector<char> lto(lto_size);
    LIBMATHDX_CHECK(commondxGetCodeLTOIR(code, lto.size(), lto.data()));
    LIBMATHDX_CHECK(commondxDestroyCode(code));

    // `lto` is an LTOIR container with NVVM/LTO inside. It contains a device function called `my_fft` with signature
    // void my_fft(void* rmem, void* smem)
    // which can be called to compute the FFT described above
    // smem must point to `CUFFTDX_TRAIT_SHARED_MEMORY_SIZE` bytes of shared memory
    // rmem is the usual register input
    //
    // In order to create a valid kernel, `lto` must be linked to the user kernel (compiled, e.g., with NVRTC) using
    // nvJitLink.

    printf("Successfully generated LTOIR, %zu bytes for FFT of size %lld\n", lto_size, size);

    /**
     * Query shared memory size, suggested elements per thread and suggested ffts per block
     */

    // How much shared memory to allocate?
    long long int shared_memory_size = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_SHARED_MEMORY_SIZE, &shared_memory_size));

    // What EPT is suggested?
    long long int ept = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_ELEMENTS_PER_THREAD, &ept));

    // How many FFTs per Block?
    long long int fpb = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_FFTS_PER_BLOCK, &fpb));

    // How many FFTs per Block is *suggested*?
    long long int sfpb = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_SUGGESTED_FFTS_PER_BLOCK, &sfpb));

    // What's the corresponding block dim (for "non suggested" FFTs per block)?
    std::array<long long int, 3> block_dim = { 0 };
    LIBMATHDX_CHECK(cufftdxGetTraitInt64s(h, CUFFTDX_TRAIT_BLOCK_DIM, block_dim.size(), block_dim.data()));

    printf("Function requires %lld B of shared memory with EPT %lld and FPB %lld / %lld, block_dim is %lld %lld %lld\n",
           shared_memory_size,
           ept,
           fpb,
           sfpb,
           block_dim.at(0),
           block_dim.at(1),
           block_dim.at(2));

    LIBMATHDX_CHECK(cufftdxDestroyDescriptor(h));
}

API reference

typedef long long int cufftdxDescriptor

A cuFFTDx descriptor.

Equivalent to using FFT = ... in cuFFTDx CUDA C++.

enum cufftdxApi_t

Type of cufftDx API.

Handling problems with input in register or in shared memory buffers.

Check cufftdx::execute method documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/methods.html#block-execute-method)

Values:

enumerator CUFFTDX_API_LMEM

Input-output is in thread-local memory. Shared memory is used as scratch for block execution. Function API is defined by its signature. Block execution: void (value_type* rmem, char* smem). Thread execution: void (value_type* rmem)

• rmem is a pointer to an array of value_type values. See CUFFTDX_TRAIT_VALUE_TYPE

• smem is a pointer to CUFFTDX_TRAIT_SHARED_MEMORY_SIZE bytes in shared memory.

The function is extern "C" and the symbol name can be queried using CUFFTDX_TRAIT_SYMBOL_NAME

enumerator CUFFTDX_API_SMEM

Input-output is in shared memory Function API is defined by its signature: void (value_type* smem)

• smem is a pointer to a shared memory array of value_type values. See CUFFTDX_TRAIT_VALUE_TYPE.

The function is extern "C" and the symbol name can be queried using CUFFTDX_TRAIT_SYMBOL_NAME

enum cufftdxType_t

Type of computation data.

Check cufftdx::Type operator documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/methods.html#block-execute-method)

Values:

enumerator CUFFTDX_TYPE_C2C

Complex-to-complex FFT

enumerator CUFFTDX_TYPE_R2C

Real-to-complex FFT

enumerator CUFFTDX_TYPE_C2R

Complex-to-real FFT

enum cufftdxDirection_t

FFT direction.

Check cufftdx::Direction operator documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/operators.html#direction-operator)

Values:

enumerator CUFFTDX_DIRECTION_FORWARD

Forward FFT

enumerator CUFFTDX_DIRECTION_INVERSE

Inverse FFT

enum cufftdxComplexLayout_t

Complex data layout.

Check cufftdx layout documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/methods.html#complex-element-layouts)

Values:

enumerator CUFFTDX_COMPLEX_LAYOUT_NATURAL

Natural complex layout (N/2+1 complex values)

enumerator CUFFTDX_COMPLEX_LAYOUT_PACKED

Packed complex layout (N/2 complex values)

enumerator CUFFTDX_COMPLEX_LAYOUT_FULL

Full complex layout (N complex values)

enum cufftdxRealMode_t

Real data mode.

Check cufftdx real data mode documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/methods.html#real-element-layouts)

Values:

enumerator CUFFTDX_REAL_MODE_NORMAL

Normal real mode (N real values)

enumerator CUFFTDX_REAL_MODE_FOLDED

Fold real mode (N/2-1 complex values)

enum cufftdxCodeType_t

Code type.

Check cufftdx code type documentation for more details (https://docs.nvidia.com/cuda/1.4.0-ea/cufftdx/api/methods.html#code-type)

Values:

enumerator CUFFTDX_CODE_TYPE_PTX

inlined-PTX implementation

enumerator CUFFTDX_CODE_TYPE_LTOIR

LTOIR implementation

enum cufftdxOperatorType_t

cufftDx operators

The set of supported cufftDx operators.

Check cufftDx description operators documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/operators.html#description-operators)

Check cufftDx execution operators documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/operators.html#execution-operators)

Values:

enumerator CUFFTDX_OPERATOR_SIZE

Operator data type: long long int. Expected content: >= 1. Operator definition: required

enumerator CUFFTDX_OPERATOR_DIRECTION

Operator data type: cufftdxDirection_t. Operator definition: required

enumerator CUFFTDX_OPERATOR_TYPE

Operator data type: cufftdxType_t. Operator definition: optional

enumerator CUFFTDX_OPERATOR_PRECISION

Operator data type: commondxPrecision_t. Operator definition: required

enumerator CUFFTDX_OPERATOR_SM

Operator data type: long long int. Expected content: 700 (Volta), 800 (Ampere), …. Operator definition: required

enumerator CUFFTDX_OPERATOR_EXECUTION

Operator data type: commondxExecution_t. Operator definition: required

enumerator CUFFTDX_OPERATOR_FFTS_PER_BLOCK

Operator data type: long long int. Expected content: >= 0. Operator definition: optional

enumerator CUFFTDX_OPERATOR_ELEMENTS_PER_THREAD

Operator data type: long long int. Expected content: >= 0. Operator definition: optional

enumerator CUFFTDX_OPERATOR_BLOCK_DIM

Operator data type: long long int * 3. Expected content: <x, y, z> block dimensions. Operator definition: optional

enumerator CUFFTDX_OPERATOR_REAL_FFT_OPTIONS

Operator data type: cufftdxComplexLayout_t followed by cufftdxRealMode_t. Operator definition: optional

enumerator CUFFTDX_OPERATOR_API

Operator data type: cufftdxApi_t. Operator definition: required

enumerator CUFFTDX_OPERATOR_CODE_TYPE

Operator data type: cufftdxCodeType_t. Operator definition: optional

enum cufftdxKnobType_t

cufftDx configuration knobs

The set of supported knobs used for accessing cufftDx operator’s performance configuration details.

Check cufftDx Execution Traits documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/traits.html#execution-traits)

Values:

enumerator CUFFTDX_KNOB_ELEMENTS_PER_THREAD

Elements per thread

enumerator CUFFTDX_KNOB_FFTS_PER_BLOCK

FFTs per block

enum cufftdxTraitType_t

cufftDx traits

The set of supported types of traits that can be accessed from finalized sources that use cufftDx.

Check cufftDx Execution Thread Traits documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/traits.html#thread-traits)

Check cufftDx Execution Block Traits documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/traits.html#block-traits)

Values:

enumerator CUFFTDX_TRAIT_VALUE_TYPE

Trait data type: commondxValueType_t. Expected content: complex type of the underlying data used to compute the FFT

enumerator CUFFTDX_TRAIT_INPUT_TYPE

Trait data type: commondxValueType_t. Expected content: type of the underlying data used as input for the FFT

enumerator CUFFTDX_TRAIT_OUTPUT_TYPE

Trait data type: commondxValueType_t. Expected content: type of the underlying data used as output for the FFT

enumerator CUFFTDX_TRAIT_IMPLICIT_TYPE_BATCHING

Trait data type: long long int

enumerator CUFFTDX_TRAIT_ELEMENTS_PER_THREAD

Trait data type: long long int. Expected content: >= 0

enumerator CUFFTDX_TRAIT_STORAGE_SIZE

Trait data type: long long int. Expected content: >= 0, in compute type elements

enumerator CUFFTDX_TRAIT_STRIDE

Trait data type: long long int. Expected content: >= 0

enumerator CUFFTDX_TRAIT_BLOCK_DIM

Trait data type: long long int * 3. Expected content: <x, y, z> block dimensions

enumerator CUFFTDX_TRAIT_SHARED_MEMORY_SIZE

Trait data type: long long int (SMEM size in bytes)

enumerator CUFFTDX_TRAIT_FFTS_PER_BLOCK

Trait data type: long long int

enumerator CUFFTDX_TRAIT_SYMBOL_NAME

Trait data type: char*

enumerator CUFFTDX_TRAIT_INPUT_LENGTH

Trait data type: long long int, in input type elements

enumerator CUFFTDX_TRAIT_OUTPUT_LENGTH

Trait data type: long long int, in output type elements

enumerator CUFFTDX_TRAIT_INPUT_ELEMENTS_PER_THREAD

Trait data type: long long int

enumerator CUFFTDX_TRAIT_OUTPUT_ELEMENTS_PER_THREAD

Trait data type: long long int

enumerator CUFFTDX_TRAIT_SUGGESTED_FFTS_PER_BLOCK

Trait data type: long long int

commondxStatusType cufftdxGetVersion(

int *major,

int *minor,

int *patch

)

Returns the major.minor.patch version of cuFFTDx.

Parameters:

• major – [out] The major version

• minor – [out] The minor version

• patch – [out] The patch version

Returns:

COMMONDX_SUCCESS

commondxStatusType cufftdxCreateDescriptor(cufftdxDescriptor *handle)

Creates a cuFFTDx descriptor.

Parameters:

handle – [inout] A pointer to a cuFFTDx descriptor handle, and as output, a valid initialized descriptor.

Returns:

COMMONDX_SUCCESS on success, or an error code

commondxStatusType cufftdxSetOptionStr(

cufftdxDescriptor handle,

commondxOption opt,

const char *value

)

Set a C-string option on a cuFFTDx descriptor.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• opt – [in] The option to set

• value – [in] The C-string to set the option to

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxSetOptionStrs(

cufftdxDescriptor handle,

commondxOption opt,

size_t count,

const char **values

)

Sets one or more C-string option on a cuFFTDx descriptor.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor .

• opt – [in] The option to set.

• count – [in] The number of options.

• values – [in] An array of count C-string values.

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxIsSupported(

cufftdxDescriptor handle,

int *value

)

Check if a given descriptor has an underlying implementation.

Returns a non-zero integer in value if handle has an underlying implementation.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• value – [out] The number of distinct sets of knobs.

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetKnobInt64Size(

cufftdxDescriptor handle,

size_t num_knobs,

cufftdxKnobType *knobs_ptr,

size_t *size

)

Returns the number of knobs for a set of knobs.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• num_knobs – [in] The number of knobs

• knobs_ptr – [in] An array of num_knobs knobs

• size – [out] The number of distinct sets of knobs.

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetKnobInt64s(

cufftdxDescriptor handle,

size_t num_knobs,

cufftdxKnobType *knobs_ptr,

size_t size,

long long int *values

)

Returns the knobs values for a set of knobs.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• num_knobs – [in] The number of knobs

• knobs_ptr – [in] A pointer to an array of num_knobs knobs

• size – [in] The number of knobs.

• values – [out] The knob values. Must be a pointer to an array of at least size knobs values (integer)

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxSetOperatorInt64(

cufftdxDescriptor handle,

cufftdxOperatorType op,

long long int value

)

Set an integer operator to a cuFFTDx descriptor.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• op – [in] The operator to set the descriptor to

• value – [in] The value to set the operator to

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxSetOperatorInt64s(

cufftdxDescriptor handle,

cufftdxOperatorType op,

size_t count,

const long long int *array

)

Set an array operator to a cuFFTDx descriptor.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• op – [in] The operator to set the descriptor to

• count – [in] The array size

• array – [in] A pointer to at least count integers, the arrat to set the descriptor to.

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetLTOIRSize(

cufftdxDescriptor handle,

size_t *lto_size

)

Get the LTOIR’s size from a cuFFTDx descriptor.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• lto_size – [out] The size (in bytes) of the LTOIR

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetLTOIR(

cufftdxDescriptor handle,

size_t size,

void *lto

)

Get the LTOIR from a cuFFTDx descriptor.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• size – [in] The LTOIR size, in bytes

• lto – [out] The LTOIR code.

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetTraitStrSize(

cufftdxDescriptor handle,

cufftdxTraitType trait,

size_t *size

)

Returns a C-string trait’s value size.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• trait – [in] The trait to query the descriptor for

• size – [out] The C-string length (including \0)

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetTraitStr(

cufftdxDescriptor handle,

cufftdxTraitType trait,

size_t size,

char *value

)

Returns a C-string trait value.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• trait – [in] The trait to query the descriptor for

• size – [in] The C-string size (including the \0)

• value – [out] As output, the C-string trait value

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetTraitInt64(

cufftdxDescriptor handle,

cufftdxTraitType trait,

long long int *value

)

Returns an integer trait.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• trait – [in] The trait to query the descriptor for

• value – [out] The trait integer value.

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetTraitInt64s(

cufftdxDescriptor handle,

cufftdxTraitType trait,

size_t count,

long long int *array

)

Returns an array of integers trait.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• trait – [in] The trait to query the descriptor for

• count – [in] The array size

• array – [out] The trait array.

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxGetTraitCommondxDataType(

cufftdxDescriptor handle,

cufftdxTraitType trait,

commondxValueType *value

)

Return a commondxValueType trait.

Parameters:

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

• trait – [in] The trait to query the descriptor for, of value commondxValueType

• value – [out] As output, the valuetype for the given input trait.

Returns:

COMMONDX_SUCCESS on success, or an error code.

const char *cufftdxOperatorTypeToStr(cufftdxOperatorType op)

Convert a cufftdxOperatorType instance to a human readable C-string.

Parameters:

op – [in] A cufftdxOperatorType instance

Returns:

A C-string

const char *cufftdxTraitTypeToStr(cufftdxTraitType op)

Convert a cufftdxTraitType instance to a human readable C-string.

Parameters:

op – [in] A cufftdxTraitType instance

Returns:

A C-string

commondxStatusType cufftdxFinalizeCode(

commondxCode code,

cufftdxDescriptor handle

)

Generate code from the cuFFTDx descriptor and stores it in code.

Parameters:

• code – [out] A commondxCode instance, output of commondxCreateCode.

• handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

Returns:

COMMONDX_SUCCESS on success, or an error code.

commondxStatusType cufftdxDestroyDescriptor(cufftdxDescriptor handle)

Destroys a cuFFTDx descriptor.

Parameters:

handle – [in] A cuFFTDx descriptor, output of cufftdxCreateDescriptor

Returns:

COMMONDX_SUCCESS on success, or an error code.

const char *cufftdxApiToStr(cufftdxApi api)

Convert an API enum to a human readable C-string.

Parameters:

api – [in] The API enum to convert

Returns:

The C-string

const char *cufftdxTypeToStr(cufftdxType type)

Convert a type enum to a human readable C-string.

Parameters:

type – [in] The type enum to convert

Returns:

The C-string

const char *cufftdxDirectionToStr(cufftdxDirection direction)

Convert a direction enum to a human readable C-string.

Parameters:

direction – [in] The direction enum to convert

Returns:

The C-string

const char *cufftdxComplexLayoutToStr(cufftdxComplexLayout layout)

Convert a complex layout enum to a human readable C-string.

Parameters:

layout – [in] The complex layout enum to convert

Returns:

The C-string

const char *cufftdxRealModeToStr(cufftdxRealMode mode)

Convert a real mode enum to a human readable C-string.

Parameters:

mode – [in] The real mode enum to convert

Returns:

The C-string

const char *cufftdxKnobTypeToStr(cufftdxKnobType knob)

Convert a knob type enum to a human readable C-string.

Parameters:

knob – [in] The knob type enum to convert

Returns:

The C-string

const char *cufftdxCodeTypeToStr(cufftdxCodeType code_type)

Convert a code type enum to a human readable C-string.

Parameters:

code_type – [in] The code type enum to convert

Returns:

The C-string