The ``export_to_c`` interface is provided by the ``JitCompiledFunction`` class. It accepts the following parameters:

* ``file_path``: The path to the directory where the header and object files will be saved.
* ``file_name``: The base name for the header and object files. The same file name will always overwrite existing files.
* ``function_prefix``: The prefix of the function symbol in the generated object file. This should be a unique identifier to avoid symbol conflicts. Users should ensure the function prefix is unique for each exported function. Defaults to the ``file_name``.

It generates the following files:

* ``{file_path}/{file_name}.h``: A C header file containing API function declarations. This header specifies the runtime function signatures in C, mirroring the original Python function interfaces.
* ``{file_path}/{file_name}.o``: A standard object file containing the compiled kernel code. You can link this object file into either a static or shared library. It includes the host entry function, fatbin data, and helper functions such as ``cuda_init`` and ``cuda_load_to_device``. Additionally, it embeds metadata for runtime loading and version verification.

Example:

.. code-block:: python

   import cutlass.cute as cute
   import cutlass.cute.cuda as cuda

   @cute.kernel
   def print_tensor_kernel(a: cute.Tensor):
       cute.printf("a: {}", a)

   @cute.jit
   def print_tensor(a: cute.Tensor, stream: cuda.CUstream):
       print_tensor_kernel(a).launch(grid=(1, 1, 1), block=(1, 1, 1), stream=stream)

   compiled_func = cute.compile(print_tensor)
   # Export compiled functions to object files and headers
   compiled_func.export_to_c(file_path="./artifacts", file_name="print_tensor_example", function_prefix="print_tensor")

Loading in Python

Integrate compiled kernels directly into your C++ executable during the build process. The generated header file supplies the necessary API for loading the module and invoking the function.

Example:

.. code-block:: cpp

   #include "print_tensor_example.h"
   #include <cuda_runtime.h>

   void run_print_tensor() {
       // Prepare tensor, the tensor declaration is in the header file
       print_tensor_Tensor_a_t tensor_a;
       tensor_a.data = nullptr; // GPU memory is set to nullptr.
       // Set dynamic shapes and strides
       tensor_a.dynamic_shapes[0] = 32;
       tensor_a.dynamic_shapes[1] = 16;
       tensor_a.dynamic_strides[0] = 16;
       
       // Create stream
       cudaStream_t stream;
       cudaStreamCreate(&stream);
       
       // Load module before calling the kernel
       print_tensor_Kernel_Module_t module;
       print_tensor_Kernel_Module_Load(&module);
       
       // Call the kernel; the kernel wrapper function is defined in the header file
       cute_dsl_print_tensor_wrapper(&module, &tensor_a, stream);
       
       // Cleanup
       print_tensor_Kernel_Module_Unload(&module);
       cudaStreamDestroy(stream);
   }

The ``print_tensor_example.h`` header file is generated by the ``export_to_c`` interface. It includes:

* The ``print_tensor_Kernel_Module_t`` type: Represents the kernel module.
* The ``print_tensor_Tensor_a_t`` type: A tensor-specific type that defines the ABI for a particular CuTe tensor.
* The ``cute_dsl_print_tensor_wrapper`` function: The user-facing entry point to invoke the kernel.

The compilation of the C++ executable requires the ``libcuda_dialect_runtime.so`` or ``libcuda_dialect_runtime_static.a`` library which is involved in ``<wheel_install_path>/lib``, along with the CUDA driver and runtime libraries, to function properly.


C++ Integration with Dynamic Loading

   // Lookup function
   CuteDSLRT_Function_t *func = nullptr;
   err = CuteDSLRT_Module_Get_Function(&func, module, "print_tensor");
   check_error(err);
   
   // Prepare arguments, matching the argument type defined in the header file
   typedef struct {
       void *data;
       int32_t dynamic_shapes[2];
       int64_t dynamic_strides[1];
   } print_tensor_Tensor_a_t;

   print_tensor_Tensor_a_t tensor_a;
   tensor_a.data = nullptr;
   tensor_a.dynamic_shapes[0] = 32;
   tensor_a.dynamic_shapes[1] = 16;
   tensor_a.dynamic_strides[0] = 16;

   // Create stream
   cudaStream_t stream;
   cudaStreamCreate(&stream);

   // Call the function; the runtime function accepts packed arguments, refer to the wrapper in the header file
   int ret;
   void* args[] = {&tensor_a, &stream, &ret};
   err = CuteDSLRT_Function_Run(func, args, 3);
   check_error(err);
   cudaStreamSynchronize(stream);
   
   // Cleanup
   CuteDSLRT_Module_Destroy(module);
   cudaStreamDestroy(stream);

* **Build host**: the machine where you run Python, ``cute.compile``, and ``export_to_c``.
* **Target**: the machine where the resulting shared library or executable will run.

When the host target is left empty (the default), the ``.o`` is compiled for the build host. When a target is requested, the ``.o`` is an ELF object for that target's triple, and the accompanying ``.h`` uses only fixed-width and opaque types, so the C ABI is identical regardless of any word-size difference between the build host and the target.

Selecting the Host Target

Once the function is compiled with a host target, export it exactly as in the native workflow; no additional arguments are required.

.. code-block:: python

   compiled.export_to_c(file_path="./artifacts", file_name="kernel",
                        function_prefix="kernel")

The resulting ``kernel.o`` is an AArch64 ELF object, and ``kernel.h`` is portable across the build host and the target.

Resolving Runtime Symbols at Link Time

Link the object with your AArch64 cross toolchain (for example ``aarch64-linux-gnu-g++``) and a sysroot that provides the CUDA headers and libraries for the target:

.. code-block:: bash

   aarch64-linux-gnu-g++ -shared -o kernel.so kernel.o \
       $(python -m cutlass.cute.export.aot_config --ldflags --target aarch64-unknown-linux-gnu) \
       $(python -m cutlass.cute.export.aot_config --libs    --target aarch64-unknown-linux-gnu)

Linking succeeds against the stub on the build host. Deploy ``kernel.so`` to the AArch64 target, where it binds against the real runtime library.

Limitations
~~~~~~~~~~~

* **AArch64 only.** Only ``aarch64-unknown-linux-gnu`` is currently supported. Other architectures fail with a "target not registered" error during code generation.
* **Not compatible with TVM FFI.** Combining ``--enable-tvm-ffi`` with ``--host-target`` raises an error; drop ``--enable-tvm-ffi`` and use the plain AOT export path described here.
* **CUDA runtime version must match.** The exported object depends on the CUDA runtime; the target's CUDA runtime/toolkit version must match the one |DSL| was built against (see :ref:`dsl_aot_object_compat`).
* **Linking is your responsibility.** You must supply your own cross toolchain and a target sysroot with the CUDA headers and libraries; the stub only resolves the |DSL| runtime symbols.

Supported Argument Types
------------------------

|DSL| supports the following argument types:

* ``cute.Tensor``
* ``cute.Shape`` / ``cute.Coord`` / ``cute.Tile`` / ``cute.IntTuple`` / ``cute.Stride``
* ``cuda.CUstream``
* ``cutlass.Int8`` / ``cutlass.Int16`` / ``cutlass.Int32`` / ``cutlass.Int64`` / ``cutlass.Boolean``
* ``cutlass.Uint8`` / ``cutlass.Uint16`` / ``cutlass.Uint32`` / ``cutlass.Uint64``
* ``cutlass.Float32`` / ``cutlass.TFloat32`` / ``cutlass.Float64`` / ``cutlass.Float16``

Note that:

1. ``cute.Tensor`` is a dynamic tensor type that only contains dynamic shapes and strides in its ABI representation. As a result, different compilations may produce different tensor ABIs. This is why declarations for each tensor type are included in the generated header file.
2. ``strides`` in ``cute.Tensor`` are determined by the ``use_32bit_strides`` compile argument. When ``use_32bit_strides`` is set to ``True``, the strides are 32-bit; when set to ``False``, they are 64-bit.
3. Currently, custom types are not supported for AOT compilation.

.. _dsl_aot_object_compat:

Object File Compatibility Issues
--------------------------------

The object file generated by |DSL| depends on the CUDA runtime library. Therefore, ensure that the version of the CUDA runtime/toolkit library matches the version used by |DSL|. Otherwise, ABI compatibility with the CUDA runtime cannot be guaranteed.

When using C++ static linking integration, compatibility is assured because the header and object files are generated together and guaranteed to match.

For C++ dynamic loading integration and Python loading, the binary file is loaded at runtime. To ensure compatibility, version information is embedded in the metadata of the generated binary file. At runtime, this version information is checked, and if it does not match the expected version, the binary file will be rejected.

Relation to Apache TVM FFI AOT
------------------------------

Apache TVM FFI AOT offers a comparable capability, enabling TVM functions to be compiled into binary files that can be loaded and executed at runtime.
For more information, see the section "Exporting Compiled Module" in :doc:`tvm_ffi_compilation`.

The primary distinction is that, when TVM FFI is enabled, |DSL| generates a dedicated wrapper function on top of the underlying CuTe ABI. This wrapper adheres to the calling conventions defined by TVM FFI.
In contrast, the CuTe ABI entry function is specified directly in the generated header file, which affects how arguments must be provided.

For instance, with the TVM FFI wrapper function, users are able to pass in arguments such as ``torch.Tensor`` directly. However, when calling the CuTe ABI entry function, arguments should be provided as ``cute.Tensor`` types.