You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by ya...@apache.org on 2026/03/27 14:19:29 UTC
(tvm-ffi) branch main updated: [OrcJIT] Add LLVM ORC JIT v2 dynamic object loader addon (#254)
This is an automated email from the ASF dual-hosted git repository.
yaxingcai pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm-ffi.git
The following commit(s) were added to refs/heads/main by this push:
new 1da5637 [OrcJIT] Add LLVM ORC JIT v2 dynamic object loader addon (#254)
1da5637 is described below
commit 1da563747d5552a4d8f6a9f2bac063426d9c1ba5
Author: Yaoyao Ding <di...@gmail.com>
AuthorDate: Fri Mar 27 10:19:23 2026 -0400
[OrcJIT] Add LLVM ORC JIT v2 dynamic object loader addon (#254)
This PR introduces a new addon package `tvm-ffi-orcjit` that enables
dynamic loading of TVM-FFI exported object files (.o) at runtime using
LLVM's ORC JIT v2 engine.
The addon provides a Python API for loading compiled object files, load
the tvm-ffi functions defined in the object files.
The API is organized around three main concepts:
- **ExecutionSession**: A JIT compilation context that manages the
lifetime of dynamic libraries.
- **DynamicLibrary**: Represents a shared library that links multiple
object files.
- **Object Files**: Compiled `.o` files containing TVM-FFI exported
functions.
## Usage Example
```python
from tvm_ffi_orcjit import create_session
# Create an execution session (JIT compilation context)
session = create_session()
# Load an object file into the session, returns a DynamicLibrary handle
dylib = session.load("example.o")
# Get and call a function from the loaded object file
add_func = dylib.get_function("simple_add")
result = add_func(1, 2)
print(f"Result: {result}") # Output: Result: 3
```
For incremental loading, you can add multiple object files to the same
session:
```python
session = create_session()
# Load first object file
dylib = session.load("math_ops.o")
add = dylib.get_function("simple_add")
# Incrementally load more object files into the same library
dylib.load("string_ops.o")
# Now both functions are accessible from the same library
concat = dylib.get_function("string_concat")
```
See the test for more example.
## TODO
- Support Windows platform (currently Linux and macOS only)
- Customize memory allocation for dynamic libraries (allow user-provided
allocators)
---------
Co-authored-by: Yaxing Cai <ca...@gmail.com>
---
.github/actions/build-orcjit-wheel/action.yml | 114 ++++
.github/workflows/ci_test.yml | 42 ++
.github/workflows/publish_orcjit_wheel.yml | 83 +++
.gitignore | 2 +
addons/tvm_ffi_orcjit/CMakeLists.txt | 178 ++++++
addons/tvm_ffi_orcjit/ORCJIT_PRIMER.md | 535 ++++++++++++++++
addons/tvm_ffi_orcjit/README.md | 261 ++++++++
.../examples/quick-start/CMakeLists.txt | 67 ++
.../tvm_ffi_orcjit/examples/quick-start/README.md | 83 +++
addons/tvm_ffi_orcjit/examples/quick-start/add.cc | 46 ++
addons/tvm_ffi_orcjit/examples/quick-start/add_c.c | 59 ++
addons/tvm_ffi_orcjit/examples/quick-start/run.py | 137 ++++
addons/tvm_ffi_orcjit/pyproject.toml | 83 +++
.../python/tvm_ffi_orcjit/__init__.py | 94 +++
.../python/tvm_ffi_orcjit/_ffi_api.py | 21 +
.../tvm_ffi_orcjit/python/tvm_ffi_orcjit/dylib.py | 85 +++
.../python/tvm_ffi_orcjit/session.py | 90 +++
addons/tvm_ffi_orcjit/src/ffi/orcjit_dylib.cc | 212 +++++++
addons/tvm_ffi_orcjit/src/ffi/orcjit_dylib.h | 129 ++++
addons/tvm_ffi_orcjit/src/ffi/orcjit_session.cc | 696 +++++++++++++++++++++
addons/tvm_ffi_orcjit/src/ffi/orcjit_session.h | 129 ++++
addons/tvm_ffi_orcjit/src/ffi/orcjit_utils.h | 90 +++
addons/tvm_ffi_orcjit/tests/CMakeLists.txt | 107 ++++
addons/tvm_ffi_orcjit/tests/README.md | 109 ++++
.../tests/sources/c/test_call_global.c | 106 ++++
.../tests/sources/c/test_ctor_dtor.c | 152 +++++
addons/tvm_ffi_orcjit/tests/sources/c/test_error.c | 31 +
addons/tvm_ffi_orcjit/tests/sources/c/test_funcs.c | 43 ++
.../tvm_ffi_orcjit/tests/sources/c/test_funcs2.c | 42 ++
.../tests/sources/c/test_funcs_conflict.c | 43 ++
.../tests/sources/c/test_link_order_base.c | 33 +
.../tests/sources/c/test_link_order_caller.c | 35 ++
addons/tvm_ffi_orcjit/tests/sources/c/test_types.c | 60 ++
.../tests/sources/cc/test_call_global.cc | 37 ++
.../tests/sources/cc/test_ctor_dtor.cc | 82 +++
.../tvm_ffi_orcjit/tests/sources/cc/test_error.cc | 26 +
.../tvm_ffi_orcjit/tests/sources/cc/test_funcs.cc | 26 +
.../tvm_ffi_orcjit/tests/sources/cc/test_funcs2.cc | 26 +
.../tests/sources/cc/test_funcs_conflict.cc | 26 +
.../tests/sources/cc/test_link_order_base.cc | 24 +
.../tests/sources/cc/test_link_order_caller.cc | 43 ++
.../tvm_ffi_orcjit/tests/sources/cc/test_types.cc | 33 +
.../tests/sources/cuda/test_funcs.cu | 61 ++
addons/tvm_ffi_orcjit/tests/test_basic.py | 475 ++++++++++++++
addons/tvm_ffi_orcjit/tests/utils.py | 232 +++++++
docs/conf.py | 3 +
python/tvm_ffi/cpp/extension.py | 300 ++++++---
tests/lint/check_file_type.py | 1 +
tests/python/test_build.py | 2 +-
tests/python/test_build_inline.py | 71 +++
tests/python/test_stl.py | 2 +-
51 files changed, 5375 insertions(+), 92 deletions(-)
diff --git a/.github/actions/build-orcjit-wheel/action.yml b/.github/actions/build-orcjit-wheel/action.yml
new file mode 100644
index 0000000..0b8844f
--- /dev/null
+++ b/.github/actions/build-orcjit-wheel/action.yml
@@ -0,0 +1,114 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+
+name: Build OrcJIT Wheel
+description: >
+ Build and test OrcJIT wheels for a given OS/architecture combination.
+ Handles LLVM caching, conda installation, and cibuildwheel execution.
+
+inputs:
+ arch:
+ description: "Target architecture (e.g., x86_64, aarch64, arm64, AMD64)"
+ required: true
+ build:
+ description: "cibuildwheel build selector (e.g., cp312-manylinux_x86_64)"
+ required: true
+ checkout_ref:
+ description: "Branch, tag, or SHA to check out before building"
+ required: true
+
+runs:
+ using: "composite"
+ steps:
+ - name: Check out source
+ uses: actions/checkout@v5
+ with:
+ ref: ${{ inputs.checkout_ref }}
+ submodules: recursive
+
+ - uses: ./.github/actions/detect-env-vars
+ id: env_vars
+
+ # ---- Cache LLVM prefix ----
+ - name: Cache LLVM
+ uses: actions/cache@v4
+ id: llvm-cache
+ with:
+ path: ${{ runner.os == 'Windows' && 'C:/opt/llvm' || '/opt/llvm' }}
+ key: llvm-22.1.0-${{ runner.os }}-${{ inputs.arch }}-v3
+
+ # ---- Install LLVM via conda (cache miss only) ----
+ - name: Setup conda
+ if: steps.llvm-cache.outputs.cache-hit != 'true'
+ uses: conda-incubator/setup-miniconda@fc2d68f6413eb2d87b895e92f8584b5b94a10167 # v3.3.0
+ continue-on-error: true
+ id: conda1
+ with:
+ miniforge-version: latest
+
+ - name: Setup conda (retry with tar.bz2)
+ if: steps.llvm-cache.outputs.cache-hit != 'true' && steps.conda1.outcome == 'failure'
+ uses: conda-incubator/setup-miniconda@fc2d68f6413eb2d87b895e92f8584b5b94a10167 # v3.3.0
+ with:
+ miniforge-version: latest
+ use-only-tar-bz2: true
+
+ - name: Create /opt/llvm (macOS)
+ if: steps.llvm-cache.outputs.cache-hit != 'true' && runner.os == 'macOS'
+ shell: bash
+ run: sudo mkdir -p /opt/llvm && sudo chown -R $(whoami) /opt/llvm
+
+ - name: Install LLVM (Unix)
+ if: steps.llvm-cache.outputs.cache-hit != 'true' && runner.os != 'Windows'
+ shell: bash -l {0}
+ run: |
+ conda create -q -p /opt/llvm -c conda-forge \
+ llvmdev=22.1.0 clangdev=22.1.0 compiler-rt=22.1.0 zlib zstd-static \
+ -y
+
+ - name: Install LLVM (Windows)
+ if: steps.llvm-cache.outputs.cache-hit != 'true' && runner.os == 'Windows'
+ shell: cmd /C call {0}
+ run: |
+ conda create -q -p C:\opt\llvm -c conda-forge llvmdev=22.1.0 zlib zstd-static -y
+
+ # ---- Build and test wheels ----
+ - name: Build and test wheels
+ uses: pypa/cibuildwheel@298ed2fb2c105540f5ed055e8a6ad78d82dd3a7e # v3.3.1
+ with:
+ package-dir: addons/tvm_ffi_orcjit
+ output-dir: wheelhouse
+ env:
+ CIBW_BUILD: ${{ inputs.build }}
+ CIBW_ARCHS_LINUX: ${{ inputs.arch }}
+ CIBW_ARCHS_MACOS: ${{ inputs.arch }}
+ CIBW_ARCHS_WINDOWS: ${{ inputs.arch }}
+ CIBW_BUILD_VERBOSITY: 1
+ CMAKE_BUILD_PARALLEL_LEVEL: ${{ steps.env_vars.outputs.cpu_count }}
+ CIBW_ENVIRONMENT: LLVM_PREFIX=/opt/llvm
+ CIBW_ENVIRONMENT_WINDOWS: LLVM_PREFIX="C:/opt/llvm"
+ CIBW_CONTAINER_ENGINE: "docker; create_args: --volume /opt/llvm:/opt/llvm"
+ CIBW_TEST_REQUIRES: pytest ninja
+ CIBW_TEST_COMMAND: >-
+ pip install --force-reinstall {project} &&
+ pytest {project}/addons/tvm_ffi_orcjit/tests -v &&
+ python {project}/addons/tvm_ffi_orcjit/examples/quick-start/run.py --lang cpp &&
+ python {project}/addons/tvm_ffi_orcjit/examples/quick-start/run.py --lang c
+ CIBW_TEST_COMMAND_WINDOWS: >-
+ pip install --force-reinstall {project} &&
+ pytest {project}/addons/tvm_ffi_orcjit/tests -v &&
+ python {project}/addons/tvm_ffi_orcjit/examples/quick-start/run.py --lang c
diff --git a/.github/workflows/ci_test.yml b/.github/workflows/ci_test.yml
index 37f00af..ec7bce3 100644
--- a/.github/workflows/ci_test.yml
+++ b/.github/workflows/ci_test.yml
@@ -35,6 +35,7 @@ jobs:
should_skip_ci_docs_only: ${{ steps.detect.outputs.should_skip_ci_docs_only }}
cpp_changed: ${{ steps.cpp_files.outputs.changed }}
cpp_files: ${{ steps.cpp_files.outputs.files }}
+ orcjit_changed: ${{ steps.orcjit_files.outputs.changed }}
steps:
- uses: actions/checkout@v5
with:
@@ -54,6 +55,13 @@ jobs:
src/ tests/ | grep -E '\.(c|cc|cpp|cxx)$' | tr '\n' ' ')
echo "files=$FILES" >> $GITHUB_OUTPUT
[ -n "$FILES" ] && echo "changed=true" >> $GITHUB_OUTPUT || echo "changed=false" >> $GITHUB_OUTPUT
+ - name: Get changed OrcJIT files
+ id: orcjit_files
+ run: |
+ FILES=$(git diff --name-only --diff-filter=ACMR origin/${{ github.base_ref }}...HEAD -- \
+ addons/tvm_ffi_orcjit/ .github/actions/build-orcjit-wheel/ | tr '\n' ' ')
+ echo "files=$FILES" >> $GITHUB_OUTPUT
+ [ -n "$FILES" ] && echo "changed=true" >> $GITHUB_OUTPUT || echo "changed=false" >> $GITHUB_OUTPUT
lint:
needs: [prepare]
@@ -199,3 +207,37 @@ jobs:
working-directory: rust
run: |
cargo test
+
+ orcjit:
+ needs: [lint, prepare]
+ if: >
+ needs.prepare.outputs.orcjit_changed == 'true' &&
+ needs.prepare.outputs.should_skip_ci_commit != 'true' &&
+ needs.prepare.outputs.should_skip_ci_docs_only != 'true'
+ name: orcjit ${{ matrix.os }} (${{ matrix.arch }})
+ runs-on: ${{ matrix.os }}
+ strategy:
+ fail-fast: false
+ matrix:
+ include:
+ - {os: ubuntu-latest, arch: x86_64, build: "cp312-manylinux_x86_64"}
+ - {os: ubuntu-24.04-arm, arch: aarch64, build: "cp312-manylinux_aarch64"}
+ - {os: macos-14, arch: arm64, build: "cp312-macosx_arm64"}
+ - {os: windows-latest, arch: AMD64, build: "cp312-win_amd64"}
+
+ steps:
+ - uses: actions/checkout@v5
+ with:
+ fetch-depth: 1
+
+ - name: Build OrcJIT wheel
+ uses: ./.github/actions/build-orcjit-wheel
+ with:
+ arch: ${{ matrix.arch }}
+ build: ${{ matrix.build }}
+ checkout_ref: ${{ github.sha }}
+
+ - uses: actions/upload-artifact@v4
+ with:
+ name: cibw-orcjit-wheels-${{ matrix.os }}-${{ matrix.arch }}-${{ strategy.job-index }}
+ path: wheelhouse/*.whl
diff --git a/.github/workflows/publish_orcjit_wheel.yml b/.github/workflows/publish_orcjit_wheel.yml
new file mode 100644
index 0000000..27bd390
--- /dev/null
+++ b/.github/workflows/publish_orcjit_wheel.yml
@@ -0,0 +1,83 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+
+name: Publish OrcJIT wheel
+
+on:
+ workflow_dispatch:
+ inputs:
+ tag:
+ description: "Tag to publish (manual run)"
+ required: true
+
+jobs:
+ build_wheels:
+ name: ${{ matrix.os }} (${{ matrix.arch }})
+ runs-on: ${{ matrix.os }}
+ strategy:
+ fail-fast: false
+ matrix:
+ include:
+ - {os: ubuntu-latest, arch: x86_64, build: "cp312-manylinux_x86_64"}
+ - {os: ubuntu-24.04-arm, arch: aarch64, build: "cp312-manylinux_aarch64"}
+ - {os: macos-14, arch: arm64, build: "cp312-macosx_arm64"}
+ - {os: windows-latest, arch: AMD64, build: "cp312-win_amd64"}
+
+ steps:
+ - name: Checkout repository (for local composite action)
+ uses: actions/checkout@v5
+ with:
+ fetch-depth: 1
+
+ - name: Build OrcJIT wheel
+ uses: ./.github/actions/build-orcjit-wheel
+ with:
+ arch: ${{ matrix.arch }}
+ build: ${{ matrix.build }}
+ checkout_ref: ${{ inputs.tag }}
+
+ - name: Upload wheels
+ uses: actions/upload-artifact@v4
+ with:
+ name: cibw-orcjit-wheels-${{ matrix.os }}-${{ matrix.arch }}-${{ strategy.job-index }}
+ path: ./wheelhouse/*.whl
+
+ upload_pypi:
+ needs: [build_wheels]
+ runs-on: ubuntu-latest
+ environment: pypi
+ permissions:
+ id-token: write
+ attestations: write
+ if: github.event_name == 'workflow_dispatch'
+ steps:
+ - uses: actions/download-artifact@v4
+ with:
+ pattern: cibw-orcjit-*
+ path: dist
+ merge-multiple: true
+
+ - name: Generate artifact attestation for wheels
+ uses: actions/attest-build-provenance@v1
+ with:
+ subject-path: dist/*
+
+ - name: Publish package distributions to PyPI
+ uses: pypa/gh-action-pypi-publish@release/v1
+ with:
+ attestations: true
+ verbose: true
diff --git a/.gitignore b/.gitignore
index c1638b4..fa91f21 100644
--- a/.gitignore
+++ b/.gitignore
@@ -122,6 +122,8 @@ celerybeat-schedule
# Allow composite actions under .github/actions despite build-* ignore pattern
!.github/actions/build-wheel-for-publish/
!.github/actions/build-wheel-for-publish/action.yml
+!.github/actions/build-orcjit-wheel/
+!.github/actions/build-orcjit-wheel/action.yml
celerybeat.pid
# SageMath parsed files
diff --git a/addons/tvm_ffi_orcjit/CMakeLists.txt b/addons/tvm_ffi_orcjit/CMakeLists.txt
new file mode 100644
index 0000000..5281238
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/CMakeLists.txt
@@ -0,0 +1,178 @@
+cmake_minimum_required(VERSION 3.20)
+project(
+ tvm_ffi_orcjit
+ VERSION 0.1.0
+ LANGUAGES C CXX
+)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
+
+# ---- LLVM_PREFIX handling ----
+# scikit-build-core overrides CMAKE_PREFIX_PATH via its init cache, so we read LLVM_PREFIX from the
+# environment instead.
+if (DEFINED ENV{LLVM_PREFIX})
+ list(APPEND CMAKE_PREFIX_PATH "$ENV{LLVM_PREFIX}")
+ if (EXISTS "$ENV{LLVM_PREFIX}/Library")
+ list(APPEND CMAKE_PREFIX_PATH "$ENV{LLVM_PREFIX}/Library")
+ endif ()
+endif ()
+
+# ---- Find packages ----
+find_package(LLVM REQUIRED CONFIG)
+message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION} in ${LLVM_DIR}")
+
+find_package(
+ Python
+ COMPONENTS Interpreter
+ REQUIRED
+)
+execute_process(
+ COMMAND "${Python_EXECUTABLE}" -m tvm_ffi.config --cmakedir
+ OUTPUT_STRIP_TRAILING_WHITESPACE
+ OUTPUT_VARIABLE tvm_ffi_ROOT
+)
+find_package(tvm_ffi CONFIG REQUIRED)
+
+# ---- Build shared library ----
+add_library(tvm_ffi_orcjit SHARED src/ffi/orcjit_session.cc src/ffi/orcjit_dylib.cc)
+set_target_properties(
+ tvm_ffi_orcjit PROPERTIES CXX_VISIBILITY_PRESET hidden VISIBILITY_INLINES_HIDDEN ON
+)
+target_include_directories(
+ tvm_ffi_orcjit PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/src ${LLVM_INCLUDE_DIRS}
+)
+separate_arguments(LLVM_DEFINITIONS_LIST NATIVE_COMMAND ${LLVM_DEFINITIONS})
+target_compile_definitions(tvm_ffi_orcjit PRIVATE ${LLVM_DEFINITIONS_LIST})
+
+# ---- Static LLVM linking via llvm-config ----
+set(_llvm_config_hints ${LLVM_TOOLS_BINARY_DIR} "${LLVM_DIR}/../../../bin")
+if (DEFINED ENV{LLVM_PREFIX})
+ list(APPEND _llvm_config_hints "$ENV{LLVM_PREFIX}/bin" "$ENV{LLVM_PREFIX}/Library/bin")
+endif ()
+find_program(
+ LLVM_CONFIG_EXE
+ NAMES llvm-config-${LLVM_VERSION_MAJOR} llvm-config
+ HINTS ${_llvm_config_hints}
+)
+if (NOT LLVM_CONFIG_EXE)
+ message(FATAL_ERROR "llvm-config not found")
+endif ()
+
+execute_process(
+ COMMAND ${LLVM_CONFIG_EXE} --link-static --libs Core OrcJIT Support native
+ OUTPUT_STRIP_TRAILING_WHITESPACE
+ OUTPUT_VARIABLE _llvm_libs
+)
+execute_process(
+ COMMAND ${LLVM_CONFIG_EXE} --link-static --ldflags
+ OUTPUT_STRIP_TRAILING_WHITESPACE
+ OUTPUT_VARIABLE _llvm_ldflags
+)
+separate_arguments(_llvm_libs_list NATIVE_COMMAND "${_llvm_libs}")
+separate_arguments(_llvm_ldflags_list NATIVE_COMMAND "${_llvm_ldflags}")
+
+# ---- zlib (static, from LLVM prefix) ----
+cmake_path(GET LLVM_DIR PARENT_PATH _llvm_prefix)
+cmake_path(GET _llvm_prefix PARENT_PATH _llvm_prefix)
+cmake_path(GET _llvm_prefix PARENT_PATH _llvm_prefix)
+set(_lib_hints "${_llvm_prefix}/lib" "${_llvm_prefix}/lib64" "${_llvm_prefix}/Library/lib")
+
+if (WIN32)
+ find_library(
+ ZLIB_STATIC
+ NAMES zlibstatic z
+ HINTS ${_lib_hints} REQUIRED
+ )
+else ()
+ find_library(
+ ZLIB_STATIC
+ NAMES libz.a
+ HINTS ${_lib_hints} REQUIRED
+ )
+endif ()
+
+# ---- zstd (static, from LLVM prefix) ----
+if (WIN32)
+ find_library(
+ ZSTD_STATIC
+ NAMES zstd_static zstd
+ HINTS ${_lib_hints} REQUIRED
+ )
+else ()
+ find_library(
+ ZSTD_STATIC
+ NAMES libzstd.a
+ HINTS ${_lib_hints} REQUIRED
+ )
+endif ()
+
+# ---- Link everything ----
+target_link_libraries(
+ tvm_ffi_orcjit PRIVATE tvm_ffi::header tvm_ffi::shared ${_llvm_ldflags_list} ${_llvm_libs_list}
+ ${ZLIB_STATIC} ${ZSTD_STATIC}
+)
+
+# LLVM system libs (ntdll/psapi on Windows, rt/dl on Linux), minus zlib/zstd (linked above).
+execute_process(
+ COMMAND ${LLVM_CONFIG_EXE} --system-libs
+ OUTPUT_STRIP_TRAILING_WHITESPACE
+ OUTPUT_VARIABLE _sys_libs
+ RESULT_VARIABLE _sys_libs_rc
+)
+if (_sys_libs_rc EQUAL 0 AND _sys_libs)
+ separate_arguments(_sys_libs_list NATIVE_COMMAND "${_sys_libs}")
+ list(FILTER _sys_libs_list EXCLUDE REGEX "(zstd|^-lz$|^z\\.lib$|xml2)")
+ target_link_libraries(tvm_ffi_orcjit PRIVATE ${_sys_libs_list})
+endif ()
+
+# ---- Platform-specific fixups ----
+if (APPLE)
+ add_custom_command(
+ TARGET tvm_ffi_orcjit
+ POST_BUILD
+ COMMAND install_name_tool -change @rpath/libc++.1.dylib /usr/lib/libc++.1.dylib
+ $<TARGET_FILE:tvm_ffi_orcjit>
+ COMMENT "Fixing libc++ rpath to use system library"
+ )
+elseif (UNIX AND NOT WIN32)
+ target_link_options(tvm_ffi_orcjit PRIVATE -static-libstdc++ -static-libgcc)
+endif ()
+
+# ---- Find and bundle liborc_rt ----
+set(ORC_RT_PATH "")
+if (DEFINED ENV{ORC_RT_PATH})
+ set(ORC_RT_PATH "$ENV{ORC_RT_PATH}")
+else ()
+ if (WIN32)
+ file(GLOB_RECURSE _orc_rt_candidates "${_llvm_prefix}/lib/clang/*/lib/orc_rt*.lib")
+ if (NOT _orc_rt_candidates)
+ file(GLOB_RECURSE _orc_rt_candidates "${_llvm_prefix}/Library/lib/clang/*/lib/orc_rt*.lib")
+ endif ()
+ else ()
+ file(GLOB_RECURSE _orc_rt_candidates "${_llvm_prefix}/lib/clang/*/lib/liborc_rt*.a")
+ endif ()
+ if (_orc_rt_candidates)
+ list(GET _orc_rt_candidates 0 ORC_RT_PATH)
+ endif ()
+endif ()
+
+if (NOT ORC_RT_PATH)
+ message(WARNING "Could not find liborc_rt. ORC runtime features will be disabled.")
+endif ()
+
+if (ORC_RT_PATH)
+ message(STATUS "Found liborc_rt: ${ORC_RT_PATH}")
+ file(COPY "${ORC_RT_PATH}" DESTINATION "${CMAKE_BINARY_DIR}")
+ install(FILES "${ORC_RT_PATH}" DESTINATION lib)
+endif ()
+
+# ---- Install ----
+install(
+ TARGETS tvm_ffi_orcjit
+ LIBRARY DESTINATION lib
+ ARCHIVE DESTINATION lib
+ RUNTIME DESTINATION lib
+)
diff --git a/addons/tvm_ffi_orcjit/ORCJIT_PRIMER.md b/addons/tvm_ffi_orcjit/ORCJIT_PRIMER.md
new file mode 100644
index 0000000..af3a7e7
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/ORCJIT_PRIMER.md
@@ -0,0 +1,535 @@
+<!--- Licensed to the Apache Software Foundation (ASF) under one -->
+<!--- or more contributor license agreements. See the NOTICE file -->
+<!--- distributed with this work for additional information -->
+<!--- regarding copyright ownership. The ASF licenses this file -->
+<!--- to you 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. -->
+
+# ORC JIT v2 Primer — Background for the `tvm_ffi_orcjit` Addon
+
+This document explains the background knowledge needed to understand the
+`tvm_ffi_orcjit` addon: object file formats, the classical linker model, and LLVM's
+ORC JIT v2 architecture. It then maps those concepts onto the addon's implementation.
+
+---
+
+## 1. Object File Formats
+
+A **compiled object file** (`.o` / `.obj`) is not an executable. It is an intermediate
+container that holds machine code, data, and metadata about unresolved references. The
+three dominant formats are:
+
+| Format | Platform | File extensions |
+| --- | --- | --- |
+| **ELF** (Executable and Linkable Format) | Linux, most Unix | `.o`, `.so`, `.elf` |
+| **Mach-O** (Mach Object) | macOS, iOS | `.o`, `.dylib`, `.macho` |
+| **COFF/PE** (Common Object File Format / Portable Executable) | Windows | `.obj`, `.dll`, `.exe` |
+
+Despite surface differences, all three share the same conceptual structure:
+
+```text
+┌─────────────────────────────────┐
+│ File Header │ magic, target arch, section count
+├─────────────────────────────────┤
+│ Section Headers │ name, type, file offset, size, flags
+├─────────────────────────────────┤
+│ .text (code) │ machine instructions
+│ .rodata (read-only data) │ string literals, constants
+│ .data (writable data) │ initialized globals
+│ .bss (zero-init data) │ uninitialized globals (no file bytes)
+│ .init_array (constructors) │ array of function pointers — C++ ctors
+│ .fini_array (destructors) │ array of function pointers — C++ dtors
+│ .eh_frame (unwind info) │ exception/stack-unwind tables
+│ … other sections … │
+├─────────────────────────────────┤
+│ Symbol Table │ name → (section, offset, binding, type)
+├─────────────────────────────────┤
+│ Relocation Tables │ (section, offset, symbol, type, addend)
+└─────────────────────────────────┘
+```
+
+### Symbols
+
+A **symbol** is a named location in the file — a function entry point, a global variable,
+or a section boundary. Symbols have:
+
+- **binding**: `LOCAL` (file-private), `GLOBAL` (externally visible), `WEAK` (override-able default)
+- **definition status**: *defined* (has an address in this file) vs. *undefined* (imported — must be
+ resolved at link time)
+
+### Relocations
+
+A **relocation** is a "fixup recipe" stored alongside the machine code. It says:
+
+> "At byte offset X in section S, patch in the address of symbol FOO, using formula R."
+
+The formula R is a **relocation type** that encodes what arithmetic to apply:
+
+| Type | Meaning |
+| --- | --- |
+| `R_X86_64_64` | Absolute 64-bit address of the symbol |
+| `R_X86_64_PC32` | Symbol address minus the patch location (PC-relative 32-bit) |
+| `R_X86_64_PLT32` | PC-relative call through the PLT (procedure linkage table) |
+| `IMAGE_REL_AMD64_ADDR32NB` (COFF) | Address relative to image base (Pointer32NB) |
+
+The linker (or JIT) processes these tables to produce the final binary.
+
+### Platform-Specific Initialization Sections
+
+C++ global constructors and destructors must run before/after `main`. Compilers encode
+them as arrays of function pointers in special sections:
+
+| Platform | Constructors | Destructors | Notes |
+| --- | --- | --- | --- |
+| ELF (Linux) | `.init_array` | `.fini_array` | Priority suffix `.init_array.NNN` (lower = earlier) |
+| ELF (legacy) | `.ctors` | `.dtors` | Older GCC style |
+| Mach-O | `__DATA,__mod_init_func` | `__DATA,__mod_term_func` | Processed by dyld |
+| COFF (Windows) | `.CRT$XCU` (default) | `.CRT$XTZ` | Suffix encodes priority via ASCII ordering |
+
+The linker or OS loader is responsible for iterating these arrays and calling each pointer
+in priority order before handing control to user code. A JIT that loads `.o` files must
+replicate this behavior itself — this is a key responsibility of `tvm_ffi_orcjit`.
+
+---
+
+## 2. The Classical Linker Model
+
+A **static linker** (`ld`, `link.exe`) takes multiple object files and libraries and
+produces a single loadable image (executable or shared library). The process is:
+
+```text
+Object files + Archives
+ │
+ ▼
+┌───────────────────┐
+│ Symbol Resolution │ Match every undefined symbol to a definition.
+│ │ Archive (.a) members pulled in on demand.
+└────────┬──────────┘
+ │
+ ▼
+┌───────────────────┐
+│ Section Merging │ All .text sections → one .text; same for .data, etc.
+│ │ Assign virtual addresses (VMA) to each section.
+└────────┬──────────┘
+ │
+ ▼
+┌───────────────────┐
+│ Relocation │ Apply each relocation fixup now that all VMAs are known.
+│ │ Patch raw bytes in the output image.
+└────────┬──────────┘
+ │
+ ▼
+┌───────────────────┐
+│ Output Emission │ Write ELF/PE/Mach-O header + sections + dynamic table.
+└───────────────────┘
+```
+
+### Dynamic Linking
+
+At **load time** (when `dlopen` / `LoadLibrary` runs), the OS dynamic linker
+(`ld.so`, `dyld`, `ntdll`) performs a *reduced* version of the above:
+
+1. Map shared library sections into the process address space.
+2. Resolve cross-library symbol references (PLT/GOT stubs).
+3. Apply load-time relocations for position-dependent code.
+4. Run `.init_array` / `__mod_init_func` constructors.
+
+### Why JIT Is Different
+
+A JIT allocates memory at arbitrary addresses *at runtime*. It cannot use precomputed
+link-time addresses. Every relocation must be re-evaluated against the JIT-allocated
+addresses. Furthermore, if the JIT loads object files incrementally (one at a time),
+symbol resolution must be deferred until all relevant objects are present.
+
+---
+
+## 3. LLVM ORC JIT v2
+
+**ORC** stands for *On Request Compilation*. LLVM ORC JIT v2 (introduced in LLVM 9,
+stabilized in LLVM 13+) is a complete redesign of LLVM's JIT infrastructure. It is
+designed to be composable, asynchronous, and correct for production use (unlike the
+older `MCJIT` which had several fundamental limitations around multi-module linking).
+
+### 3.1 Core Concepts
+
+#### ExecutionSession
+
+`llvm::orc::ExecutionSession` is the root object of any ORC JIT instance. It owns:
+
+- A set of `JITDylib`s (the symbol namespaces).
+- The dispatch mechanism for asynchronous compilation tasks.
+- The global symbol interning table (maps string → `SymbolStringPtr`).
+
+Think of it as the "JIT process" — one per logical JIT environment.
+
+#### JITDylib (JIT Dynamic Library)
+
+A `JITDylib` is a **symbol namespace** that loosely mirrors a shared library. It:
+
+- Holds a **symbol table**: `name → {flags, address}`.
+- Has a **link order** (a list of other `JITDylib`s to search for unresolved symbols).
+- Can be populated via *materialization units* (object files, LLVM IR, inline asm, etc.).
+
+Multiple `JITDylib`s can coexist in one `ExecutionSession`, enabling isolation: e.g.,
+one library per compiled kernel, sharing a common runtime library.
+
+```text
+ExecutionSession
+ ├── JITDylib "main" ← default, process symbols
+ ├── JITDylib "libA" ← user kernel A
+ │ link_order: [libA, libB, main]
+ └── JITDylib "libB" ← user kernel B (shared by A)
+```
+
+#### LLJIT
+
+`llvm::orc::LLJIT` is the high-level, batteries-included wrapper around
+`ExecutionSession`. It:
+
+- Sets up a target machine and data layout.
+- Creates a default `JITDylib` ("main") with process symbol support.
+- Configures the linking pipeline (see below).
+- Exposes `addObjectFile()`, `addIRModule()`, and `lookup()`.
+
+`LLJIT` is what `tvm_ffi_orcjit` wraps in `ORCJITExecutionSessionObj`.
+
+#### MaterializationUnit and MaterializationResponsibility
+
+A **MaterializationUnit** is a lazy producer of symbols. When a symbol is first looked
+up and not yet defined, ORC triggers its materialization unit to produce the definition
+asynchronously. This is the "on request" in ORC.
+
+An object file becomes a `StaticLibraryDefinitionGenerator` or `ObjectLayer`-level unit:
+the object is parsed, linked, and its symbols resolved only when someone asks for them.
+
+#### Layers
+
+ORC processes objects through a pipeline of **layers**, each transforming the input:
+
+```text
+addObjectFile(buffer)
+ │
+ ▼
+┌─────────────────────┐
+│ ObjectTransformLayer│ (optional) transform raw object bytes before linking
+│ e.g. strip .pdata │
+└────────┬────────────┘
+ │
+ ▼
+┌─────────────────────┐
+│ ObjectLinkingLayer │ runs JITLink to:
+│ (uses JITLink) │ • parse object format
+│ │ • resolve symbols across JITDylibs
+│ │ • apply relocations
+│ │ • allocate + write JIT memory
+└────────┬────────────┘
+ │
+ ▼
+ JITDylib symbol table updated; code is live
+```
+
+#### JITLink
+
+`JITLink` is LLVM's low-level, graph-based linker used inside `ObjectLinkingLayer`. It:
+
+1. **Parses** the object file into an in-memory `LinkGraph` (nodes = sections/atoms,
+ edges = relocations).
+2. **Runs pass pipelines** (pre-prune, post-allocation, post-fixup) — plugins can
+ inspect and modify the graph at each stage.
+3. **Allocates** JIT memory (code + data) via a `JITLinkMemoryManager`.
+4. **Resolves** relocations using the current `ExecutionSession` symbol lookup.
+5. **Finalizes** by writing machine code into the allocated pages and marking them
+ executable.
+
+The pass pipeline is where `tvm_ffi_orcjit`'s `InitFiniPlugin` does its work.
+
+### 3.2 Symbol Lookup and Resolution
+
+When `session.lookup(search_order, symbol_name)` is called:
+
+```text
+1. Check each JITDylib in search_order, in order.
+2. If symbol is defined → return its ExecutorAddr.
+3. If symbol is in a MaterializationUnit not yet realized → trigger materialization.
+4. Materialization runs JITLink for the relevant object → symbol becomes defined.
+5. Return the address.
+```
+
+Symbol names are **mangled** (C++ name mangling) or explicitly prefixed. `tvm_ffi_orcjit`
+uses the `__tvm_ffi_` prefix to namespace exported functions.
+
+### 3.3 DefinitionGenerators
+
+A `DefinitionGenerator` is a fallback attached to a `JITDylib`. When a symbol is not
+found in the dylib's own table, the generator is invoked to dynamically create a
+definition — typically by looking up a symbol in the host process or another library.
+
+`LLJIT` attaches a `DynamicLibrarySearchGenerator` ("ProcessSymbols") to the main
+dylib, which resolves symbols like `malloc`, `printf`, or any C runtime function by
+looking them up in the running process.
+
+On Windows, `tvm_ffi_orcjit` adds a custom `DLLImportDefinitionGenerator` that handles
+`__imp_XXX` import stubs that MSVC-compiled objects expect.
+
+### 3.4 Platform Support
+
+#### What an ORC Platform Is
+
+When the OS dynamic linker loads a shared library, it does more than just map bytes
+into memory — it runs constructors, registers `atexit` handlers, and sets up
+thread-local storage. An **ORC platform** is the JIT-side object that replicates this
+OS loader behavior for JIT-linked code.
+
+Concretely, an ORC platform:
+
+1. **Intercepts `__cxa_atexit`**: C++ objects with destructors register their cleanup
+ via `__cxa_atexit`. The platform installs an interposer so these registrations are
+ scoped to a `JITDylib` and can be drained when that dylib is torn down — instead of
+ running at process exit like normal.
+2. **Drives initialization**: After an object file is linked into a `JITDylib`, calling
+ `jit_->initialize(dylib)` asks the platform to run constructors (e.g., iterate
+ `__mod_init_func` on macOS, `.init_array` on ELF).
+3. **Drives deinitialization**: `jit_->deinitialize(dylib)` drains the `atexit`
+ handlers registered during step 2 and runs destructors.
+
+This requires a small **ORC runtime library** (part of LLVM's `compiler-rt`) to be
+compiled for the target and loaded into the JIT — it provides the actual
+`__cxa_atexit` interposer and initialization trampolines that the platform object
+coordinates with.
+
+The three platform objects in LLVM are:
+
+| Platform | OS | Init section driven |
+| --- | --- | --- |
+| `MachOPlatform` | macOS / iOS | `__DATA,__mod_init_func` |
+| `ELFNativePlatform` | Linux / ELF | `.init_array`, TLS |
+| `COFFPlatform` | Windows | `.CRT$XC*` init, `__cxa_atexit` interop |
+
+`ExecutorNativePlatform` is a convenience builder that auto-selects the right platform
+for the host OS and loads the ORC runtime from a given path.
+
+#### How `tvm_ffi_orcjit` Uses (or Avoids) ORC Platforms
+
+The addon takes a different approach on each platform:
+
+- **macOS**: ORC platform support is *optional*. When the caller passes an ORC runtime
+ path to `ExecutionSession`, `ExecutorNativePlatform` activates `MachOPlatform`.
+ `jit_->initialize(dylib)` and `jit_->deinitialize(dylib)` then drive `__mod_init_func`
+ and `__cxa_atexit` teardown natively. Without the path, the addon falls back to its
+ own `InitFiniPlugin`.
+- **Windows**: `COFFPlatform` is skipped entirely because it requires MSVC CRT symbols
+ (`_CxxThrowException`, RTTI vtables, iostream objects) that are not resolvable in
+ the JIT context. Instead, `InitFiniPlugin` manually handles `.CRT$XC*` / `.CRT$XT*`
+ init/fini sections.
+- **Linux**: `ELFNativePlatform` is not used. `InitFiniPlugin` handles `.init_array` /
+ `.fini_array` / `.ctors` / `.dtors` directly, without the ORC runtime.
+
+---
+
+## 4. How `tvm_ffi_orcjit` Uses ORC JIT v2
+
+With the background above, here is how the addon maps onto the concepts.
+
+### 4.1 Object Model
+
+```text
+Python / C++ API LLVM ORC v2 concept
+─────────────────────────────────────────────────────────────────
+ExecutionSession LLJIT + ExecutionSession
+DynamicLibrary JITDylib
+dylib.add("foo.o") ObjectLinkingLayer.add(buffer)
+dylib.get_function("add") ExecutionSession.lookup("__tvm_ffi_add")
+dylib.set_link_order([a, b]) JITDylib.setLinkOrder([a, b, main, ...])
+```
+
+### 4.2 Object File Loading Pipeline
+
+```text
+dylib.add("foo.o")
+ │
+ ▼ ORCJITDynamicLibraryObj::AddObjectFile()
+ │ jit_->addObjectFile(*dylib_, MemoryBuffer)
+ │
+ ▼ ObjectTransformLayer (macOS: skipped; Linux/Win: may strip .pdata)
+ │ Windows: strip .pdata/.xdata avoids JITLink COMDAT limitation
+ │ fix __ImageBase fixes Pointer32NB relocations
+ │
+ ▼ ObjectLinkingLayer (JITLink)
+ │ Parse object → LinkGraph
+ │ Run InitFiniPlugin passes:
+ │ PrePrunePasses: mark .init_array / .ctors blocks as live
+ │ PostAllocationPasses: (Windows) set __ImageBase, strip SEH sections
+ │ PostFixupPasses: extract resolved init/fini function pointers
+ │ → session.AddPendingInitializer(dylib, entry)
+ │ Resolve relocations via ExecutionSession.lookup()
+ │ Allocate JIT memory, write code, mark executable
+ │
+ ▼ JITDylib symbol table updated
+ (symbols are defined but constructors not yet run)
+```
+
+### 4.3 Symbol Lookup and Initialization
+
+```text
+dylib.get_function("add")
+ │
+ ▼ ORCJITDynamicLibraryObj::GetFunction("add")
+ │ → GetSymbol("__tvm_ffi_add")
+ │ build JITDylibSearchOrder: [this, linked dylibs, LLJIT default]
+ │ jit_->getExecutionSession().lookup(order, "__tvm_ffi_add")
+ │ ← ExecutorAddr
+ │
+ ▼ Linux/Windows: RunPendingInitializers()
+ │ Sort entries by priority
+ │ Call each function pointer (C++ ctors, .init_array entries)
+ │
+ ▼ macOS: jit_->initialize(*dylib_)
+ │ ORC MachOPlatform calls __mod_init_func pointers
+ │
+ ▼ Wrap raw function pointer as tvm_ffi::Function
+ via Function::FromPacked lambda (marshals AnyView args ↔ TVMFFIAny)
+```
+
+### 4.4 InitFiniPlugin Detail
+
+`InitFiniPlugin` is an `ObjectLinkingLayer::Plugin` — it receives callbacks during
+JITLink's pass pipeline for every object being linked.
+
+```text
+JITLink pass pipeline for each object file:
+─────────────────────────────────────────────────────────────
+PrePrunePasses
+ → InitFiniPlugin::modifyPassConfig (PrePrunePasses)
+ For each section named .init_array* / .ctors* / .fini_array* / .dtors*
+ (ELF) or __DATA,__mod_init_func (Mach-O) or .CRT$XC* (COFF):
+ Mark all blocks in that section as "keep" (live)
+ → Prevents dead-code elimination from removing constructors
+
+PostAllocationPasses (Windows only)
+ → InitFiniPlugin::modifyPassConfig (PostAllocationPasses)
+ Set __ImageBase = lowest allocated block address
+ Strip .pdata / .xdata exception handler sections
+
+PostFixupPasses
+ → InitFiniPlugin::modifyPassConfig (PostFixupPasses)
+ Iterate all blocks in init/fini sections
+ Read each 8-byte slot as an ExecutorAddr
+ Parse section name to determine: priority, is_init vs is_fini
+ Call session->AddPendingInitializer(dylib, InitFiniEntry{addr, section, priority})
+ (or AddPendingDeinitializer for dtors/fini_array)
+```
+
+### 4.5 Cross-Library Symbol Resolution
+
+```text
+libA depends on a symbol defined in libB:
+
+libA.set_link_order([libB])
+ → JITDylibSearchOrder for libA: [libA, libB, main(ProcessSymbols)]
+
+When libA's object file has an unresolved symbol "foo":
+ JITLink asks ExecutionSession.lookup([libA, libB, main], "foo")
+ → found in libB → returns libB's ExecutorAddr for "foo"
+ → relocation in libA patched with that address
+```
+
+### 4.6 Windows DLL Import Stubs
+
+Windows MSVC objects reference DLL functions through `__imp_XXX` pointer stubs (the
+Import Address Table pattern). At static link time the linker creates these stubs. In
+JIT mode there is no linker, so `DLLImportDefinitionGenerator` creates them on demand:
+
+```text
+JITLink encounters undefined symbol "__imp_malloc"
+ │
+ ▼ DLLImportDefinitionGenerator::tryToGenerate()
+ │ Search ucrtbase.dll, msvcrt.dll, then all process modules
+ │ for the real address of "malloc"
+ │
+ ▼ Allocate two JIT-memory stubs:
+ │ __imp_malloc → 8-byte slot containing &malloc (host address)
+ │ malloc → x86_64 jmp [__imp_malloc] trampoline
+ │
+ ▼ Define both symbols in the JITDylib
+ → JITLink can now apply PCRel32 reloc to __imp_malloc (stub is close in JIT memory)
+```
+
+The stubs must live in JIT-allocated memory (not at the host process address) because
+x86_64 `PCRel32` relocations can only reach ±2 GB. The host's `malloc` may be farther
+than 2 GB from the JIT allocation.
+
+---
+
+## 5. End-to-End Example
+
+```python
+import tvm_ffi_orcjit as oj
+
+# 1. Create an ExecutionSession (wraps LLJIT)
+sess = oj.ExecutionSession()
+
+# 2. Create a JITDylib
+lib = sess.create_dynamic_library()
+
+# 3. Load a compiled object file
+# → object parsed, JITLink links it, InitFiniPlugin collects ctors
+lib.add("add.o")
+
+# 4. Look up a function
+# → LLVM resolves "__tvm_ffi_add", RunPendingInitializers() fires ctors
+add = lib.get_function("add") # returns tvm_ffi.Function
+
+# 5. Call it
+result = add(3, 4) # → 7
+```
+
+The corresponding C++ side of `add.o`:
+
+```cpp
+// add.cc — compiled to add.o with clang++ -c -O2 add.cc
+#include <tvm/ffi/function.h>
+
+static tvm::ffi::Function add_impl = [](int a, int b) { return a + b; };
+
+// Exports symbol "__tvm_ffi_add" using TVMFFISafeCallType ABI
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(add, add_impl);
+```
+
+---
+
+## 6. Key Concepts Summary
+
+| Concept | What it is | Where in the addon |
+| --- | --- | --- |
+| Object file | Container of machine code, data, symbols, relocations | Input to `dylib.add()` |
+| Relocation | Recipe to patch a code address at link/JIT time | Applied by JITLink |
+| `.init_array` / `.ctors` | Array of C++ constructor pointers in ELF objects | Collected by `InitFiniPlugin` |
+| `ExecutionSession` | Root of the ORC JIT environment | `ORCJITExecutionSessionObj` |
+| `LLJIT` | High-level ORC JIT wrapper | Stored in `ORCJITExecutionSessionObj::jit_` |
+| `JITDylib` | Symbol namespace / virtual shared library | `ORCJITDynamicLibraryObj::dylib_` |
+| `JITLink` | LLVM's JIT-aware linker | Used inside `ObjectLinkingLayer` |
+| JITLink pass pipeline | Pre-prune → post-alloc → post-fixup hooks | Where `InitFiniPlugin` runs |
+| `DefinitionGenerator` | Fallback symbol provider | `DLLImportDefinitionGenerator` (Win) |
+| Link order | Search path across JITDylibs for symbol resolution | `SetLinkOrder()` |
+| `__tvm_ffi_` prefix | Namespace for TVM-FFI exported functions | Used in `GetFunction()` |
+
+---
+
+## 7. Further Reading
+
+- LLVM ORC JIT documentation: <https://llvm.org/docs/ORCv2.html>
+- JITLink design: <https://llvm.org/docs/JITLink.html>
+- ELF specification: <https://refspecs.linuxfoundation.org/elf/elf.pdf>
+- PE/COFF specification: <https://learn.microsoft.com/en-us/windows/win32/debug/pe-format>
+- Mach-O reference: <https://developer.apple.com/documentation/kernel/mach-o_file_format_reference>
+- Ian Lance Taylor's linker series (20-part blog): foundational reading on linkers
diff --git a/addons/tvm_ffi_orcjit/README.md b/addons/tvm_ffi_orcjit/README.md
new file mode 100644
index 0000000..c68d193
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/README.md
@@ -0,0 +1,261 @@
+<!--- Licensed to the Apache Software Foundation (ASF) under one -->
+<!--- or more contributor license agreements. See the NOTICE file -->
+<!--- distributed with this work for additional information -->
+<!--- regarding copyright ownership. The ASF licenses this file -->
+<!--- to you 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. -->
+
+# TVM-FFI OrcJIT
+
+A Python package that enables dynamic loading of compiled object files (`.o`)
+using LLVM ORC JIT v2, providing a flexible JIT execution environment for
+TVM-FFI exported functions.
+
+## Features
+
+- **JIT Execution**: Load and execute compiled object files at runtime using LLVM's ORC JIT v2
+- **Multiple Libraries**: Create separate dynamic libraries with independent symbol namespaces
+- **Incremental Loading**: Add multiple object files to the same library incrementally
+- **Symbol Isolation**: Different libraries can define the same symbol without conflicts
+- **Init/Fini Support**: Handles static constructors/destructors across ELF (`.init_array`/`.ctors`), Mach-O (`__mod_init_func`), and COFF (`.CRT$XC*`/`.CRT$XT*`)
+- **Cross-Platform**: Linux (x86_64, aarch64), macOS (arm64), Windows (AMD64)
+- **Multi-Compiler**: Tested with LLVM Clang, GCC, Apple Clang, MSVC, and clang-cl
+- **TVM-FFI Integration**: Seamlessly works with TVM-FFI's stable C ABI
+- **Python API**: Simple Pythonic interface for JIT compilation and execution
+
+## Supported Platforms and Compilers
+
+Object files compiled with any of the following compiler/platform combinations
+can be loaded and executed by the ORC JIT:
+
+| Platform | Compilers | C | C++ |
+| -------- | --------- | :-: | :-: |
+| Linux (x86_64, aarch64) | LLVM Clang, GCC | yes | yes |
+| macOS (arm64) | LLVM Clang, Apple Clang | yes | yes |
+| Windows (AMD64) | LLVM Clang, MSVC, clang-cl | yes | no |
+
+Windows is C-only across all compilers. C++ objects compiled with
+`TVM_FFI_DLL_EXPORT_TYPED_FUNC` use `try`/`catch` (via `TVM_FFI_SAFE_CALL_BEGIN/END`),
+which requires Itanium exception ABI symbols (`__cxa_begin_catch`,
+`__gxx_personality_v0`, etc.) that the MSVC-built host process cannot provide.
+Pure C objects using the `TVMFFISafeCallType` ABI work on all platforms.
+
+## Installation
+
+### Install from PyPI
+
+```bash
+pip install apache-tvm-ffi apache-tvm_ffi_orcjit
+```
+
+### Build from Source
+
+#### Prerequisites
+
+- Python 3.10+, CMake 3.20+, C++17 compiler
+- LLVM 22+ development libraries (`llvmdev`, `llvm-config`)
+- Static `zlib` and `zstd` libraries (in the same prefix as LLVM)
+
+#### Install LLVM via conda-forge
+
+The easiest way to get all dependencies is via conda-forge:
+
+```bash
+conda create -p /opt/llvm -c conda-forge \
+ llvmdev=22.1.0 clangdev=22.1.0 compiler-rt=22.1.0 zlib zstd-static -y
+export LLVM_PREFIX=/opt/llvm
+```
+
+On Windows:
+
+```cmd
+conda create -p C:\opt\llvm -c conda-forge llvmdev=22.1.0 zlib zstd-static -y
+set LLVM_PREFIX=C:\opt\llvm
+```
+
+#### Build and install
+
+```bash
+git clone --recursive https://github.com/apache/tvm-ffi.git
+cd tvm-ffi
+
+# Install tvm-ffi first
+pip install -e .
+
+# Build and install the orcjit addon
+cd addons/tvm_ffi_orcjit
+pip install -e .
+```
+
+The `LLVM_PREFIX` environment variable tells CMake where to find LLVM. If
+LLVM is installed in a conda env or a standard system path, CMake can
+auto-discover it and `LLVM_PREFIX` is not needed.
+
+## Usage
+
+### Basic Example
+
+```python
+from tvm_ffi_orcjit import ExecutionSession
+
+# Create an execution session
+session = ExecutionSession()
+
+# Create a dynamic library
+lib = session.create_library()
+
+# Load an object file
+lib.add("example.o")
+
+# Get and call a function
+add_func = lib.get_function("add")
+result = add_func(1, 2)
+print(f"Result: {result}") # Output: Result: 3
+```
+
+### Multiple Libraries with Symbol Isolation
+
+```python
+session = ExecutionSession()
+
+lib1 = session.create_library("lib1")
+lib2 = session.create_library("lib2")
+
+lib1.add("implementation_v1.o")
+lib2.add("implementation_v2.o")
+
+add_v1 = lib1.get_function("add")
+add_v2 = lib2.get_function("add")
+
+print(add_v1(5, 3)) # Uses implementation from lib1
+print(add_v2(5, 3)) # Uses implementation from lib2
+```
+
+### Cross-Library Linking
+
+```python
+session = ExecutionSession()
+
+base_lib = session.create_library("base")
+base_lib.add("math_ops.o")
+
+caller_lib = session.create_library("caller")
+caller_lib.set_link_order(base_lib) # Can resolve symbols from base_lib
+caller_lib.add("caller.o")
+
+result = caller_lib.get_function("call_math")(10, 20)
+```
+
+## Writing Functions for OrcJIT
+
+### C++ (Linux/macOS)
+
+```cpp
+#include <tvm/ffi/function.h>
+
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(add, [](int a, int b) {
+ return a + b;
+});
+```
+
+Compile: `clang++ -std=c++17 -fPIC -O2 -c -o example.o example.cc`
+
+### Pure C (all platforms including Windows)
+
+```c
+#include <tvm/ffi/c_api.h>
+
+TVM_FFI_DLL_EXPORT int __tvm_ffi_add(
+ void* self, const TVMFFIAny* args, int32_t num_args, TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->v_int64 = args[0].v_int64 + args[1].v_int64;
+ return 0;
+}
+```
+
+Compile: `clang -O2 -c -o example.o example.c`
+
+## How It Works
+
+- **LLJIT**: Built on LLVM's ORC JIT v2 with `ObjectLinkingLayer` (JITLink) for
+ all platforms.
+- **InitFiniPlugin**: Custom `ObjectLinkingLayer::Plugin` that collects function
+ pointers from init/fini sections (ELF `.init_array`/`.ctors`/`.fini_array`/`.dtors`,
+ Mach-O `__mod_init_func`/`__mod_term_func`, COFF `.CRT$XC*`/`.CRT$XT*`) and
+ runs them in priority order at symbol lookup / library teardown.
+- **DLL Import Stubs** (Windows): Custom `DefinitionGenerator` that resolves host
+ process symbols from all loaded DLLs and creates `__imp_*` pointer stubs in
+ JIT memory, keeping all fixups within PCRel32 range.
+- **SEH Stripping** (Windows): `ObjectTransformLayer` strips `.pdata`/`.xdata`
+ relocations from COFF objects before JITLink graph building, working around a
+ JITLink limitation with COMDAT section symbols.
+
+Please refers to [ORCJIT_PRIMER.md](./ORCJIT_PRIMER.md) to learn more about object file, linking, llvm orcjit v2, and how the addon works.
+
+## Project Structure
+
+```text
+tvm_ffi_orcjit/
+├── CMakeLists.txt # Build configuration
+├── pyproject.toml # Python package metadata
+├── src/ffi/
+│ ├── orcjit_session.cc # ExecutionSession (LLJIT setup, plugins)
+│ ├── orcjit_session.h
+│ ├── orcjit_dylib.cc # DynamicLibrary (object loading, symbol lookup)
+│ ├── orcjit_dylib.h
+│ └── orcjit_utils.h # LLVM error handling utilities
+├── python/tvm_ffi_orcjit/
+│ ├── __init__.py # Module exports and library loading
+│ ├── session.py # Python ExecutionSession wrapper
+│ └── dylib.py # Python DynamicLibrary wrapper
+├── tests/ # See tests/README.md
+└── examples/quick-start/ # Complete example with CMake
+```
+
+## CI
+
+Runs on Linux (x86_64, aarch64), macOS (arm64), Windows (AMD64) via
+`cibuildwheel`. Each platform builds test objects with multiple compilers
+and runs the full test suite. See `.github/workflows/tvm_ffi_orcjit.yml`.
+
+## Troubleshooting
+
+### "Cannot find global function" error
+
+The shared library wasn't loaded. Reinstall: `pip install --force-reinstall apache-tvm_ffi_orcjit`
+
+### "Duplicate definition of symbol" error
+
+Use separate libraries for different implementations of the same symbol.
+
+### "Symbol not found" error
+
+Ensure functions are exported with TVM-FFI macros (`TVM_FFI_DLL_EXPORT_TYPED_FUNC`
+for C++, or `__tvm_ffi_` prefix for C).
+
+### Relocation errors on Windows
+
+MSVC/clang-cl objects must be compiled with `/GS-` to disable buffer security
+checks (`__security_cookie`) which are CRT symbols the JIT cannot resolve.
+
+### LLVM version mismatch
+
+The package requires LLVM 22+. Set `LLVM_PREFIX` to the LLVM install prefix:
+
+```bash
+export LLVM_PREFIX=/path/to/llvm
+```
+
+## License
+
+Apache License 2.0
diff --git a/addons/tvm_ffi_orcjit/examples/quick-start/CMakeLists.txt b/addons/tvm_ffi_orcjit/examples/quick-start/CMakeLists.txt
new file mode 100644
index 0000000..458f7fd
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/examples/quick-start/CMakeLists.txt
@@ -0,0 +1,67 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+
+cmake_minimum_required(VERSION 3.18)
+project(tvm_ffi_orcjit_example)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
+
+# Find tvm-ffi package
+find_package(
+ Python
+ COMPONENTS Interpreter
+ REQUIRED
+)
+execute_process(
+ COMMAND "${Python_EXECUTABLE}" -m tvm_ffi.config --cmakedir
+ OUTPUT_STRIP_TRAILING_WHITESPACE
+ OUTPUT_VARIABLE tvm_ffi_ROOT
+)
+find_package(tvm_ffi CONFIG REQUIRED)
+
+# Helper: build a source file into an object and copy as <name>.o
+function (add_example_object name source)
+ get_filename_component(_ext "${source}" EXT)
+ add_library(${name}_obj OBJECT "${source}")
+ target_link_libraries(${name}_obj PRIVATE tvm_ffi::header)
+ if (_ext STREQUAL ".c" AND MSVC)
+ target_compile_options(${name}_obj PRIVATE /O2)
+ else ()
+ target_compile_options(${name}_obj PRIVATE -O2)
+ if (CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64|arm64")
+ target_compile_options(${name}_obj PRIVATE -mno-outline-atomics)
+ endif ()
+ endif ()
+ add_custom_target(
+ copy_${name} ALL
+ COMMAND ${CMAKE_COMMAND} -E copy $<TARGET_OBJECTS:${name}_obj>
+ ${CMAKE_CURRENT_SOURCE_DIR}/${name}.o
+ DEPENDS ${name}_obj
+ COMMENT "Copying ${name}.o to example directory"
+ )
+endfunction ()
+
+# C++ object — skip on Windows (C-only strategy for ORC JIT on Windows)
+if (NOT WIN32)
+ add_example_object(add add.cc)
+endif ()
+
+# Pure C object — built on all platforms
+enable_language(C)
+add_example_object(add_c add_c.c)
diff --git a/addons/tvm_ffi_orcjit/examples/quick-start/README.md b/addons/tvm_ffi_orcjit/examples/quick-start/README.md
new file mode 100644
index 0000000..41834bb
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/examples/quick-start/README.md
@@ -0,0 +1,83 @@
+<!--- Licensed to the Apache Software Foundation (ASF) under one -->
+<!--- or more contributor license agreements. See the NOTICE file -->
+<!--- distributed with this work for additional information -->
+<!--- regarding copyright ownership. The ASF licenses this file -->
+<!--- to you 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. -->
+
+# Quick Start Example
+
+Demonstrates basic usage of tvm_ffi_orcjit: compile functions to object files,
+load them into the ORC JIT at runtime, and call them from Python.
+
+## Files
+
+| File | Description |
+| ------ | ------------- |
+| `add.cc` | C++ source using `TVM_FFI_DLL_EXPORT_TYPED_FUNC` (automatic type marshaling, supports `std::string`) |
+| `add_c.c` | Pure C source using `TVMFFISafeCallType` ABI (`__tvm_ffi_` prefix). No C++ runtime needed. |
+| `run.py` | Python script that compiles, loads, and calls the functions |
+| `CMakeLists.txt` | Alternative CMake build configuration |
+
+## Prerequisites
+
+- Python 3.10+, C/C++ compiler
+- `apache-tvm-ffi` and `apache-tvm_ffi_orcjit` packages installed
+
+## Run
+
+`run.py` compiles the source files to object files using `tvm_ffi.cpp.build`,
+then loads them into the ORC JIT:
+
+```bash
+# C++ variant (Linux/macOS)
+python run.py
+
+# Pure C variant (all platforms including Windows)
+python run.py --lang c
+```
+
+## How It Works
+
+**C++ variant** (`add.cc`): Functions are exported with
+`TVM_FFI_DLL_EXPORT_TYPED_FUNC`, which wraps a typed C++ lambda/function into
+TVM-FFI's packed calling convention. Supports C++ types like `std::string`.
+
+**C variant** (`add_c.c`): Functions follow the `TVMFFISafeCallType` ABI
+directly — each function is named `__tvm_ffi_<name>` and manually packs
+arguments/results via `TVMFFIAny`. Zero C++ dependencies, works on all
+platforms including Windows with MSVC or clang-cl.
+
+**Python side** (`run.py`):
+
+```python
+from tvm_ffi_orcjit import ExecutionSession
+
+session = ExecutionSession() # Create ORC JIT session
+lib = session.create_library() # Create a JITDylib
+lib.add("add_c.o") # Load object file into JIT
+add = lib.get_function("add") # Look up symbol
+print(add(10, 20)) # Call like a normal function → 30
+```
+
+## Platform Notes
+
+| Platform | C++ (`add.o`) | C (`add_c.o`) |
+| ---------- | :-: | :-: |
+| Linux (Clang/GCC) | yes | yes |
+| macOS (Clang/Apple Clang) | yes | yes |
+| Windows (all compilers) | no | yes |
+
+C++ is not supported for ORC JIT on Windows. The `TVM_FFI_DLL_EXPORT_TYPED_FUNC`
+macro uses `try`/`catch` which requires Itanium exception ABI symbols that the
+MSVC-built host process cannot provide. Use the pure C variant on Windows.
diff --git a/addons/tvm_ffi_orcjit/examples/quick-start/add.cc b/addons/tvm_ffi_orcjit/examples/quick-start/add.cc
new file mode 100644
index 0000000..dd3c3e2
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/examples/quick-start/add.cc
@@ -0,0 +1,46 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Quick Start Example - Simple Math Functions
+ *
+ * This file demonstrates how to export C++ functions using TVM-FFI
+ * so they can be loaded dynamically at runtime with tvm_ffi_orcjit.
+ */
+
+#include <tvm/ffi/function.h>
+
+// Simple addition function
+int add_impl(int a, int b) { return a + b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(add, add_impl);
+
+// Multiplication function
+int multiply_impl(int a, int b) { return a * b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(multiply, multiply_impl);
+
+// Fibonacci function (recursive)
+int fib_impl(int n) {
+ if (n <= 1) return n;
+ return fib_impl(n - 1) + fib_impl(n - 2);
+}
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(fibonacci, fib_impl);
+
+// String concatenation example
+std::string concat_impl(std::string a, std::string b) { return a + b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(concat, concat_impl);
diff --git a/addons/tvm_ffi_orcjit/examples/quick-start/add_c.c b/addons/tvm_ffi_orcjit/examples/quick-start/add_c.c
new file mode 100644
index 0000000..198f1a8
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/examples/quick-start/add_c.c
@@ -0,0 +1,59 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Quick Start Example - Simple Math Functions (Pure C)
+ *
+ * This file demonstrates how to export C functions using TVM-FFI's
+ * C ABI (TVMFFISafeCallType) so they can be loaded dynamically at
+ * runtime with tvm_ffi_orcjit. No C++ runtime dependencies.
+ */
+#include <tvm/ffi/c_api.h>
+
+/* add: add two integers */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_add(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 + args[1].v_int64;
+ return 0;
+}
+
+/* multiply: multiply two integers */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_multiply(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 * args[1].v_int64;
+ return 0;
+}
+
+/* fibonacci: recursive fibonacci */
+static int64_t fib(int64_t n) {
+ if (n <= 1) return n;
+ return fib(n - 1) + fib(n - 2);
+}
+
+TVM_FFI_DLL_EXPORT int __tvm_ffi_fibonacci(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = fib(args[0].v_int64);
+ return 0;
+}
diff --git a/addons/tvm_ffi_orcjit/examples/quick-start/run.py b/addons/tvm_ffi_orcjit/examples/quick-start/run.py
new file mode 100755
index 0000000..3a648a1
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/examples/quick-start/run.py
@@ -0,0 +1,137 @@
+#!/usr/bin/env python3
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+
+"""Quick Start Example - Compile, load, and call functions via ORC JIT.
+
+This script demonstrates how to:
+1. Compile C/C++ source files to object files using tvm_ffi.cpp.build
+2. Load them into an ORC JIT ExecutionSession
+3. Get functions by name
+4. Call them like regular Python functions
+
+Usage:
+ python run.py # Load C++ object file (add.o)
+ python run.py --lang c # Load pure C object file (add_c.o)
+"""
+
+from __future__ import annotations
+
+import argparse
+import platform
+import sys
+from pathlib import Path
+
+import tvm_ffi.cpp
+from tvm_ffi_orcjit import ExecutionSession
+
+SCRIPT_DIR = Path(__file__).resolve().parent
+
+
+def _extra_cflags() -> list[str]:
+ """Return extra compiler flags for the current platform."""
+ machine = platform.machine()
+ if machine in ("aarch64", "arm64"):
+ return ["-mno-outline-atomics"]
+ return []
+
+
+def _build_object(lang: str) -> str:
+ """Compile the source file to a relocatable object file."""
+ if lang == "c":
+ return tvm_ffi.cpp.build(
+ name="add_c",
+ sources=[str(SCRIPT_DIR / "add_c.c")],
+ output="add_c.o",
+ extra_cflags=_extra_cflags(),
+ )
+ else:
+ return tvm_ffi.cpp.build(
+ name="add",
+ sources=[str(SCRIPT_DIR / "add.cc")],
+ output="add.o",
+ extra_cflags=_extra_cflags(),
+ )
+
+
+def _run_tests(obj_file: str, lang: str) -> None:
+ """Load object file and run all test assertions.
+
+ All JIT references (functions, lib, session) are released automatically
+ when this function returns.
+ """
+ print(f"Loading object file: {obj_file} (lang={lang})")
+
+ # Create execution session and dynamic library
+ session = ExecutionSession()
+ lib = session.create_library()
+ lib.add(obj_file)
+
+ print("Object file loaded successfully\n")
+
+ # Get and call the 'add' function
+ print("=== Testing add function ===")
+ add = lib.get_function("add")
+ result = add(10, 20)
+ print(f"add(10, 20) = {result}")
+ assert result == 30, f"Expected 30, got {result}"
+
+ # Get and call the 'multiply' function
+ print("\n=== Testing multiply function ===")
+ multiply = lib.get_function("multiply")
+ result = multiply(7, 6)
+ print(f"multiply(7, 6) = {result}")
+ assert result == 42, f"Expected 42, got {result}"
+
+ # Get and call the 'fibonacci' function
+ print("\n=== Testing fibonacci function ===")
+ fibonacci = lib.get_function("fibonacci")
+ result = fibonacci(10)
+ print(f"fibonacci(10) = {result}")
+ assert result == 55, f"Expected 55, got {result}"
+
+ if lang == "cpp":
+ # String concatenation only available in C++ variant (uses std::string)
+ print("\n=== Testing concat function ===")
+ concat = lib.get_function("concat")
+ result = concat("Hello, ", "World!")
+ print(f"concat('Hello, ', 'World!') = '{result}'")
+ assert result == "Hello, World!", f"Expected 'Hello, World!', got '{result}'"
+
+ print("\n" + "=" * 50)
+ print("All tests passed successfully!")
+ print("=" * 50)
+
+
+def main() -> int:
+ """Run the quick start example."""
+ parser = argparse.ArgumentParser(description="Quick Start Example")
+ parser.add_argument(
+ "--lang",
+ choices=["cpp", "c"],
+ default="cpp",
+ help="Language variant to load: 'cpp' for add.o (default), 'c' for add_c.o",
+ )
+ args = parser.parse_args()
+
+ obj_file = _build_object(args.lang)
+ _run_tests(obj_file, args.lang)
+ return 0
+
+
+if __name__ == "__main__":
+ sys.exit(main())
diff --git a/addons/tvm_ffi_orcjit/pyproject.toml b/addons/tvm_ffi_orcjit/pyproject.toml
new file mode 100644
index 0000000..49bab10
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/pyproject.toml
@@ -0,0 +1,83 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+
+[build-system]
+requires = ["scikit-build-core>=0.10.0", "apache-tvm-ffi"]
+build-backend = "scikit_build_core.build"
+
+[project]
+name = "apache-tvm_ffi_orcjit"
+version = "0.1.0"
+description = "Load TVM-FFI exported object files using LLVM ORC JIT v2"
+readme = "README.md"
+requires-python = ">=3.10"
+license = { text = "Apache-2.0" }
+authors = [{ name = "TVM-FFI OrcJIT Contributors" }]
+keywords = ["tvm-ffi", "llvm", "jit", "orcjit"]
+classifiers = [
+ "Development Status :: 3 - Alpha",
+ "Intended Audience :: Developers",
+ "License :: OSI Approved :: Apache Software License",
+ "Programming Language :: Python :: 3",
+ "Programming Language :: Python :: 3.10",
+ "Programming Language :: Python :: 3.11",
+ "Programming Language :: Python :: 3.12",
+ "Programming Language :: Python :: 3.13",
+ "Programming Language :: C++",
+]
+dependencies = ["apache-tvm-ffi>=0.1.0"]
+
+[project.urls]
+Homepage = "https://github.com/apache/tvm-ffi"
+Repository = "https://github.com/apache/tvm-ffi"
+
+[tool.scikit-build]
+cmake.build-type = "Release"
+wheel.py-api = "py3"
+build-dir = "build"
+build.verbose = true
+editable.rebuild = false
+editable.verbose = true
+wheel.packages = ["python/tvm_ffi_orcjit"]
+wheel.install-dir = "tvm_ffi_orcjit"
+sdist.include = [
+ "/README.md",
+ "/LICENSE",
+ "/pyproject.toml",
+ "/CMakeLists.txt",
+ "/src/**/*.h",
+ "/src/**/*.cc",
+ "/src/**/*.cpp",
+ "/python/**/*.py",
+]
+
+[tool.scikit-build.cmake.define]
+CMAKE_EXPORT_COMPILE_COMMANDS = "ON"
+
+[tool.cibuildwheel]
+# One build per platform is sufficient: the shared library does not link against
+# Python C API, so wheel.py-api = "py3" produces a version-agnostic wheel.
+build = "cp312-*"
+skip = "*-win32 *-manylinux_i686 *-musllinux*"
+
+[tool.cibuildwheel.linux]
+# LLVM_PREFIX is unnecessary when LLVM is installed in a conda env
+# (cmake auto-discovers via CONDA_PREFIX) or in a standard system path.
+repair-wheel-command = "auditwheel repair --exclude libtvm_ffi.so -w {dest_dir} {wheel}"
+
+[tool.cibuildwheel.macos]
+repair-wheel-command = "delocate-wheel --ignore-missing-dependencies --exclude libtvm_ffi.dylib -w {dest_dir} -v {wheel}"
diff --git a/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/__init__.py b/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/__init__.py
new file mode 100644
index 0000000..1a94d89
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/__init__.py
@@ -0,0 +1,94 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+
+"""TVM-FFI OrcJIT.
+
+This module provides functionality to load object files (.o) compiled with TVM-FFI
+exports using LLVM ORC JIT v2.
+
+Example:
+ >>> from tvm_ffi_orcjit import ExecutionSession
+ >>> session = ExecutionSession()
+ >>> lib = session.create_library()
+ >>> lib.add("example.o")
+ >>> func = lib.get_function("my_function")
+ >>> result = func(arg1, arg2)
+
+"""
+
+import ctypes
+import os
+import platform
+import sys
+from pathlib import Path
+
+from tvm_ffi import load_module
+
+# Determine the library name based on platform
+if platform.system() == "Windows":
+ _LIB_NAME = "tvm_ffi_orcjit.dll"
+elif platform.system() == "Darwin":
+ _LIB_NAME = "libtvm_ffi_orcjit.dylib"
+else:
+ _LIB_NAME = "libtvm_ffi_orcjit.so"
+
+# Load the orcjit extension library
+# - lib/: normal install (wheel)
+# - ../../build/: editable install (cmake build output relative to python/tvm_ffi_orcjit/)
+_LIB_PATH = [
+ Path(__file__).parent / "lib" / _LIB_NAME,
+ Path(__file__).parent.parent.parent / "build" / _LIB_NAME,
+]
+_lib_dir = None
+for path in _LIB_PATH:
+ if path.exists():
+ _ = load_module(str(path))
+ _lib_dir = path.parent
+if _lib_dir is None:
+ raise RuntimeError(
+ f"Could not find {_LIB_NAME}. "
+ f"Searched in {_LIB_PATH} and site-packages. "
+ f"Please ensure the package is installed correctly."
+ )
+
+# Explicitly initialize the library to register functions
+# This is needed because static initializers may not run when loaded via dlopen
+try:
+ # The dll search path need to be added explicitly in windows
+ if sys.platform.startswith("win32"):
+ os.add_dll_directory(str(_lib_dir))
+ # Load the library with ctypes and call the initialization function
+ c_lib = ctypes.CDLL(str(_lib_dir / _LIB_NAME), mode=ctypes.RTLD_GLOBAL)
+ init_func = c_lib.TVMFFIOrcJITInitialize
+ init_func.restype = None
+ init_func()
+except Exception as e:
+ import warnings
+
+ warnings.warn(f"Failed to explicitly initialize orcjit library: {e}")
+
+from .dylib import DynamicLibrary
+from .session import ExecutionSession
+
+__all__ = ["DynamicLibrary", "ExecutionSession"]
+
+try:
+ from importlib.metadata import version
+
+ __version__ = version("apache-tvm_ffi_orcjit")
+except Exception:
+ __version__ = "0.0.0.dev0"
diff --git a/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/_ffi_api.py b/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/_ffi_api.py
new file mode 100644
index 0000000..fb2464b
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/_ffi_api.py
@@ -0,0 +1,21 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+"""FFI APIs for orcjit."""
+
+import tvm_ffi
+
+tvm_ffi.init_ffi_api("orcjit", __name__)
diff --git a/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/dylib.py b/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/dylib.py
new file mode 100644
index 0000000..6cfda6c
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/dylib.py
@@ -0,0 +1,85 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+"""ORC JIT Dynamic Library."""
+
+from __future__ import annotations
+
+from pathlib import Path
+
+from tvm_ffi import Module
+
+
+class DynamicLibrary(Module):
+ """ORC JIT Dynamic Library (JITDylib).
+
+ Represents a collection of symbols that can be loaded from object files and linked
+ against other dynamic libraries. Supports JIT compilation and symbol resolution.
+
+ Examples
+ --------
+ >>> session = ExecutionSession()
+ >>> lib = session.create_library()
+ >>> lib.add("add.o")
+ >>> lib.add("multiply.o")
+ >>> add_func = lib.get_function("add")
+ >>> result = add_func(1, 2)
+
+ """
+
+ def add(self, object_file: str | Path) -> None:
+ """Add an object file to this dynamic library.
+
+ Parameters
+ ----------
+ object_file : str or Path
+ Path to the object file to load.
+
+ Examples
+ --------
+ >>> lib.add("add.o")
+ >>> lib.add(Path("multiply.o"))
+
+ """
+ if isinstance(object_file, Path):
+ object_file = str(object_file)
+ self.get_function("orcjit.add_object_file")(object_file)
+
+ def set_link_order(self, *libraries: DynamicLibrary) -> None:
+ """Set the link order for symbol resolution.
+
+ When resolving symbols, this library will search in the specified libraries
+ in the order provided. This replaces any previous link order.
+
+ Parameters
+ ----------
+ *libraries : DynamicLibrary
+ One or more dynamic libraries to search for symbols (in order).
+
+ Examples
+ --------
+ >>> session = ExecutionSession()
+ >>> lib_utils = session.create_library()
+ >>> lib_utils.add("utils.o")
+ >>> lib_core = session.create_library()
+ >>> lib_core.add("core.o")
+ >>> lib_main = session.create_library()
+ >>> lib_main.add("main.o")
+ >>> # main can call symbols from utils and core (utils searched first)
+ >>> lib_main.set_link_order(lib_utils, lib_core)
+
+ """
+ self.get_function("orcjit.set_link_order")(list(libraries))
diff --git a/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/session.py b/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/session.py
new file mode 100644
index 0000000..66fb1e4
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/python/tvm_ffi_orcjit/session.py
@@ -0,0 +1,90 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+"""ORC JIT Execution Session."""
+
+from __future__ import annotations
+
+import sys
+
+from tvm_ffi import Object, register_object
+
+from . import _ffi_api, _lib_dir
+from .dylib import DynamicLibrary
+
+
+def _find_orc_rt_library() -> str | None:
+ """Find the bundled liborc_rt library in the same directory as the .so/.dll."""
+ # Windows: skip ORC runtime entirely. LLVM's COFFPlatform (loaded via
+ # ExecutorNativePlatform with liborc_rt) depends on MSVC C++ runtime symbols
+ # that are not available in the JIT environment. On Windows, ORC JIT uses a
+ # C-only strategy: JIT objects are compiled as pure C (TVMFFISafeCallType ABI),
+ # avoiding all C++ runtime dependencies (magic statics, RTTI, sized delete,
+ # SEH, COMDAT). Our custom InitFiniPlugin handles .CRT$XC*/.CRT$XT* init/fini
+ # sections, and DLLImportDefinitionGenerator resolves __imp_ DLL import stubs.
+ if sys.platform == "win32":
+ return None
+ patterns = ["liborc_rt*.a"]
+ for pattern in patterns:
+ for lib_path in _lib_dir.glob(pattern):
+ return str(lib_path)
+ return None
+
+
+@register_object("orcjit.ExecutionSession")
+class ExecutionSession(Object):
+ """ORC JIT Execution Session.
+
+ Manages the LLVM ORC JIT execution environment and creates dynamic libraries (JITDylibs).
+ This is the top-level context for JIT compilation and symbol management.
+
+ Examples
+ --------
+ >>> session = ExecutionSession()
+ >>> lib = session.create_library(name="main")
+ >>> lib.add("add.o")
+ >>> add_func = lib.get_function("add")
+
+ """
+
+ def __init__(self, orc_rt_path: str | None = None) -> None:
+ """Initialize ExecutionSession.
+
+ Args:
+ orc_rt_path: Optional path to the liborc_rt library. If not provided,
+ it will be automatically discovered using clang.
+
+ """
+ if orc_rt_path is None:
+ orc_rt_path = _find_orc_rt_library()
+ if orc_rt_path is None:
+ orc_rt_path = ""
+ self.__init_handle_by_constructor__(_ffi_api.ExecutionSession, orc_rt_path) # type: ignore
+
+ def create_library(self, name: str = "") -> DynamicLibrary:
+ """Create a new dynamic library associated with this execution session.
+
+ Args:
+ name: Optional name for the library. If empty, a unique name will be generated.
+
+ Returns:
+ A new DynamicLibrary instance.
+
+ """
+ handle = _ffi_api.ExecutionSessionCreateDynamicLibrary(self, name) # type: ignore
+ lib = DynamicLibrary.__new__(DynamicLibrary)
+ lib.__move_handle_from__(handle)
+ return lib
diff --git a/addons/tvm_ffi_orcjit/src/ffi/orcjit_dylib.cc b/addons/tvm_ffi_orcjit/src/ffi/orcjit_dylib.cc
new file mode 100644
index 0000000..82c5290
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/src/ffi/orcjit_dylib.cc
@@ -0,0 +1,212 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*!
+ * \file orcjit_dylib.cc
+ * \brief LLVM ORC JIT DynamicLibrary implementation
+ */
+
+#include "orcjit_dylib.h"
+
+#include <llvm/ExecutionEngine/Orc/Core.h>
+#include <llvm/ExecutionEngine/Orc/LLJIT.h>
+#include <llvm/Object/ObjectFile.h>
+#include <llvm/Support/Error.h>
+#include <llvm/Support/MemoryBuffer.h>
+#include <tvm/ffi/c_api.h>
+#include <tvm/ffi/cast.h>
+#include <tvm/ffi/container/array.h>
+#include <tvm/ffi/error.h>
+#include <tvm/ffi/extra/c_env_api.h>
+#include <tvm/ffi/extra/module.h>
+#include <tvm/ffi/function.h>
+#include <tvm/ffi/reflection/registry.h>
+
+#include "orcjit_session.h"
+#include "orcjit_utils.h"
+
+namespace tvm {
+namespace ffi {
+namespace orcjit {
+
+ORCJITDynamicLibraryObj::ORCJITDynamicLibraryObj(ORCJITExecutionSession session,
+ llvm::orc::JITDylib* dylib, llvm::orc::LLJIT* jit,
+ String name)
+ : session_(std::move(session)), dylib_(dylib), jit_(jit), name_(std::move(name)) {
+ TVM_FFI_CHECK(dylib_ != nullptr, ValueError) << "JITDylib cannot be null";
+ TVM_FFI_CHECK(jit_ != nullptr, ValueError) << "LLJIT cannot be null";
+ if (void** ctx_addr = reinterpret_cast<void**>(GetSymbol(ffi::symbol::tvm_ffi_library_ctx))) {
+ *ctx_addr = this;
+ }
+ Module::VisitContextSymbols([this](const ffi::String& name, void* symbol) {
+ if (void** ctx_addr = reinterpret_cast<void**>(GetSymbol(name))) {
+ *ctx_addr = symbol;
+ }
+ });
+}
+
+ORCJITDynamicLibraryObj::~ORCJITDynamicLibraryObj() {
+ // See InitFiniPlugin class doc in orcjit_session.cc for the three-platform
+ // init/fini strategy (macOS native vs. Linux/Windows plugin).
+#if defined(__linux__) || defined(_WIN32)
+ session_->RunPendingDeinitializers(GetJITDylib());
+#else
+ if (auto err = jit_->deinitialize(*dylib_)) {
+ llvm::consumeError(std::move(err));
+ }
+#endif
+}
+
+void ORCJITDynamicLibraryObj::AddObjectFile(const String& path) {
+ // Read object file
+ auto buffer_or_err = llvm::MemoryBuffer::getFile(path.c_str());
+ if (!buffer_or_err) {
+ TVM_FFI_THROW(IOError) << "Failed to read object file: " << path;
+ }
+
+ // Add object file to this JITDylib
+ TVM_FFI_ORCJIT_LLVM_CALL(jit_->addObjectFile(*dylib_, std::move(*buffer_or_err)));
+}
+
+void ORCJITDynamicLibraryObj::SetLinkOrder(const std::vector<llvm::orc::JITDylib*>& dylibs) {
+ // Rebuild the link order: user-specified libraries first, then the LLJIT
+ // default link order (Main → Platform → ProcessSymbols). Preserving the
+ // default link order is essential — without ProcessSymbols, C++ objects
+ // that need host-process symbols (runtime, libtvm_ffi) would fail to link.
+ link_order_.clear();
+
+ for (auto* lib : dylibs) {
+ link_order_.emplace_back(lib, llvm::orc::JITDylibLookupFlags::MatchAllSymbols);
+ }
+ for (auto& kv : jit_->defaultLinkOrder()) {
+ link_order_.emplace_back(kv.first, kv.second);
+ }
+
+ // Set the link order in the LLVM JITDylib
+ dylib_->setLinkOrder(link_order_, false);
+}
+
+void* ORCJITDynamicLibraryObj::GetSymbol(const String& name) {
+ // Build search order: this dylib first, then all linked dylibs
+ llvm::orc::JITDylibSearchOrder search_order;
+ search_order.emplace_back(dylib_, llvm::orc::JITDylibLookupFlags::MatchAllSymbols);
+ // Append linked libraries
+ search_order.insert(search_order.end(), link_order_.begin(), link_order_.end());
+
+ // Look up symbol using the full search order
+ auto symbol_or_err =
+ jit_->getExecutionSession().lookup(search_order, jit_->mangleAndIntern(name.c_str()));
+
+ // Run pending initializers. Both paths are idempotent and support incremental
+ // loading (new object files added between GetSymbol calls).
+ // See InitFiniPlugin class doc in orcjit_session.cc for three-platform strategy.
+#if defined(__linux__) || defined(_WIN32)
+ // Linux/Windows: RunPendingInitializers drains and erases the map entry;
+ // subsequent calls are no-ops until new object files add fresh entries.
+ session_->RunPendingInitializers(GetJITDylib());
+#else
+ // macOS: LLJIT tracks an InitializedDylib set (LLJIT.h:624). First call
+ // invokes dlopen (refcount 0→1, runs initializers). Subsequent calls
+ // invoke dlupdate (no refcount bump, runs only new initializers from
+ // newly-added object files). The single deinitialize() in the destructor
+ // calls dlclose (refcount 1→0), so the count stays balanced.
+ if (auto err = jit_->initialize(*dylib_)) {
+ llvm::consumeError(std::move(err));
+ }
+#endif
+ // Convert ExecutorAddr to pointer
+ return symbol_or_err ? symbol_or_err->getAddress().toPtr<void*>() : nullptr;
+}
+
+llvm::orc::JITDylib& ORCJITDynamicLibraryObj::GetJITDylib() {
+ TVM_FFI_CHECK(dylib_ != nullptr, InternalError) << "JITDylib is null";
+ return *dylib_;
+}
+
+Optional<Function> ORCJITDynamicLibraryObj::GetFunction(const String& name) {
+ if (name == "orcjit.add_object_file") {
+ return Function::FromTyped([this](const String& path) { AddObjectFile(path); });
+ }
+ if (name == "orcjit.set_link_order") {
+ return Function::FromTyped([this](const Array<ORCJITDynamicLibrary>& libraries) {
+ std::vector<llvm::orc::JITDylib*> libs;
+ libs.reserve(libraries.size());
+ for (const ORCJITDynamicLibrary& lib : libraries) {
+ libs.push_back(&lib->GetJITDylib());
+ }
+ SetLinkOrder(libs);
+ });
+ }
+
+ // TVM-FFI exports have __tvm_ffi_ prefix
+ std::string symbol_name = symbol::tvm_ffi_symbol_prefix + std::string(name);
+
+ // Try to get the symbol - return NullOpt if not found
+ if (void* symbol = GetSymbol(symbol_name)) {
+ // Wrap C function pointer as tvm-ffi Function.
+ // Capture a strong self-ref so the module (and its JIT code pages) stays
+ // alive as long as the returned Function is alive.
+ TVMFFISafeCallType c_func = reinterpret_cast<TVMFFISafeCallType>(symbol);
+ Module self_strong_ref = GetRef<Module>(this);
+ return Function::FromPacked([c_func, self_strong_ref](PackedArgs args, Any* rv) {
+ TVM_FFI_ICHECK_LT(rv->type_index(), ffi::TypeIndex::kTVMFFIStaticObjectBegin);
+ TVM_FFI_CHECK_SAFE_CALL((*c_func)(nullptr, reinterpret_cast<const TVMFFIAny*>(args.data()),
+ args.size(), reinterpret_cast<TVMFFIAny*>(rv)));
+ });
+ }
+ return std::nullopt;
+}
+
+//-------------------------------------
+// Registration
+//-------------------------------------
+
+static void RegisterOrcJITFunctions() {
+ static bool registered = false;
+ if (registered) return;
+ registered = true;
+
+ namespace refl = tvm::ffi::reflection;
+
+ refl::ObjectDef<ORCJITExecutionSessionObj>();
+
+ refl::GlobalDef()
+ .def("orcjit.ExecutionSession",
+ [](const std::string& orc_rt_path) { return ORCJITExecutionSession(orc_rt_path); })
+ .def("orcjit.ExecutionSessionCreateDynamicLibrary",
+ [](const ORCJITExecutionSession& session, const String& name) -> Module {
+ return session->CreateDynamicLibrary(name);
+ });
+}
+
+TVM_FFI_STATIC_INIT_BLOCK() {
+ // This block may not execute when loaded via dlopen on some platforms.
+ // Call TVMFFIOrcJITInitialize() explicitly if functions are not registered.
+ RegisterOrcJITFunctions();
+}
+
+} // namespace orcjit
+} // namespace ffi
+} // namespace tvm
+
+// C API for explicit initialization
+extern "C" {
+
+TVM_FFI_DLL_EXPORT void TVMFFIOrcJITInitialize() { tvm::ffi::orcjit::RegisterOrcJITFunctions(); }
+}
diff --git a/addons/tvm_ffi_orcjit/src/ffi/orcjit_dylib.h b/addons/tvm_ffi_orcjit/src/ffi/orcjit_dylib.h
new file mode 100644
index 0000000..43bd688
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/src/ffi/orcjit_dylib.h
@@ -0,0 +1,129 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*!
+ * \file orcjit_dylib.h
+ * \brief LLVM ORC JIT DynamicLibrary (JITDylib) wrapper
+ */
+#ifndef TVM_FFI_ORCJIT_ORCJIT_DYLIB_H_
+#define TVM_FFI_ORCJIT_ORCJIT_DYLIB_H_
+
+#include <llvm/ExecutionEngine/Orc/LLJIT.h>
+#include <tvm/ffi/container/array.h>
+#include <tvm/ffi/extra/module.h>
+#include <tvm/ffi/object.h>
+#include <tvm/ffi/string.h>
+
+#include "orcjit_session.h"
+
+namespace tvm {
+namespace ffi {
+namespace orcjit {
+
+class ORCJITExecutionSession;
+
+class ORCJITDynamicLibraryObj : public ModuleObj {
+ public:
+ /*!
+ * \brief Constructor
+ * \param session The parent execution session
+ * \param dylib The LLVM JITDylib
+ * \param jit The LLJIT instance
+ * \param name The library name
+ */
+ ORCJITDynamicLibraryObj(ORCJITExecutionSession session, llvm::orc::JITDylib* dylib,
+ llvm::orc::LLJIT* jit, String name);
+
+ ~ORCJITDynamicLibraryObj();
+
+ const char* kind() const final { return "orcjit"; }
+
+ Optional<Function> GetFunction(const String& name) override;
+
+ private:
+ /*!
+ * \brief Add an object file to this library
+ * \param path Path to the object file to load
+ */
+ void AddObjectFile(const String& path);
+
+ /*!
+ * \brief Set the link order for symbol resolution
+ * \param dylibs Vector of libraries to search for symbols (in order)
+ *
+ * When resolving symbols, this library will search in the specified libraries
+ * in the order provided. This replaces any previous link order.
+ */
+ void SetLinkOrder(const std::vector<llvm::orc::JITDylib*>& dylibs);
+
+ /*!
+ * \brief Look up a symbol in this library
+ * \param name The symbol name to look up
+ * \return Pointer to the symbol, or nullptr if not found
+ */
+ void* GetSymbol(const String& name);
+
+ /*!
+ * \brief Get the underlying LLVM JITDylib
+ * \return Reference to the LLVM JITDylib
+ */
+ llvm::orc::JITDylib& GetJITDylib();
+
+ /*!
+ * \brief Get the name of this library
+ * \return The library name
+ */
+ String GetName() const { return name_; }
+
+ /*! \brief Parent execution session (for lifetime management) */
+ ORCJITExecutionSession session_;
+
+ /*! \brief The LLVM JITDylib */
+ llvm::orc::JITDylib* dylib_;
+
+ /*! \brief The LLJIT instance (for addObjectFile API) */
+ llvm::orc::LLJIT* jit_;
+
+ /*! \brief Library name */
+ String name_;
+
+ /*! \brief Link order tracking (to support incremental linking) */
+ llvm::orc::JITDylibSearchOrder link_order_;
+};
+
+/*!
+ * \brief DynamicLibrary wrapper for LLVM ORC JIT v2 JITDylib
+ *
+ * This class wraps an LLVM JITDylib and provides functionality to:
+ * - Load object files
+ * - Link against other dynamic libraries
+ * - Look up symbols
+ */
+class ORCJITDynamicLibrary : public Module {
+ public:
+ explicit ORCJITDynamicLibrary(const ObjectPtr<ORCJITDynamicLibraryObj>& ptr) : Module(ptr) {};
+ TVM_FFI_DEFINE_OBJECT_REF_METHODS_NOTNULLABLE(ORCJITDynamicLibrary, Module,
+ ORCJITDynamicLibraryObj);
+};
+
+} // namespace orcjit
+} // namespace ffi
+} // namespace tvm
+
+#endif // TVM_FFI_ORCJIT_ORCJIT_DYLIB_H_
diff --git a/addons/tvm_ffi_orcjit/src/ffi/orcjit_session.cc b/addons/tvm_ffi_orcjit/src/ffi/orcjit_session.cc
new file mode 100644
index 0000000..ed2290e
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/src/ffi/orcjit_session.cc
@@ -0,0 +1,696 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*!
+ * \file orcjit_session.cc
+ * \brief LLVM ORC JIT ExecutionSession implementation
+ */
+
+#include "orcjit_session.h"
+
+#include <llvm/ExecutionEngine/JITLink/JITLink.h>
+#include <llvm/ExecutionEngine/JITLink/x86_64.h>
+#include <llvm/ExecutionEngine/Orc/AbsoluteSymbols.h>
+#include <llvm/ExecutionEngine/Orc/LLJIT.h>
+#include <llvm/ExecutionEngine/Orc/ObjectLinkingLayer.h>
+#include <llvm/ExecutionEngine/Orc/ObjectTransformLayer.h>
+#include <llvm/ExecutionEngine/Orc/Shared/ExecutorAddress.h>
+#include <llvm/ExecutionEngine/Orc/Shared/ExecutorSymbolDef.h>
+#include <llvm/Support/DynamicLibrary.h>
+#include <llvm/Support/Error.h>
+#include <llvm/Support/TargetSelect.h>
+#include <llvm/TargetParser/SubtargetFeature.h>
+#include <tvm/ffi/cast.h>
+#include <tvm/ffi/error.h>
+#include <tvm/ffi/object.h>
+#include <tvm/ffi/reflection/registry.h>
+
+#include <cstddef>
+#include <cstring>
+
+#ifdef _WIN32
+#ifndef NOMINMAX
+#define NOMINMAX
+#endif
+#ifndef WIN32_LEAN_AND_MEAN
+#define WIN32_LEAN_AND_MEAN
+#endif
+#include <psapi.h>
+#include <windows.h>
+#endif
+
+#include "orcjit_dylib.h"
+#include "orcjit_utils.h"
+
+namespace tvm {
+namespace ffi {
+namespace orcjit {
+
+// Initialize LLVM native target (only once)
+struct LLVMInitializer {
+ LLVMInitializer() {
+ llvm::InitializeNativeTarget();
+ llvm::InitializeNativeTargetAsmPrinter();
+ llvm::InitializeNativeTargetAsmParser();
+ }
+};
+
+static LLVMInitializer llvm_initializer;
+
+// Custom ObjectLinkingLayer plugin for init/fini section handling.
+//
+// Collects function pointers from init/fini sections (.init_array, .fini_array,
+// .ctors, .dtors, .CRT$XC*, .CRT$XT*) and runs them in priority order.
+//
+// Three-platform init/fini strategy:
+//
+// macOS: MachOPlatform (via orc_rt) handles __mod_init_func/__mod_term_func
+// and __cxa_atexit natively. We delegate to jit_->initialize() and
+// deinitialize(). No InitFiniPlugin needed.
+//
+// Windows: COFFPlatform is unusable -- it requires MSVC CRT symbols
+// (_CxxThrowException, RTTI vtables, etc.) that LLVM's COFF ORC
+// runtime cannot provide (stalled for 2+ years). Our plugin handles
+// .CRT$XC*/.CRT$XT* sections instead.
+//
+// Linux: ELFNixPlatform does not handle .init_array/.fini_array correctly
+// prior to llvm/llvm-project#175981. Once a new LLVM release includes
+// that patch, we can switch Linux to ELFNixPlatform and remove the
+// plugin for that platform.
+//
+// The plugin is compiled only on Linux/Windows and can be removed per-platform
+// as LLVM's native platform support matures.
+class InitFiniPlugin : public llvm::orc::ObjectLinkingLayer::Plugin {
+ // Store a raw pointer to avoid a reference cycle:
+ // Session → LLJIT → ObjectLinkingLayer → Plugin → Session
+ // The plugin's lifetime is bounded by the ObjectLinkingLayer which is
+ // owned by LLJIT which is owned by the session, so the pointer is always valid.
+ ORCJITExecutionSessionObj* session_;
+
+ public:
+ explicit InitFiniPlugin(ORCJITExecutionSessionObj* session) : session_(session) {}
+
+ void modifyPassConfig(llvm::orc::MaterializationResponsibility& MR, llvm::jitlink::LinkGraph& G,
+ llvm::jitlink::PassConfiguration& Config) override {
+ auto& jit_dylib = MR.getTargetJITDylib();
+ // Mark all init/fini section blocks and their edge targets as live
+ // so they survive dead-stripping.
+ Config.PrePrunePasses.emplace_back([](llvm::jitlink::LinkGraph& G) {
+ for (auto& Section : G.sections()) {
+ auto section_name = Section.getName();
+ // ELF: .init_array*, .fini_array*, .ctors*, .dtors*
+ // Mach-O: __DATA,__mod_init_func, __DATA,__mod_term_func
+ // COFF: .CRT$XC* (ctors), .CRT$XT* (dtors)
+ if (section_name.starts_with(".init_array") || section_name.starts_with(".fini_array") ||
+ section_name.starts_with(".ctors") || section_name.starts_with(".dtors") ||
+ section_name == "__DATA,__mod_init_func" || section_name == "__DATA,__mod_term_func" ||
+ section_name.starts_with(".CRT$XC") || section_name.starts_with(".CRT$XT")) {
+ for (auto* Block : Section.blocks()) {
+ bool has_live_sym = false;
+ for (auto* Sym : G.defined_symbols()) {
+ if (&Sym->getBlock() == Block) {
+ Sym->setLive(true);
+ has_live_sym = true;
+ }
+ }
+ // MSVC may emit .CRT$XC* blocks with data but no symbol table
+ // entries (static variables in __declspec(allocate) sections).
+ // Add an anonymous symbol so the block survives dead-stripping.
+ if (!has_live_sym) {
+ G.addAnonymousSymbol(*Block, 0, Block->getSize(), false, true).setLive(true);
+ }
+ for (auto& Edge : Block->edges()) {
+ Edge.getTarget().setLive(true);
+ }
+ }
+ }
+ }
+ return llvm::Error::success();
+ });
+#ifdef _WIN32
+ // Without COFFPlatform, __ImageBase (used by IMAGE_REL_AMD64_ADDR32NB /
+ // Pointer32NB relocations) defaults to 0. This causes all Pointer32 fixups
+ // to overflow since JIT addresses don't fit in 32 bits.
+ //
+ // Fix: set __ImageBase to the lowest block address in the graph after
+ // allocation. This makes all intra-graph Pointer32NB offsets small.
+ //
+ // Also strip .pdata/.xdata edges: SEH unwind data references external
+ // handlers (e.g., __CxxFrameHandler3) in DLLs that may be >4GB from
+ // __ImageBase. Since we don't call RtlAddFunctionTable, SEH data is
+ // unused anyway.
+ Config.PostAllocationPasses.emplace_back([](llvm::jitlink::LinkGraph& G) {
+ // Set __ImageBase to the lowest allocated block address.
+ auto ImageBaseName = G.intern("__ImageBase");
+ llvm::jitlink::Symbol* ImageBase = nullptr;
+ for (auto* Sym : G.external_symbols()) {
+ if (Sym->getName() == ImageBaseName) {
+ ImageBase = Sym;
+ break;
+ }
+ }
+ if (ImageBase) {
+ llvm::orc::ExecutorAddr BaseAddr;
+ for (auto* B : G.blocks()) {
+ if (!BaseAddr || B->getAddress() < BaseAddr) {
+ BaseAddr = B->getAddress();
+ }
+ }
+ ImageBase->getAddressable().setAddress(BaseAddr);
+ }
+ // Strip .pdata/.xdata edges: external handlers may be >4GB from
+ // __ImageBase, and we don't register SEH data anyway.
+ for (auto& Sec : G.sections()) {
+ if (Sec.getName() == ".pdata" || Sec.getName().starts_with(".xdata")) {
+ for (auto* B : Sec.blocks()) {
+ while (!B->edges_empty()) {
+ B->removeEdge(B->edges().begin());
+ }
+ }
+ }
+ }
+ return llvm::Error::success();
+ });
+#endif
+ // After fixups, read resolved function pointers from all init/fini data sections.
+ // Handles ELF (.init_array, .ctors, .fini_array, .dtors),
+ // Mach-O (__DATA,__mod_init_func, __DATA,__mod_term_func),
+ // and COFF (.CRT$XC*, .CRT$XT*) section conventions.
+ Config.PostFixupPasses.emplace_back([this, &jit_dylib](llvm::jitlink::LinkGraph& G) {
+ using Entry = ORCJITExecutionSessionObj::InitFiniEntry;
+ for (auto& Sec : G.sections()) {
+ auto section_name = Sec.getName();
+
+ // --- ELF sections ---
+ bool is_init_array = section_name.starts_with(".init_array");
+ bool is_ctors = section_name.starts_with(".ctors");
+ bool is_fini_array = section_name.starts_with(".fini_array");
+ bool is_dtors = section_name.starts_with(".dtors");
+ // --- Mach-O sections ---
+ bool is_mod_init = (section_name == "__DATA,__mod_init_func");
+ bool is_mod_term = (section_name == "__DATA,__mod_term_func");
+ // --- COFF sections ---
+ bool is_crt_xc = section_name.starts_with(".CRT$XC");
+ bool is_crt_xt = section_name.starts_with(".CRT$XT");
+
+ if (!is_init_array && !is_ctors && !is_fini_array && !is_dtors && !is_mod_init &&
+ !is_mod_term && !is_crt_xc && !is_crt_xt)
+ continue;
+
+ int priority = 0;
+ Entry::Section sec;
+ bool is_init;
+ // ELF default priority for sections without a numeric suffix is 65535.
+ // Lower priority numbers run first for .init_array; .fini_array and .ctors
+ // negate so that higher-numbered entries run first (reverse order).
+ if (is_init_array) {
+ if (section_name.consume_front(".init_array.")) {
+ section_name.getAsInteger(10, priority);
+ } else {
+ priority = 65535;
+ }
+ sec = Entry::Section::kInitArray;
+ is_init = true;
+ } else if (is_ctors) {
+ if (section_name.consume_front(".ctors.") && !section_name.getAsInteger(10, priority)) {
+ priority = -priority;
+ }
+ sec = Entry::Section::kCtors;
+ is_init = true;
+ } else if (is_fini_array) {
+ if (section_name.consume_front(".fini_array.") &&
+ !section_name.getAsInteger(10, priority)) {
+ priority = -priority;
+ } else {
+ priority = -65535;
+ }
+ sec = Entry::Section::kFiniArray;
+ is_init = false;
+ } else if (is_dtors) {
+ if (section_name.consume_front(".dtors.")) {
+ section_name.getAsInteger(10, priority);
+ }
+ sec = Entry::Section::kDtors;
+ is_init = false;
+ } else if (is_mod_init) {
+ // Mach-O __mod_init_func: no priority system, treated as init_array
+ sec = Entry::Section::kInitArray;
+ is_init = true;
+ } else if (is_mod_term) {
+ // Mach-O __mod_term_func: no priority system, treated as fini_array
+ sec = Entry::Section::kFiniArray;
+ is_init = false;
+ } else if (is_crt_xc) {
+ // COFF .CRT$XC[suffix]: C++ constructors, sorted alphabetically by suffix.
+ // Convert suffix to integer priority that preserves alphabetical ordering.
+ // E.g., .CRT$XCA → 'A'*100000=6500000, .CRT$XCU → 'U'*100000=8500000,
+ // .CRT$XCT00200 → 'T'*100000+200=8400200
+ sec = Entry::Section::kInitArray;
+ is_init = true;
+ auto suffix = section_name.substr(7); // after ".CRT$XC"
+ if (!suffix.empty()) {
+ priority = static_cast<int>(suffix[0]) * 100000;
+ if (suffix.size() > 1) {
+ int num = 0;
+ suffix.substr(1).getAsInteger(10, num);
+ priority += num;
+ }
+ }
+ } else {
+ // COFF .CRT$XT[suffix]: C++ destructors, same suffix-to-priority scheme.
+ sec = Entry::Section::kFiniArray;
+ is_init = false;
+ auto suffix = section_name.substr(7); // after ".CRT$XT"
+ if (!suffix.empty()) {
+ priority = static_cast<int>(suffix[0]) * 100000;
+ if (suffix.size() > 1) {
+ int num = 0;
+ suffix.substr(1).getAsInteger(10, num);
+ priority += num;
+ }
+ }
+ }
+
+ for (auto* Block : Sec.blocks()) {
+ auto Content = Block->getContent();
+ size_t PtrSize = G.getPointerSize();
+ for (size_t Offset = 0; Offset + PtrSize <= Content.size(); Offset += PtrSize) {
+ uint64_t FnAddr = 0;
+ memcpy(&FnAddr, Content.data() + Offset, PtrSize);
+ if (FnAddr != 0) {
+ Entry entry{llvm::orc::ExecutorAddr(FnAddr), sec, priority};
+ if (is_init) {
+ session_->AddPendingInitializer(&jit_dylib, entry);
+ } else {
+ session_->AddPendingDeinitializer(&jit_dylib, entry);
+ }
+ }
+ }
+ }
+ }
+ return llvm::Error::success();
+ });
+ }
+
+ llvm::Error notifyFailed(llvm::orc::MaterializationResponsibility& MR) override {
+ return llvm::Error::success();
+ }
+
+ llvm::Error notifyRemovingResources(llvm::orc::JITDylib& JD, llvm::orc::ResourceKey K) override {
+ return llvm::Error::success();
+ }
+
+ void notifyTransferringResources(llvm::orc::JITDylib& JD, llvm::orc::ResourceKey DstKey,
+ llvm::orc::ResourceKey SrcKey) override {}
+};
+
+#ifdef _WIN32
+/*!
+ * \brief Custom definition generator for Windows DLL import symbols.
+ *
+ * On Windows with the MSVC ABI, COFF objects reference DLL-imported symbols
+ * via __imp_XXX pointer stubs and direct calls. Without COFFPlatform (which
+ * we skip due to MSVC CRT dependency issues), JITLink cannot resolve these.
+ *
+ * For each resolved symbol, this generator creates a JIT-allocated LinkGraph
+ * containing:
+ * - __imp_XXX: a GOT-like pointer entry holding the real address
+ * - XXX: a PLT-like jump stub (`jmp [__imp_XXX]`) for direct calls
+ *
+ * By allocating stubs in JIT memory (rather than using absoluteSymbols at
+ * host-process addresses), all PCRel32 fixups from JIT code reach safely.
+ *
+ * Symbol search order:
+ * 1. Specific MSVC runtime DLLs (vcruntime140, ucrtbase, msvcp140)
+ * 2. All loaded process modules (EnumProcessModules)
+ * 3. LLVM's SearchForAddressOfSymbol
+ */
+class DLLImportDefinitionGenerator : public llvm::orc::DefinitionGenerator {
+ llvm::orc::ExecutionSession& ES_;
+ llvm::orc::ObjectLinkingLayer& L_;
+
+ static void* FindInProcessModules(const std::string& Name) {
+ // Try specific runtime DLLs first, then tvm_ffi.dll (loaded by Python),
+ // then all process modules, then LLVM's search.
+ static const char* kRuntimeDLLs[] = {
+ "vcruntime140.dll",
+ "vcruntime140_1.dll",
+ "ucrtbase.dll",
+ "msvcp140.dll",
+ };
+ // NOTE: We intentionally do not call FreeLibrary() here. These runtime DLLs
+ // (vcruntime140, ucrtbase, etc.) are already loaded by the process and will
+ // remain loaded for its lifetime. LoadLibraryA merely increments the refcount;
+ // the extra refcount is harmless and avoids the overhead of balancing
+ // Get/FreeLibrary for every symbol lookup.
+ for (const char* dll : kRuntimeDLLs) {
+ if (HMODULE hMod = LoadLibraryA(dll)) {
+ if (auto addr = GetProcAddress(hMod, Name.c_str())) {
+ return reinterpret_cast<void*>(addr);
+ }
+ }
+ }
+ // Also check tvm_ffi.dll (host process symbol provider)
+ if (HMODULE hTvmFfi = GetModuleHandleA("tvm_ffi.dll")) {
+ if (auto addr = GetProcAddress(hTvmFfi, Name.c_str())) {
+ return reinterpret_cast<void*>(addr);
+ }
+ }
+ HMODULE hMods[1024];
+ DWORD cbNeeded;
+ if (EnumProcessModules(GetCurrentProcess(), hMods, sizeof(hMods), &cbNeeded)) {
+ DWORD count = cbNeeded / sizeof(HMODULE);
+ if (count > 1024) count = 1024;
+ for (DWORD i = 0; i < count; ++i) {
+ if (auto addr = GetProcAddress(hMods[i], Name.c_str())) {
+ return reinterpret_cast<void*>(addr);
+ }
+ }
+ }
+ if (void* addr = llvm::sys::DynamicLibrary::SearchForAddressOfSymbol(Name)) {
+ return addr;
+ }
+ return nullptr;
+ }
+
+ public:
+ DLLImportDefinitionGenerator(llvm::orc::ExecutionSession& ES, llvm::orc::ObjectLinkingLayer& L)
+ : ES_(ES), L_(L) {}
+
+ llvm::Error tryToGenerate(llvm::orc::LookupState& LS, llvm::orc::LookupKind K,
+ llvm::orc::JITDylib& JD, llvm::orc::JITDylibLookupFlags JDLookupFlags,
+ const llvm::orc::SymbolLookupSet& LookupSet) override {
+ // Step 1: Collect unique base names (strip __imp_ prefix) and resolve addresses.
+ llvm::DenseMap<llvm::orc::SymbolStringPtr, llvm::orc::ExecutorAddr> Resolved;
+ for (auto& [Name, Flags] : LookupSet) {
+ llvm::StringRef NameStr = *Name;
+ std::string BaseName =
+ NameStr.starts_with("__imp_") ? NameStr.drop_front(6).str() : NameStr.str();
+ if (BaseName == "__ImageBase") continue;
+ auto InternedBase = ES_.intern(BaseName);
+ if (Resolved.count(InternedBase)) continue;
+ void* Addr = FindInProcessModules(BaseName);
+ if (Addr) {
+ Resolved[InternedBase] = llvm::orc::ExecutorAddr::fromPtr(Addr);
+ }
+ }
+ if (Resolved.empty()) return llvm::Error::success();
+
+ // Step 2: Build a LinkGraph with __imp_ pointers and PLT jump stubs.
+ auto G = std::make_unique<llvm::jitlink::LinkGraph>(
+ "<DLL_IMPORT_STUBS>", ES_.getSymbolStringPool(), ES_.getTargetTriple(),
+ llvm::SubtargetFeatures(), llvm::jitlink::getGenericEdgeKindName);
+ auto Prot = static_cast<llvm::orc::MemProt>(static_cast<unsigned>(llvm::orc::MemProt::Read) |
+ static_cast<unsigned>(llvm::orc::MemProt::Exec));
+ auto& Sec = G->createSection("__dll_stubs", Prot);
+
+ for (auto& [InternedName, Addr] : Resolved) {
+ // Absolute symbol at the real address (local to this graph)
+ auto& Target = G->addAbsoluteSymbol(G->intern(("__real_" + *InternedName).str()), Addr,
+ G->getPointerSize(), llvm::jitlink::Linkage::Strong,
+ llvm::jitlink::Scope::Local, false);
+ // __imp_XXX pointer (GOT-like entry)
+ auto& Ptr = llvm::jitlink::x86_64::createAnonymousPointer(*G, Sec, &Target);
+ Ptr.setName(G->intern(("__imp_" + *InternedName).str()));
+ Ptr.setLinkage(llvm::jitlink::Linkage::Strong);
+ Ptr.setScope(llvm::jitlink::Scope::Default);
+ // XXX jump stub (PLT-like entry) for direct calls
+ auto& StubBlock = llvm::jitlink::x86_64::createPointerJumpStubBlock(*G, Sec, Ptr);
+ G->addDefinedSymbol(StubBlock, 0, *InternedName, StubBlock.getSize(),
+ llvm::jitlink::Linkage::Strong, llvm::jitlink::Scope::Default, true,
+ false);
+ }
+ return L_.add(JD, std::move(G));
+ }
+};
+#endif // _WIN32
+
+ORCJITExecutionSessionObj::ORCJITExecutionSessionObj(const std::string& orc_rt_path)
+ : jit_(nullptr) {
+ // Helper: force JITLink's ObjectLinkingLayer on platforms where
+ // the default RTDyldObjectLinkingLayer won't work.
+ //
+ // macOS: MachOPlatform (via ExecutorNativePlatform) requires ObjectLinkingLayer.
+ //
+ // Windows: LLJIT defaults to RTDyldObjectLinkingLayer for COFF x86_64
+ // (see LLJIT.cpp, LLJITBuilderState::prepareForConstruction). We need
+ // ObjectLinkingLayer because:
+ // 1. Our InitFiniPlugin inherits ObjectLinkingLayer::Plugin — RTDyld has
+ // no plugin API, so the static_cast<ObjectLinkingLayer&> would crash.
+ // 2. We skip the ORC runtime on Windows (COFFPlatform requires MSVC CRT
+ // symbols like _CxxThrowException, RTTI vtables, iostream objects that
+ // are not resolvable in the JIT), so we handle .CRT$XC*/.CRT$XT*
+ // init/fini sections ourselves via the plugin.
+ //
+ // Linux: LLJIT already defaults to ObjectLinkingLayer for ELF, no override needed.
+ auto setup_builder = [](llvm::orc::LLJITBuilder& builder) {
+#if defined(__APPLE__) || defined(_WIN32)
+ builder.setObjectLinkingLayerCreator(
+ [](llvm::orc::ExecutionSession& ES)
+ -> llvm::Expected<std::unique_ptr<llvm::orc::ObjectLayer>> {
+ return std::make_unique<llvm::orc::ObjectLinkingLayer>(ES);
+ });
+#endif
+#ifdef _WIN32
+ // Override ProcessSymbols setup to NOT add the default
+ // EPCDynamicLibrarySearchGenerator. That generator resolves symbols to
+ // absolute host-process addresses, which causes PCRel32 overflow when
+ // JIT code calls into DLLs >2GB away. Our DLLImportDefinitionGenerator
+ // (added after construction) wraps every resolved address in a
+ // JIT-allocated PLT stub, keeping all fixups in range.
+ builder.setProcessSymbolsJITDylibSetup(
+ [](llvm::orc::LLJIT& J) -> llvm::Expected<llvm::orc::JITDylibSP> {
+ return &J.getExecutionSession().createBareJITDylib("<Process Symbols>");
+ });
+#endif
+ (void)builder;
+ };
+
+ if (!orc_rt_path.empty()) {
+ auto builder = llvm::orc::LLJITBuilder();
+ builder.setPlatformSetUp(llvm::orc::ExecutorNativePlatform(orc_rt_path));
+ setup_builder(builder);
+ jit_ = TVM_FFI_ORCJIT_LLVM_CALL(builder.create());
+ } else {
+ auto builder = llvm::orc::LLJITBuilder();
+ setup_builder(builder);
+ jit_ = TVM_FFI_ORCJIT_LLVM_CALL(builder.create());
+ }
+#ifdef _WIN32
+ // Strip .pdata/.xdata relocations from COFF objects before JITLink graph building.
+ // clang-cl puts static functions in COMDAT sections, and .pdata SEH unwind data
+ // has relocations targeting COMDAT leader symbols. JITLink's COFFLinkGraphBuilder
+ // doesn't register COMDAT leaders in its symbol table when the second COMDAT symbol
+ // is CLASS_STATIC (not CLASS_EXTERNAL), causing "Could not find symbol" errors.
+ // We already strip .pdata/.xdata edges in a PostAllocationPass; this moves the
+ // stripping earlier to prevent the graph builder error.
+ // Strip .pdata/.xdata relocations using raw COFF binary manipulation.
+ // We avoid llvm/Object/COFF.h because windows.h (included transitively by
+ // LLJIT.h) defines IMAGE_* macros that conflict with LLVM's COFF enums.
+ jit_->getObjTransformLayer().setTransform(
+ [](std::unique_ptr<llvm::MemoryBuffer> Buf)
+ -> llvm::Expected<std::unique_ptr<llvm::MemoryBuffer>> {
+ const char* Data = Buf->getBufferStart();
+ size_t Size = Buf->getBufferSize();
+ if (Size < 20) return std::move(Buf);
+
+ // Parse COFF header (regular or bigobj format)
+ uint16_t w0, w1;
+ std::memcpy(&w0, Data, 2);
+ std::memcpy(&w1, Data + 2, 2);
+ bool bigobj = (w0 == 0 && w1 == 0xFFFF);
+
+ uint16_t machine;
+ uint32_t num_sections, ptr_to_symtab, num_symbols;
+ size_t sec_hdr_start, sym_entry_size;
+ if (bigobj) {
+ if (Size < 56) return std::move(Buf);
+ std::memcpy(&machine, Data + 6, 2);
+ std::memcpy(&num_sections, Data + 44, 4);
+ std::memcpy(&ptr_to_symtab, Data + 48, 4);
+ std::memcpy(&num_symbols, Data + 52, 4);
+ sec_hdr_start = 56;
+ sym_entry_size = 20;
+ } else {
+ machine = w0;
+ uint16_t ns, opt_hdr_size;
+ std::memcpy(&ns, Data + 2, 2);
+ std::memcpy(&opt_hdr_size, Data + 16, 2);
+ std::memcpy(&ptr_to_symtab, Data + 8, 4);
+ std::memcpy(&num_symbols, Data + 12, 4);
+ num_sections = ns;
+ sec_hdr_start = 20 + opt_hdr_size;
+ sym_entry_size = 18;
+ }
+ if (machine != 0x8664) return std::move(Buf);
+
+ // String table follows the symbol table
+ size_t strtab_start = ptr_to_symtab + static_cast<size_t>(num_symbols) * sym_entry_size;
+
+ // Resolve a section name (inline 8-byte or "/offset" string table ref)
+ constexpr size_t kSecHdrSize = 40;
+ auto resolve_name = [&](size_t hdr_off) -> llvm::StringRef {
+ const char* raw = Data + hdr_off;
+ if (raw[0] == '/' && raw[1] >= '0' && raw[1] <= '9') {
+ uint32_t offset = 0;
+ for (int j = 1; j < 8 && raw[j] >= '0' && raw[j] <= '9'; ++j)
+ offset = offset * 10 + (raw[j] - '0');
+ size_t pos = strtab_start + offset;
+ if (pos < Size) {
+ size_t len = 0;
+ while (pos + len < Size && Data[pos + len]) ++len;
+ return {Data + pos, len};
+ }
+ }
+ size_t len = 0;
+ while (len < 8 && raw[len]) ++len;
+ return {raw, len};
+ };
+
+ // Collect section header offsets needing relocation stripping
+ llvm::SmallVector<size_t, 8> strip_offsets;
+ for (uint32_t i = 0; i < num_sections; ++i) {
+ size_t off = sec_hdr_start + i * kSecHdrSize;
+ if (off + kSecHdrSize > Size) break;
+ auto name = resolve_name(off);
+ if (name.starts_with(".pdata") || name.starts_with(".xdata")) {
+ uint16_t num_relocs;
+ std::memcpy(&num_relocs, Data + off + 32, 2);
+ if (num_relocs > 0) strip_offsets.push_back(off);
+ }
+ }
+ if (strip_offsets.empty()) return std::move(Buf);
+
+ // Create mutable copy, zero out PointerToRelocations and NumberOfRelocations
+ llvm::SmallVector<char> MutableBuf(Data, Data + Size);
+ for (auto off : strip_offsets) {
+ std::memset(&MutableBuf[off + 24], 0, 4); // PointerToRelocations
+ std::memset(&MutableBuf[off + 32], 0, 2); // NumberOfRelocations
+ }
+ return llvm::MemoryBuffer::getMemBufferCopy(
+ llvm::StringRef(MutableBuf.data(), MutableBuf.size()), Buf->getBufferIdentifier());
+ });
+#endif
+#if defined(__linux__) || defined(_WIN32)
+ // Linux/Windows: use our custom InitFiniPlugin for init/fini section
+ // collection and priority-ordered execution. See InitFiniPlugin class
+ // documentation for the three-platform init/fini strategy.
+ auto& objlayer = jit_->getObjLinkingLayer();
+ static_cast<llvm::orc::ObjectLinkingLayer&>(objlayer).addPlugin(
+ std::make_unique<InitFiniPlugin>(this));
+#endif
+#ifdef _WIN32
+ // On Windows, the default process-symbol generator only searches the main
+ // exe module via GetProcAddress(GetModuleHandle(NULL), ...). Add a
+ // comprehensive generator that searches all loaded DLLs (vcruntime140,
+ // ucrtbase, tvm_ffi, etc.) and creates __imp_* pointer stubs.
+ if (auto PSG = jit_->getProcessSymbolsJITDylib()) {
+ auto& ObjLayer = static_cast<llvm::orc::ObjectLinkingLayer&>(jit_->getObjLinkingLayer());
+ PSG->addGenerator(
+ std::make_unique<DLLImportDefinitionGenerator>(jit_->getExecutionSession(), ObjLayer));
+ }
+#endif
+}
+
+ORCJITExecutionSession::ORCJITExecutionSession(const std::string& orc_rt_path) {
+ ObjectPtr<ORCJITExecutionSessionObj> obj = make_object<ORCJITExecutionSessionObj>(orc_rt_path);
+ data_ = std::move(obj);
+}
+
+ORCJITDynamicLibrary ORCJITExecutionSessionObj::CreateDynamicLibrary(const String& name) {
+ TVM_FFI_CHECK(jit_ != nullptr, InternalError) << "ExecutionSession not initialized";
+
+ // Generate name if not provided
+ String lib_name = name;
+ if (lib_name.empty()) {
+ std::ostringstream oss;
+ oss << "dylib_" << dylib_counter_++;
+ lib_name = oss.str();
+ }
+
+ llvm::orc::JITDylib& jit_dylib =
+ TVM_FFI_ORCJIT_LLVM_CALL(jit_->getExecutionSession().createJITDylib(lib_name.c_str()));
+ // Use the LLJIT's default link order (Main → Platform → ProcessSymbols).
+ // This provides host process symbols via the ProcessSymbols JITDylib's generator,
+ // while ensuring the platform's __cxa_atexit interposer (in PlatformJD) takes
+ // precedence — so __cxa_atexit handlers are managed by the platform and can be
+ // drained per-JITDylib via __lljit_run_atexits at teardown.
+ for (auto& kv : jit_->defaultLinkOrder()) {
+ jit_dylib.addToLinkOrder(*kv.first, kv.second);
+ }
+
+ auto dylib = ORCJITDynamicLibrary(make_object<ORCJITDynamicLibraryObj>(
+ GetRef<ORCJITExecutionSession>(this), &jit_dylib, jit_.get(), lib_name));
+
+ return dylib;
+}
+
+llvm::orc::ExecutionSession& ORCJITExecutionSessionObj::GetLLVMExecutionSession() {
+ TVM_FFI_CHECK(jit_ != nullptr, InternalError) << "ExecutionSession not initialized";
+ return jit_->getExecutionSession();
+}
+
+llvm::orc::LLJIT& ORCJITExecutionSessionObj::GetLLJIT() {
+ TVM_FFI_CHECK(jit_ != nullptr, InternalError) << "ExecutionSession not initialized";
+ return *jit_;
+}
+
+using CtorDtor = void (*)();
+
+void ORCJITExecutionSessionObj::RunPendingInitializers(llvm::orc::JITDylib& jit_dylib) {
+ auto it = pending_initializers_.find(&jit_dylib);
+ if (it != pending_initializers_.end()) {
+ llvm::sort(it->second, [](const InitFiniEntry& a, const InitFiniEntry& b) {
+ if (a.section != b.section) return static_cast<int>(a.section) < static_cast<int>(b.section);
+ return a.priority < b.priority;
+ });
+ for (const auto& entry : it->second) {
+ entry.address.toPtr<CtorDtor>()();
+ }
+ pending_initializers_.erase(it);
+ }
+}
+
+void ORCJITExecutionSessionObj::RunPendingDeinitializers(llvm::orc::JITDylib& jit_dylib) {
+ auto it = pending_deinitializers_.find(&jit_dylib);
+ if (it != pending_deinitializers_.end()) {
+ llvm::sort(it->second, [](const InitFiniEntry& a, const InitFiniEntry& b) {
+ if (a.section != b.section) return static_cast<int>(a.section) < static_cast<int>(b.section);
+ return a.priority < b.priority;
+ });
+ for (const auto& entry : it->second) {
+ entry.address.toPtr<CtorDtor>()();
+ }
+ pending_deinitializers_.erase(it);
+ }
+}
+
+void ORCJITExecutionSessionObj::AddPendingInitializer(llvm::orc::JITDylib* jit_dylib,
+ const InitFiniEntry& entry) {
+ pending_initializers_[jit_dylib].push_back(entry);
+}
+
+void ORCJITExecutionSessionObj::AddPendingDeinitializer(llvm::orc::JITDylib* jit_dylib,
+ const InitFiniEntry& entry) {
+ pending_deinitializers_[jit_dylib].push_back(entry);
+}
+
+} // namespace orcjit
+} // namespace ffi
+} // namespace tvm
diff --git a/addons/tvm_ffi_orcjit/src/ffi/orcjit_session.h b/addons/tvm_ffi_orcjit/src/ffi/orcjit_session.h
new file mode 100644
index 0000000..cacffcd
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/src/ffi/orcjit_session.h
@@ -0,0 +1,129 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*!
+ * \file orcjit_session.h
+ * \brief LLVM ORC JIT ExecutionSession wrapper
+ */
+#ifndef TVM_FFI_ORCJIT_ORCJIT_SESSION_H_
+#define TVM_FFI_ORCJIT_ORCJIT_SESSION_H_
+
+#include <llvm/ExecutionEngine/Orc/Core.h>
+#include <llvm/ExecutionEngine/Orc/LLJIT.h>
+#include <llvm/ExecutionEngine/Orc/Shared/ExecutorAddress.h>
+#include <tvm/ffi/object.h>
+#include <tvm/ffi/string.h>
+
+#include <atomic>
+#include <string>
+#include <unordered_map>
+
+namespace tvm {
+namespace ffi {
+namespace orcjit {
+
+// Forward declaration
+class ORCJITDynamicLibrary;
+
+/*!
+ * \brief ExecutionSession object for LLVM ORC JIT v2
+ *
+ * This class manages the lifetime of an LLVM ExecutionSession and provides
+ * functionality to create and manage multiple JITDylibs (DynamicLibraries).
+ */
+class ORCJITExecutionSessionObj : public Object {
+ public:
+ /*!
+ * \brief Default constructor (for make_object)
+ */
+ explicit ORCJITExecutionSessionObj(const std::string& orc_rt_path = "");
+
+ /*!
+ * \brief Create a new DynamicLibrary (JITDylib) in this session
+ * \param name Optional name for the library (for debugging)
+ * \return The created dynamic library instance
+ */
+ ORCJITDynamicLibrary CreateDynamicLibrary(const String& name);
+
+ /*!
+ * \brief Get the underlying LLVM ExecutionSession
+ * \return Reference to the LLVM ExecutionSession
+ */
+ llvm::orc::ExecutionSession& GetLLVMExecutionSession();
+
+ /*!
+ * \brief Get the underlying LLJIT instance
+ * \return Reference to the LLJIT instance
+ */
+ llvm::orc::LLJIT& GetLLJIT();
+
+ static constexpr bool _type_mutable = true;
+ TVM_FFI_DECLARE_OBJECT_INFO_FINAL("orcjit.ExecutionSession", ORCJITExecutionSessionObj, Object);
+
+ struct InitFiniEntry {
+ enum class Section {
+ kInitArray = 0,
+ kCtors = 1,
+ kDtors = 2,
+ kFiniArray = 3,
+ };
+ llvm::orc::ExecutorAddr address;
+ Section section;
+ int priority;
+ };
+
+ void RunPendingInitializers(llvm::orc::JITDylib& jit_dylib);
+ void RunPendingDeinitializers(llvm::orc::JITDylib& jit_dylib);
+
+ void AddPendingInitializer(llvm::orc::JITDylib* jd, const InitFiniEntry& entry);
+ void AddPendingDeinitializer(llvm::orc::JITDylib* jd, const InitFiniEntry& entry);
+
+ private:
+ /*! \brief The LLVM ORC JIT instance */
+ std::unique_ptr<llvm::orc::LLJIT> jit_;
+
+ /*! \brief Counter for auto-generating library names */
+ std::atomic<int> dylib_counter_{0};
+
+ std::unordered_map<llvm::orc::JITDylib*, std::vector<InitFiniEntry>> pending_initializers_;
+ std::unordered_map<llvm::orc::JITDylib*, std::vector<InitFiniEntry>> pending_deinitializers_;
+};
+
+/*!
+ * \brief Reference wrapper for ORCJITExecutionSessionObj
+ *
+ * A reference wrapper serves as a reference-counted pointer to the session object.
+ */
+class ORCJITExecutionSession : public ObjectRef {
+ public:
+ /*!
+ * \brief Create a new ExecutionSession
+ * \return The created execution session instance
+ */
+ explicit ORCJITExecutionSession(const std::string& orc_rt_path = "");
+ // Required: define object reference methods
+ TVM_FFI_DEFINE_OBJECT_REF_METHODS_NOTNULLABLE(ORCJITExecutionSession, ObjectRef,
+ ORCJITExecutionSessionObj);
+};
+
+} // namespace orcjit
+} // namespace ffi
+} // namespace tvm
+
+#endif // TVM_FFI_ORCJIT_ORCJIT_SESSION_H_
diff --git a/addons/tvm_ffi_orcjit/src/ffi/orcjit_utils.h b/addons/tvm_ffi_orcjit/src/ffi/orcjit_utils.h
new file mode 100644
index 0000000..a8d52bc
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/src/ffi/orcjit_utils.h
@@ -0,0 +1,90 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*!
+ * \file orcjit_utils.h
+ * \brief LLVM ORC JIT utility macros and helpers.
+ *
+ * Provides TVM_FFI_ORCJIT_LLVM_CALL, a macro that checks an LLVM
+ * Error or Expected<T> result and throws a TVM FFI InternalError
+ * on failure.
+ */
+#ifndef TVM_FFI_ORCJIT_ORCJIT_UTILS_H_
+#define TVM_FFI_ORCJIT_ORCJIT_UTILS_H_
+
+#include <llvm/ExecutionEngine/Orc/LLJIT.h>
+#include <llvm/Support/Error.h>
+#include <tvm/ffi/string.h>
+
+#include <string>
+
+namespace tvm {
+namespace ffi {
+namespace orcjit {
+namespace detail {
+
+/*! \brief Check an llvm::Error and throw on failure. */
+inline void CallLLVM(llvm::Error err, const char* expr, const char* file, int line) {
+ if (err) {
+ std::string err_msg;
+ llvm::handleAllErrors(std::move(err),
+ [&](const llvm::ErrorInfoBase& eib) { err_msg = eib.message(); });
+ TVM_FFI_THROW(InternalError) << file << ":" << line << ": " << expr << " failed: " << err_msg;
+ }
+}
+
+/*! \brief Check an llvm::Expected<T> and return the value or throw on failure. */
+template <typename T>
+inline T CallLLVM(llvm::Expected<T> value_or_err, const char* expr, const char* file, int line) {
+ if (value_or_err) return std::move(*value_or_err);
+
+ std::string err_msg;
+ llvm::handleAllErrors(std::move(value_or_err.takeError()),
+ [&](const llvm::ErrorInfoBase& eib) { err_msg = eib.message(); });
+ TVM_FFI_THROW(InternalError) << file << ":" << line << ": " << expr << " failed: " << err_msg;
+}
+
+/*! \brief Check an llvm::Expected<T&> and return the reference or throw on failure. */
+template <typename T>
+inline T& CallLLVM(llvm::Expected<T&> value_or_err, const char* expr, const char* file, int line) {
+ if (value_or_err) return *value_or_err;
+
+ std::string err_msg;
+ llvm::handleAllErrors(std::move(value_or_err.takeError()),
+ [&](const llvm::ErrorInfoBase& eib) { err_msg = eib.message(); });
+ TVM_FFI_THROW(InternalError) << file << ":" << line << ": " << expr << " failed: " << err_msg;
+}
+
+} // namespace detail
+} // namespace orcjit
+} // namespace ffi
+} // namespace tvm
+
+/*!
+ * \brief Check an LLVM Error or Expected<T> and throw on failure.
+ *
+ * Usage:
+ * TVM_FFI_ORCJIT_LLVM_CALL(builder.create());
+ * auto jit = TVM_FFI_ORCJIT_LLVM_CALL(builder.create());
+ * TVM_FFI_ORCJIT_LLVM_CALL(jit->addObjectFile(...));
+ */
+#define TVM_FFI_ORCJIT_LLVM_CALL(expr) \
+ ::tvm::ffi::orcjit::detail::CallLLVM((expr), #expr, __FILE__, __LINE__)
+
+#endif // TVM_FFI_ORCJIT_ORCJIT_UTILS_H_
diff --git a/addons/tvm_ffi_orcjit/tests/CMakeLists.txt b/addons/tvm_ffi_orcjit/tests/CMakeLists.txt
new file mode 100644
index 0000000..9259e1e
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/CMakeLists.txt
@@ -0,0 +1,107 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+
+cmake_minimum_required(VERSION 3.18)
+project(tvm_ffi_orcjit_tests)
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
+
+set(TEST_OBJ_INSTALL_SUFFIX
+ ""
+ CACHE STRING "Suffix for install subdirectory (e.g., -gcc)"
+)
+
+if (CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64|arm64")
+ add_compile_options(-mno-outline-atomics)
+endif ()
+
+# Find tvm-ffi package
+find_package(
+ Python
+ COMPONENTS Interpreter
+ REQUIRED
+)
+execute_process(
+ COMMAND "${Python_EXECUTABLE}" -m tvm_ffi.config --cmakedir
+ OUTPUT_STRIP_TRAILING_WHITESPACE
+ OUTPUT_VARIABLE tvm_ffi_ROOT
+)
+find_package(tvm_ffi CONFIG REQUIRED)
+
+# Helper: build a source file into an object and install as <subdir>/<name>.o Derives the target
+# name and install subdirectory from the source path.
+function (add_test_object source)
+ get_filename_component(_name "${source}" NAME_WE)
+ get_filename_component(_ext "${source}" EXT)
+ get_filename_component(_dir "${source}" DIRECTORY)
+ get_filename_component(_subdir "${_dir}" NAME)
+ set(target_name "${_subdir}_${_name}")
+ if (_ext STREQUAL ".c")
+ set(lang C)
+ else ()
+ set(lang CXX)
+ endif ()
+ add_library(${target_name}_obj OBJECT "${source}")
+ target_link_libraries(${target_name}_obj PRIVATE tvm_ffi::header)
+ if (lang STREQUAL "C" AND MSVC)
+ # /GS- disables security cookie checks (__security_cookie) which are CRT symbols the ORC JIT
+ # cannot resolve.
+ target_compile_options(${target_name}_obj PRIVATE /O2 /GS-)
+ else ()
+ target_compile_options(${target_name}_obj PRIVATE -O2)
+ endif ()
+ install(
+ FILES $<TARGET_OBJECTS:${target_name}_obj>
+ DESTINATION ${CMAKE_CURRENT_SOURCE_DIR}/${_subdir}${TEST_OBJ_INSTALL_SUFFIX}
+ RENAME ${_name}.o
+ )
+endfunction ()
+
+# C++ object files — skip on Windows (C-only strategy for ORC JIT on Windows)
+if (NOT WIN32)
+ add_test_object(sources/cc/test_funcs.cc)
+ add_test_object(sources/cc/test_funcs2.cc)
+ add_test_object(sources/cc/test_funcs_conflict.cc)
+ add_test_object(sources/cc/test_ctor_dtor.cc)
+ add_test_object(sources/cc/test_call_global.cc)
+ add_test_object(sources/cc/test_types.cc)
+ add_test_object(sources/cc/test_link_order_base.cc)
+ add_test_object(sources/cc/test_link_order_caller.cc)
+ add_test_object(sources/cc/test_error.cc)
+endif ()
+
+# Pure C object files — built on all platforms (no C++ runtime deps)
+enablelanguage(C)
+add_test_object(sources/c/test_funcs.c)
+add_test_object(sources/c/test_funcs2.c)
+add_test_object(sources/c/test_funcs_conflict.c)
+add_test_object(sources/c/test_call_global.c)
+add_test_object(sources/c/test_types.c)
+add_test_object(sources/c/test_link_order_base.c)
+add_test_object(sources/c/test_link_order_caller.c)
+add_test_object(sources/c/test_error.c)
+add_test_object(sources/c/test_ctor_dtor.c)
+
+# CUDA object files — optional
+find_package(CUDAToolkit)
+if (CUDAToolkit_FOUND)
+ enablelanguage(CUDA)
+ message(STATUS "CUDA found: ${CUDAToolkit_VERSION}")
+ add_test_object(sources/cuda/test_funcs.cu)
+endif ()
diff --git a/addons/tvm_ffi_orcjit/tests/README.md b/addons/tvm_ffi_orcjit/tests/README.md
new file mode 100644
index 0000000..0ab9a2d
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/README.md
@@ -0,0 +1,109 @@
+<!--- Licensed to the Apache Software Foundation (ASF) under one -->
+<!--- or more contributor license agreements. See the NOTICE file -->
+<!--- distributed with this work for additional information -->
+<!--- regarding copyright ownership. The ASF licenses this file -->
+<!--- to you 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. -->
+
+# TVM-FFI-OrcJIT Tests
+
+## Quick Start
+
+Run pytest directly — test objects are auto-built by `utils.build_test_objects()`
+into a temporary directory using `tvm_ffi.cpp.build`:
+
+```bash
+pytest tests/ -v
+```
+
+Run the quick-start example:
+
+```bash
+python examples/quick-start/run.py --lang c
+```
+
+### Alternative: cmake-based build
+
+A `CMakeLists.txt` is also provided for cmake-based workflows.
+See `TEST_OBJ_INSTALL_SUFFIX` for multi-compiler builds.
+
+## Tested Configurations
+
+Tests are parametrized over compiler variants. Each platform builds objects
+with multiple compilers; `test_basic.py` auto-discovers which are available.
+
+| Platform | Variant | Subdir | Compiler | Languages |
+| -------- | ------- | ------ | -------- | --------- |
+| Linux, macOS | default | `c/`, `cc/` | LLVM Clang | C, C++ |
+| Linux | gcc | `c-gcc/`, `cc-gcc/` | GCC | C, C++ |
+| macOS | appleclang | `c-appleclang/`, `cc-appleclang/` | Apple Clang (`/usr/bin/clang`) | C, C++ |
+| Windows | default | `c/` | LLVM Clang | C only |
+| Windows | msvc | `c-msvc/` | MSVC (`cl`) | C only |
+| Windows | clang-cl | `c-clang-cl/` | clang-cl | C only |
+
+All Windows variants are C-only. C++ objects use `try`/`catch`
+(via `TVM_FFI_SAFE_CALL_BEGIN/END`) which requires Itanium exception ABI
+symbols that the MSVC-built host process cannot resolve.
+
+### Platform-specific compiler flags
+
+| Flag | Applies to | Reason |
+| ---- | ---------- | ------ |
+| `/O2 /GS-` | MSVC, clang-cl | `/GS-` disables buffer security checks (`__security_cookie`) which are CRT symbols the ORC JIT cannot resolve |
+| `-O2` | All others | Standard optimization |
+| `-mno-outline-atomics` | aarch64 | Avoids libgcc outline atomics helper calls that the JIT cannot resolve |
+
+## Test Structure
+
+```text
+tests/
+ sources/
+ c/ C source files (.c)
+ cc/ C++ source files (.cc)
+ cuda/ CUDA source files (optional)
+ utils.py Build helpers (compile sources to objects in /tmp)
+ test_basic.py Python test cases
+ CMakeLists.txt Alternative cmake build
+```
+
+Built objects are placed in `{tempdir}/tvm_ffi_orcjit_tests/` with per-compiler
+variant subdirectories (c/, cc/, c-gcc/, etc.).
+
+### Source files
+
+| Source | Description |
+| ------ | ----------- |
+| `test_funcs` | Basic arithmetic (add, multiply) |
+| `test_funcs2` | More arithmetic (subtract, divide) |
+| `test_funcs_conflict` | Symbol conflict testing (duplicate `add`) |
+| `test_call_global` | Callbacks into Python-registered global functions |
+| `test_types` | Zero-arg, multi-arg, float, void return types |
+| `test_link_order_base` / `test_link_order_caller` | Cross-library symbol resolution |
+| `test_error` | Error propagation from JIT'd code |
+| `test_ctor_dtor` | Constructor/destructor and init/fini sections |
+
+### Constructor/destructor test coverage
+
+`test_ctor_dtor` verifies that the ORC JIT correctly handles platform-specific
+static initialization and finalization:
+
+| Platform | Init mechanism | Fini mechanism |
+| -------- | -------------- | -------------- |
+| Linux (ELF) | `.init_array` (with priorities), `.ctors` (reversed) | `.dtors` |
+| macOS (Mach-O) | `.init_array` (with priorities), `__mod_init_func` | `.fini_array` (with priorities) |
+| Windows (COFF) | `.CRT$XCA`..`XCU` (alphabetical order) | `.CRT$XTA`..`XTZ` (alphabetical order) |
+
+## CI
+
+The CI workflow (`.github/workflows/tvm_ffi_orcjit.yml`) runs `pytest` and the
+quick-start example via `cibuildwheel`.
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_call_global.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_call_global.c
new file mode 100644
index 0000000..21a35ea
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_call_global.c
@@ -0,0 +1,106 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Pure C test: look up a host-registered global function by name and call it.
+ * Demonstrates JIT code calling back into the host process via the TVM FFI C API.
+ */
+#include <tvm/ffi/c_api.h>
+
+/*
+ * test_call_global_add: look up "test_host_add" global function and call it
+ * with two integer arguments. Returns the result from the host function.
+ */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_call_global_add(void* self, const TVMFFIAny* args,
+ int32_t num_args, TVMFFIAny* result) {
+ int ret_code = 0;
+ TVMFFIByteArray func_name;
+ TVMFFIObjectHandle func_handle = NULL;
+ TVMFFIAny call_args[2];
+ TVMFFIAny call_result;
+
+ /* Look up the host-registered global function */
+ func_name.data = "test_host_add";
+ func_name.size = 13;
+ ret_code = TVMFFIFunctionGetGlobal(&func_name, &func_handle);
+ if (ret_code != 0) return ret_code;
+
+ /* Prepare arguments: pass through the two integer args */
+ call_args[0].type_index = kTVMFFIInt;
+ call_args[0].zero_padding = 0;
+ call_args[0].v_int64 = args[0].v_int64;
+ call_args[1].type_index = kTVMFFIInt;
+ call_args[1].zero_padding = 0;
+ call_args[1].v_int64 = args[1].v_int64;
+
+ /* Call the global function */
+ call_result.type_index = kTVMFFINone;
+ call_result.zero_padding = 0;
+ call_result.v_int64 = 0;
+ ret_code = TVMFFIFunctionCall(func_handle, call_args, 2, &call_result);
+
+ /* Release the function handle */
+ TVMFFIObjectDecRef((TVMFFIObject*)func_handle);
+
+ if (ret_code != 0) return ret_code;
+
+ /* Forward the result */
+ result->type_index = call_result.type_index;
+ result->zero_padding = call_result.zero_padding;
+ result->v_int64 = call_result.v_int64;
+ return 0;
+}
+
+/*
+ * test_call_global_mul: look up "test_host_multiply" global function and call it.
+ */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_call_global_mul(void* self, const TVMFFIAny* args,
+ int32_t num_args, TVMFFIAny* result) {
+ int ret_code = 0;
+ TVMFFIByteArray func_name;
+ TVMFFIObjectHandle func_handle = NULL;
+ TVMFFIAny call_args[2];
+ TVMFFIAny call_result;
+
+ func_name.data = "test_host_multiply";
+ func_name.size = 18;
+ ret_code = TVMFFIFunctionGetGlobal(&func_name, &func_handle);
+ if (ret_code != 0) return ret_code;
+
+ call_args[0].type_index = kTVMFFIInt;
+ call_args[0].zero_padding = 0;
+ call_args[0].v_int64 = args[0].v_int64;
+ call_args[1].type_index = kTVMFFIInt;
+ call_args[1].zero_padding = 0;
+ call_args[1].v_int64 = args[1].v_int64;
+
+ call_result.type_index = kTVMFFINone;
+ call_result.zero_padding = 0;
+ call_result.v_int64 = 0;
+ ret_code = TVMFFIFunctionCall(func_handle, call_args, 2, &call_result);
+
+ TVMFFIObjectDecRef((TVMFFIObject*)func_handle);
+
+ if (ret_code != 0) return ret_code;
+
+ result->type_index = call_result.type_index;
+ result->zero_padding = call_result.zero_padding;
+ result->v_int64 = call_result.v_int64;
+ return 0;
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_ctor_dtor.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_ctor_dtor.c
new file mode 100644
index 0000000..5af1ea6
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_ctor_dtor.c
@@ -0,0 +1,152 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Pure C version of the constructor/destructor test.
+ * Uses TVMFFIFunctionGetGlobal + TVMFFIFunctionCall to call the host
+ * "append_log" function from constructors and destructors.
+ *
+ * Platform-specific init/deinit mechanisms tested:
+ * ELF → .init_array/.fini_array (via __attribute__) + .ctors/.dtors
+ * Mach-O → __DATA,__mod_init_func (via __attribute__) + explicit section
+ * COFF → .CRT$XC* sections (init) + .CRT$XT* sections (term)
+ */
+#include <tvm/ffi/c_api.h>
+
+static void puts_log(const char* msg) {
+ TVMFFIByteArray func_name;
+ TVMFFIObjectHandle func_handle = NULL;
+ TVMFFIAny call_args[1];
+ TVMFFIAny call_result;
+
+ func_name.data = "append_log";
+ func_name.size = 10;
+ if (TVMFFIFunctionGetGlobal(&func_name, &func_handle) != 0) return;
+
+ call_args[0].type_index = kTVMFFIRawStr;
+ call_args[0].zero_padding = 0;
+ call_args[0].v_c_str = msg;
+
+ call_result.type_index = kTVMFFINone;
+ call_result.zero_padding = 0;
+ call_result.v_int64 = 0;
+ TVMFFIFunctionCall(func_handle, call_args, 1, &call_result);
+ TVMFFIObjectDecRef((TVMFFIObject*)func_handle);
+}
+
+typedef void (*ctor_t)(void);
+typedef void (*dtor_t)(void);
+
+/* =========================================================================
+ * GCC / Clang (ELF + Mach-O): __attribute__((constructor/destructor))
+ * ========================================================================= */
+#ifndef _MSC_VER
+
+__attribute__((constructor)) void init_array(void) { puts_log("<init_array>"); }
+__attribute__((constructor(101))) void init_array_101(void) { puts_log("<init_array.101>"); }
+__attribute__((constructor(102))) void init_array_102(void) { puts_log("<init_array.102>"); }
+__attribute__((constructor(103))) void init_array_103(void) { puts_log("<init_array.103>"); }
+
+__attribute__((destructor)) void fini_array(void) { puts_log("<fini_array>"); }
+__attribute__((destructor(101))) void fini_array_101(void) { puts_log("<fini_array.101>"); }
+__attribute__((destructor(102))) void fini_array_102(void) { puts_log("<fini_array.102>"); }
+__attribute__((destructor(103))) void fini_array_103(void) { puts_log("<fini_array.103>"); }
+
+/* ELF-specific: explicit .ctors/.dtors section placements with priorities.
+ * .ctors priorities are reversed: lower number = later execution. */
+#ifdef __ELF__
+static void ctors(void) { puts_log("<ctors>"); }
+__attribute__((section(".ctors"), used)) static ctor_t ctors_ptr = ctors;
+
+static void ctors_101(void) { puts_log("<ctors.101>"); }
+__attribute__((section(".ctors.101"), used)) static ctor_t ctors_1_ptr = ctors_101;
+
+static void ctors_102(void) { puts_log("<ctors.102>"); }
+__attribute__((section(".ctors.102"), used)) static ctor_t ctors_2_ptr = ctors_102;
+
+static void ctors_103(void) { puts_log("<ctors.103>"); }
+__attribute__((section(".ctors.103"), used)) static ctor_t ctors_3_ptr = ctors_103;
+
+static void dtors(void) { puts_log("<dtors>"); }
+__attribute__((section(".dtors"), used)) static dtor_t dtors_ptr = dtors;
+
+static void dtors_101(void) { puts_log("<dtors.101>"); }
+__attribute__((section(".dtors.101"), used)) static dtor_t dtors_1_ptr = dtors_101;
+
+static void dtors_102(void) { puts_log("<dtors.102>"); }
+__attribute__((section(".dtors.102"), used)) static dtor_t dtors_2_ptr = dtors_102;
+
+static void dtors_103(void) { puts_log("<dtors.103>"); }
+__attribute__((section(".dtors.103"), used)) static dtor_t dtors_3_ptr = dtors_103;
+#endif /* __ELF__ */
+
+/* Mach-O-specific: explicit __DATA,__mod_init_func section placement. */
+#ifdef __APPLE__
+static void mod_init_func(void) { puts_log("<mod_init_func>"); }
+__attribute__((section("__DATA,__mod_init_func"), used)) static ctor_t mod_init_ptr = mod_init_func;
+#endif /* __APPLE__ */
+
+#endif /* !_MSC_VER */
+
+/* =========================================================================
+ * MSVC (COFF): .CRT$XC* sections + atexit()
+ * Subsections run in alphabetical order: XCA < XCB < XCC < XCU.
+ *
+ * Variables in these sections MUST NOT be static: both MSVC and clang-cl
+ * may optimize away unreferenced static variables even with __declspec(allocate).
+ * External linkage + unique names ensures the section data is emitted.
+ * ========================================================================= */
+#ifdef _MSC_VER
+
+#pragma section(".CRT$XCA", read)
+#pragma section(".CRT$XCB", read)
+#pragma section(".CRT$XCC", read)
+#pragma section(".CRT$XCU", read)
+#pragma section(".CRT$XTA", read)
+#pragma section(".CRT$XTZ", read)
+
+static void __cdecl crt_init_a(void) { puts_log("<crt.XCA>"); }
+__declspec(allocate(".CRT$XCA")) ctor_t __tvm_test_crt_init_a = crt_init_a;
+
+static void __cdecl crt_init_b(void) { puts_log("<crt.XCB>"); }
+__declspec(allocate(".CRT$XCB")) ctor_t __tvm_test_crt_init_b = crt_init_b;
+
+static void __cdecl crt_init_c(void) { puts_log("<crt.XCC>"); }
+__declspec(allocate(".CRT$XCC")) ctor_t __tvm_test_crt_init_c = crt_init_c;
+
+static void __cdecl crt_init_u(void) { puts_log("<crt.XCU>"); }
+__declspec(allocate(".CRT$XCU")) ctor_t __tvm_test_crt_init_u = crt_init_u;
+
+static void __cdecl crt_term_a(void) { puts_log("<crt.XTA>"); }
+__declspec(allocate(".CRT$XTA")) dtor_t __tvm_test_crt_term_a = crt_term_a;
+
+static void __cdecl crt_term_z(void) { puts_log("<crt.XTZ>"); }
+__declspec(allocate(".CRT$XTZ")) dtor_t __tvm_test_crt_term_z = crt_term_z;
+
+#endif /* _MSC_VER */
+
+/* main: callable entry point */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_main(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ puts_log("<main>");
+ result->type_index = kTVMFFINone;
+ result->zero_padding = 0;
+ result->v_int64 = 0;
+ return 0;
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_error.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_error.c
new file mode 100644
index 0000000..7d414b7
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_error.c
@@ -0,0 +1,31 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Error propagation test: JIT function signals error via
+ * TVMFFIErrorSetRaisedFromCStr, which should surface as a Python exception.
+ */
+#include <tvm/ffi/c_api.h>
+
+/* test_error: always raises a ValueError */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_error(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ TVMFFIErrorSetRaisedFromCStr("ValueError", "test error");
+ return -1;
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs.c
new file mode 100644
index 0000000..d8a2d26
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs.c
@@ -0,0 +1,43 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Pure C test functions using the TVMFFISafeCallType ABI directly.
+ * No C++ features: no exceptions, RTTI, static variables, or std library.
+ * This isolates the JIT infrastructure from MSVC C++ runtime issues.
+ */
+#include <tvm/ffi/c_api.h>
+
+/* test_add: add two integers */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_add(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 + args[1].v_int64;
+ return 0;
+}
+
+/* test_multiply: multiply two integers */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_multiply(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 * args[1].v_int64;
+ return 0;
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs2.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs2.c
new file mode 100644
index 0000000..ea470ed
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs2.c
@@ -0,0 +1,42 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Pure C test functions (subtract, divide) using the TVMFFISafeCallType ABI directly.
+ * C version of test_funcs2.cc — no C++ runtime dependencies.
+ */
+#include <tvm/ffi/c_api.h>
+
+/* test_subtract: subtract two integers */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_subtract(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 - args[1].v_int64;
+ return 0;
+}
+
+/* test_divide: divide two integers */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_divide(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 / args[1].v_int64;
+ return 0;
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs_conflict.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs_conflict.c
new file mode 100644
index 0000000..8ca5231
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_funcs_conflict.c
@@ -0,0 +1,43 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Pure C conflicting test functions using the TVMFFISafeCallType ABI directly.
+ * C version of test_funcs_conflict.cc — same symbol names as test_funcs.c
+ * but different implementations to test symbol conflict handling.
+ */
+#include <tvm/ffi/c_api.h>
+
+/* test_add: conflicting add — adds 1000 to the result */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_add(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 + args[1].v_int64 + 1000;
+ return 0;
+}
+
+/* test_multiply: conflicting multiply — doubles the result */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_multiply(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 * args[1].v_int64 * 2;
+ return 0;
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_link_order_base.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_link_order_base.c
new file mode 100644
index 0000000..21f26df
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_link_order_base.c
@@ -0,0 +1,33 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Base library for cross-library linking test.
+ * Exports helper_add which is called by test_link_order_caller.c.
+ */
+#include <tvm/ffi/c_api.h>
+
+/* helper_add: add two integers — used as a dependency by the caller library */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_helper_add(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 + args[1].v_int64;
+ return 0;
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_link_order_caller.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_link_order_caller.c
new file mode 100644
index 0000000..215a267
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_link_order_caller.c
@@ -0,0 +1,35 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Caller library for cross-library linking test.
+ * References __tvm_ffi_helper_add from test_link_order_base.c via extern
+ * declaration, and exports cross_lib_add which forwards to it.
+ */
+#include <tvm/ffi/c_api.h>
+
+/* Declare external symbol from the base library */
+extern int __tvm_ffi_helper_add(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result);
+
+/* cross_lib_add: forwards to helper_add in the base library */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_cross_lib_add(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ return __tvm_ffi_helper_add(self, args, num_args, result);
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/c/test_types.c b/addons/tvm_ffi_orcjit/tests/sources/c/test_types.c
new file mode 100644
index 0000000..c0b45cf
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/c/test_types.c
@@ -0,0 +1,60 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements. See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership. The ASF licenses this file
+ * to you 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.
+ */
+
+/*
+ * Pure C test functions exercising type variety: zero-arg, four-arg,
+ * float, and void return types via the TVMFFISafeCallType ABI.
+ */
+#include <tvm/ffi/c_api.h>
+
+/* test_zero_arg: ignores arguments, returns integer 42 */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_zero_arg(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = 42;
+ return 0;
+}
+
+/* test_four_args: sum of four integer arguments */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_four_args(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 + args[1].v_int64 + args[2].v_int64 + args[3].v_int64;
+ return 0;
+}
+
+/* test_float_multiply: multiply two doubles */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_float_multiply(void* self, const TVMFFIAny* args,
+ int32_t num_args, TVMFFIAny* result) {
+ result->type_index = kTVMFFIFloat;
+ result->zero_padding = 0;
+ result->v_float64 = args[0].v_float64 * args[1].v_float64;
+ return 0;
+}
+
+/* test_void_function: does nothing, returns None */
+TVM_FFI_DLL_EXPORT int __tvm_ffi_test_void_function(void* self, const TVMFFIAny* args,
+ int32_t num_args, TVMFFIAny* result) {
+ result->type_index = kTVMFFINone;
+ result->zero_padding = 0;
+ result->v_int64 = 0;
+ return 0;
+}
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_call_global.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_call_global.cc
new file mode 100644
index 0000000..1ec47e2
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_call_global.cc
@@ -0,0 +1,37 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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.
+
+// C++ test: look up a host-registered global function by name and call it.
+// Demonstrates JIT code calling back into the host process via the TVM FFI C++ API.
+
+#include <tvm/ffi/function.h>
+
+using namespace tvm::ffi;
+
+// Look up "test_host_add" global function and call it with two integer arguments.
+int test_call_global_add_impl(int a, int b) {
+ Function host_add = Function::GetGlobalRequired("test_host_add");
+ return host_add(a, b).cast<int>();
+}
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_call_global_add, test_call_global_add_impl);
+
+// Look up "test_host_multiply" global function and call it.
+int test_call_global_mul_impl(int a, int b) {
+ Function host_mul = Function::GetGlobalRequired("test_host_multiply");
+ return host_mul(a, b).cast<int>();
+}
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_call_global_mul, test_call_global_mul_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_ctor_dtor.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_ctor_dtor.cc
new file mode 100644
index 0000000..35af0af
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_ctor_dtor.cc
@@ -0,0 +1,82 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 <tvm/ffi/function.h>
+
+using namespace tvm::ffi;
+
+#define PUTS_LOG(msg) \
+ auto append_log_func = Function::GetGlobalRequired("append_log"); \
+ append_log_func(msg);
+
+using ctor_t = void (*)();
+using dtor_t = void (*)();
+
+// __attribute__((constructor/destructor)) works on both ELF and Mach-O:
+// ELF → .init_array / .fini_array
+// Mach-O → __DATA,__mod_init_func / __cxa_atexit
+__attribute__((constructor)) void init_array() { PUTS_LOG("<init_array>"); }
+
+__attribute__((constructor(101))) void init_array_101() { PUTS_LOG("<init_array.101>"); }
+
+__attribute__((constructor(102))) void init_array_102() { PUTS_LOG("<init_array.102>"); }
+__attribute__((constructor(103))) void init_array_103() { PUTS_LOG("<init_array.103>"); }
+
+__attribute__((destructor)) void fini_array() { PUTS_LOG("<fini_array>"); }
+
+__attribute__((destructor(101))) void fini_array_101() { PUTS_LOG("<fini_array.101>"); }
+
+__attribute__((destructor(102))) void fini_array_102() { PUTS_LOG("<fini_array.102>"); }
+
+__attribute__((destructor(103))) void fini_array_103() { PUTS_LOG("<fini_array.103>"); }
+
+// ELF-specific: explicit .ctors/.dtors section placements with priorities.
+// These sections don't exist on Mach-O or COFF.
+#ifdef __ELF__
+static void ctors() { PUTS_LOG("<ctors>"); }
+__attribute__((section(".ctors"), used)) static ctor_t ctors_ptr = ctors;
+
+static void ctors_101() { PUTS_LOG("<ctors.101>"); }
+__attribute__((section(".ctors.101"), used)) static ctor_t ctors_1_ptr = ctors_101;
+
+static void ctors_102() { PUTS_LOG("<ctors.102>"); }
+__attribute__((section(".ctors.102"), used)) static ctor_t ctors_2_ptr = ctors_102;
+
+static void ctors_103() { PUTS_LOG("<ctors.103>"); }
+__attribute__((section(".ctors.103"), used)) static ctor_t ctors_3_ptr = ctors_103;
+
+static void dtors() { PUTS_LOG("<dtors>"); }
+__attribute__((section(".dtors"), used)) static dtor_t dtors_ptr = dtors;
+
+static void dtors_101() { PUTS_LOG("<dtors.101>"); }
+__attribute__((section(".dtors.101"), used)) static dtor_t dtors_1_ptr = dtors_101;
+
+static void dtors_102() { PUTS_LOG("<dtors.102>"); }
+__attribute__((section(".dtors.102"), used)) static dtor_t dtors_2_ptr = dtors_102;
+
+static void dtors_103() { PUTS_LOG("<dtors.103>"); }
+__attribute__((section(".dtors.103"), used)) static dtor_t dtors_3_ptr = dtors_103;
+#endif // __ELF__
+
+// Mach-O-specific: explicit __DATA,__mod_init_func section placement.
+#ifdef __APPLE__
+static void mod_init_func() { PUTS_LOG("<mod_init_func>"); }
+__attribute__((section("__DATA,__mod_init_func"), used)) static ctor_t mod_init_ptr = mod_init_func;
+#endif // __APPLE__
+
+void main_impl() { PUTS_LOG("<main>"); }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(main, main_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_error.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_error.cc
new file mode 100644
index 0000000..a12d3d2
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_error.cc
@@ -0,0 +1,26 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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.
+
+// Error propagation test (C++ version): throws std::runtime_error
+// which TVM_FFI_SAFE_CALL_END catches and converts to an InternalError.
+
+#include <tvm/ffi/function.h>
+
+#include <stdexcept>
+
+int test_error_impl() { throw std::runtime_error("test error"); }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_error, test_error_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs.cc
new file mode 100644
index 0000000..94aecfb
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs.cc
@@ -0,0 +1,26 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 <tvm/ffi/function.h>
+
+// Simple addition function
+int test_add_impl(int a, int b) { return a + b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_add, test_add_impl);
+
+// Multiplication function
+int test_multiply_impl(int a, int b) { return a * b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_multiply, test_multiply_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs2.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs2.cc
new file mode 100644
index 0000000..d786bc3
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs2.cc
@@ -0,0 +1,26 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 <tvm/ffi/function.h>
+
+// Subtraction function
+int test_subtract_impl(int a, int b) { return a - b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_subtract, test_subtract_impl);
+
+// Division function
+int test_divide_impl(int a, int b) { return a / b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_divide, test_divide_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs_conflict.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs_conflict.cc
new file mode 100644
index 0000000..09d35a8
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_funcs_conflict.cc
@@ -0,0 +1,26 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 <tvm/ffi/function.h>
+
+// Conflicting test_add function - different implementation
+int test_add_conflict_impl(int a, int b) { return a + b + 1000; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_add, test_add_conflict_impl);
+
+// Conflicting test_multiply function - different implementation
+int test_multiply_conflict_impl(int a, int b) { return a * b * 2; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_multiply, test_multiply_conflict_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_link_order_base.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_link_order_base.cc
new file mode 100644
index 0000000..c878c6e
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_link_order_base.cc
@@ -0,0 +1,24 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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.
+
+// Base library for cross-library linking test (C++ version).
+// Exports helper_add which is called by test_link_order_caller.cc.
+
+#include <tvm/ffi/function.h>
+
+int helper_add_impl(int a, int b) { return a + b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(helper_add, helper_add_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_link_order_caller.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_link_order_caller.cc
new file mode 100644
index 0000000..85dde57
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_link_order_caller.cc
@@ -0,0 +1,43 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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.
+
+// Caller library for cross-library linking test (C++ version).
+// References __tvm_ffi_helper_add from test_link_order_base.cc and
+// exports cross_lib_add which forwards to it.
+
+#include <tvm/ffi/c_api.h>
+#include <tvm/ffi/function.h>
+
+extern "C" int __tvm_ffi_helper_add(void* self, const TVMFFIAny* args, int32_t num_args,
+ TVMFFIAny* result);
+
+int cross_lib_add_impl(int a, int b) {
+ TVMFFIAny call_args[2];
+ TVMFFIAny call_result;
+ call_args[0].type_index = kTVMFFIInt;
+ call_args[0].zero_padding = 0;
+ call_args[0].v_int64 = a;
+ call_args[1].type_index = kTVMFFIInt;
+ call_args[1].zero_padding = 0;
+ call_args[1].v_int64 = b;
+ call_result.type_index = kTVMFFINone;
+ call_result.zero_padding = 0;
+ call_result.v_int64 = 0;
+ __tvm_ffi_helper_add(nullptr, call_args, 2, &call_result);
+ return static_cast<int>(call_result.v_int64);
+}
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(cross_lib_add, cross_lib_add_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cc/test_types.cc b/addons/tvm_ffi_orcjit/tests/sources/cc/test_types.cc
new file mode 100644
index 0000000..0056520
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cc/test_types.cc
@@ -0,0 +1,33 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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.
+
+// C++ test functions exercising type variety: zero-arg, four-arg,
+// float, and void return types.
+
+#include <tvm/ffi/function.h>
+
+int test_zero_arg_impl() { return 42; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_zero_arg, test_zero_arg_impl);
+
+int test_four_args_impl(int a, int b, int c, int d) { return a + b + c + d; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_four_args, test_four_args_impl);
+
+double test_float_multiply_impl(double a, double b) { return a * b; }
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_float_multiply, test_float_multiply_impl);
+
+void test_void_function_impl() {}
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_void_function, test_void_function_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/sources/cuda/test_funcs.cu b/addons/tvm_ffi_orcjit/tests/sources/cuda/test_funcs.cu
new file mode 100644
index 0000000..fda0d7c
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/sources/cuda/test_funcs.cu
@@ -0,0 +1,61 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 <tvm/ffi/function.h>
+
+#include <cstdio>
+
+// Simple addition function
+__global__ void test_add_kernel(int* a, int* b, int* c) { *c = *a + *b; }
+int test_add_impl(int a, int b) {
+ int c;
+ int *d_a, *d_b, *d_c;
+ cudaMalloc(&d_a, sizeof(int));
+ cudaMalloc(&d_b, sizeof(int));
+ cudaMalloc(&d_c, sizeof(int));
+ cudaMemcpy(d_a, &a, sizeof(int), cudaMemcpyHostToDevice);
+ cudaMemcpy(d_b, &b, sizeof(int), cudaMemcpyHostToDevice);
+ test_add_kernel<<<1, 1>>>(d_a, d_b, d_c);
+ cudaError_t err = cudaGetLastError();
+ if (err != cudaSuccess) printf("Kernel launch error: %s\n", cudaGetErrorString(err));
+ cudaMemcpy(&c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
+ cudaFree(d_a);
+ cudaFree(d_b);
+ cudaFree(d_c);
+ return c;
+}
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_add, test_add_impl);
+
+// Multiplication function
+
+__global__ void test_multiply_kernel(int* a, int* b, int* c) { *c = *a * *b; }
+int test_multiply_impl(int a, int b) {
+ int c;
+ int *d_a, *d_b, *d_c;
+ cudaMalloc(&d_a, sizeof(int));
+ cudaMalloc(&d_b, sizeof(int));
+ cudaMalloc(&d_c, sizeof(int));
+ cudaMemcpy(d_a, &a, sizeof(int), cudaMemcpyHostToDevice);
+ cudaMemcpy(d_b, &b, sizeof(int), cudaMemcpyHostToDevice);
+ test_multiply_kernel<<<1, 1>>>(d_a, d_b, d_c);
+ cudaMemcpy(&c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
+ cudaFree(d_a);
+ cudaFree(d_b);
+ cudaFree(d_c);
+ return c;
+}
+TVM_FFI_DLL_EXPORT_TYPED_FUNC(test_multiply, test_multiply_impl);
diff --git a/addons/tvm_ffi_orcjit/tests/test_basic.py b/addons/tvm_ffi_orcjit/tests/test_basic.py
new file mode 100644
index 0000000..ca5ba18
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/test_basic.py
@@ -0,0 +1,475 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+"""Basic tests for tvm_ffi_orcjit functionality."""
+
+from __future__ import annotations
+
+import sys
+import tempfile
+
+import pytest
+import tvm_ffi
+from tvm_ffi_orcjit import ExecutionSession
+from tvm_ffi_orcjit.dylib import DynamicLibrary
+from utils import build_test_objects
+
+# ---------------------------------------------------------------------------
+# Helpers
+# ---------------------------------------------------------------------------
+
+OBJ_DIR = build_test_objects()
+
+_KNOWN_SUBDIRS = [
+ "cc",
+ "c", # default (LLVM clang)
+ "cc-gcc",
+ "c-gcc", # GCC (Linux)
+ "cc-appleclang",
+ "c-appleclang", # Apple Clang (macOS)
+ "c-msvc", # MSVC (Windows, C only)
+ "c-clang-cl", # clang-cl (Windows, C only)
+]
+
+
+def obj(name: str) -> str:
+ """Return path to a pre-built test object file, or skip if missing."""
+ path = OBJ_DIR / f"{name}.o"
+ if not path.exists():
+ pytest.skip(f"{path.name} not found (not built)")
+ return str(path)
+
+
+def make_lib(
+ *obj_names: str, session: ExecutionSession | None = None, name: str = ""
+) -> tuple[ExecutionSession, DynamicLibrary]:
+ """Create a library and load one or more object files into it."""
+ if session is None:
+ session = ExecutionSession()
+ lib = session.create_library(name)
+ for o in obj_names:
+ lib.add(obj(o))
+ return session, lib
+
+
+# ---------------------------------------------------------------------------
+# Variants: C and C++ sources live in separate subdirectories (c/, cc/).
+# Function names are identical; only the object file path differs.
+# Tests are parametrized over both variants where the logic is identical.
+# ---------------------------------------------------------------------------
+
+
+class Variant:
+ """Describes a C or C++ test variant (object file path mapping)."""
+
+ def __init__(self, subdir: str) -> None:
+ self.subdir = subdir # "cc" for C++, "c" for C
+
+ def funcs_obj(self) -> str:
+ """Return path prefix for test_funcs object."""
+ return f"{self.subdir}/test_funcs"
+
+ def funcs2_obj(self) -> str:
+ """Return path prefix for test_funcs2 object."""
+ return f"{self.subdir}/test_funcs2"
+
+ def conflict_obj(self) -> str:
+ """Return path prefix for test_funcs_conflict object."""
+ return f"{self.subdir}/test_funcs_conflict"
+
+ def call_global_obj(self) -> str:
+ """Return path prefix for test_call_global object."""
+ return f"{self.subdir}/test_call_global"
+
+ def types_obj(self) -> str:
+ """Return path prefix for test_types object."""
+ return f"{self.subdir}/test_types"
+
+ def link_order_base_obj(self) -> str:
+ """Return path prefix for test_link_order_base object."""
+ return f"{self.subdir}/test_link_order_base"
+
+ def link_order_caller_obj(self) -> str:
+ """Return path prefix for test_link_order_caller object."""
+ return f"{self.subdir}/test_link_order_caller"
+
+ def error_obj(self) -> str:
+ """Return path prefix for test_error object."""
+ return f"{self.subdir}/test_error"
+
+ def ctor_dtor_obj(self) -> str:
+ """Return path prefix for test_ctor_dtor object."""
+ return f"{self.subdir}/test_ctor_dtor"
+
+ def fn(self, base_name: str) -> str:
+ """Return the function name for a given base name."""
+ return base_name
+
+ def __repr__(self) -> str:
+ parts = self.subdir.split("-", 1)
+ lang = "C++" if parts[0] == "cc" else "C"
+ return f"{lang}-{parts[1]}" if len(parts) > 1 else lang
+
+
+def _discover_variants() -> list[Variant]:
+ return [Variant(s) for s in _KNOWN_SUBDIRS if (OBJ_DIR / s / "test_funcs.o").exists()]
+
+
+_all_variants = _discover_variants()
+_cpp_only = [v for v in _all_variants if v.subdir.startswith("cc")]
+
+
+def _variant_id(v: Variant) -> str:
+ return repr(v)
+
+
+# ---------------------------------------------------------------------------
+# Session / library creation
+# ---------------------------------------------------------------------------
+
+
+def test_create_session() -> None:
+ """Create an ExecutionSession successfully."""
+ session = ExecutionSession()
+ assert session is not None
+
+
+def test_create_library() -> None:
+ """Create a DynamicLibrary from an ExecutionSession."""
+ session = ExecutionSession()
+ lib = session.create_library()
+ assert lib is not None
+
+
+def test_multiple_libraries() -> None:
+ """Create multiple named libraries in one session."""
+ session = ExecutionSession()
+ lib1 = session.create_library("lib1")
+ lib2 = session.create_library("lib2")
+ assert lib1 is not None
+ assert lib2 is not None
+
+
+# ---------------------------------------------------------------------------
+# Load & execute — parametrized over C / C++
+# ---------------------------------------------------------------------------
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_load_and_execute(v: Variant) -> None:
+ """Load test_funcs and verify add/multiply."""
+ _, lib = make_lib(v.funcs_obj())
+ assert lib.get_function(v.fn("test_add"))(10, 20) == 30
+ assert lib.get_function(v.fn("test_multiply"))(7, 6) == 42
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_load_and_execute_second_set(v: Variant) -> None:
+ """Load test_funcs2 and verify subtract/divide."""
+ _, lib = make_lib(v.funcs2_obj())
+ assert lib.get_function(v.fn("test_subtract"))(10, 3) == 7
+ assert lib.get_function(v.fn("test_divide"))(20, 4) == 5
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_function_not_found(v: Variant) -> None:
+ """Raise AttributeError for a missing function name."""
+ _, lib = make_lib(v.funcs_obj())
+ with pytest.raises(AttributeError, match="Module has no function"):
+ lib.get_function("nonexistent_function")
+
+
+# ---------------------------------------------------------------------------
+# Multi-object / multi-library — parametrized over C / C++
+# ---------------------------------------------------------------------------
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_gradually_add_objects(v: Variant) -> None:
+ """Add multiple objects incrementally and verify all functions work."""
+ _session, lib = make_lib(v.funcs_obj())
+
+ add_func = lib.get_function(v.fn("test_add"))
+ mul_func = lib.get_function(v.fn("test_multiply"))
+ assert add_func(5, 3) == 8
+ assert mul_func(4, 5) == 20
+
+ lib.add(obj(v.funcs2_obj()))
+ assert lib.get_function(v.fn("test_subtract"))(10, 3) == 7
+ assert lib.get_function(v.fn("test_divide"))(20, 4) == 5
+
+ # First object's functions still work
+ assert add_func(10, 20) == 30
+ assert mul_func(7, 6) == 42
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_two_separate_libraries(v: Variant) -> None:
+ """Separate libraries expose only their own functions."""
+ session = ExecutionSession()
+ _, lib1 = make_lib(v.funcs_obj(), session=session, name="lib1")
+ _, lib2 = make_lib(v.funcs2_obj(), session=session, name="lib2")
+
+ assert lib1.get_function(v.fn("test_add"))(5, 3) == 8
+ assert lib1.get_function(v.fn("test_multiply"))(4, 5) == 20
+ assert lib2.get_function(v.fn("test_subtract"))(10, 3) == 7
+ assert lib2.get_function(v.fn("test_divide"))(20, 4) == 5
+
+ with pytest.raises(AttributeError, match="Module has no function"):
+ lib1.get_function(v.fn("test_subtract"))
+ with pytest.raises(AttributeError, match="Module has no function"):
+ lib2.get_function(v.fn("test_add"))
+
+
+# ---------------------------------------------------------------------------
+# Symbol conflicts — parametrized over C / C++
+# ---------------------------------------------------------------------------
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_symbol_conflict_same_library(v: Variant) -> None:
+ """Duplicate symbol in the same library raises an error."""
+ _, lib = make_lib(v.funcs_obj())
+ assert lib.get_function(v.fn("test_add"))(10, 20) == 30
+ with pytest.raises(Exception):
+ lib.add(obj(v.conflict_obj()))
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_symbol_conflict_different_libraries(v: Variant) -> None:
+ """Same symbol in different libraries resolves independently."""
+ session = ExecutionSession()
+ _, lib1 = make_lib(v.funcs_obj(), session=session, name="lib1")
+ _, lib2 = make_lib(v.conflict_obj(), session=session, name="lib2")
+
+ assert lib1.get_function(v.fn("test_add"))(10, 20) == 30
+ assert lib2.get_function(v.fn("test_add"))(10, 20) == 1030
+ assert lib1.get_function(v.fn("test_multiply"))(5, 6) == 30
+ assert lib2.get_function(v.fn("test_multiply"))(5, 6) == 60
+
+
+# ---------------------------------------------------------------------------
+# Global function callbacks — parametrized over C / C++
+# ---------------------------------------------------------------------------
+
+
+@pytest.fixture()
+def _register_host_functions() -> None:
+ """Register host add/multiply functions for JIT code to call."""
+
+ @tvm_ffi.register_global_func("test_host_add", override=True)
+ def _host_add(a: int, b: int) -> int:
+ return a + b
+
+ @tvm_ffi.register_global_func("test_host_multiply", override=True)
+ def _host_mul(a: int, b: int) -> int:
+ return a * b
+
+
+@pytest.mark.usefixtures("_register_host_functions")
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_call_global(v: Variant) -> None:
+ """JIT code calls back into Python-registered global functions."""
+ _, lib = make_lib(v.call_global_obj())
+ add_func = lib.get_function(v.fn("test_call_global_add"))
+ assert add_func(10, 20) == 30
+ assert add_func(100, 200) == 300
+
+ mul_func = lib.get_function(v.fn("test_call_global_mul"))
+ assert mul_func(7, 6) == 42
+ assert mul_func(11, 11) == 121
+
+
+# ---------------------------------------------------------------------------
+# Error handling — pure Python (Group 1)
+# ---------------------------------------------------------------------------
+
+
+def test_empty_library() -> None:
+ """get_function on an empty library raises AttributeError."""
+ session = ExecutionSession()
+ lib = session.create_library("empty")
+ with pytest.raises(AttributeError, match="Module has no function"):
+ lib.get_function("nonexistent")
+
+
+def test_invalid_object_file_path() -> None:
+ """Adding a nonexistent object file raises an error."""
+ session = ExecutionSession()
+ lib = session.create_library("bad_path")
+ with pytest.raises(Exception):
+ lib.add("/nonexistent_path/does_not_exist.o")
+
+
+def test_invalid_object_file_content() -> None:
+ """Adding a file with garbage content raises an error."""
+ with tempfile.NamedTemporaryFile(suffix=".o", delete=False) as f:
+ f.write(b"this is not a valid object file")
+ f.flush()
+ session = ExecutionSession()
+ lib = session.create_library("bad_content")
+ with pytest.raises(Exception):
+ lib.add(f.name)
+
+
+def test_multiple_independent_sessions() -> None:
+ """Two independent sessions don't interfere with each other."""
+ session1 = ExecutionSession()
+ session2 = ExecutionSession()
+ _, lib1 = make_lib(_all_variants[0].funcs_obj(), session=session1)
+ _, lib2 = make_lib(_all_variants[0].funcs2_obj(), session=session2)
+
+ assert lib1.get_function("test_add")(5, 3) == 8
+ assert lib2.get_function("test_subtract")(10, 3) == 7
+
+ # Each session's library doesn't see the other's functions
+ with pytest.raises(AttributeError, match="Module has no function"):
+ lib1.get_function("test_subtract")
+ with pytest.raises(AttributeError, match="Module has no function"):
+ lib2.get_function("test_add")
+
+
+# ---------------------------------------------------------------------------
+# Type variety — parametrized over C / C++ (Group 2)
+# ---------------------------------------------------------------------------
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_zero_arg_function(v: Variant) -> None:
+ """Zero-arg function returns constant 42."""
+ _, lib = make_lib(v.types_obj())
+ assert lib.get_function(v.fn("test_zero_arg"))() == 42
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_four_arg_function(v: Variant) -> None:
+ """Four integer arguments summed."""
+ _, lib = make_lib(v.types_obj())
+ assert lib.get_function(v.fn("test_four_args"))(1, 2, 3, 4) == 10
+ assert lib.get_function(v.fn("test_four_args"))(100, 200, 300, 400) == 1000
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_float_function(v: Variant) -> None:
+ """Float multiply returns approximate result."""
+ _, lib = make_lib(v.types_obj())
+ result = lib.get_function(v.fn("test_float_multiply"))(3.14, 2.0)
+ assert result == pytest.approx(6.28)
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_void_function(v: Variant) -> None:
+ """Void function returns None."""
+ _, lib = make_lib(v.types_obj())
+ result = lib.get_function(v.fn("test_void_function"))()
+ assert result is None
+
+
+# ---------------------------------------------------------------------------
+# Advanced — cross-library linking and error propagation (Group 3)
+# ---------------------------------------------------------------------------
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_set_link_order(v: Variant) -> None:
+ """Cross-library symbol resolution via set_link_order."""
+ session = ExecutionSession()
+ # Base library exports helper_add
+ lib_base = session.create_library("base")
+ lib_base.add(obj(v.link_order_base_obj()))
+ # Caller library references helper_add from base
+ lib_caller = session.create_library("caller")
+ lib_caller.set_link_order(lib_base)
+ lib_caller.add(obj(v.link_order_caller_obj()))
+
+ cross_add = lib_caller.get_function(v.fn("cross_lib_add"))
+ assert cross_add(10, 20) == 30
+ assert cross_add(100, 200) == 300
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_error_propagation(v: Variant) -> None:
+ """JIT function that signals an error raises a Python exception."""
+ _, lib = make_lib(v.error_obj())
+ with pytest.raises(Exception, match="test error"):
+ lib.get_function(v.fn("test_error"))()
+
+
+# ---------------------------------------------------------------------------
+# CUDA (optional)
+# ---------------------------------------------------------------------------
+
+
+def test_load_and_execute_cuda_function() -> None:
+ """Load and execute CUDA-compiled test objects."""
+ _, lib = make_lib("cuda/test_funcs")
+ assert lib.get_function("test_add")(10, 20) == 30
+ assert lib.get_function("test_multiply")(7, 6) == 42
+
+
+# ---------------------------------------------------------------------------
+# Constructor / destructor (parametrized over C / C++)
+# ---------------------------------------------------------------------------
+
+
+@pytest.mark.parametrize("v", _all_variants, ids=_variant_id)
+def test_ctor_dtor(v: Variant) -> None:
+ """Verify constructor/destructor ordering across platforms."""
+ log = ""
+
+ @tvm_ffi.register_global_func("append_log", override=True)
+ def _append_ctor_log(x: str) -> None:
+ nonlocal log
+ log += x
+
+ _, lib = make_lib(v.ctor_dtor_obj())
+ lib.get_function(v.fn("main"))()
+ del lib
+
+ main_idx = log.index("<main>")
+ pre = log[:main_idx]
+ post = log[main_idx:]
+
+ if sys.platform == "win32":
+ # Windows (all compilers): COFF .CRT$XC* constructors + .CRT$XT* terminators.
+ # All Windows compilers (MSVC, clang-cl, and LLVM Clang targeting MSVC ABI)
+ # define _MSC_VER, so the source uses #pragma section / __declspec(allocate).
+ assert "<crt.XCA>" in pre, f"CRT initializers not found in log: {log!r}"
+ assert pre.index("<crt.XCA>") < pre.index("<crt.XCB>")
+ assert pre.index("<crt.XCB>") < pre.index("<crt.XCC>")
+ assert pre.index("<crt.XCC>") < pre.index("<crt.XCU>")
+ assert "<crt.XTA>" in post, f"CRT terminators not found in log: {log!r}"
+ assert post.index("<crt.XTA>") < post.index("<crt.XTZ>")
+ elif sys.platform == "linux":
+ # ELF: init_array (priority order) + .ctors (reversed priority)
+ assert pre.index("<init_array.101>") < pre.index("<init_array.102>")
+ assert pre.index("<init_array.102>") < pre.index("<init_array.103>")
+ assert pre.index("<init_array.103>") < pre.index("<init_array>")
+ assert pre.index("<ctors.103>") < pre.index("<ctors.102>")
+ assert pre.index("<ctors.102>") < pre.index("<ctors.101>")
+ assert pre.index("<ctors.101>") < pre.index("<ctors>")
+ assert "<dtors>" in post
+ elif sys.platform == "darwin":
+ # Mach-O: init_array (priority order) + explicit __mod_init_func
+ assert pre.index("<init_array.101>") < pre.index("<init_array.102>")
+ assert pre.index("<init_array.102>") < pre.index("<init_array.103>")
+ assert pre.index("<init_array.103>") < pre.index("<init_array>")
+ assert "<mod_init_func>" in pre
+ assert post.index("<fini_array>") < post.index("<fini_array.103>")
+ assert post.index("<fini_array.103>") < post.index("<fini_array.102>")
+ assert post.index("<fini_array.102>") < post.index("<fini_array.101>")
+ assert "<ctors>" not in log
+ assert "<dtors>" not in log
diff --git a/addons/tvm_ffi_orcjit/tests/utils.py b/addons/tvm_ffi_orcjit/tests/utils.py
new file mode 100644
index 0000000..e1fbc96
--- /dev/null
+++ b/addons/tvm_ffi_orcjit/tests/utils.py
@@ -0,0 +1,232 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you 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.
+"""Build test object files for all available compiler variants.
+
+Uses ``tvm_ffi.cpp.build`` to compile C/C++ test sources to relocatable
+object files. Detects platform and available compilers:
+
+- Linux: LLVM Clang (default) + GCC
+- macOS: LLVM Clang (default) + Apple Clang
+- Windows: LLVM Clang (default) + MSVC + clang-cl
+
+Objects are placed under a temporary directory and skipped if they already
+exist. LLVM clang is found via ``$LLVM_PREFIX/bin`` or ``$PATH``; no
+hardcoded install path is assumed.
+"""
+
+from __future__ import annotations
+
+import os
+import platform
+import shutil
+import tempfile
+from pathlib import Path
+
+import tvm_ffi.cpp
+
+SOURCES_DIR = Path(__file__).resolve().parent / "sources"
+SOURCES_C = SOURCES_DIR / "c"
+SOURCES_CC = SOURCES_DIR / "cc"
+SOURCES_CUDA = SOURCES_DIR / "cuda"
+
+_DEFAULT_OUT_DIR = Path(tempfile.gettempdir()) / "tvm_ffi_orcjit_tests"
+
+
+# ---------------------------------------------------------------------------
+# Helpers
+# ---------------------------------------------------------------------------
+
+
+def _extra_cflags() -> list[str]:
+ machine = platform.machine()
+ if machine in ("aarch64", "arm64"):
+ return ["-mno-outline-atomics"]
+ return []
+
+
+def _build_objects(
+ src_dir: Path,
+ out_dir: Path,
+ *,
+ ext_glob: str,
+ extra_cflags: list[str],
+) -> None:
+ """Compile all sources in *src_dir* to object files in *out_dir*."""
+ out_dir.mkdir(parents=True, exist_ok=True)
+
+ for src in sorted(src_dir.glob(ext_glob)):
+ dest = out_dir / f"{src.stem}.o"
+ if dest.exists():
+ continue
+ build_dir = out_dir / f".build_{src.stem}"
+ obj_path = tvm_ffi.cpp.build(
+ name=src.stem,
+ sources=[str(src)],
+ output=f"{src.stem}.o",
+ extra_cflags=extra_cflags,
+ build_directory=str(build_dir),
+ )
+ shutil.copy2(obj_path, dest)
+
+
+def _build_variant(
+ name: str,
+ *,
+ cc: str | None,
+ cxx: str | None,
+ extra_cflags: list[str],
+ c_outdir: Path,
+ cc_outdir: Path | None,
+) -> None:
+ """Build all test objects for one compiler variant."""
+ saved_cc = os.environ.get("CC")
+ saved_cxx = os.environ.get("CXX")
+ try:
+ if cc:
+ os.environ["CC"] = cc
+ if cxx:
+ os.environ["CXX"] = cxx
+
+ if cc:
+ _build_objects(SOURCES_C, c_outdir, ext_glob="*.c", extra_cflags=extra_cflags)
+ if cxx and cc_outdir:
+ _build_objects(SOURCES_CC, cc_outdir, ext_glob="*.cc", extra_cflags=extra_cflags)
+ finally:
+ if saved_cc is None:
+ os.environ.pop("CC", None)
+ else:
+ os.environ["CC"] = saved_cc
+ if saved_cxx is None:
+ os.environ.pop("CXX", None)
+ else:
+ os.environ["CXX"] = saved_cxx
+
+
+# ---------------------------------------------------------------------------
+# Compiler detection
+# ---------------------------------------------------------------------------
+
+
+def _find_llvm_clang() -> tuple[str, str] | None:
+ """Find LLVM clang/clang++ via LLVM_PREFIX or PATH.
+
+ Returns (clang, clang++) paths, or None if not found.
+ """
+ llvm_prefix = os.environ.get("LLVM_PREFIX")
+ if llvm_prefix:
+ p = Path(llvm_prefix)
+ bin_dir = p / "Library" / "bin" if (p / "Library" / "bin").exists() else p / "bin"
+ cc = bin_dir / ("clang.exe" if platform.system() == "Windows" else "clang")
+ if cc.exists():
+ cxx = bin_dir / ("clang++.exe" if platform.system() == "Windows" else "clang++")
+ return str(cc), str(cxx)
+
+ cc = shutil.which("clang")
+ if cc:
+ cxx = shutil.which("clang++")
+ return cc, cxx or cc
+
+ return None
+
+
+# ---------------------------------------------------------------------------
+# Public API
+# ---------------------------------------------------------------------------
+
+
+def build_test_objects(out_dir: Path | None = None) -> Path:
+ """Build test objects for all available compiler variants.
+
+ Objects are placed under *out_dir* (default:
+ ``{tempdir}/tvm_ffi_orcjit_tests``). Already-built objects are
+ skipped. Returns the output directory.
+ """
+ if out_dir is None:
+ out_dir = _DEFAULT_OUT_DIR
+ out_dir.mkdir(parents=True, exist_ok=True)
+
+ system = platform.system()
+ extra = _extra_cflags()
+
+ if system in ("Linux", "Darwin"):
+ llvm = _find_llvm_clang()
+ if llvm:
+ clang, clangxx = llvm
+ _build_variant(
+ "LLVM Clang",
+ cc=clang,
+ cxx=clangxx,
+ extra_cflags=extra,
+ c_outdir=out_dir / "c",
+ cc_outdir=out_dir / "cc",
+ )
+ if system == "Linux" and shutil.which("gcc"):
+ _build_variant(
+ "GCC",
+ cc="gcc",
+ cxx="g++",
+ extra_cflags=extra,
+ c_outdir=out_dir / "c-gcc",
+ cc_outdir=out_dir / "cc-gcc",
+ )
+ if system == "Darwin" and Path("/usr/bin/clang").exists():
+ _build_variant(
+ "Apple Clang",
+ cc="/usr/bin/clang",
+ cxx="/usr/bin/clang++",
+ extra_cflags=extra,
+ c_outdir=out_dir / "c-appleclang",
+ cc_outdir=out_dir / "cc-appleclang",
+ )
+
+ elif system == "Windows":
+ llvm = _find_llvm_clang()
+ if llvm:
+ clang_cl = Path(llvm[0]).parent / "clang-cl.exe"
+ if clang_cl.exists():
+ _build_variant(
+ "LLVM clang-cl",
+ cc=str(clang_cl),
+ cxx=None,
+ extra_cflags=["/GS-"],
+ c_outdir=out_dir / "c",
+ cc_outdir=None,
+ )
+ if shutil.which("cl"):
+ _build_variant(
+ "MSVC",
+ cc="cl",
+ cxx=None,
+ extra_cflags=["/GS-"],
+ c_outdir=out_dir / "c-msvc",
+ cc_outdir=None,
+ )
+ if shutil.which("clang-cl"):
+ _build_variant(
+ "clang-cl",
+ cc="clang-cl",
+ cxx=None,
+ extra_cflags=["/GS-"],
+ c_outdir=out_dir / "c-clang-cl",
+ cc_outdir=None,
+ )
+
+ # CUDA (platform-independent, uses nvcc)
+ if shutil.which("nvcc"):
+ _build_objects(SOURCES_CUDA, out_dir / "cuda", ext_glob="*.cu", extra_cflags=[])
+
+ return out_dir
diff --git a/docs/conf.py b/docs/conf.py
index b8b1021..61ffb60 100644
--- a/docs/conf.py
+++ b/docs/conf.py
@@ -482,3 +482,6 @@ html_extra_path = ["reference/rust/generated"] if build_rust_docs else []
html_css_files = ["custom.css"]
+
+
+show_warning_types = True
diff --git a/python/tvm_ffi/cpp/extension.py b/python/tvm_ffi/cpp/extension.py
index af30348..a7b3514 100644
--- a/python/tvm_ffi/cpp/extension.py
+++ b/python/tvm_ffi/cpp/extension.py
@@ -338,24 +338,38 @@ def _generate_ninja_build( # noqa: PLR0915, PLR0912
extra_cuda_cflags: Sequence[str],
extra_ldflags: Sequence[str],
extra_include_paths: Sequence[str],
- cpp_files: Sequence[str],
- cuda_files: Sequence[str],
+ sources: Sequence[str],
embed_cubin: Mapping[str, bytes] | None = None,
backend: str | None = None,
+ output: str | None = None,
) -> str:
"""Generate the content of build.ninja for building the module."""
+ # Determine output format from extension
+ if output is not None:
+ out_ext = Path(output).suffix.lower()
+ object_mode = out_ext in (".o", ".obj")
+ output_name = output
+ else:
+ object_mode = False
+ output_name = f"{name}{'.dll' if IS_WINDOWS else '.so'}"
+ has_cuda_sources = any(Path(s).suffix.lower() == ".cu" for s in sources)
with_hip = backend == "hip"
with_cuda = backend == "cuda"
- with_backend = with_hip or with_cuda
+ with_backend = with_hip or with_cuda or has_cuda_sources
+ if has_cuda_sources and not (with_hip or with_cuda):
+ # Auto-detect backend from available GPU
+ detected = _resolve_gpu_backend(None)
+ with_hip = detected == "hip"
+ with_cuda = detected == "cuda"
default_include_paths = [find_include_path(), find_dlpack_include_path()]
tvm_ffi_lib = Path(find_libtvm_ffi())
tvm_ffi_lib_path = str(tvm_ffi_lib.parent)
tvm_ffi_lib_name = tvm_ffi_lib.stem
if IS_WINDOWS:
- default_cflags = [
- "/std:c++17",
- "/MD",
+ default_cflags = ["/O2", "/MD"]
+ default_cxxflags = ["/std:c++17", "/MD", "/EHsc"]
+ _win_warnings = [
"/wd4819",
"/wd4251",
"/wd4244",
@@ -366,8 +380,9 @@ def _generate_ninja_build( # noqa: PLR0915, PLR0912
"/wd4624",
"/wd4067",
"/wd4068",
- "/EHsc",
]
+ default_cflags += _win_warnings
+ default_cxxflags += _win_warnings
default_cuda_cflags = ["-Xcompiler", "/std:c++17", "/O2"]
default_ldflags = [
"/DLL",
@@ -375,7 +390,8 @@ def _generate_ninja_build( # noqa: PLR0915, PLR0912
f"{tvm_ffi_lib_name}.lib",
]
else:
- default_cflags = ["-std=c++17", "-fPIC", "-O2"]
+ default_cflags = ["-fPIC", "-O2"]
+ default_cxxflags = ["-std=c++17", "-fPIC", "-O2"]
default_cuda_cflags = ["-std=c++17", "-O2"]
default_ldflags = ["-shared", f"-L{tvm_ffi_lib_path}", "-ltvm_ffi"]
@@ -396,7 +412,9 @@ def _generate_ninja_build( # noqa: PLR0915, PLR0912
"-lcudart", # cuda runtime library
]
- cflags = default_cflags + [flag.strip() for flag in extra_cflags]
+ extra_cflags_list = [flag.strip() for flag in extra_cflags]
+ cflags = default_cflags + extra_cflags_list
+ cxxflags = default_cxxflags + extra_cflags_list
cuda_cflags = default_cuda_cflags + [flag.strip() for flag in extra_cuda_cflags]
ldflags = default_ldflags + [flag.strip() for flag in extra_ldflags]
include_paths = default_include_paths + [
@@ -405,14 +423,20 @@ def _generate_ninja_build( # noqa: PLR0915, PLR0912
# append include paths
for path in include_paths:
- cflags.append("-I{}".format(path.replace(":", "$:")))
- cuda_cflags.append("-I{}".format(path.replace(":", "$:")))
+ inc = "-I{}".format(path.replace(":", "$:"))
+ cflags.append(inc)
+ cxxflags.append(inc)
+ cuda_cflags.append(inc)
- # flags
+ # Classify sources by extension to determine which rules are needed
+ with_c = any(Path(s).suffix.lower() == ".c" for s in sources)
ninja: list[str] = []
ninja.append("ninja_required_version = 1.3")
ninja.append("cxx = {}".format(os.environ.get("CXX", "cl" if IS_WINDOWS else "c++")))
- ninja.append("cflags = {}".format(" ".join(cflags)))
+ ninja.append("cxxflags = {}".format(" ".join(cxxflags)))
+ if with_c:
+ ninja.append("cc = {}".format(os.environ.get("CC", "cl" if IS_WINDOWS else "cc")))
+ ninja.append("cflags = {}".format(" ".join(cflags)))
if with_backend:
if with_hip:
ninja.append("nvcc = {}".format(str(Path(_find_rocm_home()) / "bin" / "hipcc")))
@@ -425,14 +449,25 @@ def _generate_ninja_build( # noqa: PLR0915, PLR0912
ninja.append("")
ninja.append("rule compile")
if IS_WINDOWS:
- ninja.append(" command = $cxx /showIncludes $cflags -c $in /Fo$out")
+ ninja.append(" command = $cxx /showIncludes $cxxflags -c $in /Fo$out")
ninja.append(" deps = msvc")
else:
ninja.append(" depfile = $out.d")
ninja.append(" deps = gcc")
- ninja.append(" command = $cxx -MMD -MF $out.d $cflags -c $in -o $out")
+ ninja.append(" command = $cxx -MMD -MF $out.d $cxxflags -c $in -o $out")
ninja.append("")
+ if with_c:
+ ninja.append("rule c_compile")
+ if IS_WINDOWS:
+ ninja.append(" command = $cc /showIncludes $cflags -c $in /Fo$out")
+ ninja.append(" deps = msvc")
+ else:
+ ninja.append(" depfile = $out.d")
+ ninja.append(" deps = gcc")
+ ninja.append(" command = $cc -MMD -MF $out.d $cflags -c $in -o $out")
+ ninja.append("")
+
if with_backend:
ninja.append("rule compile_cuda")
ninja.append(" depfile = $out.d")
@@ -447,38 +482,62 @@ def _generate_ninja_build( # noqa: PLR0915, PLR0912
# Add rules for object merging and cubin embedding (Unix only)
if not IS_WINDOWS:
- if embed_cubin:
- ninja.append("rule merge_objects")
- ninja.append(" command = ld -r -o $out $in")
- ninja.append("")
+ ninja.append("rule merge_objects")
+ ninja.append(" command = ld -r -o $out $in")
+ ninja.append("")
+ if embed_cubin:
ninja.append("rule embed_cubin")
ninja.append(
f" command = {sys.executable} -m tvm_ffi.utils.embed_cubin --output-obj $out --input-obj $in --cubin $cubin --name $name"
)
ninja.append("")
- ninja.append("rule link")
- if IS_WINDOWS:
- ninja.append(" command = $cxx $in /link $ldflags /out:$out")
- else:
- ninja.append(" command = $cxx $in $ldflags -o $out")
- ninja.append("")
+ if not object_mode:
+ ninja.append("rule link")
+ if IS_WINDOWS:
+ ninja.append(" command = $cxx $in /link $ldflags /out:$out")
+ else:
+ ninja.append(" command = $cxx $in $ldflags -o $out")
+ ninja.append("")
- # build targets
+ # build targets — dispatch by file extension
obj_files: list[str] = []
- for i, cpp_path in enumerate(sorted(cpp_files)):
- obj_name = f"cpp_{i}.o"
- ninja.append("build {}: compile {}".format(obj_name, cpp_path.replace(":", "$:")))
- obj_files.append(obj_name)
-
- for i, cuda_path in enumerate(sorted(cuda_files)):
- obj_name = f"cuda_{i}.o"
- ninja.append("build {}: compile_cuda {}".format(obj_name, cuda_path.replace(":", "$:")))
- obj_files.append(obj_name)
-
- # Use appropriate extension based on platform
- ext = ".dll" if IS_WINDOWS else ".so"
+ c_idx = cpp_idx = cuda_idx = 0
+ for src in sorted(sources):
+ ext = Path(src).suffix.lower()
+ escaped = src.replace(":", "$:")
+ if ext in (".o", ".obj"):
+ # Pre-compiled object file: pass directly to linker
+ obj_files.append(escaped)
+ elif ext == ".c":
+ obj_name = f"c_{c_idx}.o"
+ ninja.append(f"build {obj_name}: c_compile {escaped}")
+ obj_files.append(obj_name)
+ c_idx += 1
+ elif ext == ".cu":
+ obj_name = f"cuda_{cuda_idx}.o"
+ ninja.append(f"build {obj_name}: compile_cuda {escaped}")
+ obj_files.append(obj_name)
+ cuda_idx += 1
+ else:
+ # .cc, .cpp, .cxx — default to C++ compilation
+ obj_name = f"cpp_{cpp_idx}.o"
+ ninja.append(f"build {obj_name}: compile {escaped}")
+ obj_files.append(obj_name)
+ cpp_idx += 1
+
+ if object_mode:
+ # Object-only output: merge all object files into the target.
+ if not IS_WINDOWS:
+ ninja.append(f"build {output_name}: merge_objects {' '.join(obj_files)}")
+ ninja.append("")
+ ninja.append(f"default {output_name}")
+ else:
+ # Windows: no ld -r available; default to the first intermediate object
+ ninja.append(f"default {obj_files[0]}")
+ ninja.append("")
+ return "\n".join(ninja)
# For Unix systems with embed_cubin, use a 3-step process:
# 1. Merge all object files into a unified object file
@@ -507,16 +566,16 @@ def _generate_ninja_build( # noqa: PLR0915, PLR0912
current_obj = next_obj
# Step 3: Link the final object file
- ninja.append(f"build {name}{ext}: link {current_obj}")
+ ninja.append(f"build {output_name}: link {current_obj}")
ninja.append("")
else:
- # Original behavior: directly link object files (for Windows or no cubin embedding)
+ # Directly link object files (for Windows or no cubin embedding)
link_files_str = " ".join(obj_files)
- ninja.append(f"build {name}{ext}: link {link_files_str}")
+ ninja.append(f"build {output_name}: link {link_files_str}")
ninja.append("")
# default target
- ninja.append(f"default {name}{ext}")
+ ninja.append(f"default {output_name}")
ninja.append("")
return "\n".join(ninja)
@@ -595,8 +654,7 @@ def _str_seq2list(seq: Sequence[str] | str | None) -> list[str]:
def _build_impl( # noqa: PLR0913
name: str,
- cpp_files: Sequence[str] | str | None,
- cuda_files: Sequence[str] | str | None,
+ sources: Sequence[str] | str | None,
extra_cflags: Sequence[str] | None,
extra_cuda_cflags: Sequence[str] | None,
extra_ldflags: Sequence[str] | None,
@@ -605,16 +663,15 @@ def _build_impl( # noqa: PLR0913
need_lock: bool = True,
embed_cubin: Mapping[str, bytes] | None = None,
backend: str | None = None,
+ output: str | None = None,
) -> str:
"""Real implementation of build function."""
# need to resolve the path to make it unique
- cpp_path_list = [str(Path(p).resolve()) for p in _str_seq2list(cpp_files)]
- cuda_path_list = [str(Path(p).resolve()) for p in _str_seq2list(cuda_files)]
- with_cpp = bool(cpp_path_list)
- with_backend = bool(cuda_path_list)
- assert with_cpp or with_backend, "Either cpp_files or cuda_files must be provided."
+ source_path_list = [str(Path(p).resolve()) for p in _str_seq2list(sources)]
+ assert source_path_list, "sources must be provided."
- resolved_backend = _resolve_gpu_backend(backend) if with_backend else None
+ has_cuda = any(Path(p).suffix.lower() == ".cu" for p in source_path_list)
+ resolved_backend = _resolve_gpu_backend(backend) if has_cuda else None
extra_ldflags_list = list(extra_ldflags) if extra_ldflags is not None else []
extra_cflags_list = list(extra_cflags) if extra_cflags is not None else []
extra_cuda_cflags_list = list(extra_cuda_cflags) if extra_cuda_cflags is not None else []
@@ -626,8 +683,8 @@ def _build_impl( # noqa: PLR0913
source_hash: str = _hash_sources(
None,
None,
- cpp_path_list,
- cuda_path_list,
+ source_path_list,
+ None,
{},
extra_cflags_list,
extra_cuda_cflags_list,
@@ -658,10 +715,10 @@ def _build_impl( # noqa: PLR0913
extra_cuda_cflags=extra_cuda_cflags_list,
extra_ldflags=extra_ldflags_list,
extra_include_paths=extra_include_paths_list,
- cpp_files=cpp_path_list,
- cuda_files=cuda_path_list,
+ sources=source_path_list,
embed_cubin=embed_cubin,
backend=resolved_backend,
+ output=output,
)
# may not hold lock when build_directory is specified, prevent deadlock
@@ -670,9 +727,26 @@ def _build_impl( # noqa: PLR0913
_maybe_write(str(build_dir / "build.ninja"), ninja_source)
# build the module
build_ninja(str(build_dir))
- # Use appropriate extension based on platform
- ext = ".dll" if IS_WINDOWS else ".so"
- return str((build_dir / f"{name}{ext}").resolve())
+ # Determine the output filename (mirrors _generate_ninja_build logic)
+ if output is not None:
+ out_ext = Path(output).suffix.lower()
+ object_mode = out_ext in (".o", ".obj")
+ output_name = Path(output).name
+ else:
+ object_mode = False
+ output_name = f"{name}{'.dll' if IS_WINDOWS else '.so'}"
+ if object_mode and IS_WINDOWS:
+ # Windows has no ld -r; the actual target is the first intermediate object.
+ # The name must match _generate_ninja_build: c_0.o / cpp_0.o / cuda_0.o.
+ first_ext = Path(sorted(source_path_list)[0]).suffix.lower() if source_path_list else ""
+ if first_ext == ".c":
+ obj_name = "c_0.o"
+ elif first_ext == ".cu":
+ obj_name = "cuda_0.o"
+ else:
+ obj_name = "cpp_0.o"
+ return str((build_dir / obj_name).resolve())
+ return str((build_dir / output_name).resolve())
def build_inline( # noqa: PLR0913
@@ -688,12 +762,15 @@ def build_inline( # noqa: PLR0913
build_directory: str | None = None,
embed_cubin: Mapping[str, bytes] | None = None,
backend: str | None = None,
+ output: str | None = None,
) -> str:
"""Compile and build a C++/CUDA module from inline source code.
- This function compiles the given C++ and/or CUDA source code into a shared library. Both ``cpp_sources`` and
- ``cuda_sources`` are compiled to an object file, and then linked together into a shared library. It's possible to only
- provide cpp_sources or cuda_sources. The path to the compiled shared library is returned.
+ This function compiles the given C++ and/or CUDA source code into a shared library or object file.
+ Both ``cpp_sources`` and ``cuda_sources`` are compiled to an object file. When ``output`` is
+ ``None`` (the default) or has a shared-library extension (``.so``, ``.dll``), object files are
+ linked into a shared library. When ``output`` has an object-file extension (``.o``, ``.obj``),
+ linking is skipped and the path to the object file is returned directly.
The ``functions`` parameter is used to specify which functions in the source code should be exported to the tvm ffi
module. It can be a mapping, a sequence, or a single string. When a mapping is given, the keys are the names of the
@@ -766,10 +843,16 @@ def build_inline( # noqa: PLR0913
The GPU backend to use. It can be "cuda" or "hip".
If not specified, the backend will be automatically determined based on the available GPU and the provided source code.
+ output
+ Output filename that determines the build type from its extension. When ``None``
+ (the default), builds a shared library (``.so`` on Unix, ``.dll`` on Windows).
+ Use an object-file extension (e.g., ``"hello.o"``) to skip linking and produce
+ a relocatable object file. The file is placed in the build directory.
+
Returns
-------
- lib_path: str
- The path to the built shared library.
+ path: str
+ The path to the built shared library or object file.
Example
-------
@@ -874,10 +957,14 @@ def build_inline( # noqa: PLR0913
if with_backend:
_maybe_write(cuda_file, cuda_source)
+ src_files = []
+ if with_cpp:
+ src_files.append(cpp_file)
+ if with_backend:
+ src_files.append(cuda_file)
return _build_impl(
name=name,
- cpp_files=[cpp_file] if with_cpp else [],
- cuda_files=[cuda_file] if with_backend else [],
+ sources=src_files,
extra_cflags=extra_cflags_list,
extra_cuda_cflags=extra_cuda_cflags_list,
extra_ldflags=extra_ldflags_list,
@@ -886,6 +973,7 @@ def build_inline( # noqa: PLR0913
need_lock=False, # already hold the lock
embed_cubin=embed_cubin,
backend=backend,
+ output=output,
)
@@ -1049,9 +1137,10 @@ def load_inline( # noqa: PLR0913
)
-def build(
+def build( # noqa: PLR0913
name: str,
*,
+ sources: Sequence[str] | str | None = None,
cpp_files: Sequence[str] | str | None = None,
cuda_files: Sequence[str] | str | None = None,
extra_cflags: Sequence[str] | None = None,
@@ -1060,12 +1149,20 @@ def build(
extra_include_paths: Sequence[str] | None = None,
build_directory: str | None = None,
backend: str | None = None,
+ output: str | None = None,
) -> str:
- """Compile and build a C++/CUDA module from source files.
+ """Compile and build a C/C++/CUDA module from source files.
+
+ This function compiles the given C, C++, and/or CUDA source files into a shared library or
+ object file. The compiler is selected automatically based on file extension:
- This function compiles the given C++ and/or CUDA source files into a shared library. Both ``cpp_files`` and
- ``cuda_files`` are compiled to object files, and then linked together into a shared library. It's possible to only
- provide cpp_files or cuda_files. The path to the compiled shared library is returned.
+ - ``.c`` — compiled with the C compiler (``$CC``)
+ - ``.cc``, ``.cpp``, ``.cxx`` — compiled with the C++ compiler (``$CXX``)
+ - ``.o``, ``.obj`` — pre-compiled objects, passed directly to the linker
+
+ When ``output`` is ``None`` (the default) or has a shared-library extension, object files are
+ linked into a shared library. When ``output`` has an object-file extension (``.o``, ``.obj``),
+ linking is skipped and the path to the object file is returned.
Note that this function does not automatically export functions to the tvm ffi module. You need to
manually use the TVM FFI export macros (e.g., ``TVM_FFI_DLL_EXPORT_TYPED_FUNC``) in your source files to export
@@ -1084,23 +1181,35 @@ def build(
The default cache directory can be specified via the `TVM_FFI_CACHE_DIR` environment variable. If not specified,
the default cache directory is ``~/.cache/tvm-ffi``.
+ The C compiler is controlled by the ``$CC`` environment variable (default: ``cc`` on Unix, ``cl`` on Windows).
+ The C++ compiler is controlled by the ``$CXX`` environment variable (default: ``c++`` on Unix, ``cl`` on Windows).
+
Parameters
----------
name
The name of the tvm ffi module.
+ sources
+ Source files to compile. The compiler is auto-detected from the file extension:
+
+ - ``.c`` → C compiler (``$CC``)
+ - ``.cc``, ``.cpp``, ``.cxx`` → C++ compiler (``$CXX``)
+ - ``.cu`` → CUDA/HIP compiler (``nvcc`` or ``hipcc``)
+ - ``.o``, ``.obj`` → pre-compiled objects, passed directly to the linker
+
+ It can be a list of file paths or a single file path.
cpp_files
- The C++ source files to compile. It can be a list of file paths or a single file path. Both absolute and
- relative paths are supported.
+ Alias for ``sources``, kept for backward compatibility.
cuda_files
- The CUDA source files to compile. It can be a list of file paths or a single file path. Both absolute and
- relative paths are supported.
+ Alias for ``sources``, kept for backward compatibility.
extra_cflags
- The extra compiler flags for C++ compilation.
- The default flags are:
+ Extra compiler flags applied to both C and C++ compilation.
+ The C++ default flags are:
- On Linux/macOS: ['-std=c++17', '-fPIC', '-O2']
- On Windows: ['/std:c++17', '/MD', '/O2']
+ The C default flags omit ``-std=c++17`` and ``/EHsc``.
+
extra_cuda_cflags
The extra compiler flags for CUDA compilation.
The default flags are:
@@ -1127,10 +1236,16 @@ def build(
The GPU backend to use. It can be "cuda" or "hip".
If not specified, the backend will be automatically determined based on the available GPU and the provided source code.
+ output
+ Output filename that determines the build type from its extension. When ``None``
+ (the default), builds a shared library (``.so`` on Unix, ``.dll`` on Windows).
+ Use an object-file extension (e.g., ``"my_ops.o"``) to skip linking and produce
+ a relocatable object file. The file is placed in the build directory.
+
Returns
-------
- lib_path: str
- The path to the built shared library.
+ path: str
+ The path to the built shared library or object file.
Example
-------
@@ -1167,7 +1282,7 @@ def build(
# compile the cpp source file and get the library path
lib_path: str = tvm_ffi.cpp.build(
name="my_ops",
- cpp_files="my_ops.cpp",
+ sources="my_ops.cpp",
)
# load the module
@@ -1180,10 +1295,11 @@ def build(
torch.testing.assert_close(x + 1, y)
"""
+ # Merge sources, cpp_files, and cuda_files (backward compat aliases)
+ merged = _str_seq2list(sources) + _str_seq2list(cpp_files) + _str_seq2list(cuda_files)
return _build_impl(
name=name,
- cpp_files=cpp_files,
- cuda_files=cuda_files,
+ sources=merged or None,
extra_cflags=extra_cflags,
extra_cuda_cflags=extra_cuda_cflags,
extra_ldflags=extra_ldflags,
@@ -1191,12 +1307,14 @@ def build(
build_directory=build_directory,
need_lock=True,
backend=backend,
+ output=output,
)
-def load(
+def load( # noqa: PLR0913
name: str,
*,
+ sources: Sequence[str] | str | None = None,
cpp_files: Sequence[str] | str | None = None,
cuda_files: Sequence[str] | str | None = None,
extra_cflags: Sequence[str] | None = None,
@@ -1207,11 +1325,10 @@ def load(
keep_module_alive: bool = True,
backend: str | None = None,
) -> Module:
- """Compile, build and load a C++/CUDA module from source files.
+ """Compile, build and load a C/C++/CUDA module from source files.
- This function compiles the given C++ and/or CUDA source files into a shared library and loads it as a tvm ffi
- module. Both ``cpp_files`` and ``cuda_files`` are compiled to object files, and then linked together into a shared
- library. It's possible to only provide cpp_files or cuda_files.
+ This function compiles the given source files into a shared library and loads it as a tvm ffi
+ module. The compiler is selected automatically based on file extension.
Note that this function does not automatically export functions to the tvm ffi module. You need to
manually use the TVM FFI export macros (e.g., :c:macro:`TVM_FFI_DLL_EXPORT_TYPED_FUNC`) in your source files to export
@@ -1234,12 +1351,14 @@ def load(
----------
name
The name of the tvm ffi module.
+ sources
+ Source files to compile. The compiler is auto-detected from the file extension:
+ ``.c`` → C, ``.cc``/``.cpp``/``.cxx`` → C++, ``.cu`` → CUDA/HIP,
+ ``.o``/``.obj`` → linker passthrough. It can be a list of file paths or a single file path.
cpp_files
- The C++ source files to compile. It can be a list of file paths or a single file path. Both absolute and
- relative paths are supported.
+ Alias for ``sources``, kept for backward compatibility.
cuda_files
- The CUDA source files to compile. It can be a list of file paths or a single file path. Both absolute and
- relative paths are supported.
+ Alias for ``sources``, kept for backward compatibility.
extra_cflags
The extra compiler flags for C++ compilation.
The default flags are:
@@ -1321,7 +1440,7 @@ def load(
# compile the cpp source file and load the module
mod: Module = tvm_ffi.cpp.load(
name="my_ops",
- cpp_files="my_ops.cpp",
+ sources="my_ops.cpp",
)
# use the function from the loaded module
@@ -1334,6 +1453,7 @@ def load(
return load_module(
build(
name=name,
+ sources=sources,
cpp_files=cpp_files,
cuda_files=cuda_files,
extra_cflags=extra_cflags,
diff --git a/tests/lint/check_file_type.py b/tests/lint/check_file_type.py
index 9849274..f43e917 100644
--- a/tests/lint/check_file_type.py
+++ b/tests/lint/check_file_type.py
@@ -49,6 +49,7 @@ ALLOW_EXTENSION = {
"cu",
"cuh",
"bat",
+ "ps1",
# configurations
"mk",
"in",
diff --git a/tests/python/test_build.py b/tests/python/test_build.py
index 8844616..6cf717b 100644
--- a/tests/python/test_build.py
+++ b/tests/python/test_build.py
@@ -28,7 +28,7 @@ def test_build_cpp() -> None:
cpp_path = pathlib.Path(__file__).parent.resolve() / "test_build.cc"
output_lib_path = tvm_ffi.cpp.build(
name="hello",
- cpp_files=[str(cpp_path)],
+ sources=[str(cpp_path)],
)
mod: Module = tvm_ffi.load_module(output_lib_path)
diff --git a/tests/python/test_build_inline.py b/tests/python/test_build_inline.py
index fa0e6c9..e754e16 100644
--- a/tests/python/test_build_inline.py
+++ b/tests/python/test_build_inline.py
@@ -14,6 +14,9 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+import tempfile
+from pathlib import Path
+
import numpy
import pytest
import tvm_ffi.cpp
@@ -48,5 +51,73 @@ def test_build_inline_cpp() -> None:
numpy.testing.assert_equal(x + 1, y)
+def test_build_c_file() -> None:
+ """Test building a C source file to .o and linking into .so."""
+ c_source = r"""
+#include <tvm/ffi/c_api.h>
+
+TVM_FFI_DLL_EXPORT int __tvm_ffi_c_add(
+ void* self, const TVMFFIAny* args, int32_t num_args, TVMFFIAny* result) {
+ result->type_index = kTVMFFIInt;
+ result->zero_padding = 0;
+ result->v_int64 = args[0].v_int64 + args[1].v_int64;
+ return 0;
+}
+"""
+ with tempfile.TemporaryDirectory() as tmpdir:
+ c_file = str(Path(tmpdir) / "test_c.c")
+ Path(c_file).write_text(c_source)
+
+ obj_path = tvm_ffi.cpp.build(
+ name="test_c",
+ sources=[c_file],
+ output="test_c.o",
+ )
+ assert obj_path.endswith(".o")
+
+ lib_path = tvm_ffi.cpp.build(
+ name="test_c_lib",
+ sources=[obj_path],
+ )
+
+ mod: Module = tvm_ffi.load_module(lib_path)
+ assert mod.c_add(10, 20) == 30
+
+
+def test_build_inline_object_output() -> None:
+ """Test building a .o then linking it into a .so and calling the function."""
+ obj_path = tvm_ffi.cpp.build_inline(
+ name="hello_obj",
+ cpp_sources=r"""
+ void add_one_cpu(tvm::ffi::TensorView x, tvm::ffi::TensorView y) {
+ TVM_FFI_ICHECK(x.ndim() == 1) << "x must be a 1D tensor";
+ DLDataType f32_dtype{kDLFloat, 32, 1};
+ TVM_FFI_ICHECK(x.dtype() == f32_dtype) << "x must be a float tensor";
+ TVM_FFI_ICHECK(y.ndim() == 1) << "y must be a 1D tensor";
+ TVM_FFI_ICHECK(y.dtype() == f32_dtype) << "y must be a float tensor";
+ TVM_FFI_ICHECK(x.size(0) == y.size(0)) << "x and y must have the same shape";
+ for (int i = 0; i < x.size(0); ++i) {
+ static_cast<float*>(y.data_ptr())[i] = static_cast<float*>(x.data_ptr())[i] + 1;
+ }
+ }
+ """,
+ functions=["add_one_cpu"],
+ output="hello_obj.o",
+ )
+ assert obj_path.endswith(".o")
+
+ # Link the object file into a shared library
+ lib_path = tvm_ffi.cpp.build(
+ name="hello_from_obj",
+ sources=[obj_path],
+ )
+
+ mod: Module = tvm_ffi.load_module(lib_path)
+ x = numpy.array([1, 2, 3, 4, 5], dtype=numpy.float32)
+ y = numpy.empty_like(x)
+ mod.add_one_cpu(x, y)
+ numpy.testing.assert_equal(x + 1, y)
+
+
if __name__ == "__main__":
pytest.main([__file__])
diff --git a/tests/python/test_stl.py b/tests/python/test_stl.py
index 3ba70b8..2468a84 100644
--- a/tests/python/test_stl.py
+++ b/tests/python/test_stl.py
@@ -25,7 +25,7 @@ def test_stl() -> None:
cpp_path = pathlib.Path(__file__).parent.resolve() / "cpp_src" / "test_stl.cc"
output_lib_path = tvm_ffi.cpp.build(
name="test_stl",
- cpp_files=[str(cpp_path)],
+ sources=[str(cpp_path)],
)
mod: Module = tvm_ffi.load_module(output_lib_path)