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 examplecufftdxCreateDescriptor()orcublasdxCreateDescriptor()).cufftdxDescriptor desc; cufftdxCreateDescriptor(&desc);
Set operators — Use
<lib>SetOperatorInt64and<lib>SetOperatorInt64sto 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>GetTraitStrSizeon 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:
Create a
commondxCodehandle withcommondxCreateCode().commondxCode code; commondxCreateCode(&code);
Set code options such as the target SM with
commondxSetCodeOptionInt64().commondxSetCodeOptionInt64(code, COMMONDX_OPTION_TARGET_SM, 800);
Finalize the code with
<lib>FinalizeCode(code, handle)for the domain. For the cuBLASDx pipeline path, usecublasdxFinalizeDeviceFunctions()instead.cufftdxFinalizeCode(code, desc);
Extract the generated LTOIR with
commondxGetCodeLTOIRSize()andcommondxGetCodeLTOIR(). For multiple LTOIR blobs, usecommondxGetCodeNumLTOIRs(),commondxGetCodeLTOIRSizes(), andcommondxGetCodeLTOIRs().size_t ltoir_size = 0; commondxGetCodeLTOIRSize(code, <oir_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
Creating a code handle with
commondxCreateCode()Setting options (e.g. target SM) with
commondxSetCodeOptionInt64()and related APIsGenerate the code with
cublasdxFinalizeCode()and extract it withcommondxGetCodeLTOIR()and related APIs.Query properties of the code (e.g. ISA version) with
commondxGetCodeOptionInt64()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()APIsCreate opaque tensors using
cublasdxCreateTensor()andcublasdxFinalizeTensors().cublasdxFinalizeTensors()can be done on a set of tensors.Tensor traits can then be queried using
cublasdxGetTensorTraitInt64()APIs. This must be done after callingcublasdxFinalizeTensors().
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 usingcublasdxFinalizeTensors()before callingcublasdxCreateDeviceFunction().Device function traits can then be queried using
cublasdxGetDeviceFunctionTraitStrSize()andcublasdxGetDeviceFunctionTraitStr()APIs. This must be done after callingcublasdxCreateDeviceFunction().
Generate the device function code
Creating a code handle with
commondxCreateCode()Setting options (e.g. target SM) with
commondxSetCodeOptionInt64()APIsGenerate the code with
cublasdxFinalizeDeviceFunctions()and extract it withcommondxGetCodeLTOIR()APIsQuery properties of the code (e.g. ISA version) with
commondxGetCodeOptionInt64()Destroy the code handle with
commondxDestroyCode()
Destroy the device functions, tensors, and descriptor with their corresponding destroy APIs (
cublasdxDestroyDeviceFunction(),cublasdxDestroyTensor(), andcublasdxDestroyDescriptor()).
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 (seecublasdxTensorType_t) are trivial and are passed by value.When calling the device function from CUDA C++ or from non-CUDA C++ (e.g. LLVM):
Determine the device function symbol name (e.g. cublasdx_execute_8745486) using the
CUBLASDX_DEVICE_FUNCTION_TRAIT_SYMBOLtrait. The function is extern “C” and can be called as is.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; }.
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 withcommondxCreateCode(), destroy withcommondxDestroyCode().Code options —
commondxSetCodeOptionInt64(),commondxSetCodeOptionInt64s(),commondxSetCodeOptionStr(),commondxSetCodeOptionStrs(). Options includeCOMMONDX_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, andCOMMONDX_OPTION_TARGET_CODE. Query options back withcommondxGetCodeOptionInt64()andcommondxGetCodeOptionsInt64s().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 viacommondxSetCodeOptionInt64s()with (sm, modifier).
Using the generated code (LTOIR)#
Compile the kernel — Use NVRTC or the CUDA compiler with
-dltoand--relocatable-device-code=trueso 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 withcuModuleGetFunction().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.