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_FUNCTION_ABI — COMMONDX_FUNCTION_ABI_NATIVE (the default) or COMMONDX_FUNCTION_ABI_TILE.

• 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).

• CUFFTDX_OPERATOR_TRANSPOSITION — cufftdxTransposition_t (optional, experimental). Affects how much shared memory per FFT the implementation uses; defaults to CUFFTDX_TRANSPOSITION_ANY.

• CUFFTDX_OPERATOR_THREAD_LINEARIZATION — CUFFTDX_THREAD_LINEARIZATION_DISABLED (default) or CUFFTDX_THREAD_LINEARIZATION_ENABLED (optional, experimental, PTX-only).

• CUFFTDX_OPERATOR_THREAD_LAYOUT — cufftdxThreadLayout_t (optional, experimental, block execution only).

Options

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

Heuristics

Use cufftdxGetKnobInt64Size() and cufftdxGetKnobInt64s() (knob type CUFFTDX_KNOB_ELEMENTS_PER_THREAD, CUFFTDX_KNOB_FFTS_PER_BLOCK or CUFFTDX_KNOB_TRANSPOSITION) to get valid EPT/FPB/transposition values. Then set one and call cufftdxIsSupported() to check validity. Implementations sharing an EPT/FPB pair differ only by transposition, hence by shared memory usage.

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

The CUFFTDX_API_LMEM API keeps input and output in thread-local memory. Block execution additionally uses shared memory as scratch.

Execution	Native function ABI	Tile function ABI
Block	void symbol(value_type* rmem, void* scratch_smem)	Same as Native
Thread	void symbol(value_type* rmem)	Same as Native

rmem is an inout operand. scratch_smem points to CUFFTDX_TRAIT_SHARED_MEMORY_SIZE bytes. The Native caller provides this storage; for Tile, the compiler supplies it through the final parameter.

SMEM API

The CUFFTDX_API_SMEM API keeps input and output in shared memory.

Native function ABI	Tile function ABI
void symbol(value_type* data)	Same as Native

data is an inout shared-memory operand. It holds both the operand and the FFT’s scratch, so it must always be allocated as max(CUFFTDX_TRAIT_SHARED_MEMORY_SIZE, 2 * sizeof(precision) * CUFFTDX_OPERATOR_SIZE * CUFFTDX_TRAIT_FFTS_PER_BLOCK) bytes. The second term is a full complex array, which bounds both the input and the output for every type and layout. Which term is the largest depends on the FFT size, the precision, the transposition and the real FFT layout, so compute the maximum rather than assuming one of them.

Supported FFT sizes

libmathdx supports a broad range of non-power-of-two FFT sizes, including prime sizes. All supported sizes use the same API, including sizes that require a workspace in the cuFFTDx CUDA C++ API. libmathdx embeds the Bluestein workspace in the generated code, so callers do not allocate or pass a workspace. These FFTs can be tuned using the knob APIs.

General size support is available for block C2C and normal-mode R2C/C2R FFTs. Thread FFTs and folded real mode support a more restricted set of sizes. Supported sizes also depend on precision and architecture; see the cuFFTDx supported-size tables for the limits, including those for FFTs requiring workspace. Use cufftdxIsSupported to check support for a particular FFT description.

Note

cuFFTDx FP16 FFTs at some non-power-of-two sizes can overflow during intermediate calculations even when the final result fits in FP16. This behavior also applies to the cuFFTDx FFTs generated by libmathdx. Scale inputs or use higher precision if needed.

Device functions

Set CUFFTDX_OPERATOR_DEVICE_FUNCTION with cufftdxSetOperatorInt64() to select CUFFTDX_DEVICE_FUNCTION_EXECUTE, CUFFTDX_DEVICE_FUNCTION_FFT_ID, CUFFTDX_DEVICE_FUNCTION_BATCH_ID, or CUFFTDX_DEVICE_FUNCTION_IS_THREAD_ACTIVE. Combine these flags with bitwise OR to compile multiple functions from one descriptor. Execution is the default when the operator is omitted. Zero and unknown bits are invalid.

Use the descriptor APIs:

1. Select the functions before querying traits or compiling, since those operations lock the descriptor.

2. Query each symbol using cufftdxGetTraitStrSize()/cufftdxGetTraitStr() with CUFFTDX_TRAIT_SYMBOL_EXECUTE_NAME, CUFFTDX_TRAIT_SYMBOL_FFT_ID_NAME, CUFFTDX_TRAIT_SYMBOL_BATCH_ID_NAME, or CUFFTDX_TRAIT_SYMBOL_IS_THREAD_ACTIVE_NAME. Querying a function-specific symbol trait for an unselected function returns COMMONDX_INVALID_VALUE.

3. Compile the selected functions together with cufftdxFinalizeCode(), or compile multiple descriptors together with cufftdxFinalizeCodes(). Batched descriptor finalization is unavailable when the cuFFT LTO EA dependency is enabled.

4. Destroy the descriptor with cufftdxDestroyDescriptor().

CUFFTDX_TRAIT_SYMBOL_NAME returns the execution symbol even when execution is not selected. CUFFTDX_TRAIT_SYMBOL_EXECUTE_NAME requires execution to be selected, like the other function-specific traits. COMMONDX_OPTION_SYMBOL_NAME sets the exact execution symbol and does not affect helper symbols.

For example, after configuring an FFT descriptor named fft and a code object named code:

cufftdxSetOperatorInt64(fft, CUFFTDX_OPERATOR_DEVICE_FUNCTION,
                       CUFFTDX_DEVICE_FUNCTION_EXECUTE | CUFFTDX_DEVICE_FUNCTION_FFT_ID |
                       CUFFTDX_DEVICE_FUNCTION_BATCH_ID | CUFFTDX_DEVICE_FUNCTION_IS_THREAD_ACTIVE);
cufftdxFinalizeCode(code, fft);

See cufftdxDeviceFunctionType_t for the device-function signatures documented in the header.

Helpers follow the descriptor’s thread layout and linearization. Use IS_THREAD_ACTIVE to guard loads/stores; keep FFT execution outside the guard so all block threads participate. See the cuFFTDx Working Group Trait documentation for more information.

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;

    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()));
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h,
                                            CUFFTDX_OPERATOR_DEVICE_FUNCTION,
                                            CUFFTDX_DEVICE_FUNCTION_EXECUTE | CUFFTDX_DEVICE_FUNCTION_FFT_ID |
                                                CUFFTDX_DEVICE_FUNCTION_BATCH_ID |
                                                CUFFTDX_DEVICE_FUNCTION_IS_THREAD_ACTIVE));

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

    // Let cuFFTDx choose a valid configuration for the requested size.
    long long int ept = 0, bpb = 0;
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_ELEMENTS_PER_THREAD, &ept));
    LIBMATHDX_CHECK(cufftdxGetTraitInt64(h, CUFFTDX_TRAIT_FFTS_PER_BLOCK, &bpb));

    printf(
        "Successfully generated LTOIR, %zu bytes for FFT of size %d (%lld elements per thread, %lld 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));

    // Query the execution and indexing helper symbols.
    const auto get_symbol = [h](cufftdxTraitType trait) {
        size_t str_size = 0;
        LIBMATHDX_CHECK(cufftdxGetTraitStrSize(h, trait, &str_size));
        std::string name(str_size, '\0');
        LIBMATHDX_CHECK(cufftdxGetTraitStr(h, trait, str_size, name.data()));
        return name;
    };
    const auto symbol_name = get_symbol(CUFFTDX_TRAIT_SYMBOL_NAME);
    const auto fft_id_name = get_symbol(CUFFTDX_TRAIT_SYMBOL_FFT_ID_NAME);
    const auto batch_id_name = get_symbol(CUFFTDX_TRAIT_SYMBOL_BATCH_ID_NAME);
    const auto is_thread_active_name = get_symbol(CUFFTDX_TRAIT_SYMBOL_IS_THREAD_ACTIVE_NAME);

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

    const auto total_size = size * bpb;

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

    // function from cufftdx descriptor
    extern "C" __device__ void fft_function(float2* rmem, float2* smem);
    extern "C" __device__ void fft_thread_id(unsigned int*);
    extern "C" __device__ void fft_batch_id(unsigned int*);
    extern "C" __device__ void fft_is_thread_active(int*);

    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
        unsigned int fft_id, thread_id;
        int active;
        fft_batch_id(&fft_id);
        fft_thread_id(&thread_id);
        fft_is_thread_active(&active);
        unsigned int batch_offset = fft_id * size;

        // copy input to rmem
        if (active) {
            for (unsigned int index = thread_id, i = 0; index < size; index += stride, ++i) {
                rmem[i] = in[batch_offset + index];
            }
        }
    
        // call JIT compiled FFT device function
        fft_function(rmem, smem);

        // copy output to out
        if (active) {
            for (unsigned int index = thread_id, i = 0; index < size; index += stride, ++i) {
                out[batch_offset + index] = rmem[i];
            }
        }
    }
    )";

    std::string final_kernel = strprintf(kernel_template,
                                         symbol_name.data(),
                                         fft_id_name.data(),
                                         batch_id_name.data(),
                                         is_thread_active_name.data(),
                                         size,
                                         (int)storage_size,
                                         (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"));
    CU_CHECK(cuFuncSetAttribute(
        kernel, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, static_cast<int>(shared_memory_size)));

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

    // Each all-ones FFT has a DC value of size and zero elsewhere, within FP32 roundoff.
    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));

        const float error = std::abs(out[i] - expected);
        if (!std::isfinite(error) || error > 1e-5f * size) {
            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, and for each of them the transpositions it supports
     *
     * Implementations sharing an EPT differ only by transposition, i.e. by how much shared memory
     * they use per FFT: none, `sizeof(precision) * size` bytes (CUFFTDX_TRANSPOSITION_BY_COMPONENT)
     * or twice that (CUFFTDX_TRANSPOSITION_BY_COMPLEX).
     */

    std::array<cufftdxKnobType_t, 2> knobs = { CUFFTDX_KNOB_ELEMENTS_PER_THREAD, CUFFTDX_KNOB_TRANSPOSITION };
    size_t num_values = 0;
    LIBMATHDX_CHECK(cufftdxGetKnobInt64Size(h, knobs.size(), knobs.data(), &num_values));
    std::vector<long long int> values(num_values, 0);
    LIBMATHDX_CHECK(cufftdxGetKnobInt64s(h, knobs.size(), knobs.data(), values.size(), values.data()));

    for (size_t i = 0; i < values.size(); i += knobs.size()) {
        printf("Valid ept: %lld with transposition %s\n",
               values.at(i),
               cufftdxTranspositionToStr(static_cast<cufftdxTransposition>(values.at(i + 1))));
    }

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

    if (values.empty()) {
        printf("No valid EPTs!");
        LIBMATHDX_CHECK(cufftdxDestroyDescriptor(h));
        return 1;
    }
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_ELEMENTS_PER_THREAD, values.at(0)));
    LIBMATHDX_CHECK(cufftdxSetOperatorInt64(h, CUFFTDX_OPERATOR_TRANSPOSITION, values.at(1)));

    /**
     * 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) {
        LIBMATHDX_CHECK(cufftdxDestroyDescriptor(h));
        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, void* scratch_smem). Thread execution: void (value_type* rmem)

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

• scratch_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 cufftdxTransposition_t

Transposition strategy.

Selects, among the implementations of an FFT, the one using a given amount of shared memory per FFT. S below is sizeof(precision) * CUFFTDX_OPERATOR_SIZE.

Check cufftdx transposition documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/operators.html)

Values:

enumerator CUFFTDX_TRANSPOSITION_ANY

Any transposition, i.e. no constraint on the implementation

enumerator CUFFTDX_TRANSPOSITION_NONE

Does not transpose data through shared memory

enumerator CUFFTDX_TRANSPOSITION_BY_COMPONENT

Transposes real and imaginary components separately through shared memory

enumerator CUFFTDX_TRANSPOSITION_BY_COMPLEX

Transposes complex values as units through shared memory

enumerator CUFFTDX_TRANSPOSITION_BY_COMPONENT_PADDED

Transposes components through a padded shared-memory layout to reduce bank conflicts

enum cufftdxThreadLayout_t

CUDA thread layout.

Selects whether the FFT-thread or FFT-batch index is contiguous in the physical CUDA thread order for block execution.

Check cufftdx thread layout documentation for more details (https://docs.nvidia.com/cuda/cufftdx/api/operators.html)

Values:

enumerator CUFFTDX_THREAD_LAYOUT_REGULAR

FFT-thread index is contiguous and the FFT-batch index is strided

enumerator CUFFTDX_THREAD_LAYOUT_TRANSPOSED

FFT-batch index is contiguous and the FFT-thread index is strided

enum cufftdxThreadLinearization_t

Enable or disable experimental thread linearization (PTX-only).

Values:

enumerator CUFFTDX_THREAD_LINEARIZATION_DISABLED

enumerator CUFFTDX_THREAD_LINEARIZATION_ENABLED

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

enumerator CUFFTDX_OPERATOR_FUNCTION_ABI

Operator data type: commondxFunctionAbi_t. Operator definition: optional. Defaults to COMMONDX_FUNCTION_ABI_NATIVE.

enumerator CUFFTDX_OPERATOR_THREAD_LINEARIZATION

Operator data type: cufftdxThreadLinearization_t (experimental, PTX-only). Disabled by default. Operator definition: optional

enumerator CUFFTDX_OPERATOR_TRANSPOSITION

Operator data type: cufftdxTransposition_t (experimental). Operator definition: optional. Defaults to CUFFTDX_TRANSPOSITION_ANY.

enumerator CUFFTDX_OPERATOR_THREAD_LAYOUT

Operator data type: cufftdxThreadLayout_t (experimental). Operator definition: optional. Defaults to CUFFTDX_THREAD_LAYOUT_REGULAR.

enumerator CUFFTDX_OPERATOR_DEVICE_FUNCTION

Operator data type: cufftdxDeviceFunctionType_t. Bitmask selecting one or more device functions generated by this descriptor. Operator definition: optional. Defaults to CUFFTDX_DEVICE_FUNCTION_EXECUTE.

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

enumerator CUFFTDX_KNOB_TRANSPOSITION

Transposition, as a cufftdxTransposition_t. Implementations sharing an elements-per-thread and FFTs-per-block pair differ only by this knob.

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). The shared memory the FFT computation itself requires, which with CUFFTDX_API_SMEM is not on its own the amount to allocate.

enumerator CUFFTDX_TRAIT_FFTS_PER_BLOCK

Trait data type: long long int

enumerator CUFFTDX_TRAIT_SYMBOL_NAME

Trait data type: char*. Execution symbol, even when execution is not selected.

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

enumerator CUFFTDX_TRAIT_SYMBOL_EXECUTE_NAME

Trait data type: char*. Symbol of the selected execute function. Returns COMMONDX_INVALID_VALUE when execution is not selected.

enumerator CUFFTDX_TRAIT_SYMBOL_FFT_ID_NAME

Trait data type: char*. Symbol of the selected fft_id function.

enumerator CUFFTDX_TRAIT_SYMBOL_BATCH_ID_NAME

Trait data type: char*. Symbol of the selected batch_id function.

enumerator CUFFTDX_TRAIT_SYMBOL_IS_THREAD_ACTIVE_NAME

Trait data type: char*. Symbol of the selected is_thread_active function.

enum cufftdxDeviceFunctionType_t

Device functions selected by CUFFTDX_OPERATOR_DEVICE_FUNCTION. Combine flags with bitwise OR before querying traits or compiling; helpers follow the thread layout and linearization. For thread execution, both IDs are zero and IS_THREAD_ACTIVE is one.

Values:

enumerator CUFFTDX_DEVICE_FUNCTION_EXECUTE

Execute the FFT with the descriptor’s API and function ABI. See cufftdxApi.

enumerator CUFFTDX_DEVICE_FUNCTION_FFT_ID

Thread index within an FFT. Device signature: void symbol(unsigned int* out).

enumerator CUFFTDX_DEVICE_FUNCTION_BATCH_ID

Explicit batch index within a block, excluding blockIdx and implicit type batching. Device signature: void symbol(unsigned int* out).

enumerator CUFFTDX_DEVICE_FUNCTION_IS_THREAD_ACTIVE

Thread participation in FFT input/output. Device signature: void symbol(int* out), zero or one. Guard loads/stores with this predicate; all block threads must still call execute.

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.

COMMONDX_OPTION_SYMBOL_NAME sets the exact execution symbol; helper symbols are unaffected.

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 cufftdxFinalizeCodes(

commondxCode code,

size_t count,

const cufftdxDescriptor *handles,

)

Generate code from multiple cuFFTDx descriptors and store it in code.

Parameters:

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

• count – [in] The number of descriptors.

• handles – [in] An array of cuFFTDx descriptors.

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

const char *cufftdxTranspositionToStr(

cufftdxTransposition transposition,

)

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

Parameters:

transposition – [in] The transposition enum to convert

Returns:

The C-string

const char *cufftdxThreadLayoutToStr(

cufftdxThreadLayout thread_layout,

)

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

Parameters:

thread_layout – [in] The thread layout enum to convert

Returns:

The C-string

const char *cufftdxDeviceFunctionTypeToStr(

cufftdxDeviceFunctionType type,

)

Return the function type’s enum name.