compiled_add = cute.compile(add, a_torch, b_torch, options="--enable-tvm-ffi")

Note that CuTe's convention is to write the stride order for dimensions from left to right,
where a lower order number means higher priority. In the context of the ``make_fake_compact_tensor`` API,
for shape ``(2, 3, 4)`` and stride order ``(0, 1, 2)``, the stride is ``(1, 2, 6)``.
This is commonly known as column-major order. If you want to create a fake tensor with compact row-major order,
you should explicitly pass in ``stride_order=tuple(reversed(range(len(shape))))``
to ``make_fake_compact_tensor``. Alternatively, you can always precisely control the
stride via the ``stride`` argument in the ``make_fake_tensor`` API.


``cute.Tensor`` adapter for TVM FFI
-----------------------------------

To adapt the ``cute.Tensor`` to the TVM FFI function, you can use the ``cute.runtime.from_dlpack`` function with the
``enable_tvm_ffi=True`` option or the environment variable ``CUTE_DSL_ENABLE_TVM_FFI=1``. For example:

.. code-block:: python

   def example_from_dlpack():
      a_cute = cute.runtime.from_dlpack(a_torch, enable_tvm_ffi=True).mark_layout_dynamic()
      b_cute = cute.runtime.from_dlpack(b_torch, enable_tvm_ffi=True).mark_layout_dynamic()

      compiled_add_one(a_cute, b_cute)

Note that because the ``cute.runtime.from_dlpack`` function performs an explicit DLPack conversion, it is less efficient than passing the ``torch.Tensor`` directly.
You can also use ``cute.Tensor`` as an argument hint for ``cute.compile``.

.. code-block:: python

   compiled_add_one = cute.compile(add_one, a_cute, b_cute, options="--enable-tvm-ffi")


Working with torch Tensors
--------------------------

As you may have noticed in the examples above, TVM FFI-compiled functions can
directly accept ``torch.Tensor`` objects (and other DLPack-compatible tensors) as inputs.
The resulting functions add minimal overhead, enabling faster eager invocations
thanks to the optimized calling path.

Working with Streams
--------------------

In many cases, a CuTe kernel needs to run on a specific CUDA stream.
|DSL| provides two ways to work with streams through TVM FFI.
The first is to pass the stream explicitly as an argument.
The following example demonstrates this approach; the function accepts ``torch.cuda.Stream``,
``CUstream`` or any stream class that implements the CUDA stream protocol.

.. code-block:: python

   import cutlass.cute as cute
   import torch
   from cuda.bindings.driver import CUstream

   @cute.kernel
   def device_add_one(a: cute.Tensor, b: cute.Tensor):
      threads_per_block = 128
      cta_x_, _, _ = cute.arch.block_idx()
      tid_x, _, _ = cute.arch.thread_idx()
      tid = cta_x_ * threads_per_block + tid_x
      if tid < a.shape[0]:
         b[tid] = a[tid] + 1.0

   @cute.jit
   def add_one_with_stream(a: cute.Tensor, b: cute.Tensor, stream: CUstream):
      n = a.shape[0]
      threads_per_block = 128
      blocks = (n + threads_per_block - 1) // threads_per_block
      device_add_one(a, b).launch(
         grid=(blocks, 1, 1),
         block=(threads_per_block, 1, 1),
         stream=stream,
      )

   def example_add_one_with_stream():
      n = cute.sym_int()
      a_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
      b_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
      # Fake stream is a placeholder for stream argument
      stream = cute.runtime.make_fake_stream()
      compiled_add_one = cute.compile(
         add_one_with_stream, a_cute, b_cute, stream, options="--enable-tvm-ffi"
      )
      a_torch = torch.arange(10, dtype=torch.float32, device="cuda")
      b_torch = torch.empty(10, dtype=torch.float32, device="cuda")
      torch_stream = torch.cuda.current_stream()
      compiled_add_one(a_torch, b_torch, torch_stream)
      torch_stream.synchronize()
      print("result of b_torch after compiled_add_one(a_torch, b_torch, torch_stream)")
      print(b_torch)

Using Environment Stream

Named tuples are also supported and help logically group related arguments together.
The example below shows how to use named tuples as arguments. Under the hood, named tuples
are passed as unnamed tuples at the ABI level. When errors occur, the function signature in
error messages will display unnamed tuple arguments.
Ensure that the compile-time CuTe named tuple type definition has the same fields
as the runtime PyTorch named tuple.
Currently, users need to explicitly unpack the named tuple outside of conditionals and then
use the unpacked variables inside the conditionals.

.. code-block:: python

   from typing import NamedTuple
   from cutlass import cute
   import torch

   class CuteNamedTuple(NamedTuple):
      a: cute.Tensor
      b: cute.Tensor
      c: cute.Float32 = cute.Float32(1)

      def __new_from_mlir_values__(self, values):
         return CuteNamedTuple(*values)

   class TorchNamedTuple(NamedTuple):
      a: torch.Tensor
      b: torch.Tensor
      c: float = 1

   @cute.kernel
   def device_add_one_named_tuple(value: CuteNamedTuple):
      tid = cute.arch.block_idx()[0] * 128 + cute.arch.thread_idx()[0]
      # need to unpack namedtuple outside conditionals
      a = value.a
      b = value.b
      c = value.c
      if tid < a.shape[0]:
         b[tid] = a[tid] + c

   @cute.jit
   def add_one_with_named_tuple(value: CuteNamedTuple):
      n = value.a.shape[0]
      threads_per_block = 128
      blocks = (n + threads_per_block - 1) // threads_per_block
      device_add_one_named_tuple(value).launch(grid=(blocks, 1, 1), block=(threads_per_block, 1, 1))

   def example_add_one_with_named_tuple():
      n = cute.sym_int()
      a_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
      b_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))

      compiled_add_one = cute.compile(
         add_one_with_named_tuple, CuteNamedTuple(a=a_cute, b=b_cute),
         options="--enable-tvm-ffi"
      )
      a_torch = torch.arange(10, dtype=torch.float32, device="cuda")
      b_torch = torch.empty(10, dtype=torch.float32, device="cuda")
      compiled_add_one(TorchNamedTuple(a=a_torch, b=b_torch))
      print("result of b_torch")
      print(b_torch)

   example_add_one_with_named_tuple()


Supported types
---------------

The TVM FFI function supports the following |DSL|-specific types as arguments:

- ``cute.Tensor``
- ``cutlass.Boolean``, ``cutlass.Int8``, ``cutlass.Int16``, ``cutlass.Int32``, ``cutlass.Int64``, ``cutlass.Uint8``, ``cutlass.Uint16``, ``cutlass.Uint32``, ``cutlass.Uint64``, ``cutlass.Float32``,  ``cutlass.Float64``
- ``cute.Shape``, ``cute.Stride``, ``cute.Coord``, ``cute.Tile``, ``cute.IntTuple``

.. list-table::
   :header-rows: 1
   :widths: 30 70

   * - Compile-time type
     - Call-time type
   * - ``cute.Pointer``
     - ``ctypes.c_void_p`` or a class that implements ``__tvm_ffi_opaque_ptr__`` protocol.
   * - ``cute.runtime.FakeTensor``
     - ``torch.Tensor`` and other DLPack-compatible tensors.
   * - Scalar types (e.g. ``cutlass.Boolean``, ``cutlass.Int32``)
     - Python scalars (e.g. True, 123).
   * - CuTe algebra types (e.g. ``cute.Shape``, ``cute.Stride``)
     - ``tvm_ffi.Shape`` or python tuple of ints.
   * - CUDA stream ``cuda.CUstream``
     - A stream class that implements the CUDA stream protocol (e.g. ``torch.cuda.Stream``, ``cuda.CUstream``).
   * - Tuple of types (e.g. ``Tuple[cute.Tensor, cute.Tensor, cutlass.Int32]``)
     - Python tuple of corresponding call-time types.

Error handling
--------------

TVM FFI functions will enable validation of arguments to make sure they match the expected type
and value constraints declared by the user. These checks are compiled into the function, run very fast,
and have no observable overhead during function invocation. Each of those errors will translate
into a proper Python exception that can be caught and handled. The example below shows some
example error cases that can be checked:

.. code-block:: python

   def example_constraint_checks():
      n = cute.sym_int(divisibility=16)
      # assume align to 16 bytes (4 int32), both should share same shape variable n
      a_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,), assumed_align=16)
      b_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,), assumed_align=16)
      compiled_add_one = cute.compile(add_one, a_cute, b_cute, options="--enable-tvm-ffi")
      a = torch.zeros(128, dtype=torch.float32, device="cuda")
      b = torch.zeros(128, dtype=torch.float32, device="cuda")

      try:
         # raises type mismatch error because we expect a and b to be float32
         compiled_add_one(a, 1)
      except TypeError as e:
         # Mismatched type on argument #1 when calling:
         # `add_one(a: Tensor([n0], float32), b: Tensor([n0], float32))`,
         # expected Tensor
         print(f"TypeError: {e}")

      try:
         # raises shape mismatch error because we expect both a and b have shap [n]
         compiled_add_one(a, b[:126])
      except ValueError as e:
         # Mismatched b.shape[0] on argument #1 when calling:
         # `add_one(a: Tensor([n0], float32), b: Tensor([n0], float32))`,
         # expected to match a.shape[0]
         print(f"ValueError: {e}")

      try:
         # triggers divisibility mismatch error because 126 is not divisible by 16
         compiled_add_one(a[:126], b[:126])
      except ValueError as e:
         # Invalid a.shape[0] on argument #0 when calling:
         # `add_one(a: Tensor([n0], float32), b: Tensor([n0], float32)`,
         # expected to be divisible by 16
         print(f"ValueError: {e}")

      try:
         a = torch.zeros(129, dtype=torch.float32, device="cuda")
         b = torch.zeros(129, dtype=torch.float32, device="cuda")
         # triggers data alignment mismatch error because x and y are not aligned to 16 bytes
         compiled_add_one(a[1:], b[1:])
      except ValueError as e:
         # raises: Misaligned Tensor data on argument #0 when calling:
         # `add_one(a: Tensor([n0], float32), b: Tensor([n0], float32)`,
         # expected data alignment=16 bytes
         print(f"ValueError: {e}")

Any CUDA errors encountered will also be automatically converted into Python exceptions by the TVM FFI function.

.. code-block:: python

   @cute.jit
   def add_one_invalid_launch(a: cute.Tensor, b: cute.Tensor):
      # Intentionally exceed the maximum block dimension (1024 threads) so the
      # CUDA runtime reports an invalid configuration error.
      device_add_one(a, b).launch(grid=(1, 1, 1), block=(4096, 1, 1))

   def example_error_cuda_error():
      a_torch = torch.zeros((10,), dtype=torch.float32, device="cuda")
      b_torch = torch.zeros((10,), dtype=torch.float32, device="cuda")

      a_cute = cute.runtime.from_dlpack(a_torch, enable_tvm_ffi=True)
      b_cute = cute.runtime.from_dlpack(b_torch, enable_tvm_ffi=True)
      compiled_add_one_invalid_launch = cute.compile(
         add_one_invalid_launch, a_cute, b_cute, options="--enable-tvm-ffi"
      )

      try:
         compiled_add_one_invalid_launch(a_torch, b_torch)
      except RuntimeError as e:
         # raises RuntimeError: CUDA Error: cudaErrorInvalidValue
         print(f"RuntimeError: {e}")


Working with Devices
--------------------
TVM FFI-compiled functions naturally work across GPU devices.
The device index of the first input GPU tensor determines the kernel's device context.
The TVM FFI function calls ``cudaSetDevice`` to set the correct device
before launching the kernel based on that tensor's device index.
For advanced scenarios that pass raw pointers instead of tensors, you should call
``cudaSetDevice`` explicitly through the CUDA Python API.

Exporting Compiled Module
-------------------------

The TVM FFI function supports exporting the compiled module to an object file
for further use. For example:

.. code-block:: python

   import subprocess
   import cutlass.cute as cute

   def example_add_one_export():
      n = cute.sym_int()
      a_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
      b_cute = cute.runtime.make_fake_compact_tensor(cute.Float32, (n,))
      # compile the kernel with "--enable-tvm-ffi" option and example input tensors
      compiled_add_one = cute.compile(add_one, a_cute, b_cute, options="--enable-tvm-ffi")
      # export the compiled module to object file
      compiled_add_one.export_to_c("./add_one.o", function_name="add_one")
      # obtain necessary runtime libs for loading the shared library
      runtime_libs = cute.runtime.find_runtime_libraries(enable_tvm_ffi=True)
      # compile the object file to a shared library
      cmd = ["gcc", "-shared", "-o", "./add_one.so", "./add_one.o", *runtime_libs]
      print(cmd)
      subprocess.run(cmd, check=True)
      print(f"Successfully created shared library: ./add_one.so")

Then you can load back the exported module and use it in different ways:

.. code-block:: python

   import torch
   from cutlass import cute

   def example_load_module_add_one():
      mod = cute.runtime.load_module("./add_one.so", enable_tvm_ffi=True)
      a_torch = torch.arange(10, dtype=torch.float32, device="cuda")
      b_torch = torch.empty(10, dtype=torch.float32, device="cuda")
      mod.add_one(a_torch, b_torch)
      print("result of b_torch after mod.add_one(a_torch, b_torch)")
      print(b_torch)


The exported object file exposes the function symbol ``__tvm_ffi_add_one`` that is
compatible with TVM FFI and can be used in various frameworks and programming languages.
You can either build a shared library and load it back, or link the object file directly
into your application and invoke the function via the ``InvokeExternC`` mechanism in TVM FFI.
For more information, see the `quick start guide <https://tvm.apache.org/ffi/get_started/quickstart>`_
in the official documentation.

When you build your own libraries, make sure you link against the necessary runtime libraries.
You can use ``cute.runtime.find_runtime_libraries(enable_tvm_ffi=True)`` to get the path to these libraries.
``cute.runtime.load_module(path, enable_tvm_ffi=True)`` will load these libraries automatically before loading
an exported module. You can also manually load these libraries in advanced use cases.

For low-level cute ABI AOT compilation support without TVM FFI, you can refer to :doc:`ahead_of_time_compilation`.


Keyword Arguments and Defaults

  a_torch = torch.arange(10, dtype=torch.float32, device="cuda")
  b_torch = torch.empty(10, dtype=torch.float32, device="cuda")

  mod = cute.runtime.load_module("./add_constant.so")
  try:
     mod.add_constant(a_torch, b_torch)
  except Exception as e:
     # Raises a missing arguments error because kwargs and default information are lost
     print(e)
  # We rewrap the function to regain argument and kwargs support.
  # Alternatively, use the TVM FFI wrapper generator to generate a kwargs wrapper function.
  from tvm_ffi.utils import kwargs_wrapper
  # arg_defaults are aligned to the end of the argument list
  wrapped_func = kwargs_wrapper.make_kwargs_wrapper(
     mod.add_constant, arg_names=["a", "b", "offset"], arg_defaults=(1,)
  )
  wrapped_func(a_torch, b_torch)
  print("result of b_torch after wrapped_func(a_torch, b_torch)")
  print(b_torch)
  # You can also use the signature of the original function
  # to generate a kwargs wrapper function. Make sure to exclude
  # arguments that are not included in the runtime,
  # such as 'self', constexpr, and env stream arguments.
  wrapped_func = kwargs_wrapper.make_kwargs_wrapper_from_signature(
     mod.add_constant, signature=inspect.signature(add_constant),
     exclude_arg_names=["self"]
  )
  wrapped_func(a_torch, b_torch, offset=4)
  print("result of b_torch after wrapped_func(a_torch, b_torch, offset=4)")
  print(b_torch)