How to use libmathdx?#

This guide documents the libmathdx C API for use with the shipped headers, library binary, and the examples. libmathdx JIT-compiles MathDx operations (cuFFTDx, cuBLASDx, cuSolverDx, cuRANDDx) and returns device code (LTOIR or PTX) ready to be linked and executed.

libmathdx C API reference (per-library operators, traits, and examples):

MathDx API documentation (C++ API concepts, operators, and traits; the libmathdx C API mirrors these):

Using libmathdx from CMake#

The installed libmathdx ships a CMake package config under <prefix>/lib/cmake/libmathdx/, so a downstream CMake project can locate it with find_package and link against an imported target:

find_package(libmathdx CONFIG REQUIRED)
target_link_libraries(my_app PRIVATE libmathdx::libmathdx)         # shared
# ...or, for the static archive:
target_link_libraries(my_app PRIVATE libmathdx::libmathdx_static)  # static

Both targets carry the public include directory and the transitive CUDA::nvrtc dependency; the static target additionally re-exposes system link deps (libdl on UNIX) so consumers’ final link step resolves cleanly. If libmathdx is installed in a non-standard prefix, point CMake at it with -Dlibmathdx_ROOT=<prefix> or -DCMAKE_PREFIX_PATH=<prefix> at configure time.

Common workflow (all domains)#

The following steps apply to all libmathdx domains (cuFFTDx, cuBLASDx, cuSolverDx, cuRANDDx), with domain-specific details in the respective API sections. In the steps below, <lib> stands for the library-specific prefix (cufftdx, cublasdx, cusolverdx, or curanddx).

Create descriptor — Call <lib>CreateDescriptor(&handle) for the domain (for example cufftdxCreateDescriptor() or cublasdxCreateDescriptor()).
```
cufftdxDescriptor desc;
cufftdxCreateDescriptor(&desc);
```

Set operators — Use <lib>SetOperatorInt64 and <lib>SetOperatorInt64s to define the problem: sizes, precision, API, execution model, target SM, and related properties.

cufftdxSetOperatorInt64(desc, CUFFTDX_OPERATOR_SIZE, 128);
cufftdxSetOperatorInt64(desc, CUFFTDX_OPERATOR_PRECISION, COMMONDX_PRECISION_F32);
cufftdxSetOperatorInt64(desc, CUFFTDX_OPERATOR_SM, 800);

Set descriptor options — Optionally set the generated symbol name with <lib>SetOptionStr(handle, COMMONDX_OPTION_SYMBOL_NAME, "my_symbol").
```
cufftdxSetOptionStr(desc, COMMONDX_OPTION_SYMBOL_NAME, "my_fft");
```
Query traits — Use <lib>GetTraitInt64, <lib>GetTraitInt64s, <lib>GetTraitStr, and <lib>GetTraitStrSize on the descriptor to get shared memory size, block dimensions, symbol name, and other launch data.
```
int64_t block_dim[3];
cufftdxGetTraitInt64s(desc, CUFFTDX_TRAIT_BLOCK_DIM, block_dim);
```
Generate code and get LTOIR — The code handle creation, code options, finalization, and LTOIR extraction are typically done together:
1. Create a commondxCode handle with commondxCreateCode().
```
commondxCode code;
commondxCreateCode(&code);
```
2. Set code options such as the target SM with commondxSetCodeOptionInt64().
```
commondxSetCodeOptionInt64(code, COMMONDX_OPTION_TARGET_SM, 800);
```
3. Finalize the code with <lib>FinalizeCode(code, handle) for the domain. For the cuBLASDx pipeline path, use cublasdxFinalizeDeviceFunctions() instead.
```
cufftdxFinalizeCode(code, desc);
```
4. Extract the generated LTOIR with commondxGetCodeLTOIRSize() and commondxGetCodeLTOIR(). For multiple LTOIR blobs, use commondxGetCodeNumLTOIRs(), commondxGetCodeLTOIRSizes(), and commondxGetCodeLTOIRs().
```
size_t ltoir_size = 0;
commondxGetCodeLTOIRSize(code, &ltoir_size);
commondxGetCodeLTOIR(code, ltoir_size, ltoir_buffer);
```
Link and run — Link the LTOIR with the application kernel (for example an NVRTC-compiled wrapper) via nvJitLink with input type NVJITLINK_INPUT_LTOIR. Then load it with the CUDA driver and launch the kernel that calls the generated device function.
```
nvJitLinkAddData(linker, NVJITLINK_INPUT_LTOIR, ltoir_buffer, ltoir_size,
                 "libmathdx.ltoir");
```
Destroy — Destroy the code handle and descriptor when they are no longer needed.
```
commondxDestroyCode(code);
cufftdxDestroyDescriptor(desc);
```

Steps 3 and 4 can be done in any order. Operators should not be changed after querying traits or finalizing code.

Pointer APIs#

In this mode, the (LTOIR) device function API is implicitly defined by the API operator such as cublasdxApi_t, cufftdxApi_t and cusolverdxApi_t.

Taking cuBLASDx as an example (See cuBLASDx “pointer API” example), the API works like this:

Create a descriptor using cublasdxCreateDescriptor()
Populate the descriptor with operators using cublasdxSetOperatorInt64() and related APIs.
Query traits using cublasdxGetTraitInt64() and related APIs.
Generate the LTOIR by
1. Creating a code handle with commondxCreateCode()
2. Setting options (e.g. target SM) with commondxSetCodeOptionInt64() and related APIs
3. Generate the code with cublasdxFinalizeCode() and extract it with commondxGetCodeLTOIR() and related APIs.
4. Query properties of the code (e.g. ISA version) with commondxGetCodeOptionInt64()
5. Destroy the code handle with commondxDestroyCode()
Destroy the descriptor with cublasdxDestroyDescriptor()

Step 3 and 4 can be done in any order. However, note that operators cannot be set after step 3 and/or 4.

Opaque Tensors APIs#

In this mode, the (LTOIR) device function API is defined by the input and output tensor specified when creating the device function through cublasdxCreateDeviceFunction()

This is an extension to the previous workflow (See cuBLASDx “tensor API” example).

Create a descriptor using cublasdxCreateDescriptor()
Populate the descriptor with operators using cublasdxSetOperatorInt64() APIs.
Query traits using cublasdxGetTraitInt64() APIs
Create opaque tensors using cublasdxCreateTensor() and cublasdxFinalizeTensors(). cublasdxFinalizeTensors() can be done on a set of tensors.
1. Tensor traits can then be queried using cublasdxGetTensorTraitInt64() APIs. This must be done after calling cublasdxFinalizeTensors().
Define the device function to be generated using cublasdxCreateDeviceFunction(). Device functions are defined as operations (e.g. copy, execution, etc) on a set of input and output tensors. Tensors must be finalized using cublasdxFinalizeTensors() before calling cublasdxCreateDeviceFunction().
1. Device function traits can then be queried using cublasdxGetDeviceFunctionTraitStrSize() and cublasdxGetDeviceFunctionTraitStr() APIs. This must be done after calling cublasdxCreateDeviceFunction().
Generate the device function code
1. Creating a code handle with commondxCreateCode()
2. Setting options (e.g. target SM) with commondxSetCodeOptionInt64() APIs
3. Generate the code with cublasdxFinalizeDeviceFunctions() and extract it with commondxGetCodeLTOIR() APIs
4. Query properties of the code (e.g. ISA version) with commondxGetCodeOptionInt64()
5. Destroy the code handle with commondxDestroyCode()
Destroy the device functions, tensors, and descriptor with their corresponding destroy APIs (cublasdxDestroyDeviceFunction(), cublasdxDestroyTensor(), and cublasdxDestroyDescriptor()).

How to call the device function (API & ABI)?#

For pointer APIs,

Determine the device function symbol name using the *_SYMBOL_NAME (e.g. CUBLASDX_TRAIT_SYMBOL_NAME) trait. The function is extern “C” and can be called as is.
Determine the device function API from the specific Dx lib implementation. Scalars (real and complex) and integers are passed by pointers. Pointers are passed by value.
The function can be forward-declared and called as-is.

For tensor APIs,

Determine the device function API from cublasdxDeviceFunctionType_t. Scalars (real and complex) are passed by pointers, all tensor arguments (see cublasdxTensorType_t) are trivial and are passed by value.
When calling the device function from CUDA C++ or from non-CUDA C++ (e.g. LLVM):
1. Determine the device function symbol name (e.g. cublasdx_execute_8745486) using the CUBLASDX_DEVICE_FUNCTION_TRAIT_SYMBOL trait. The function is extern “C” and can be called as is.
2. Global memory tensor names are struct libmathdx_gmem_tensor { void* ptr; long long int ld; } with ld the leading dimension. Shared and register memory tensors are struct libmathdx_tensor { void* ptr; }.
3. Forward-declare the function using its symbol name and each of the argument name, and call the function when needed.

Non CUDA C++ users must adhere to the calling convention as defined in the Itanium ABI. See also this.

Error handling#

All API return values are of type commondxStatusType. Check for COMMONDX_SUCCESS. On failure (e.g. COMMONDX_COMPILATION_ERROR), use commondxGetLastErrorStrSize() and commondxGetLastErrorStr() to retrieve the last error message.

Common API (libcommondx.h)#

Code handle — commondxCode (opaque). Create with commondxCreateCode(), destroy with commondxDestroyCode().
Code options — commondxSetCodeOptionInt64(), commondxSetCodeOptionInt64s(), commondxSetCodeOptionStr(), commondxSetCodeOptionStrs(). Options include COMMONDX_OPTION_TARGET_SM (e.g. 800), COMMONDX_OPTION_SYMBOL_NAME, COMMONDX_OPTION_CODE_CONTAINER (commondxCodeContainer_t), COMMONDX_OPTION_CODE_ISA (commondxCodeType_t), COMMONDX_OPTION_EXTRA_NVRTC_ARGS, and COMMONDX_OPTION_TARGET_CODE. Query options back with commondxGetCodeOptionInt64() and commondxGetCodeOptionsInt64s().
LTOIR — commondxGetCodeLTOIRSize(), commondxGetCodeLTOIR(). For multiple LTOIRs: commondxGetCodeNumLTOIRs(), commondxGetCodeLTOIRSizes(), commondxGetCodeLTOIRs(). For PTX output: commondxGetCodePTXSize(), commondxGetCodePTX().
Code container — commondxCodeContainer (COMMONDX_CODE_CONTAINER_LTOIR, COMMONDX_CODE_CONTAINER_FATBIN, COMMONDX_CODE_CONTAINER_PTX). commondxCodeType (COMMONDX_CODE_TYPE_LTO, COMMONDX_CODE_TYPE_PTX).
Precision and execution — commondxPrecision (e.g. COMMONDX_PRECISION_F32, COMMONDX_PRECISION_F16, COMMONDX_PRECISION_I8), commondxExecution (COMMONDX_EXECUTION_BLOCK, COMMONDX_EXECUTION_THREAD).
Architecture — Target SM is an integer (e.g. 700, 800, 890, 900). Optional modifier commondxArchModifier_t (COMMONDX_ARCH_MODIFIER_GENERIC, COMMONDX_ARCH_MODIFIER_ARCH_SPECIFIC, COMMONDX_ARCH_MODIFIER_FAMILY_SPECIFIC) can be set via commondxSetCodeOptionInt64s() with (sm, modifier).

Using the generated code (LTOIR)#

Compile the kernel — Use NVRTC or the CUDA compiler with -dlto and --relocatable-device-code=true so the application kernel (or wrapper that calls the generated device function) is compiled to LTOIR.
Link — Link the application LTOIR and the libmathdx LTOIR with nvJitLink (input type NVJITLINK_INPUT_LTOIR) for the target SM to produce cubin or PTX.
Load — Load the linked result with cuModuleLoadDataEx() and obtain the kernel or device function with cuModuleGetFunction().
Launch — Allocate device memory and launch using the block dimensions and shared memory size from the traits of the descriptor.

The shipped examples provide a helper that compiles a string kernel with NVRTC and links with LTOIR via nvJitLink; this pattern or an equivalent can be reused. For the exact API calls and order for each feature, see the full example source on the cuFFTDx, cuBLASDx, cuSolverDx, and cuRANDDx pages.