diff --git a/.github/actions/build-wheel-for-publish/action.yml b/.github/actions/build-wheel-for-publish/action.yml index 9039669e486f..6333856f83dc 100644 --- a/.github/actions/build-wheel-for-publish/action.yml +++ b/.github/actions/build-wheel-for-publish/action.yml @@ -28,7 +28,7 @@ inputs: description: "cibuildwheel build selector (e.g., cp310-manylinux_x86_64)" required: true cmake_defines: - description: "Feature CMake defines that vary by wheel (e.g. -DUSE_METAL=ON on macOS)" + description: "Feature CMake defines that vary by wheel (e.g. -DTVM_TIRX_USE_METAL=ON on macOS)" required: false default: "" include_cuda_runtime: @@ -122,22 +122,22 @@ runs: # which build without a container. CIBW_CONTAINER_ENGINE: "docker; create_args: --volume /opt/llvm:/opt/llvm:ro" # Each per-OS block carries the toolchain location (conda LLVM prefix + - # llvm-config path), the universal static-link policy (USE_LLVM --link-static + # llvm-config path), the universal static-link policy (TVM_TIRX_USE_LLVM --link-static # and ZLIB_USE_STATIC_LIBS), and the CUDA sidecar. Feature defines that vary by - # wheel (e.g. USE_METAL on macOS) come from the matrix via inputs.cmake_defines. + # wheel (e.g. TVM_TIRX_USE_METAL on macOS) come from the matrix via inputs.cmake_defines. # One explicit block per build platform (macOS / Linux / Windows); cibuildwheel # env overrides replace rather than merge, so there is no shared base block. CIBW_ENVIRONMENT_MACOS: >- CMAKE_PREFIX_PATH="/opt/llvm" - CMAKE_ARGS="-DUSE_LLVM='/opt/llvm/bin/llvm-config --link-static' -DZLIB_USE_STATIC_LIBS=ON -DCMAKE_PREFIX_PATH=/opt/llvm ${{ inputs.cmake_defines }}" + CMAKE_ARGS="-DTVM_TIRX_USE_LLVM='/opt/llvm/bin/llvm-config --link-static' -DZLIB_USE_STATIC_LIBS=ON -DCMAKE_PREFIX_PATH=/opt/llvm ${{ inputs.cmake_defines }}" CIBW_ENVIRONMENT_LINUX: >- CMAKE_PREFIX_PATH="/opt/llvm" LIBRARY_PATH="/opt/llvm/lib" - CMAKE_ARGS="-DUSE_LLVM='/opt/llvm/bin/llvm-config --link-static' -DZLIB_USE_STATIC_LIBS=ON -DCMAKE_PREFIX_PATH=/opt/llvm ${{ inputs.cmake_defines }} -DTVM_PACKAGE_EXTRA_LIBS=/project/build-wheel-cuda/lib/libtvm_runtime_cuda.so" + CMAKE_ARGS="-DTVM_TIRX_USE_LLVM='/opt/llvm/bin/llvm-config --link-static' -DZLIB_USE_STATIC_LIBS=ON -DCMAKE_PREFIX_PATH=/opt/llvm ${{ inputs.cmake_defines }} -DTVM_TIRX_PACKAGE_EXTRA_LIBS=/project/build-wheel-cuda/lib/libtvm_runtime_cuda.so" CIBW_ENVIRONMENT_WINDOWS: >- CMAKE_PREFIX_PATH="C:/opt/llvm/Library" PATH="C:/opt/llvm/Library/bin;$PATH" - CMAKE_ARGS="-DUSE_LLVM='C:/opt/llvm/Library/bin/llvm-config.exe --link-static' -DZLIB_USE_STATIC_LIBS=ON -DCMAKE_PREFIX_PATH=C:/opt/llvm/Library ${{ inputs.cmake_defines }} -DTVM_PACKAGE_EXTRA_LIBS=${{ env.TVM_CUDA_EXTRA_LIB }}" + CMAKE_ARGS="-DTVM_TIRX_USE_LLVM='C:/opt/llvm/Library/bin/llvm-config.exe --link-static' -DZLIB_USE_STATIC_LIBS=ON -DCMAKE_PREFIX_PATH=C:/opt/llvm/Library ${{ inputs.cmake_defines }} -DTVM_TIRX_PACKAGE_EXTRA_LIBS=${{ env.TVM_CUDA_EXTRA_LIB }}" # Turns the wheel-specific assertions in # tirx/tests/python/all-platform-minimal-test/test_validate_runtime_library.py # ON (they skip unless these are set). Every published wheel is LLVM-enabled, diff --git a/.github/actions/setup/action.yml b/.github/actions/setup/action.yml index 19bcedbcd2b3..20e1266e51c7 100644 --- a/.github/actions/setup/action.yml +++ b/.github/actions/setup/action.yml @@ -44,4 +44,4 @@ runs: - name: Install tvm-ffi pip package shell: bash -l {0} run: | - pip install -v ./3rdparty/tvm-ffi + pip install -v ./tirx/3rdparty/tvm-ffi diff --git a/.github/workflows/main.yml b/.github/workflows/main.yml index 8603fdc95f39..929932900854 100644 --- a/.github/workflows/main.yml +++ b/.github/workflows/main.yml @@ -57,7 +57,7 @@ jobs: run: | set -euxo pipefail pip install scikit-build-core - export CMAKE_ARGS="-DUSE_LLVM=ON -DUSE_METAL=ON -DBUILD_TESTING=OFF" + export CMAKE_ARGS="-DTVM_TIRX_USE_LLVM=ON -DTVM_TIRX_USE_METAL=ON -DTVM_TIRX_BUILD_TESTS=OFF" pip wheel --no-deps -w dist ./tirx -v export PIP_FIND_LINKS="$(pwd)/dist" export CMAKE_ARGS="-DBUILD_TESTING=OFF" @@ -122,7 +122,7 @@ jobs: CMAKE_BUILD_PARALLEL_LEVEL: ${{ steps.env_vars.outputs.cpu_count }} run: | pip install scikit-build-core || exit /b 1 - set CMAKE_ARGS=-DUSE_LLVM=ON -DBUILD_TESTING=OFF + set CMAKE_ARGS=-DTVM_TIRX_USE_LLVM=ON -DTVM_TIRX_BUILD_TESTS=OFF pip wheel --no-deps -w dist ./tirx -v || exit /b 1 set PIP_FIND_LINKS=%CD%\dist set CMAKE_ARGS=-DBUILD_TESTING=OFF diff --git a/.github/workflows/publish_wheel.yml b/.github/workflows/publish_wheel.yml index 45fcf9c59834..c43c4727a9fa 100644 --- a/.github/workflows/publish_wheel.yml +++ b/.github/workflows/publish_wheel.yml @@ -142,7 +142,7 @@ jobs: build: cp310-macosx_arm64 include_cuda_runtime: "0" artifact_suffix: macos-arm64 - cmake_defines: "-DUSE_METAL=ON" + cmake_defines: "-DTVM_TIRX_USE_METAL=ON" - name: "Windows AMD64 wheel with CUDA runtime" os: windows-2022 arch: AMD64 diff --git a/.gitmodules b/.gitmodules index a7cbc9d0191d..8a2a3ef81d6f 100644 --- a/.gitmodules +++ b/.gitmodules @@ -11,5 +11,5 @@ path = 3rdparty/libflash_attn url = https://github.com/tlc-pack/libflash_attn [submodule "3rdparty/tvm-ffi"] - path = 3rdparty/tvm-ffi + path = tirx/3rdparty/tvm-ffi url = https://github.com/apache/tvm-ffi diff --git a/CMakeLists.txt b/CMakeLists.txt index 4bd18b92a06b..df54344dec9a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -34,20 +34,59 @@ if(TVM_BUILD_PYTHON_MODULE) endif() if(USE_CLML_GRAPH_EXECUTOR) if(USE_CLML_GRAPH_EXECUTOR STREQUAL "ON") - set(USE_OPENCL "${CMAKE_CURRENT_SOURCE_DIR}/clml") + set(TVM_TIRX_USE_OPENCL "${CMAKE_CURRENT_SOURCE_DIR}/clml") else() - set(USE_OPENCL "${USE_CLML_GRAPH_EXECUTOR}") + set(TVM_TIRX_USE_OPENCL "${USE_CLML_GRAPH_EXECUTOR}") endif() endif() if(TVM_USE_INSTALLED_TIRX) find_package(Python COMPONENTS Interpreter REQUIRED) if(NOT tvm_tirx_DIR) - execute_process(COMMAND "${Python_EXECUTABLE}" -c - "import importlib.util; from pathlib import Path; print(Path(importlib.util.find_spec('tvm_tirx').origin).parent / 'lib/cmake/tvm_tirx')" + execute_process(COMMAND "${Python_EXECUTABLE}" -m tvm_tirx.config --cmakedir OUTPUT_VARIABLE tvm_tirx_DIR OUTPUT_STRIP_TRAILING_WHITESPACE COMMAND_ERROR_IS_FATAL ANY) endif() find_package(tvm_tirx CONFIG REQUIRED) + if(NOT TARGET tvm_ffi::shared) + if(NOT tvm_ffi_DIR) + execute_process(COMMAND "${Python_EXECUTABLE}" -m tvm_ffi.config --cmakedir + OUTPUT_VARIABLE tvm_ffi_DIR OUTPUT_STRIP_TRAILING_WHITESPACE + COMMAND_ERROR_IS_FATAL ANY) + endif() + find_package(tvm_ffi CONFIG REQUIRED) + endif() + # Parent extensions use installed sidecars privately. They are not part of + # the public TIRx package target interface. + get_target_property(_tirx_runtime_location tvm_runtime::shared IMPORTED_LOCATION) + get_filename_component(_tirx_libdir "${_tirx_runtime_location}" DIRECTORY) + foreach(_backend cuda rocm opencl vulkan metal) + find_library(_tirx_backend_library NAMES tvm_runtime_${_backend} + PATHS "${_tirx_libdir}" NO_DEFAULT_PATH NO_CACHE) + if(_tirx_backend_library AND NOT TARGET tvm_runtime_${_backend}) + add_library(tvm_runtime_${_backend} UNKNOWN IMPORTED) + set_target_properties(tvm_runtime_${_backend} PROPERTIES + IMPORTED_LOCATION "${_tirx_backend_library}") + target_link_libraries(tvm_runtime_${_backend} INTERFACE tvm_runtime::shared tvm_ffi::shared) + endif() + unset(_tirx_backend_library) + endforeach() + foreach(_backend CUDA ROCM) + if(NOT DEFINED TVM_TIRX_USE_${_backend}) + string(TOLOWER "${_backend}" _backend_lower) + set(TVM_TIRX_USE_${_backend} OFF) + if(TARGET tvm_runtime_${_backend_lower}) + set(TVM_TIRX_USE_${_backend} ON) + endif() + endif() + endforeach() + if(TARGET tvm_runtime_opencl) + if(WIN32) + target_compile_definitions(tvm_runtime_opencl INTERFACE "TVM_OPENCL_DLL=__declspec(dllimport)") + else() + target_compile_definitions(tvm_runtime_opencl INTERFACE + "TVM_OPENCL_DLL=__attribute__((visibility(\"default\")))") + endif() + endif() if(NOT TVM_TIRX_VERSION STREQUAL "0.1.0a0") message(FATAL_ERROR "TVM requires apache-tvm-tirx 0.1.0a0") endif() @@ -55,9 +94,31 @@ else() if(NOT TVM_OPENCL_HEADERS_ROOT AND EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/3rdparty/OpenCL-Headers/CL/cl.h") set(TVM_OPENCL_HEADERS_ROOT "${CMAKE_CURRENT_SOURCE_DIR}/3rdparty/OpenCL-Headers") endif() + # Keep source-built core libraries beside the parent libraries. + set(_tirx_output_defaults "") + foreach(_kind ARCHIVE LIBRARY RUNTIME) + if(NOT DEFINED CMAKE_${_kind}_OUTPUT_DIRECTORY) + set(CMAKE_${_kind}_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib") + list(APPEND _tirx_output_defaults ${_kind}) + endif() + endforeach() add_subdirectory(tirx) + if(TARGET tirx_cpptest) + set_target_properties(tirx_cpptest PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}") + foreach(_config IN LISTS CMAKE_CONFIGURATION_TYPES) + string(TOUPPER "${_config}" _config_upper) + set_target_properties(tirx_cpptest PROPERTIES + RUNTIME_OUTPUT_DIRECTORY_${_config_upper} "${CMAKE_BINARY_DIR}") + endforeach() + endif() + foreach(_kind IN LISTS _tirx_output_defaults) + unset(CMAKE_${_kind}_OUTPUT_DIRECTORY) + endforeach() endif() -foreach(_utility Utils Summary Linker Library FindCUDA FindROCM FindOpenCL CCache) +foreach(_utility Utils Summary Linker Library CCache) + include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/${_utility}.cmake") +endforeach() +foreach(_utility FindCUDA FindROCM FindOpenCL) include("${TVM_TIRX_CMAKE_DIR}/utils/${_utility}.cmake") endforeach() foreach(_collective NCCL RCCL NVSHMEM) @@ -66,25 +127,20 @@ endforeach() if(NOT TARGET tvm_ffi_header) add_library(tvm_ffi_header ALIAS tvm_ffi::header) endif() -# Keep existing integration modules' target names usable with an installed core. -foreach(_library tvm_runtime tvm_runtime_static tvm_runtime_cuda tvm_runtime_rocm tvm_runtime_opencl tvm_runtime_vulkan tvm_runtime_metal) - if(NOT TARGET ${_library} AND TARGET tvm_tirx::${_library}) - add_library(${_library} ALIAS tvm_tirx::${_library}) - endif() -endforeach() -set(BUILD_FOR_ANDROID "${TVM_TIRX_BUILD_FOR_ANDROID}") +# Existing parent integration modules use the unqualified runtime target. +if(NOT TARGET tvm_runtime) + add_library(tvm_runtime ALIAS tvm_runtime::shared) +endif() +set(BUILD_FOR_ANDROID "${ANDROID}") set(CMAKE_CXX_STANDARD 20) set(CMAKE_CXX_STANDARD_REQUIRED ON) set(CMAKE_CUDA_STANDARD 20) set(CMAKE_CUDA_STANDARD_REQUIRED ON) set(CMAKE_POSITION_INDEPENDENT_CODE ON) set(CMAKE_EXPORT_COMPILE_COMMANDS ON) -tvm_option(USE_CUDA "Build with CUDA" "${TVM_TIRX_USE_CUDA}") tvm_option(USE_NCCL "Build with NCCL" OFF) -tvm_option(USE_ROCM "Build with ROCM" "${TVM_TIRX_USE_ROCM}") tvm_option(USE_RCCL "Build with RCCL" OFF) tvm_option(USE_RPC "Build with RPC" ON) -tvm_option(USE_THREADS "Build with thread support" ON) tvm_option(TVM_LOG_BEFORE_THROW "Whether log before throw, for debugging purposes" OFF) tvm_option(USE_MSVC_MT "Build with MT" OFF) tvm_option(INSTALL_DEV "Install compiler infrastructure" OFF) @@ -173,7 +229,9 @@ tvm_file_glob(GLOB_RECURSE COMPILER_SRCS add_library(tvm_runtime_extra SHARED) add_library(tvm_runtime_extra_defs INTERFACE) target_link_libraries(tvm_runtime_extra_defs INTERFACE - tvm_tirx::tvm_runtime tvm_tirx::private_headers tvm_ffi::header) + tvm_runtime::shared tvm_ffi::shared) +target_include_directories(tvm_runtime_extra_defs SYSTEM INTERFACE + "${CMAKE_CURRENT_SOURCE_DIR}/tirx/3rdparty/compiler-rt") target_compile_definitions(tvm_runtime_extra_defs INTERFACE TVM_RUNTIME_EXTRA_EXPORTS TVM_FFI_EXPORTS) if(HIDE_PRIVATE_SYMBOLS AND NOT MSVC) target_compile_options(tvm_runtime_extra_defs INTERFACE -fvisibility=hidden) @@ -188,11 +246,11 @@ if(USE_RPC) target_link_libraries(tvm_runtime_extra PRIVATE ws2_32) endif() endif() -if(USE_CUDA) - if(NOT TARGET tvm_tirx::tvm_runtime_cuda) +if(TVM_TIRX_USE_CUDA) + if(NOT TARGET tvm_runtime_cuda) message(FATAL_ERROR "Parent CUDA support requires a CUDA-enabled core") endif() - find_cuda(${USE_CUDA} ${USE_CUDNN}) + find_cuda(${TVM_TIRX_USE_CUDA} ${USE_CUDNN}) enable_language(CUDA) if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES) set(CMAKE_CUDA_ARCHITECTURES native) @@ -201,11 +259,15 @@ if(USE_CUDA) tvm_file_glob(GLOB _vm_cuda_sources src/runtime/vm/cuda/*.cc) target_sources(tvm_runtime_extra PRIVATE ${_vm_cuda_sources}) endif() -if(USE_ROCM) - if(NOT TARGET tvm_tirx::tvm_runtime_rocm) +if(TVM_TIRX_USE_ROCM) + if(NOT TARGET tvm_runtime_rocm) message(FATAL_ERROR "Parent ROCm support requires a ROCm-enabled core") endif() - find_rocm(${USE_ROCM}) + if("${TVM_TIRX_USE_ROCM}" MATCHES ${IS_TRUE_PATTERN} AND IS_DIRECTORY "${TVM_TIRX_ROCM_PATH}") + find_rocm("${TVM_TIRX_ROCM_PATH}") + else() + find_rocm(${TVM_TIRX_USE_ROCM}) + endif() include_directories(SYSTEM ${ROCM_INCLUDE_DIRS}) add_compile_definitions(__HIP_PLATFORM_HCC__=1 __HIP_PLATFORM_AMD__=1) endif() @@ -227,8 +289,13 @@ include(cmake/modules/contrib/NNAPI.cmake) include(cmake/modules/contrib/ExampleNPU.cmake) include(cmake/modules/contrib/vllm.cmake) if(USE_CLML_GRAPH_EXECUTOR) - target_link_libraries(tvm_clml_objs PRIVATE tvm_tirx::tvm_runtime_opencl) - target_link_libraries(tvm_runtime_extra PRIVATE tvm_tirx::tvm_runtime_opencl) + target_link_libraries(tvm_clml_objs PRIVATE tvm_runtime_opencl) + target_link_libraries(tvm_runtime_extra PRIVATE tvm_runtime_opencl) +endif() +if(TARGET tvm_clml_objs) + # CLML shares OpenCL implementation types; these are source inputs, not an + # installed core interface. The parent sdist carries their bounded closure. + target_include_directories(tvm_clml_objs PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/tirx/src") endif() # Retain the upstream runtime-extra assembly module and its target composition. include(cmake/modules/RuntimeExtra.cmake) @@ -238,8 +305,10 @@ if(HIDE_PRIVATE_SYMBOLS AND NOT MSVC) target_compile_options(tvm_compiler PRIVATE -fvisibility=hidden) endif() target_compile_definitions(tvm_compiler PRIVATE TVM_EXPORTS TVM_FFI_EXPORTS) -target_link_libraries(tvm_compiler PUBLIC tvm_tirx::tvm_tirx tvm_runtime_extra) -target_link_libraries(tvm_compiler PRIVATE tvm_tirx::private_headers ${TVM_LINKER_LIBS}) +target_link_libraries(tvm_compiler PUBLIC tvm_tirx::shared tvm_ffi::shared tvm_runtime_extra) +target_link_libraries(tvm_compiler PRIVATE ${TVM_LINKER_LIBS}) +target_include_directories(tvm_compiler SYSTEM PRIVATE + "${CMAKE_CURRENT_SOURCE_DIR}/tirx/3rdparty/compiler-rt") foreach(_target tvm_compiler tvm_runtime_extra) target_include_directories(${_target} PUBLIC "$" @@ -274,11 +343,11 @@ if(TARGET cpptest AND TARGET tirx_cpptest) add_dependencies(cpptest tirx_cpptest) endif() add_custom_target(runtime DEPENDS tvm_runtime_extra) -add_dependencies(runtime tvm_tirx::tvm_runtime) +add_dependencies(runtime tvm_runtime::shared) set(TVM_RUNTIME_BACKEND_LIBS tvm_runtime_extra) foreach(_backend cuda opencl vulkan metal rocm) - if(TARGET tvm_tirx::tvm_runtime_${_backend}) - list(APPEND TVM_RUNTIME_BACKEND_LIBS tvm_tirx::tvm_runtime_${_backend}) + if(TARGET tvm_runtime_${_backend}) + list(APPEND TVM_RUNTIME_BACKEND_LIBS tvm_runtime_${_backend}) endif() endforeach() if(USE_CPP_RPC) @@ -290,6 +359,24 @@ endif() if(USE_IOS_RPC) add_subdirectory(apps/ios_rpc) endif() +# The parent owns installation of the core when it embeds the source build. +if(NOT TVM_USE_INSTALLED_TIRX) + set(_core_install_targets tvm_tirx tvm_runtime tvm_runtime_static) + foreach(_backend cuda opencl vulkan metal rocm) + if(TARGET tvm_runtime_${_backend}) + list(APPEND _core_install_targets tvm_runtime_${_backend}) + endif() + endforeach() + install(TARGETS ${_core_install_targets} + ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} + LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} + RUNTIME DESTINATION ${CMAKE_INSTALL_LIBDIR}) + install(DIRECTORY tirx/include/ DESTINATION include FILES_MATCHING PATTERN "*.h") + install(FILES tirx/cmake/tvm_tirx-config.cmake tirx/python/tvm_tirx/config.py + DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/tvm_tirx") + install(DIRECTORY tirx/cmake/utils/ + DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/tvm_tirx/utils" FILES_MATCHING PATTERN "Find*.cmake") +endif() install(TARGETS tvm_compiler tvm_runtime_extra ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} diff --git a/apps/android_rpc/app/src/main/jni/Android.mk b/apps/android_rpc/app/src/main/jni/Android.mk index f38d54224a71..d5a7ee042fac 100644 --- a/apps/android_rpc/app/src/main/jni/Android.mk +++ b/apps/android_rpc/app/src/main/jni/Android.mk @@ -37,8 +37,8 @@ LOCAL_SRC_FILES := org_apache_tvm_native_c_api.cc LOCAL_LDFLAGS := -L$(SYSROOT)/usr/lib/ -llog LOCAL_C_INCLUDES := $(ROOT_PATH)/include $(ROOT_PATH)/tirx/include \ - $(ROOT_PATH)/3rdparty/tvm-ffi/include \ - $(ROOT_PATH)/3rdparty/tvm-ffi/3rdparty/dlpack/include \ + $(ROOT_PATH)/tirx/3rdparty/tvm-ffi/include \ + $(ROOT_PATH)/tirx/3rdparty/tvm-ffi/3rdparty/dlpack/include \ $(ROOT_PATH)/3rdparty/OpenCL-Headers LOCAL_MODULE = tvm4j_runtime_packed diff --git a/apps/android_rpc/app/src/main/jni/tvm_runtime.h b/apps/android_rpc/app/src/main/jni/tvm_runtime.h index 60638071c60c..3db3f9756c7a 100644 --- a/apps/android_rpc/app/src/main/jni/tvm_runtime.h +++ b/apps/android_rpc/app/src/main/jni/tvm_runtime.h @@ -32,17 +32,19 @@ #define TVM_LOG_CUSTOMIZE 1 #define TVM_FFI_USE_LIBBACKTRACE 0 -#include "../3rdparty/tvm-ffi/src/ffi/backtrace.cc" -#include "../3rdparty/tvm-ffi/src/ffi/container.cc" -#include "../3rdparty/tvm-ffi/src/ffi/dtype.cc" -#include "../3rdparty/tvm-ffi/src/ffi/error.cc" -#include "../3rdparty/tvm-ffi/src/ffi/extra/library_module.cc" -#include "../3rdparty/tvm-ffi/src/ffi/extra/library_module_dynamic_lib.cc" -#include "../3rdparty/tvm-ffi/src/ffi/extra/library_module_system_lib.cc" -#include "../3rdparty/tvm-ffi/src/ffi/extra/module.cc" -#include "../3rdparty/tvm-ffi/src/ffi/function.cc" -#include "../3rdparty/tvm-ffi/src/ffi/object.cc" -#include "../3rdparty/tvm-ffi/src/ffi/tensor.cc" +// FFI registration precedes the runtime implementation. +// clang-format off +#include "../tirx/3rdparty/tvm-ffi/src/ffi/backtrace.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/container.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/dtype.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/error.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/extra/library_module.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/extra/library_module_dynamic_lib.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/extra/library_module_system_lib.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/extra/module.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/function.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/object.cc" +#include "../tirx/3rdparty/tvm-ffi/src/ffi/tensor.cc" #include "../src/runtime/registry.cc" #include "../src/runtime/rpc/rpc_channel.cc" #include "../src/runtime/rpc/rpc_endpoint.cc" @@ -63,6 +65,7 @@ #include "../tirx/src/runtime/threading_backend.cc" #include "../tirx/src/runtime/timer.cc" #include "../tirx/src/runtime/workspace_pool.cc" +// clang-format on #ifdef TVM_OPENCL_RUNTIME #include "../src/runtime/opencl/opencl_device_api.cc" diff --git a/apps/cpp_rpc/CMakeLists.txt b/apps/cpp_rpc/CMakeLists.txt index 5f98d8ffbbb6..dcc7515572e1 100644 --- a/apps/cpp_rpc/CMakeLists.txt +++ b/apps/cpp_rpc/CMakeLists.txt @@ -32,7 +32,7 @@ if (OS) endif() endif() -if(USE_OPENCL) +if(TARGET tvm_runtime_opencl) if (ANDROID_ABI) if(DEFINED ENV{ANDROID_NDK_MAJOR}) if($ENV{ANDROID_NDK_MAJOR} VERSION_LESS "23") @@ -47,7 +47,7 @@ target_include_directories( PUBLIC "../../include" ) -target_link_libraries(tvm_rpc PUBLIC tvm_ffi_header) +target_link_libraries(tvm_rpc PUBLIC tvm_ffi::shared) if(NOT MSVC AND NOT APPLE) list(APPEND TVM_RPC_LINKER_LIBS tvm_runtime "-Wl,--no-as-needed" ${TVM_RUNTIME_BACKEND_LIBS} "-Wl,--as-needed") diff --git a/apps/cpp_rpc/README.md b/apps/cpp_rpc/README.md index a4557a3fd367..513d5d34a7ba 100644 --- a/apps/cpp_rpc/README.md +++ b/apps/cpp_rpc/README.md @@ -37,9 +37,9 @@ This folder contains a simple recipe to make RPC server in c++. # Path to the desired C++ cross compiler set(CMAKE_CXX_COMPILER /path/to/cross/compiler/executable) ``` -- If you need to build cpp_rpc with OpenCL support, specify variable `USE_OPENCL` in the config: +- If you need to build cpp_rpc with OpenCL support, specify variable `TVM_TIRX_USE_OPENCL` in the config: ``` - set(USE_OPENCL ON) + set(TVM_TIRX_USE_OPENCL ON) ``` In this case [OpenCL-wrapper](../../src/runtime/opencl/opencl_wrapper) or OpenCL installed to your system will be used. When OpenCL-wrapper is used, it will dynamically load OpenCL library on the device. @@ -47,7 +47,7 @@ This folder contains a simple recipe to make RPC server in c++. If linking against a custom device OpenCL library is needed, in the config specify the path to the OpenCL SDK containing the include/CL headers and lib/ or lib64/libOpenCL.so: ``` - set(USE_OPENCL /path/to/opencl-sdk) + set(TVM_TIRX_USE_OPENCL /path/to/opencl-sdk) ``` - From within the configured tvm build directory, compile `tvm_runtime` and the `tvm_rpc` server: diff --git a/apps/ios_rpc/tvmrpc/TVMRuntime.mm b/apps/ios_rpc/tvmrpc/TVMRuntime.mm index f4866daa4d1e..593423da2464 100644 --- a/apps/ios_rpc/tvmrpc/TVMRuntime.mm +++ b/apps/ios_rpc/tvmrpc/TVMRuntime.mm @@ -33,7 +33,7 @@ #if defined(USE_CUSTOM_DSO_LOADER) && USE_CUSTOM_DSO_LOADER == 1 // internal TVM header to achieve Library class -#include <../../../3rdparty/tvm-ffi/src/ffi/extra/library_module.h> +#include <../../../tirx/3rdparty/tvm-ffi/src/ffi/extra/library_module.h> #include #endif diff --git a/ci/wheel/manylinux_build_libtvm_runtime_cuda.sh b/ci/wheel/manylinux_build_libtvm_runtime_cuda.sh index 8a5f33bf461f..9c18ef81f5db 100755 --- a/ci/wheel/manylinux_build_libtvm_runtime_cuda.sh +++ b/ci/wheel/manylinux_build_libtvm_runtime_cuda.sh @@ -51,11 +51,10 @@ rm -rf "${build_dir}" # modern CMake fills in a default so configure does not fail without it. cmake -S "${repo_root}/tirx" -B "${build_dir}" \ -DCMAKE_BUILD_TYPE=Release \ - -DBUILD_TESTING=OFF \ - -DTVM_BUILD_PYTHON_MODULE=ON \ - -DUSE_CUDA=/usr/local/cuda \ - -DUSE_LLVM=OFF \ - -DUSE_NVTX=OFF \ + -DTVM_TIRX_BUILD_TESTS=OFF \ + -DTVM_TIRX_BUILD_PYTHON_MODULE=ON \ + -DTVM_TIRX_USE_CUDA=/usr/local/cuda \ + -DTVM_TIRX_USE_LLVM=OFF \ -DCMAKE_CUDA_COMPILER=/usr/local/cuda/bin/nvcc cmake --build "${build_dir}" --target tvm_runtime tvm_runtime_cuda --parallel "${parallel}" diff --git a/ci/wheel/windows_build_libtvm_runtime_cuda.bat b/ci/wheel/windows_build_libtvm_runtime_cuda.bat index 45e62fef105d..79a826292ff4 100644 --- a/ci/wheel/windows_build_libtvm_runtime_cuda.bat +++ b/ci/wheel/windows_build_libtvm_runtime_cuda.bat @@ -90,7 +90,7 @@ set "cuda_root_fwd=%cuda_root:\=/%" set "nvcc_fwd=%nvcc_exe:\=/%" call "%vcvars%" || exit /b 1 -cmake -S "%repo_root_fwd%/tirx" -B "%build_dir_fwd%" -G Ninja -DCMAKE_BUILD_TYPE=Release -DBUILD_TESTING=OFF -DTVM_BUILD_PYTHON_MODULE=ON -DUSE_CUDA="%cuda_root_fwd%" -DUSE_LLVM=OFF -DUSE_NVTX=OFF -DCMAKE_CUDA_COMPILER="%nvcc_fwd%" -DCMAKE_CUDA_FLAGS="-allow-unsupported-compiler" || exit /b 1 +cmake -S "%repo_root_fwd%/tirx" -B "%build_dir_fwd%" -G Ninja -DCMAKE_BUILD_TYPE=Release -DTVM_TIRX_BUILD_TESTS=OFF -DTVM_TIRX_BUILD_PYTHON_MODULE=ON -DTVM_TIRX_USE_CUDA="%cuda_root_fwd%" -DTVM_TIRX_USE_LLVM=OFF -DCMAKE_CUDA_COMPILER="%nvcc_fwd%" -DCMAKE_CUDA_FLAGS="-allow-unsupported-compiler" || exit /b 1 cmake --build "%build_dir_fwd%" --target tvm_runtime tvm_runtime_cuda --config Release || exit /b 1 if not exist "%build_dir%\lib\tvm_runtime_cuda.dll" ( echo tvm_runtime_cuda.dll was not produced & exit /b 1 ) diff --git a/cmake/config.cmake b/cmake/config.cmake index 93713f4820b7..1a0f4c6161e9 100644 --- a/cmake/config.cmake +++ b/cmake/config.cmake @@ -46,7 +46,7 @@ # - ON: enable CUDA with CMake's auto search # - OFF: disable CUDA # - /path/to/cuda: use specific path to CUDA toolkit -set(USE_CUDA OFF) +set(TVM_TIRX_USE_CUDA OFF) # Whether to enable NCCL support: # - ON: enable NCCL with CMake's auto search @@ -54,7 +54,7 @@ set(USE_CUDA OFF) # - /path/to/nccl: use specific path to nccl set(USE_NCCL OFF) -# Whether to enable NVTX support (must have USE_CUDA enabled): +# Whether to enable NVTX support (must have TVM_TIRX_USE_CUDA enabled): # - ON: enable NCCL with CMake's auto search # - OFF: disable NCCL set(USE_NVTX OFF) @@ -65,7 +65,7 @@ set(USE_NVTX OFF) # - ON: enable ROCM with CMake's auto search # - OFF: disable ROCM # - /path/to/rocm: use specific path to rocm -set(USE_ROCM OFF) +set(TVM_TIRX_USE_ROCM OFF) # Whether to enable RCCL support: # - ON: enable RCCL with CMake's auto search @@ -80,15 +80,15 @@ set(USE_RCCL OFF) # time and trigger dynamic search and loading of OpenCL in runtime # - OFF: disable OpenCL # - /path/to/opencl-sdk: use specific path to opencl-sdk -set(USE_OPENCL OFF) +set(TVM_TIRX_USE_OPENCL OFF) # Whether to allow OPENCL cl_mem access to host # cl_mem will be allocated with CL_MEM_ALLOC_HOST_PTR # OpenCLWorkspace->GetHostPtr API returns the host accessible pointer -set(USE_OPENCL_ENABLE_HOST_PTR OFF) +set(TVM_TIRX_USE_OPENCL_ENABLE_HOST_PTR OFF) # Whether enable Metal runtime -set(USE_METAL OFF) +set(TVM_TIRX_USE_METAL OFF) # Whether enable Vulkan runtime # @@ -96,7 +96,7 @@ set(USE_METAL OFF) # - ON: enable Vulkan with CMake's auto search # - OFF: disable vulkan # - /path/to/vulkan-sdk: use specific path to vulkan-sdk -set(USE_VULKAN OFF) +set(TVM_TIRX_USE_VULKAN OFF) # Whether to use spirv-tools and SPIRV-Headers from Khronos GitHub or GitLab. # @@ -104,10 +104,10 @@ set(USE_VULKAN OFF) # - OFF: not to use # - /path/to/install: path to your khronis spirv-tools and SPIRV-Headers installation directory # -set(USE_KHRONOS_SPIRV OFF) +set(TVM_TIRX_USE_KHRONOS_SPIRV OFF) # whether enable SPIRV_KHR_DOT_PRODUCT -set(USE_SPIRV_KHR_INTEGER_DOT_PRODUCT OFF) +set(TVM_TIRX_USE_SPIRV_KHR_INTEGER_DOT_PRODUCT OFF) # Whether enable RPC runtime set(USE_RPC ON) @@ -129,7 +129,7 @@ set(USE_IOS_RPC OFF) # - OFF: disable llvm, note this will disable CPU codegen # which is needed for most cases # - /path/to/llvm-config: enable specific LLVM when multiple llvm-dev is available. -set(USE_LLVM OFF) +set(TVM_TIRX_USE_LLVM OFF) #--------------------------------------------- # Contrib libraries @@ -174,7 +174,7 @@ set(USE_AMX OFF) # Whether use OpenMP thread pool, choices: gnu, intel # Note: "gnu" uses gomp library, "intel" uses iomp5 library -set(USE_OPENMP none) +set(TVM_TIRX_USE_OPENMP OFF) # Whether use contrib.random in runtime set(USE_RANDOM ON) @@ -275,7 +275,7 @@ set(TVM_FFI_BACKTRACE_ON_SEGFAULT ON) set(USE_GTEST AUTO) # Enable using CUTLASS as a BYOC backend -# Need to have USE_CUDA=ON +# Need to have TVM_TIRX_USE_CUDA=ON set(USE_CUTLASS OFF) # Whether to enable FlashInfer or not. @@ -304,4 +304,4 @@ SET(CMAKE_VS_PLATFORM_NAME_DEFAULT "x64") SET(CMAKE_VS_PLATFORM_TOOLSET_HOST_ARCHITECTURE "x64") # Enable Qualcomm OpenCL extension support -set(USE_OPENCL_EXTN_QCOM OFF) +set(TVM_TIRX_USE_OPENCL_EXTN_QCOM OFF) diff --git a/cmake/modules/RuntimeExtra.cmake b/cmake/modules/RuntimeExtra.cmake index 473c0872d3ce..b2d726cdb59f 100644 --- a/cmake/modules/RuntimeExtra.cmake +++ b/cmake/modules/RuntimeExtra.cmake @@ -29,7 +29,7 @@ target_link_libraries(tvm_disco_distributed_objs PRIVATE tvm_runtime_extra_defs) target_link_libraries(tvm_runtime_extra PRIVATE tvm_disco_distributed_objs) # NCCL / cuda_ipc — requires CUDA + NCCL. -if(USE_CUDA AND USE_NCCL) +if(TVM_TIRX_USE_CUDA AND USE_NCCL) find_nccl(${USE_NCCL}) include_directories(SYSTEM ${NCCL_INCLUDE_DIR}) tvm_file_glob(GLOB _nccl_srcs src/runtime/extra/disco/nccl/*.cc src/runtime/extra/disco/cuda_ipc/*.cc 3rdparty/tensorrt_llm/*.cu) @@ -41,7 +41,7 @@ if(USE_CUDA AND USE_NCCL) endif() # NVSHMEM. -if(USE_CUDA AND USE_NVSHMEM) +if(TVM_TIRX_USE_CUDA AND USE_NVSHMEM) find_nvshmem(${USE_NVSHMEM}) if(NOT NVSHMEM_FOUND) message(FATAL_ERROR "Cannot find NVSHMEM, USE_NVSHMEM=" ${USE_NVSHMEM}) @@ -66,7 +66,7 @@ if(USE_CUDA AND USE_NVSHMEM) endif() # RCCL. -if(USE_ROCM AND USE_RCCL) +if(TVM_TIRX_USE_ROCM AND USE_RCCL) find_rccl(${USE_RCCL}) include_directories(SYSTEM ${RCCL_INCLUDE_DIR}) tvm_file_glob(GLOB _rccl_srcs src/runtime/extra/disco/nccl/*.cc) @@ -79,17 +79,17 @@ endif() target_link_libraries(tvm_runtime_extra PUBLIC tvm_runtime) # If disco/cuda_ipc is included, link the CUDA DSO. -if(USE_CUDA) +if(TVM_TIRX_USE_CUDA) target_link_libraries(tvm_runtime_extra PUBLIC tvm_runtime_cuda) target_link_libraries(tvm_runtime_extra PRIVATE ${CUDA_CUDART_LIBRARY} ${CUDA_CUDA_LIBRARY}) endif() -if(USE_ROCM) +if(TVM_TIRX_USE_ROCM) target_link_libraries(tvm_runtime_extra PRIVATE ${ROCM_HIPHCC_LIBRARY}) endif() # CUTLASS fpA_intB_gemm and flash_attn are separate shared libs. -if(USE_CUDA AND USE_CUTLASS) +if(TVM_TIRX_USE_CUDA AND USE_CUTLASS) target_link_libraries(tvm_runtime_extra PRIVATE fpA_intB_gemm fpA_intB_gemm_tvm) target_link_libraries(tvm_runtime_extra PRIVATE -Wl,--no-as-needed flash_attn) endif() diff --git a/cmake/modules/contrib/CLML.cmake b/cmake/modules/contrib/CLML.cmake index edcb9b717b2f..e085fdcbda46 100644 --- a/cmake/modules/contrib/CLML.cmake +++ b/cmake/modules/contrib/CLML.cmake @@ -80,5 +80,5 @@ if(USE_CLML_GRAPH_EXECUTOR) add_definitions(-DTVM_GRAPH_EXECUTOR_CLML) message(STATUS "Enable OpenCL as fallback to CLML") - set(USE_OPENCL ${CLML_PATH}) + set(TVM_TIRX_USE_OPENCL ${CLML_PATH}) endif() diff --git a/cmake/modules/contrib/CUDA.cmake b/cmake/modules/contrib/CUDA.cmake index c4f6d6879155..91994d337c69 100644 --- a/cmake/modules/contrib/CUDA.cmake +++ b/cmake/modules/contrib/CUDA.cmake @@ -15,10 +15,10 @@ # specific language governing permissions and limitations # under the License. -# Contrib sources gated by USE_CUDA go into libtvm_runtime_extra. +# Contrib sources gated by TVM_TIRX_USE_CUDA go into libtvm_runtime_extra. # See the RuntimeExtra assembly block in CMakeLists.txt. -if(USE_CUDA AND USE_CUDNN) +if(TVM_TIRX_USE_CUDA AND USE_CUDNN) message(STATUS "Build with cuDNN support") include_directories(SYSTEM ${CUDA_CUDNN_INCLUDE_DIRS}) tvm_file_glob(GLOB CUDNN_RELAX_CONTRIB_SRC src/relax/backend/contrib/cudnn/*.cc) @@ -27,9 +27,9 @@ if(USE_CUDA AND USE_CUDNN) add_library(tvm_cudnn_objs OBJECT ${CONTRIB_CUDNN_SRCS}) target_link_libraries(tvm_cudnn_objs PRIVATE tvm_runtime_extra_defs) target_link_libraries(tvm_runtime_extra PRIVATE tvm_cudnn_objs ${CUDA_CUDNN_LIBRARY}) -endif(USE_CUDA AND USE_CUDNN) +endif(TVM_TIRX_USE_CUDA AND USE_CUDNN) -if(USE_CUDA AND USE_CUDNN_FRONTEND) +if(TVM_TIRX_USE_CUDA AND USE_CUDNN_FRONTEND) message(STATUS "Build with cuDNN Frontend support") if(IS_DIRECTORY ${USE_CUDNN_FRONTEND}) find_file(CUDNN_FRONTEND_HEADER cudnn_frontend.h HINTS ${USE_CUDNN_FRONTEND}/include) @@ -45,9 +45,9 @@ if(USE_CUDA AND USE_CUDNN_FRONTEND) add_library(tvm_cudnn_frontend_objs OBJECT ${CONTRIB_CUDNN_FRONTEND_SRCS}) target_link_libraries(tvm_cudnn_frontend_objs PRIVATE tvm_runtime_extra_defs) target_link_libraries(tvm_runtime_extra PRIVATE tvm_cudnn_frontend_objs) -endif(USE_CUDA AND USE_CUDNN_FRONTEND) +endif(TVM_TIRX_USE_CUDA AND USE_CUDNN_FRONTEND) -if(USE_CUDA AND USE_CUBLAS) +if(TVM_TIRX_USE_CUDA AND USE_CUBLAS) message(STATUS "Build with cuBLAS support") tvm_file_glob(GLOB CUBLAS_CONTRIB_SRC src/relax/backend/contrib/cublas/*.cc) list(APPEND COMPILER_SRCS ${CUBLAS_CONTRIB_SRC}) @@ -58,9 +58,9 @@ if(USE_CUDA AND USE_CUBLAS) if(NOT CUDA_CUBLASLT_LIBRARY STREQUAL "CUDA_CUBLASLT_LIBRARY-NOTFOUND") target_link_libraries(tvm_runtime_extra PRIVATE ${CUDA_CUBLASLT_LIBRARY}) endif() -endif(USE_CUDA AND USE_CUBLAS) +endif(TVM_TIRX_USE_CUDA AND USE_CUBLAS) -if(USE_CUDA AND USE_THRUST) +if(TVM_TIRX_USE_CUDA AND USE_THRUST) message(STATUS "Build with Thrust support") tvm_file_glob(GLOB CONTRIB_THRUST_SRC src/runtime/extra/contrib/thrust/*.cu) add_library(tvm_thrust_objs OBJECT ${CONTRIB_THRUST_SRC}) @@ -71,9 +71,9 @@ if(USE_CUDA AND USE_THRUST) target_link_libraries(tvm_thrust_objs PRIVATE CCCL::Thrust) endif() target_link_libraries(tvm_runtime_extra PRIVATE tvm_thrust_objs) -endif(USE_CUDA AND USE_THRUST) +endif(TVM_TIRX_USE_CUDA AND USE_THRUST) -if(USE_CUDA AND USE_CURAND) +if(TVM_TIRX_USE_CUDA AND USE_CURAND) message(STATUS "Build with cuRAND support") message(STATUS "${CUDA_CURAND_LIBRARY}") tvm_file_glob(GLOB CONTRIB_CURAND_SRC_CC src/runtime/extra/contrib/curand/*.cc) @@ -81,4 +81,4 @@ if(USE_CUDA AND USE_CURAND) add_library(tvm_curand_objs OBJECT ${CONTRIB_CURAND_SRC_CC} ${CONTRIB_CURAND_SRC_CU}) target_link_libraries(tvm_curand_objs PRIVATE tvm_runtime_extra_defs) target_link_libraries(tvm_runtime_extra PRIVATE tvm_curand_objs ${CUDA_CURAND_LIBRARY}) -endif(USE_CUDA AND USE_CURAND) +endif(TVM_TIRX_USE_CUDA AND USE_CURAND) diff --git a/cmake/modules/contrib/CUTLASS.cmake b/cmake/modules/contrib/CUTLASS.cmake index ac0958eb96b1..1c287f36cb87 100644 --- a/cmake/modules/contrib/CUTLASS.cmake +++ b/cmake/modules/contrib/CUTLASS.cmake @@ -15,7 +15,7 @@ # specific language governing permissions and limitations # under the License. -if(USE_CUDA AND USE_CUTLASS) +if(TVM_TIRX_USE_CUDA AND USE_CUTLASS) tvm_file_glob(GLOB CUTLASS_CONTRIB_SRC src/relax/backend/contrib/cutlass/*.cc ) diff --git a/cmake/modules/contrib/ROCM.cmake b/cmake/modules/contrib/ROCM.cmake index 37ee066099fc..24c1b4e69808 100644 --- a/cmake/modules/contrib/ROCM.cmake +++ b/cmake/modules/contrib/ROCM.cmake @@ -16,7 +16,7 @@ # under the License. # HIPBLAS contrib goes into libtvm_runtime_extra. -if(USE_ROCM AND USE_HIPBLAS) +if(TVM_TIRX_USE_ROCM AND USE_HIPBLAS) message(STATUS "Build with HIPBLAS support") tvm_file_glob(GLOB HIPBLAS_CONTRIB_SRC src/relax/backend/contrib/hipblas/*.cc) list(APPEND COMPILER_SRCS ${HIPBLAS_CONTRIB_SRC}) @@ -27,9 +27,9 @@ if(USE_ROCM AND USE_HIPBLAS) if(NOT ROCM_HIPBLASLT_LIBRARY STREQUAL "ROCM_HIPBLASLT_LIBRARY-NOTFOUND") target_link_libraries(tvm_runtime_extra PRIVATE ${ROCM_HIPBLASLT_LIBRARY}) endif() -endif(USE_ROCM AND USE_HIPBLAS) +endif(TVM_TIRX_USE_ROCM AND USE_HIPBLAS) -if(USE_ROCM AND USE_THRUST) +if(TVM_TIRX_USE_ROCM AND USE_THRUST) message(STATUS "Build with rocThrust support") # We need to override CXX to hipcc. This is required by rocthrust if(${CMAKE_CXX_COMPILER} MATCHES "hipcc$") @@ -44,4 +44,4 @@ if(USE_ROCM AND USE_THRUST) add_library(tvm_rocthrust_objs OBJECT src/runtime/extra/contrib/thrust/thrust.cu) target_link_libraries(tvm_rocthrust_objs PRIVATE tvm_runtime_extra_defs) target_link_libraries(tvm_runtime_extra PRIVATE tvm_rocthrust_objs roc::rocthrust) -endif(USE_ROCM AND USE_THRUST) +endif(TVM_TIRX_USE_ROCM AND USE_THRUST) diff --git a/tirx/cmake/utils/CCache.cmake b/cmake/utils/CCache.cmake similarity index 100% rename from tirx/cmake/utils/CCache.cmake rename to cmake/utils/CCache.cmake diff --git a/cmake/utils/Library.cmake b/cmake/utils/Library.cmake new file mode 100644 index 000000000000..4b4d9f660590 --- /dev/null +++ b/cmake/utils/Library.cmake @@ -0,0 +1,86 @@ +# 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. + +# Helpers for configuring library targets. + +# Hide symbols from static archives linked into a shared library. This prevents +# downstream libraries from accidentally binding to private archive symbols such +# as those from libstdc++_nonshared.a. +function(tvm_hide_static_linked_lib_symbols target_name) + if(HIDE_PRIVATE_SYMBOLS + AND CMAKE_SYSTEM_NAME MATCHES "Linux|Android|FreeBSD|NetBSD|OpenBSD" + AND CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang" + ) + target_link_options(${target_name} PRIVATE "-Wl,--exclude-libs,ALL") + endif() +endfunction() + +####################################################### +# tvm_configure_target_library(target_name [RUNTIME_MODULE]) +# +# Configure a TVM library target. The target always gets a relative rpath +# ($ORIGIN / @loader_path) so that sibling shared libraries in the same +# directory resolve each other regardless of the install location (e.g. inside a +# Python wheel). +# +# For shared libraries and runtime modules, hide symbols that come from linked +# static archives so downstream libraries cannot accidentally bind to private +# archive symbols such as those from libstdc++_nonshared.a. +# +# With the RUNTIME_MODULE option -- used for the optional runtime backend +# libraries (tvm_runtime_cuda, tvm_runtime_vulkan, ...) -- the target is also +# emitted into the build "lib" directory and installed; when building the Python +# module it is additionally installed into the package "lib" directory. Targets +# that manage their own output directory / install rules (tvm_compiler, +# tvm_runtime, ...) omit the option and take only the rpath. +# +# No-op if the target does not exist. +function(tvm_configure_target_library target_name) + if(NOT TARGET ${target_name}) + return() + endif() + cmake_parse_arguments(ARG "RUNTIME_MODULE" "" "" ${ARGN}) + + if(APPLE) + set_target_properties(${target_name} PROPERTIES + BUILD_RPATH "@loader_path" + INSTALL_RPATH "@loader_path" + ) + elseif(UNIX) + set_target_properties(${target_name} PROPERTIES + BUILD_RPATH "\$ORIGIN" + INSTALL_RPATH "\$ORIGIN" + ) + endif() + + get_target_property(TVM_TARGET_TYPE ${target_name} TYPE) + if(TVM_TARGET_TYPE STREQUAL "SHARED_LIBRARY" OR TVM_TARGET_TYPE STREQUAL "MODULE_LIBRARY") + tvm_hide_static_linked_lib_symbols(${target_name}) + endif() + + if(ARG_RUNTIME_MODULE) + set_target_properties(${target_name} PROPERTIES + LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" + RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" + ARCHIVE_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" + ) + install(TARGETS ${target_name} DESTINATION lib${LIB_SUFFIX}) + if(TVM_BUILD_PYTHON_MODULE) + install(TARGETS ${target_name} DESTINATION "lib") + endif() + endif() +endfunction() diff --git a/tirx/cmake/utils/Linker.cmake b/cmake/utils/Linker.cmake similarity index 100% rename from tirx/cmake/utils/Linker.cmake rename to cmake/utils/Linker.cmake diff --git a/tirx/cmake/utils/Summary.cmake b/cmake/utils/Summary.cmake similarity index 100% rename from tirx/cmake/utils/Summary.cmake rename to cmake/utils/Summary.cmake diff --git a/tirx/cmake/utils/Utils.cmake b/cmake/utils/Utils.cmake similarity index 100% rename from tirx/cmake/utils/Utils.cmake rename to cmake/utils/Utils.cmake diff --git a/docs/README.md b/docs/README.md index 44621b4e07ff..777f122da517 100644 --- a/docs/README.md +++ b/docs/README.md @@ -47,7 +47,7 @@ This folder contains the source of TVM's documentation, hosted at https://tvm.ap ```bash export TVM_HOME=/path-to-tvm export TVM_LIBRARY_PATH=$TVM_HOME/build - pip install --target=$TVM_HOME/python $TVM_HOME/3rdparty/tvm-ffi + pip install --target=$TVM_HOME/python $TVM_HOME/tirx/3rdparty/tvm-ffi export PYTHONPATH=$TVM_HOME/python:$PYTHONPATH ``` diff --git a/docs/arch/runtime.rst b/docs/arch/runtime.rst index 575451d8fe91..c736464dbdff 100644 --- a/docs/arch/runtime.rst +++ b/docs/arch/runtime.rst @@ -48,7 +48,7 @@ function call whose caller and callee may be in different languages. The following code block provides an example in C++ -.. _PackedFunc: https://github.com/apache/tvm/blob/main/3rdparty/tvm-ffi/include/tvm/ffi/function.h +.. _PackedFunc: https://github.com/apache/tvm/blob/main/tirx/3rdparty/tvm-ffi/include/tvm/ffi/function.h .. code:: c @@ -138,7 +138,7 @@ which allows us to embed the PackedFunc into any languages. Besides python, so f `java`_ and `javascript`_. This philosophy of embedded API is very like Lua, except that we don't have a new language but use C++. -.. _minimum C API: https://github.com/apache/tvm/blob/main/3rdparty/tvm-ffi/include/tvm/ffi/c_api.h +.. _minimum C API: https://github.com/apache/tvm/blob/main/tirx/3rdparty/tvm-ffi/include/tvm/ffi/c_api.h .. _java: https://github.com/apache/tvm/tree/main/jvm .. _javascript: https://github.com/apache/tvm/tree/main/web @@ -169,7 +169,7 @@ TVM defines the compiled object as `Module`_. The user can get the compiled function from Module as PackedFunc. The generated compiled code can dynamically get function from Module in runtime. It caches the function handle in the first call and reuses in subsequent calls. We use this to link device code and callback into any PackedFunc(e.g., python) from generated code. -.. _Module: https://github.com/apache/tvm/blob/main/3rdparty/tvm-ffi/include/tvm/ffi/extra/module.h +.. _Module: https://github.com/apache/tvm/blob/main/tirx/3rdparty/tvm-ffi/include/tvm/ffi/extra/module.h The ModuleNode is an abstract class that can be implemented by each type of device. So far we support modules for CUDA, Metal, OpenCL and loading dynamic shared libraries. This abstraction makes introduction @@ -257,18 +257,18 @@ Each argument in PackedFunc contains a union value `TVMValue`_ and a type code. This design allows the dynamically typed language to convert to the corresponding type directly, and statically typed language to do runtime type checking during conversion. -.. _TVMValue: https://github.com/apache/tvm/blob/main/3rdparty/tvm-ffi/include/tvm/ffi/c_api.h +.. _TVMValue: https://github.com/apache/tvm/blob/main/tirx/3rdparty/tvm-ffi/include/tvm/ffi/c_api.h The relevant files are - `function.h`_ for C++ PackedFunc API - `c_api.h`_ for C API. -.. _function.h: https://github.com/apache/tvm/blob/main/3rdparty/tvm-ffi/include/tvm/ffi/function.h -.. _c_api.h: https://github.com/apache/tvm/blob/main/3rdparty/tvm-ffi/include/tvm/ffi/c_api.h +.. _function.h: https://github.com/apache/tvm/blob/main/tirx/3rdparty/tvm-ffi/include/tvm/ffi/function.h +.. _c_api.h: https://github.com/apache/tvm/blob/main/tirx/3rdparty/tvm-ffi/include/tvm/ffi/c_api.h To support extension types, we used a registry system to register type related information, like support of any -in C++. See the ``tvm-ffi`` subproject under ``3rdparty/tvm-ffi/`` for more details on the FFI type system. +in C++. See the ``tvm-ffi`` subproject under ``tirx/3rdparty/tvm-ffi/`` for more details on the FFI type system. Runtime-Specific Information diff --git a/docs/conf.py b/docs/conf.py index 361356a95e5f..1d4ec578f79a 100644 --- a/docs/conf.py +++ b/docs/conf.py @@ -365,7 +365,7 @@ def jupyter_notebook(script_blocks, gallery_conf, target_dir, real_func): # The Apache trademark/copyright footer is rendered through sphinx-book-theme's # ``extra_footer`` hook (see footer_html below). This mirrors how the tvm-ffi docs -# (3rdparty/tvm-ffi/docs/conf.py) preserve the ASF menu under the book theme. +# (tirx/3rdparty/tvm-ffi/docs/conf.py) preserve the ASF menu under the book theme. footer_dropdown = { "name": "ASF", "items": [ diff --git a/docs/how_to/tutorials/bring_your_own_codegen.py b/docs/how_to/tutorials/bring_your_own_codegen.py index de1cebdbba6b..50a57b5d5e67 100644 --- a/docs/how_to/tutorials/bring_your_own_codegen.py +++ b/docs/how_to/tutorials/bring_your_own_codegen.py @@ -43,7 +43,7 @@ **Prerequisites**: the example NPU sections need TVM built with ``USE_EXAMPLE_NPU_CODEGEN=ON`` and ``USE_EXAMPLE_NPU_RUNTIME=ON``; the TensorRT sections need ``USE_TENSORRT_CODEGEN=ON``, ``USE_TENSORRT_RUNTIME=ON`` and -``USE_CUDA=ON`` plus a CUDA GPU and a matching TensorRT install (from NVIDIA's +``TVM_TIRX_USE_CUDA=ON`` plus a CUDA GPU and a matching TensorRT install (from NVIDIA's ``pip install tensorrt`` packages or the TensorRT archive); the final deployment section also needs PyTorch. Each section degrades gracefully when its backend is unavailable. diff --git a/docs/how_to/tutorials/cross_compilation_and_rpc.py b/docs/how_to/tutorials/cross_compilation_and_rpc.py index f218582b8246..4096749e26cb 100644 --- a/docs/how_to/tutorials/cross_compilation_and_rpc.py +++ b/docs/how_to/tutorials/cross_compilation_and_rpc.py @@ -278,7 +278,7 @@ # set(CMAKE_FIND_ROOT_PATH_MODE_PACKAGE ONLY) # # Build the runtime from the TVM repository root. Enable target-specific -# options such as ``USE_OPENCL`` or vendor runtime support in ``config.cmake``. +# options such as ``TVM_TIRX_USE_OPENCL`` or vendor runtime support in ``config.cmake``. # Build any target-specific runtime libraries that your deployment needs, such # as ``tvm_runtime_opencl`` for OpenCL. # @@ -288,12 +288,12 @@ # cd cross_build # cp ../cmake/config.cmake ./ # -# # Enable other options as needed, e.g. USE_OPENCL or vendor runtimes. -# sed -i "s|USE_LLVM.*)|USE_LLVM OFF)|" config.cmake +# # Enable other options as needed, e.g. TVM_TIRX_USE_OPENCL or vendor runtimes. +# sed -i "s|TVM_TIRX_USE_LLVM.*)|TVM_TIRX_USE_LLVM OFF)|" config.cmake # # cmake -DCMAKE_TOOLCHAIN_FILE=/YYY/aarch64-linux-gnu.cmake -DCMAKE_BUILD_TYPE=Release .. # cmake --build . --target runtime -j -# # Optional example when USE_OPENCL is enabled: +# # Optional example when TVM_TIRX_USE_OPENCL is enabled: # # cmake --build . --target tvm_runtime_opencl -j # cd .. # @@ -406,7 +406,7 @@ # # mkdir -p build && cd build # cp ../cmake/config.cmake . -# sed -i "s/USE_OPENCL OFF/USE_OPENCL ON/" config.cmake +# sed -i "s/TVM_TIRX_USE_OPENCL OFF/TVM_TIRX_USE_OPENCL ON/" config.cmake # cmake .. && cmake --build . --parallel $(nproc) # # The following function shows how we run an OpenCL kernel remotely diff --git a/docs/install/from_source.rst b/docs/install/from_source.rst index 1234c7c00676..6d63f9a06d54 100644 --- a/docs/install/from_source.rst +++ b/docs/install/from_source.rst @@ -116,14 +116,14 @@ We want to specifically tweak the following flags by appending them to the end o echo "set(CMAKE_BUILD_TYPE RelWithDebInfo)" >> config.cmake # LLVM is a must dependency for compiler end - echo "set(USE_LLVM \"llvm-config --ignore-libllvm --link-static\")" >> config.cmake + echo "set(TVM_TIRX_USE_LLVM \"llvm-config --ignore-libllvm --link-static\")" >> config.cmake echo "set(HIDE_PRIVATE_SYMBOLS ON)" >> config.cmake # GPU SDKs, turn on if needed - echo "set(USE_CUDA OFF)" >> config.cmake - echo "set(USE_METAL OFF)" >> config.cmake - echo "set(USE_VULKAN OFF)" >> config.cmake - echo "set(USE_OPENCL OFF)" >> config.cmake + echo "set(TVM_TIRX_USE_CUDA OFF)" >> config.cmake + echo "set(TVM_TIRX_USE_METAL OFF)" >> config.cmake + echo "set(TVM_TIRX_USE_VULKAN OFF)" >> config.cmake + echo "set(TVM_TIRX_USE_OPENCL OFF)" >> config.cmake # cuBLAS, cuDNN, cutlass support, turn on if needed echo "set(USE_CUBLAS OFF)" >> config.cmake @@ -158,7 +158,7 @@ Therefore, after we finish the build, we need to install the tvm-ffi package. .. code-block:: bash - cd 3rdparty/tvm-ffi; pip install .; cd .. + cd tirx/3rdparty/tvm-ffi; pip install .; cd .. Leaving the build environment ``tvm-build-venv``, there are two ways to install the successful build into your environment: @@ -168,7 +168,7 @@ Leaving the build environment ``tvm-build-venv``, there are two ways to install .. code-block:: bash export TVM_HOME=/path-to-tvm - pip install --target=$TVM_HOME/python $TVM_HOME/3rdparty/tvm-ffi + pip install --target=$TVM_HOME/python $TVM_HOME/tirx/3rdparty/tvm-ffi export PYTHONPATH=$TVM_HOME/python:$PYTHONPATH - Install via pip local project @@ -332,7 +332,7 @@ Building ROCm support Currently, ROCm is supported only on linux, so all the instructions are written with linux in mind. -- Set ``set(USE_ROCM ON)``, set ROCM_PATH to the correct path. +- Set ``set(TVM_TIRX_USE_ROCM ON)``, set TVM_TIRX_ROCM_PATH to the correct path. - You need to first install HIP runtime from ROCm. Make sure the installation system has ROCm installed in it. - Install LLVM (>= 15), and LLD, make sure ``ld.lld`` is available via command line. diff --git a/docs/install/pypi.rst b/docs/install/pypi.rst index 48c25bf5d2c0..d351d16b457e 100644 --- a/docs/install/pypi.rst +++ b/docs/install/pypi.rst @@ -43,7 +43,7 @@ Install the CUDA extra in the same environment as TVM when you need this path: This extra installs Python-side CUDA bindings only. It does not make the PyPI wheel a CUDA-enabled TVM build, and it does not install NVIDIA drivers or a CUDA toolkit. If you need CUDA support in TVM itself, build TVM from source with -``USE_CUDA=ON``. +``TVM_TIRX_USE_CUDA=ON``. For more details on installing the TIRx compiler and optional kernel library, visit the :doc:`TIRx installation ` page. If you need to diff --git a/docs/tirx/install.rst b/docs/tirx/install.rst index 335a10efb076..95e917b52323 100644 --- a/docs/tirx/install.rst +++ b/docs/tirx/install.rst @@ -47,7 +47,7 @@ The wheel is enough to inspect and author TIRx IR. Some CUDA workflows also need NVIDIA's Python CUDA bindings, available through ``apache-tvm[cuda]``. That extra does not enable CUDA in the TVM library and does not install a driver or toolkit. To compile and run the CUDA examples, use a TVM build configured -with ``USE_CUDA=ON``; see :ref:`install TVM from source `. +with ``TVM_TIRX_USE_CUDA=ON``; see :ref:`install TVM from source `. Verify: @@ -88,3 +88,85 @@ list, and consult the `tirx-kernels README `_ and ``reference-dependencies.json`` in that repository for current optional requirements. + +Standalone CMake build +---------------------- + +The core can be built independently from the repository checkout. Initialize +its bundled FFI dependency, then use native CMake options (CMake 3.26 or newer):: + + git submodule update --init --recursive tirx/3rdparty/tvm-ffi + cmake -S tirx -B tirx/build -G Ninja -DCMAKE_BUILD_TYPE=Release \ + -DTVM_TIRX_USE_LLVM=ON -DTVM_TIRX_BUILD_TESTS=ON + cmake --build tirx/build + ctest --test-dir tirx/build + cmake --install tirx/build --prefix /path/to/install + +Core feature options use the ``TVM_TIRX_USE_`` prefix. SDK selectors such as +``TVM_TIRX_USE_CUDA`` and ``TVM_TIRX_USE_LLVM`` also accept an SDK path or +``llvm-config`` command. Compiler caches, linkers, and toolchain flags use +standard CMake settings. The standalone core does not read ``config.cmake``. +C++ tests are off by default; enabling them supplies GTest in the test subtree. +Installation includes public headers, libraries, and CMake package targets. + +For a Python core package build, pass these options through the ``CMAKE_ARGS`` +environment variable. Quote values containing spaces inside the variable:: + + CMAKE_ARGS='-DTVM_TIRX_USE_LLVM="llvm-config --link-static"' \ + python -m pip install -v ./tirx + +The repository's native build script also accepts ``CMAKE_ARGS`` for its TVM +configuration, including parent-only options:: + + mkdir -p build + CMAKE_ARGS='-DTVM_TIRX_USE_LLVM=ON -DUSE_RPC=ON' \ + python tests/scripts/task_build.py --build-dir build + +The native build script applies these settings only to TVM configuration; +its separate source-FFI Python package build keeps its own defaults. Direct +``cmake`` invocations use command-line options as shown above, since CMake +itself does not read ``CMAKE_ARGS``. + +A parent project can use ``add_subdirectory(path/to/tirx)`` and link to +``tvm_tirx::shared`` or ``tvm_runtime::shared``. This builds the bundled FFI +source targets and leaves tests, packaging, and installation to the parent. +TIRx registers installation rules only when it is the top-level project. +An embedding parent owns the installation rules for the libraries, public +headers, and package files it chooses to include. + +Installed consumers +------------------- + +An ordinary downstream compiler links TIRx and the separate FFI package +explicitly. Install both Python packages into the environment used by CMake:: + + python -m pip install apache-tvm-tirx apache-tvm-ffi + +Query each package's CMake directory:: + + python -m tvm_tirx.config --cmakedir + python -m tvm_ffi.config --cmakedir + +Pass those directories as ``tvm_tirx_DIR`` and ``tvm_ffi_DIR`` when configuring +the consumer. Its CMakeLists.txt uses both packages directly: + +.. code-block:: cmake + + find_package(tvm_tirx CONFIG REQUIRED) + find_package(tvm_ffi CONFIG REQUIRED) + add_executable(app main.cc) + target_link_libraries(app PRIVATE tvm_tirx::shared tvm_ffi::shared) + +The public TIRx targets are ``tvm_tirx::header``, ``tvm_tirx::shared``, and +``tvm_runtime::shared``. The package config describes the normal Python build; +it does not reconstruct build options, expose static or backend targets, or +find/link FFI implicitly. Enabled backend libraries are still installed and +loaded by the runtime. Consumers of an internally built static runtime supply +its dependencies, including OpenMP when enabled, themselves. + +TIRx installs only its own public headers, libraries, and package helpers; +it does not copy FFI libraries, headers, or CMake files. For a native prefix +installation, use its ``lib/cmake/tvm_tirx`` directory (or the configured +library directory). The handwritten config queries that exact prefix, so a +separately installed TIRx wheel cannot override the selected installation. +``Python_EXECUTABLE`` selects the Python interpreter used for discovery. diff --git a/jvm/native/linux-x86_64/pom.xml b/jvm/native/linux-x86_64/pom.xml index 31e120bc58cd..ce6e98e7f83a 100644 --- a/jvm/native/linux-x86_64/pom.xml +++ b/jvm/native/linux-x86_64/pom.xml @@ -118,7 +118,7 @@ under the License. -I../../../include - -I../../../3rdparty/tvm-ffi/include + -I../../../tirx/3rdparty/tvm-ffi/include -I${JAVA_HOME}/include -I${JAVA_HOME}/include/linux ${cflags} diff --git a/jvm/native/osx-x86_64/pom.xml b/jvm/native/osx-x86_64/pom.xml index 4f9d70d60dc1..5d6c3dcf7503 100644 --- a/jvm/native/osx-x86_64/pom.xml +++ b/jvm/native/osx-x86_64/pom.xml @@ -119,7 +119,7 @@ under the License. -I../../../include - -I../../../3rdparty/tvm-ffi/include + -I../../../tirx/3rdparty/tvm-ffi/include -I${JAVA_HOME}/include -I${JAVA_HOME}/include/darwin ${cflags} diff --git a/pyproject.toml b/pyproject.toml index cd3a9547bf29..cff0e474b3fb 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -111,6 +111,15 @@ sdist.include = [ "/README.md", "/apps/**", "/web/**", + # Private build inputs remain in the source distribution, outside the core's + # installed public-header interface. + "/tirx/3rdparty/compiler-rt/builtin_fp16.h", + "/tirx/src/backend/opencl/memory_scope.h", + "/tirx/src/backend/opencl/runtime/opencl_common.h", + "/tirx/src/backend/opencl/runtime/texture.h", + "/tirx/src/runtime/metadata.h", + "/tirx/src/runtime/pack_args.h", + "/tirx/src/runtime/workspace_pool.h", "/3rdparty/OpenCL-Headers/**", "/3rdparty/cutlass/**", "/3rdparty/cutlass_fpA_intB_gemm/**", @@ -132,7 +141,7 @@ logging.level = "INFO" [tool.scikit-build.cmake.define] TVM_BUILD_PYTHON_MODULE = "ON" TVM_USE_INSTALLED_TIRX = "ON" -USE_CUDA = "OFF" +TVM_TIRX_USE_CUDA = "OFF" BUILD_TESTING = "OFF" [tool.setuptools_scm] diff --git a/tests/scripts/task_build.py b/tests/scripts/task_build.py index f7bfd558e6f1..617e8b1a1b0b 100755 --- a/tests/scripts/task_build.py +++ b/tests/scripts/task_build.py @@ -19,6 +19,7 @@ import logging import multiprocessing import os +import shlex import shutil import subprocess import sys @@ -31,7 +32,9 @@ format="[%(filename)s:%(lineno)d %(levelname)s] %(message)s", level=logging.INFO ) - parser = argparse.ArgumentParser(description="List pytest nodeids for a folder") + parser = argparse.ArgumentParser( + description="Build TVM with optional settings from the CMAKE_ARGS environment variable" + ) parser.add_argument("--sccache-bucket", required=False, help="sccache bucket name") parser.add_argument("--sccache-region", required=False, help="sccache region") parser.add_argument("--build-dir", default="build", help="build folder") @@ -55,6 +58,8 @@ build_dir = Path(os.getcwd()) / args.build_dir build_dir = build_dir.relative_to(REPO_ROOT) env = os.environ.copy() + # CMAKE_ARGS configures TVM, not the separate source-FFI Python package. + cmake_args = shlex.split(env.pop("CMAKE_ARGS", "")) env["PYTHONPATH"] = os.pathsep.join( [ str(REPO_ROOT / "tirx" / "python"), @@ -86,7 +91,8 @@ # Build from this checkout's pinned submodule once, then transfer the exact # Python/native package together with the TVM libraries to test workers. ffi_commit = subprocess.check_output( - ["git", "-C", str(REPO_ROOT / "3rdparty" / "tvm-ffi"), "rev-parse", "HEAD"], text=True + ["git", "-C", str(REPO_ROOT / "tirx" / "3rdparty" / "tvm-ffi"), "rev-parse", "HEAD"], + text=True, ).strip() ffi_identity = f"{ffi_commit} {sys.implementation.cache_tag}" ffi_marker = REPO_ROOT / "python" / "tvm_ffi" / ".source-commit" @@ -102,7 +108,7 @@ "--no-deps", "--reinstall", "--target=python", - "./3rdparty/tvm-ffi/", + "./tirx/3rdparty/tvm-ffi/", ], cwd=REPO_ROOT, env=env, @@ -125,9 +131,12 @@ available_cpus = nproc // executors num_cpus = max(available_cpus, 1) + if "tirx_cpptest" in args.cmake_target: + cmake_args.append("-DTVM_TIRX_BUILD_TESTS=ON") + build_type = "Debug" if args.debug else "RelWithDebInfo" subprocess.run( - ["cmake", "-GNinja", f"-DCMAKE_BUILD_TYPE={build_type}", ".."], + ["cmake", "-GNinja", f"-DCMAKE_BUILD_TYPE={build_type}", *cmake_args, ".."], cwd=build_dir, env=env, check=True, diff --git a/tests/scripts/task_config_build_arm.sh b/tests/scripts/task_config_build_arm.sh index b58b21db34c7..60778f5c11eb 100755 --- a/tests/scripts/task_config_build_arm.sh +++ b/tests/scripts/task_config_build_arm.sh @@ -25,7 +25,7 @@ cp ../cmake/config.cmake . echo set\(USE_SORT ON\) >> config.cmake echo set\(USE_RPC ON\) >> config.cmake -echo set\(USE_LLVM llvm-config-17\) >> config.cmake +echo set\(TVM_TIRX_USE_LLVM llvm-config-17\) >> config.cmake echo set\(CMAKE_CXX_FLAGS -Werror\) >> config.cmake echo set\(USE_CCACHE OFF\) >> config.cmake echo set\(SUMMARIZE ON\) >> config.cmake diff --git a/tests/scripts/task_config_build_cpu.sh b/tests/scripts/task_config_build_cpu.sh index e9d528be11b9..6a8586db5b17 100755 --- a/tests/scripts/task_config_build_cpu.sh +++ b/tests/scripts/task_config_build_cpu.sh @@ -27,7 +27,7 @@ echo set\(USE_SORT ON\) >> config.cmake echo set\(USE_DNNL ON\) >> config.cmake echo set\(USE_EXAMPLE_NPU_CODEGEN ON\) >> config.cmake echo set\(USE_EXAMPLE_NPU_RUNTIME ON\) >> config.cmake -echo set\(USE_LLVM \"/usr/bin/llvm-config-17 --link-static\"\) >> config.cmake +echo set\(TVM_TIRX_USE_LLVM \"/usr/bin/llvm-config-17 --link-static\"\) >> config.cmake echo set\(CMAKE_CXX_FLAGS \"-Wno-error=range-loop-construct -Wno-error=comment\"\) >> config.cmake echo set\(HIDE_PRIVATE_SYMBOLS ON\) >> config.cmake echo set\(USE_TENSORFLOW_PATH \"/tensorflow\"\) >> config.cmake diff --git a/tests/scripts/task_config_build_gpu.sh b/tests/scripts/task_config_build_gpu.sh index a16847023a61..2f929a9e203b 100755 --- a/tests/scripts/task_config_build_gpu.sh +++ b/tests/scripts/task_config_build_gpu.sh @@ -25,11 +25,11 @@ cp ../cmake/config.cmake . echo set\(USE_CUBLAS ON\) >> config.cmake echo set\(USE_CUDNN ON\) >> config.cmake -echo set\(USE_CUDA ON\) >> config.cmake -echo set\(USE_VULKAN ON\) >> config.cmake -echo set\(USE_OPENCL ON\) >> config.cmake -echo set\(USE_OPENCL_GTEST \"/googletest\"\) >> config.cmake -echo set\(USE_LLVM \"/usr/bin/llvm-config-15 --link-static\"\) >> config.cmake +echo set\(TVM_TIRX_USE_CUDA ON\) >> config.cmake +echo set\(TVM_TIRX_USE_VULKAN ON\) >> config.cmake +echo set\(TVM_TIRX_USE_OPENCL ON\) >> config.cmake +echo set\(TVM_TIRX_USE_OPENCL_GTEST \"/googletest\"\) >> config.cmake +echo set\(TVM_TIRX_USE_LLVM \"/usr/bin/llvm-config-15 --link-static\"\) >> config.cmake echo set\(USE_RPC ON\) >> config.cmake echo set\(USE_SORT ON\) >> config.cmake echo set\(USE_BLAS openblas\) >> config.cmake diff --git a/tests/scripts/task_config_build_gpu_other.sh b/tests/scripts/task_config_build_gpu_other.sh index 0d5d2030c47d..de5b94b5dc13 100755 --- a/tests/scripts/task_config_build_gpu_other.sh +++ b/tests/scripts/task_config_build_gpu_other.sh @@ -25,8 +25,8 @@ mkdir -p "$BUILD_DIR" cd "$BUILD_DIR" cp ../cmake/config.cmake . -echo set\(USE_OPENCL ON\) >> config.cmake -echo set\(USE_ROCM ON\) >> config.cmake +echo set\(TVM_TIRX_USE_OPENCL ON\) >> config.cmake +echo set\(TVM_TIRX_USE_ROCM ON\) >> config.cmake echo set\(CMAKE_CXX_FLAGS -Werror\) >> config.cmake echo set\(USE_CCACHE OFF\) >> config.cmake echo set\(SUMMARIZE ON\) >> config.cmake diff --git a/tests/scripts/task_config_build_wasm.sh b/tests/scripts/task_config_build_wasm.sh index 11526082e481..a73464ef2145 100755 --- a/tests/scripts/task_config_build_wasm.sh +++ b/tests/scripts/task_config_build_wasm.sh @@ -24,7 +24,7 @@ cd "$BUILD_DIR" cp ../cmake/config.cmake . echo set\(USE_SORT ON\) >> config.cmake -echo set\(USE_LLVM llvm-config-15\) >> config.cmake +echo set\(TVM_TIRX_USE_LLVM llvm-config-15\) >> config.cmake echo set\(CMAKE_CXX_FLAGS -Werror\) >> config.cmake echo set\(HIDE_PRIVATE_SYMBOLS ON\) >> config.cmake echo set\(USE_CCACHE OFF\) >> config.cmake diff --git a/tests/scripts/task_jvm_build.sh b/tests/scripts/task_jvm_build.sh index 468fea4ec6be..60c0c725cb51 100755 --- a/tests/scripts/task_jvm_build.sh +++ b/tests/scripts/task_jvm_build.sh @@ -35,7 +35,7 @@ ROOTDIR="$(cd "$(dirname "$0")/../.." && pwd)" TVM_BUILD_PATH="${TVM_BUILD_PATH:-${ROOTDIR}/build}" TVM_BUILD_PATH="$(realpath "${TVM_BUILD_PATH}")" -DLPACK_PATH="${DLPACK_PATH:-${ROOTDIR}/3rdparty/tvm-ffi/3rdparty/dlpack}" +DLPACK_PATH="${DLPACK_PATH:-${ROOTDIR}/tirx/3rdparty/tvm-ffi/3rdparty/dlpack}" INCLUDE_FLAGS="-I${ROOTDIR}/include -I${DLPACK_PATH}/include" PKG_CFLAGS="-Wall -O3 ${INCLUDE_FLAGS} -fPIC" diff --git a/tests/scripts/task_python_docs.sh b/tests/scripts/task_python_docs.sh index 0a4b88df9654..57544bee2bcc 100755 --- a/tests/scripts/task_python_docs.sh +++ b/tests/scripts/task_python_docs.sh @@ -48,7 +48,7 @@ sphinx_precheck() { echo "PreCheck sphinx doc generation WARNINGS.." # setup tvm-ffi into python folder - uv pip install -v --target=python ./3rdparty/tvm-ffi/ + uv pip install -v --target=python ./tirx/3rdparty/tvm-ffi/ pushd docs make clean @@ -130,7 +130,7 @@ find . -type f -path "*.log" | xargs rm -f find . -type f -path "*.pyc" | xargs rm -f # setup tvm-ffi into python folder -uv pip install -v --target=python ./3rdparty/tvm-ffi/ +uv pip install -v --target=python ./tirx/3rdparty/tvm-ffi/ # LLVMModule JIT execution, needed by the tutorials this build runs. uv pip install --no-deps apache-tvm-ffi-orcjit==0.1.2 diff --git a/tests/scripts/task_web_wasm.sh b/tests/scripts/task_web_wasm.sh index aed91a371bc5..f901d13fee74 100755 --- a/tests/scripts/task_web_wasm.sh +++ b/tests/scripts/task_web_wasm.sh @@ -21,7 +21,7 @@ set -euxo pipefail export PYTHONPATH="$(pwd)/tirx/python:$(pwd)/python" # setup tvm-ffi into python folder -uv pip install -v --target=python ./3rdparty/tvm-ffi/ +uv pip install -v --target=python ./tirx/3rdparty/tvm-ffi/ rm -rf .emscripten_cache cd web diff --git a/3rdparty/tvm-ffi b/tirx/3rdparty/tvm-ffi similarity index 100% rename from 3rdparty/tvm-ffi rename to tirx/3rdparty/tvm-ffi diff --git a/tirx/CMakeLists.txt b/tirx/CMakeLists.txt index 0b508905eea1..510d708c5bbe 100644 --- a/tirx/CMakeLists.txt +++ b/tirx/CMakeLists.txt @@ -15,244 +15,132 @@ # specific language governing permissions and limitations # under the License. -cmake_minimum_required(VERSION 3.18) -project(tvm_tirx C CXX) -include(GNUInstallDirs) -if(TVM_BUILD_PYTHON_MODULE) - set(CMAKE_INSTALL_LIBDIR lib) +cmake_minimum_required(VERSION 3.26) +if(CMAKE_CURRENT_SOURCE_DIR STREQUAL CMAKE_SOURCE_DIR AND NOT DEFINED CMAKE_EXPORT_COMPILE_COMMANDS) + set(CMAKE_EXPORT_COMPILE_COMMANDS ON CACHE BOOL "Export compilation commands") endif() +project(tvm_tirx LANGUAGES C CXX) + +# ######### Common options ########## -# The extracted package has an independent version; parent SCM tags do not apply. set(TVM_TIRX_VERSION "0.1.0a0") -set(TVM_VERSION "${TVM_TIRX_VERSION}") -add_compile_definitions(TVM_VERSION="${TVM_VERSION}") - -# Utility functions -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/Utils.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/Summary.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/Linker.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/Library.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/FindCUDA.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/FindOpenCL.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/FindVulkan.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/FindLLVM.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/FindROCM.cmake") - -if(EXISTS ${CMAKE_BINARY_DIR}/config.cmake) - include(${CMAKE_BINARY_DIR}/config.cmake) -else() - if(EXISTS ${CMAKE_SOURCE_DIR}/config.cmake) - include(${CMAKE_SOURCE_DIR}/config.cmake) - endif() +if(PROJECT_IS_TOP_LEVEL) + option(TVM_TIRX_BUILD_TESTS "Build C++ tests" OFF) + option(TVM_TIRX_BUILD_PYTHON_MODULE "Build the Python package with scikit-build-core" OFF) endif() +option(TVM_TIRX_USE_THREADS "Build with thread support" ON) +option(TVM_TIRX_USE_CUSTOM_LOGGING "Use user-defined logging implementations" OFF) +set(TVM_TIRX_USE_OPENMP OFF CACHE STRING "OpenMP thread pool implementation (OFF, gnu, intel)") +set_property(CACHE TVM_TIRX_USE_OPENMP PROPERTY STRINGS OFF gnu intel) -# NOTE: do not modify this file to change option values. -# You can create a config.cmake at build folder -# and add set(OPTION VALUE) to override these build options. -# Alernatively, use cmake -DOPTION=VALUE through command-line. -tvm_option(USE_CUDA "Build with CUDA" OFF) -tvm_option(USE_OPENCL "Build with OpenCL" OFF) -tvm_option(USE_OPENCL_ENABLE_HOST_PTR "Enable OpenCL memory object access to host" OFF) -tvm_option(USE_OPENCL_GTEST "Path to OpenCL specific gtest version for runtime cpp tests." /path/to/opencl/gtest) -tvm_option(USE_VULKAN "Build with Vulkan" OFF) +# ######### LLVM ########## +set(TVM_TIRX_USE_LLVM OFF CACHE STRING "Build with LLVM (OFF, ON, or llvm-config command)") -# Whether to use spirv-tools and SPIRV-Headers from Khronos GitHub or GitLab. -# -# Possible values: -# - OFF: not to use -# - /path/to/install: path to your khronis spirv-tools and SPIRV-Headers installation directory -# -tvm_option(USE_KHRONOS_SPIRV "Whether to use spirv-tools and SPIRV-Headers from Khronos GitHub or GitLab" OFF) -tvm_option(USE_SPIRV_KHR_INTEGER_DOT_PRODUCT "whether enable SPIRV_KHR_DOT_PRODUCT" OFF) -tvm_option(USE_METAL "Build with Metal" OFF) -tvm_option(USE_ROCM "Build with ROCM" OFF) -tvm_option(ROCM_PATH "The path to rocm" /opt/rocm) -tvm_option(USE_HEXAGON_QHL "Enable Hexagon QHL compiler lowering (requires external runtime wrappers)" OFF) - -tvm_option(USE_THREADS "Build with thread support" ON) -tvm_option(USE_LLVM "Build with LLVM, can be set to specific llvm-config path" OFF) -tvm_option(USE_OPENMP "Build with OpenMP thread pool implementation" OFF) -tvm_option(TVM_LOG_BEFORE_THROW "Whether log before throw, for debugging purposes" OFF) -tvm_option(USE_RTTI "Build with RTTI" ON) -tvm_option(USE_MSVC_MT "Build with MT" OFF) -tvm_option(INSTALL_DEV "Install compiler infrastructure" OFF) -tvm_option(HIDE_PRIVATE_SYMBOLS "Compile with -fvisibility=hidden." ON) -tvm_option(INDEX_DEFAULT_I64 "Defaults the index datatype to int64" ON) -tvm_option(USE_GTEST "Use GoogleTest for C++ sanity tests" AUTO) -tvm_option(USE_CUSTOM_LOGGING "Use user-defined custom logging, tvm::runtime::detail::LogFatalImpl and tvm::runtime::detail::LogMessageImpl must be implemented" OFF) -tvm_option(USE_ALTERNATIVE_LINKER "Use 'mold' or 'lld' if found when invoking compiler to link artifact" AUTO) -tvm_option(USE_CCACHE "Use ccache if found when invoking compiler" AUTO) - -# 3rdparty libraries -tvm_option(COMPILER_RT_PATH "Path to COMPILER-RT" "${CMAKE_CURRENT_SOURCE_DIR}/3rdparty/compiler-rt") -# Contrib library options -tvm_option(USE_Z3 "Build with Z3 SMT solver support" AUTO) -tvm_option(USE_NVTX "Build with NVTX" OFF) -tvm_option(SUMMARIZE "Print CMake option summary after configuring" OFF) - -# Python package options -tvm_option(TVM_BUILD_PYTHON_MODULE "Build Python module with scikit-build-core" OFF) - -# include directories -include_directories(${CMAKE_INCLUDE_PATH}) -include_directories("include") -include_directories(SYSTEM ${COMPILER_RT_PATH}) - -# initial variables -set(TVM_LINKER_LIBS "") -set(TVM_RUNTIME_LINKER_LIBS "") -set(TVM_RUNTIME_BACKEND_LIBS "") -# Use one external FFI registry for both the core and its parent consumers. -find_package(Python COMPONENTS Interpreter REQUIRED) -if(NOT tvm_ffi_DIR) - execute_process(COMMAND "${Python_EXECUTABLE}" -m tvm_ffi.config --cmakedir - OUTPUT_VARIABLE tvm_ffi_DIR OUTPUT_STRIP_TRAILING_WHITESPACE - COMMAND_ERROR_IS_FATAL ANY) -endif() -if(NOT TARGET tvm_ffi::shared) - find_package(tvm_ffi CONFIG REQUIRED) - foreach(_ffi_target header shared static) - if(TARGET tvm_ffi::${_ffi_target}) - set_target_properties(tvm_ffi::${_ffi_target} PROPERTIES IMPORTED_GLOBAL TRUE) - endif() - endforeach() -endif() -if(NOT TARGET tvm_ffi_header) - add_library(tvm_ffi_header ALIAS tvm_ffi::header) -endif() +# ######### Z3 ########## + +set(TVM_TIRX_USE_Z3 AUTO CACHE STRING "Build with Z3 SMT solver support (OFF, ON, AUTO)") +set_property(CACHE TVM_TIRX_USE_Z3 PROPERTY STRINGS OFF ON AUTO) + +# ######### CUDA ########## -# Check if this is being run on its own or as a subdirectory for another project -# If we update to CMake 2.21+, we can use PROJECT_IS_TOP_LEVEL instead -get_directory_property(IS_SUBPROJECT PARENT_DIRECTORY) +set(TVM_TIRX_USE_CUDA OFF CACHE STRING "Build with CUDA (OFF, ON, or SDK path)") -if(NOT IS_SUBPROJECT AND NOT DEFINED "${CMAKE_EXPORT_COMPILE_COMMANDS}") - # If not set manually, change the default to ON - set(CMAKE_EXPORT_COMPILE_COMMANDS ON) +# ######### Other device runtimes ########## + +# OpenCL +set(TVM_TIRX_USE_OPENCL OFF CACHE STRING "Build with OpenCL (OFF, ON, or SDK path)") +option(TVM_TIRX_USE_OPENCL_ENABLE_HOST_PTR "Enable OpenCL memory object access to host" OFF) +option(TVM_TIRX_USE_OPENCL_EXTN_QCOM "Enable Qualcomm OpenCL extensions" OFF) + +# Vulkan +set(TVM_TIRX_USE_VULKAN OFF CACHE STRING "Build with Vulkan (OFF, ON, or SDK path)") +set(TVM_TIRX_USE_KHRONOS_SPIRV OFF CACHE STRING "Use Khronos SPIR-V tools (OFF or install path)") +option(TVM_TIRX_USE_SPIRV_KHR_INTEGER_DOT_PRODUCT "Enable SPIR-V integer dot product" OFF) + +# ROCm +set(TVM_TIRX_ROCM_PATH /opt/rocm CACHE PATH "Default ROCm SDK location") +set(TVM_TIRX_USE_ROCM OFF CACHE STRING "Build with ROCm (OFF, ON, or SDK path)") + +# Metal +option(TVM_TIRX_USE_METAL "Build with Metal" OFF) + +# ######### Dependencies ########## + +# Use the same source-built FFI targets throughout this build graph. +add_subdirectory(3rdparty/tvm-ffi) +include(cmake/utils/Library.cmake) +include(cmake/utils/FindLLVM.cmake) +include(cmake/utils/FindCUDA.cmake) +include(cmake/utils/FindOpenCL.cmake) +include(cmake/utils/FindVulkan.cmake) +include(cmake/utils/FindROCM.cmake) +set(IS_FALSE_PATTERN "^[Oo][Ff][Ff]$|^0$|^[Ff][Aa][Ll][Ss][Ee]$|^[Nn][Oo]$|^[Nn][Oo][Tt][Ff][Oo][Uu][Nn][Dd]$|.*-[Nn][Oo][Tt][Ff][Oo][Uu][Nn][Dd]$|^$") +set(IS_TRUE_PATTERN "^[Oo][Nn]$|^[1-9][0-9]*$|^[Tt][Rr][Uu][Ee]$|^[Yy][Ee][Ss]$|^[Yy]$") +if(TVM_TIRX_USE_THREADS) + find_package(Threads REQUIRED) endif() -if(TVM_LOG_BEFORE_THROW) - # log error before throw as - # when system have issues with stack trace - add_definitions(-DTVM_LOG_BEFORE_THROW=1) +# ######### TIRx build settings ########## + +# Apply these after the FFI subdirectory so its own settings remain independent. +set(CMAKE_POSITION_INDEPENDENT_CODE ON) +set(CMAKE_CXX_STANDARD 20) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CUDA_STANDARD 20) +set(CMAKE_CUDA_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_VISIBILITY_PRESET hidden) +set(CMAKE_VISIBILITY_INLINES_HIDDEN ON) +if(PROJECT_IS_TOP_LEVEL) + include(GNUInstallDirs) + if(TVM_TIRX_BUILD_PYTHON_MODULE) + set(CMAKE_INSTALL_LIBDIR lib) + endif() endif() +add_compile_definitions(TVM_VERSION="${TVM_TIRX_VERSION}") +include_directories("${CMAKE_CURRENT_SOURCE_DIR}/3rdparty/compiler-rt") -# Generic compilation options if(MSVC) - add_definitions(-DWIN32_LEAN_AND_MEAN) - add_definitions(-D_CRT_SECURE_NO_WARNINGS) - add_definitions(-D_SCL_SECURE_NO_WARNINGS) - add_definitions(-D_ENABLE_EXTENDED_ALIGNED_STORAGE) - add_definitions(-DNOMINMAX) - # regeneration does not work well with msbuild custom rules. - set(CMAKE_SUPPRESS_REGENERATION ON) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /EHsc") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /MP") - add_compile_options(/bigobj) - - # Use standard-conforming two-phase name resolution for templates. - # This minimizes the differences between g++/clang builds on Linux, - # and MSVC builds on Windows. - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /permissive-") - - # MSVC already errors on undefined symbols, no additional flag needed. - set(TVM_NO_UNDEFINED_SYMBOLS "") - - if(USE_MSVC_MT) - foreach(flag_var - CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_DEBUG CMAKE_CXX_FLAGS_RELEASE - CMAKE_CXX_FLAGS_MINSIZEREL CMAKE_CXX_FLAGS_RELWITHDEBINFO) - if(${flag_var} MATCHES "/MD") - string(REGEX REPLACE "/MD" "/MT" ${flag_var} "${${flag_var}}") - endif(${flag_var} MATCHES "/MD") - endforeach(flag_var) - # Static linking. CMake behavior changed in 3.15 making this necessary. - add_compile_options(/MT) - endif() - # Disable common MSVC warnings - # Integer conversion warnings(e.g. int64 to int) - add_compile_options(/wd4244) - add_compile_options(/wd4267) - # Signed unsigned constant comparison - add_compile_options(/wd4018) - # Aligned alloc may not met(need c++17) - add_compile_options(/wd4316) - # unreferenced local variables(usually in exception catch) - add_compile_options(/wd4101) - # always inline keyword not necessary - add_compile_options(/wd4180) - # DLL interface warning in c++ - add_compile_options(/wd4251) - # destructor was implicitly defined as deleted - add_compile_options(/wd4624) - # unary minus operator applied to unsigned type, result still unsigned - add_compile_options(/wd4146) - # 'inline': used more than once - add_compile_options(/wd4141) - # unknown pragma - add_compile_options(/wd4068) -else(MSVC) - set(WARNING_FLAG -Wall) - if ("${CMAKE_BUILD_TYPE}" STREQUAL "Debug") - message(STATUS "Build in Debug mode") - set(CMAKE_C_FLAGS "-O0 -g ${WARNING_FLAG} -fPIC ${CMAKE_C_FLAGS}") - set(CMAKE_CXX_FLAGS "-O0 -g ${WARNING_FLAG} -fPIC ${CMAKE_CXX_FLAGS}") - set(CMAKE_CUDA_FLAGS "-O0 -g -Xcompiler=-Wall -Xcompiler=-fPIC ${CMAKE_CUDA_FLAGS}") - else() - set(CMAKE_C_FLAGS "-O2 ${WARNING_FLAG} -fPIC ${CMAKE_C_FLAGS}") - set(CMAKE_CXX_FLAGS "-O2 ${WARNING_FLAG} -fPIC ${CMAKE_CXX_FLAGS}") - set(CMAKE_CUDA_FLAGS "-O2 -Xcompiler=-Wall -Xcompiler=-fPIC ${CMAKE_CUDA_FLAGS}") - set(TVM_VISIBILITY_FLAG "") - if (HIDE_PRIVATE_SYMBOLS) - message(STATUS "Hide private symbols...") - set(TVM_VISIBILITY_FLAG "-fvisibility=hidden") - endif(HIDE_PRIVATE_SYMBOLS) - endif () - if (CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang" AND - CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 7.0) - set(CMAKE_CXX_FLAGS "-faligned-new ${CMAKE_CXX_FLAGS}") - endif() + add_compile_definitions(WIN32_LEAN_AND_MEAN _CRT_SECURE_NO_WARNINGS + _SCL_SECURE_NO_WARNINGS _ENABLE_EXTENDED_ALIGNED_STORAGE NOMINMAX) + add_compile_options(/EHsc /MP /bigobj /permissive- + /wd4244 /wd4267 /wd4018 /wd4316 /wd4101 /wd4180 /wd4251 /wd4624 /wd4146 /wd4141 /wd4068) +else() + add_compile_options(-Wall) +endif() - # ld option to warn if symbols are undefined (e.g. libtvm_runtime.so - # using symbols only present in libtvm.so). Not needed for MSVC, - # since this is already the default there. - if(${CMAKE_SYSTEM_NAME} MATCHES "Darwin" OR ${CMAKE_SYSTEM_NAME} MATCHES "iOS") - set(TVM_NO_UNDEFINED_SYMBOLS "-Wl,-undefined,error") - else() - set(TVM_NO_UNDEFINED_SYMBOLS "-Wl,--no-undefined") - endif() - message(STATUS "Forbidding undefined symbols in shared library, using ${TVM_NO_UNDEFINED_SYMBOLS} on platform ${CMAKE_SYSTEM_NAME}") - - # Detect if we're compiling for Android. - set(TEST_FOR_ANDROID_CXX - "#ifndef __ANDROID__" - "#error" - "#endif" - "int main() {}") - set(TEST_FOR_ANDROID_DIR - "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeTmp") - set(TEST_FOR_ANDROID_FILE "${TEST_FOR_ANDROID_DIR}/test_for_android.cc") - string(REPLACE ";" "\n" TEST_FOR_ANDROID_CXX_TEXT "${TEST_FOR_ANDROID_CXX}") - file(WRITE "${TEST_FOR_ANDROID_FILE}" "${TEST_FOR_ANDROID_CXX_TEXT}") - try_compile(BUILD_FOR_ANDROID "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}" - "${TEST_FOR_ANDROID_FILE}") - file(REMOVE "${TEST_FOR_ANDROID_FILE}") - if(BUILD_FOR_ANDROID) - message(STATUS "Building for Android") +# Keep libraries together while honoring an embedding parent's output layout. +foreach(_kind ARCHIVE LIBRARY RUNTIME) + if(NOT DEFINED CMAKE_${_kind}_OUTPUT_DIRECTORY) + set(CMAKE_${_kind}_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/lib") endif() -endif(MSVC) + foreach(_config IN LISTS CMAKE_CONFIGURATION_TYPES) + string(TOUPPER "${_config}" _config) + set(CMAKE_${_kind}_OUTPUT_DIRECTORY_${_config} "${CMAKE_${_kind}_OUTPUT_DIRECTORY}") + endforeach() +endforeach() + +# ######### Target: `tvm_tirx_header` ########## -list(APPEND TVM_RUNTIME_LINKER_LIBS ${CMAKE_DL_LIBS}) +add_library(tvm_tirx_header INTERFACE) +add_library(tvm_tirx::header ALIAS tvm_tirx_header) +target_compile_features(tvm_tirx_header INTERFACE cxx_std_20) +target_compile_definitions(tvm_tirx_header INTERFACE TVM_INDEX_DEFAULT_I64=1) +target_include_directories(tvm_tirx_header INTERFACE + "$" + "$") +target_link_libraries(tvm_tirx_header INTERFACE tvm_ffi::header) -# add source group -tvm_file_glob(GLOB_RECURSE GROUP_SOURCE "src/*.cc") -tvm_file_glob(GLOB_RECURSE GROUP_INCLUDE "src/*.h" "include/*.h") -assign_source_group("Source" ${GROUP_SOURCE}) -assign_source_group("Include" ${GROUP_INCLUDE}) +# ######### Sources and backends ########## +set(TVM_LINKER_LIBS "") +set(TVM_RUNTIME_LINKER_LIBS "${CMAKE_DL_LIBS}") +set(TVM_RUNTIME_BACKEND_LIBS "") +set(BUILD_FOR_ANDROID "${ANDROID}") # Source file lists # Core compiler sources, extended by the backend modules. -tvm_file_glob(GLOB_RECURSE COMPILER_SRCS +file(GLOB_RECURSE COMPILER_SRCS CONFIGURE_DEPENDS src/ir/*.cc src/sym/*.cc src/tirx/*.cc @@ -278,7 +166,7 @@ tvm_file_glob(GLOB_RECURSE COMPILER_SRCS src/script/printer/script_printer.cc ) -tvm_file_glob(GLOB CODEGEN_SRCS +file(GLOB CODEGEN_SRCS CONFIGURE_DEPENDS src/target/*.cc src/codegen/c/*.cc src/codegen/*.cc @@ -307,302 +195,134 @@ tvm_file_glob(GLOB CODEGEN_SRCS list(APPEND COMPILER_SRCS ${CODEGEN_SRCS}) # Core runtime sources, extended by the backend modules. -tvm_file_glob(GLOB RUNTIME_SRCS +file(GLOB RUNTIME_SRCS CONFIGURE_DEPENDS src/runtime/*.cc ) set(TVM_RUNTIME_EXT_OBJS "") -# Package runtime rules -if(NOT USE_RTTI) -endif() - -if(INDEX_DEFAULT_I64) - add_definitions(-DTVM_INDEX_DEFAULT_I64=1) -endif() - -if(USE_KALLOC_ALIGNMENT) - message(STATUS "Build Alloc alignment set to ${USE_KALLOC_ALIGNMENT}") - add_definitions(-DTVM_KALLOC_ALIGNMENT=${USE_KALLOC_ALIGNMENT}) -endif(USE_KALLOC_ALIGNMENT) - -# Caches the build. -# Note that ccache-3.x doesn't support nvcc well, so CUDA kernels may never hit the cache and still -# need to be re-compiled every time. Using ccache 4.0+ can resolve this issue. -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/CCache.cmake") - -include(CheckCXXCompilerFlag) -set(CMAKE_CXX_STANDARD_REQUIRED ON) -set(CMAKE_CXX_STANDARD 20) -set(CMAKE_CUDA_STANDARD_REQUIRED ON) -set(CMAKE_CUDA_STANDARD 20) - -# Module rules -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/CUDA.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/OpenCL.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/OpenMP.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/Vulkan.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/Metal.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/ROCM.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/LLVM.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/contrib/Z3.cmake") -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/Git.cmake") - -include(GNUInstallDirs) - -if(USE_THREADS) - message(STATUS "Build with thread support...") - set(CMAKE_THREAD_PREFER_PTHREAD TRUE) - set(THREADS_PREFER_PTHREAD_FLAG TRUE) - find_package(Threads REQUIRED) -endif() +# The core runtime's thread pool implementation supplies its link dependencies. +include(cmake/modules/OpenMP.cmake) -set(TVM_IS_DEBUG_BUILD OFF) -if(CMAKE_BUILD_TYPE STREQUAL "Debug" OR CMAKE_BUILD_TYPE STREQUAL "RelWithDebInfo" OR CMAKE_CXX_FLAGS MATCHES "-g") - set(TVM_IS_DEBUG_BUILD ON) - set(FILE_PREFIX_MAP_FLAG "-ffile-prefix-map=..=${CMAKE_CURRENT_SOURCE_DIR}") - CHECK_CXX_COMPILER_FLAG("${FILE_PREFIX_MAP_FLAG}" FILE_PREFIX_MAP_SUPPORTED) -endif() +# ######### Targets: `tvm_runtime`, `tvm_runtime_static` ########## -# Core runtime objects are shared by the dynamic and static libraries. add_library(tvm_runtime_objs OBJECT ${RUNTIME_SRCS}) -target_link_libraries(tvm_runtime_objs PUBLIC tvm_ffi_header) -if(TVM_VISIBILITY_FLAG) - target_compile_options(tvm_runtime_objs PRIVATE "${TVM_VISIBILITY_FLAG}") -endif() +target_link_libraries(tvm_runtime_objs PUBLIC tvm_tirx_header) if(MSVC) - # Runtime objects export only runtime/FFI symbols, never compiler symbols. target_compile_definitions(tvm_runtime_objs PRIVATE TVM_RUNTIME_EXPORTS TVM_FFI_EXPORTS) endif() -# Loaded with RTLD_GLOBAL so downstream runtime modules can resolve its symbols. -add_library(tvm_runtime SHARED - $ - ${TVM_RUNTIME_EXT_OBJS} -) -target_link_libraries(tvm_runtime PUBLIC tvm_ffi::shared) -set_property(TARGET tvm_runtime APPEND PROPERTY LINK_OPTIONS "${TVM_NO_UNDEFINED_SYMBOLS}") -set_property(TARGET tvm_runtime APPEND PROPERTY LINK_OPTIONS "${TVM_VISIBILITY_FLAG}") - -# Static consumers must retain registration objects from both archives, with -# tvm_ffi::static before tvm_runtime_static so FFI type attributes initialize first: -# -Wl,--whole-archive tvm_ffi::static tvm_runtime_static -Wl,--no-whole-archive -# Use the corresponding archive-retention flags on other platforms. -set(_ffi_static_dependency tvm_ffi::shared) -if(TARGET tvm_ffi::static) - set(_ffi_static_dependency tvm_ffi::static) -endif() -add_library(tvm_runtime_static STATIC - $ - ${TVM_RUNTIME_EXT_OBJS} -) -target_link_libraries(tvm_runtime_static PUBLIC ${_ffi_static_dependency}) +add_library(tvm_runtime SHARED $ ${TVM_RUNTIME_EXT_OBJS}) +add_library(tvm_runtime::shared ALIAS tvm_runtime) +target_link_libraries(tvm_runtime PUBLIC tvm_tirx_header tvm_ffi::shared) +add_library(tvm_runtime_static STATIC $ ${TVM_RUNTIME_EXT_OBJS}) +target_link_libraries(tvm_runtime_static PUBLIC tvm_tirx_header tvm_ffi::static) +# Static consumers retain registration objects from both archives, with +# tvm_ffi::static before tvm_runtime_static in their whole-archive link group. if(MSVC) - # Static consumers reference ordinary symbols rather than DLL import stubs. target_compile_definitions(tvm_runtime_static INTERFACE - TVM_RUNTIME_DLL= TVM_RUNTIME_DLL_EXPORT=) - if(TARGET tvm_ffi::static) - target_compile_definitions(tvm_runtime_static INTERFACE TVM_FFI_DLL= TVM_FFI_DLL_EXPORT=) - endif() + TVM_RUNTIME_DLL= TVM_RUNTIME_DLL_EXPORT= TVM_FFI_DLL= TVM_FFI_DLL_EXPORT=) endif() - -set(_runtime_targets tvm_runtime tvm_runtime_static) - -foreach(runtime_target IN LISTS _runtime_targets) - target_include_directories(${runtime_target} PUBLIC "$" "$") - target_link_libraries(${runtime_target} PRIVATE ${TVM_RUNTIME_LINKER_LIBS}) - if(USE_THREADS) - target_link_libraries(${runtime_target} PUBLIC Threads::Threads) +foreach(_target tvm_runtime tvm_runtime_static) + target_link_libraries(${_target} PRIVATE ${TVM_RUNTIME_LINKER_LIBS}) + if(TVM_TIRX_USE_THREADS) + target_link_libraries(${_target} PUBLIC Threads::Threads) endif() endforeach() -set_target_properties(${_runtime_targets} PROPERTIES - ARCHIVE_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" - LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" - RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" -) -if(MSVC) - foreach(config_type Release RELEASE Debug DEBUG RelWithDebInfo RELWITHDEBINFO MinSizeRel MINSIZEREL) - set_target_properties(${_runtime_targets} PROPERTIES - ARCHIVE_OUTPUT_DIRECTORY_${config_type} "${CMAKE_BINARY_DIR}/lib" - LIBRARY_OUTPUT_DIRECTORY_${config_type} "${CMAKE_BINARY_DIR}/lib" - RUNTIME_OUTPUT_DIRECTORY_${config_type} "${CMAKE_BINARY_DIR}/lib" - ) - endforeach() -endif() +# ######### Device runtimes and compiler backends ########## -if(TVM_IS_DEBUG_BUILD AND FILE_PREFIX_MAP_SUPPORTED) - target_compile_options(tvm_runtime PRIVATE $<$:${FILE_PREFIX_MAP_FLAG}>) - target_compile_options(tvm_runtime_objs PRIVATE $<$:${FILE_PREFIX_MAP_FLAG}>) -endif() +# The core runtime exists before device runtimes link to it. Compiler backend +# contributions are collected before the compiler object target is created. +foreach(_module CUDA OpenCL Vulkan Metal ROCM LLVM contrib/Z3 Git) + include(cmake/modules/${_module}.cmake) +endforeach() +# ######### Target: `tvm_tirx` ########## -tvm_ffi_add_apple_dsymutil(tvm_runtime) -tvm_configure_target_library(tvm_runtime) -if(NOT IS_SUBPROJECT) - add_custom_target(runtime DEPENDS ${_runtime_targets}) -endif() -install(TARGETS tvm_runtime DESTINATION lib${LIB_SUFFIX}) -# Keep wheel contents unchanged; the archive is available to native consumers. -if(TARGET tvm_runtime_static AND NOT TVM_BUILD_PYTHON_MODULE) - install(TARGETS tvm_runtime_static DESTINATION lib${LIB_SUFFIX}) -endif() add_library(tvm_tirx_objs OBJECT ${COMPILER_SRCS}) -target_link_libraries(tvm_tirx_objs PUBLIC tvm_ffi_header) +target_link_libraries(tvm_tirx_objs PUBLIC tvm_tirx_header) if(TARGET tvm_llvm_header) - target_link_libraries(tvm_tirx_objs PUBLIC tvm_llvm_header) -endif() -if(TVM_VISIBILITY_FLAG) - target_compile_options(tvm_tirx_objs PRIVATE "${TVM_VISIBILITY_FLAG}") + target_link_libraries(tvm_tirx_objs PRIVATE tvm_llvm_header) endif() if(MSVC) - # Compiler objects import runtime symbols from the separate runtime DLL. target_compile_definitions(tvm_tirx_objs PRIVATE TVM_TIRX_EXPORTS TVM_FFI_EXPORTS) endif() +add_library(tvm_tirx SHARED $) +add_library(tvm_tirx::shared ALIAS tvm_tirx) +target_link_libraries(tvm_tirx PUBLIC tvm_runtime tvm_tirx_header) +target_link_libraries(tvm_tirx PRIVATE ${TVM_LINKER_LIBS} ${TVM_RUNTIME_LINKER_LIBS}) -# Compiler library: built from compiler-only sources. -# Links against tvm_runtime (and tvm_ffi::shared transitively). Loaded with -# RTLD_LOCAL so compiler-internal symbols don't leak into the global namespace. -add_library(tvm_tirx SHARED - $ -) -target_link_libraries(tvm_tirx PUBLIC tvm_runtime tvm_ffi::shared) -target_include_directories(tvm_tirx PUBLIC "$" "$") -set_property(TARGET tvm_tirx APPEND PROPERTY LINK_OPTIONS "${TVM_NO_UNDEFINED_SYMBOLS}") -set_property(TARGET tvm_tirx APPEND PROPERTY LINK_OPTIONS "${TVM_VISIBILITY_FLAG}") - -# Work around a GNU ld (binutils) relaxation bug that miscompiles -# R_X86_64_GOTPCRELX relocations inside very large statically-linked archives. -# When the full LLVM static libraries are linked into libtvm_tirx.so, the -# library is large enough that ld can relax an indirect GOT call (LLVM built -# with -fno-plt emits these) into a direct call with an incorrect displacement. -# The call then targets read-only data instead of the intended function and -# crashes at runtime with a SIGSEGV inside llvm::X86Subtarget during code -# generation. Disabling linker relaxation keeps the GOT-indirect sequences and -# avoids the miscompilation; it is harmless when LLVM is linked dynamically. -# See binutils bug ld/25754. -if(CMAKE_SYSTEM_NAME STREQUAL "Linux" AND NOT "${USE_LLVM}" MATCHES ${IS_FALSE_PATTERN}) - set_property(TARGET tvm_tirx APPEND PROPERTY LINK_OPTIONS "-Wl,--no-relax") - # LLVM's --system-libs reports -lxml2, but TVM calls no libxml2 symbols. The - # manylinux toolchain does not default to --as-needed, so set it explicitly: - # the unused libxml2 is then not recorded as a NEEDED dependency, so auditwheel - # does not vendor it (and its deps) into the wheel. - set_property(TARGET tvm_tirx APPEND PROPERTY LINK_OPTIONS "-Wl,--as-needed") +# Avoid GNU ld relaxation of LLVM's GOT calls into incorrect direct calls. +if(CMAKE_SYSTEM_NAME STREQUAL "Linux" AND NOT "${TVM_TIRX_USE_LLVM}" MATCHES ${IS_FALSE_PATTERN}) + target_link_options(tvm_tirx PRIVATE "-Wl,--no-relax" "-Wl,--as-needed") endif() +include(cmake/modules/Logging.cmake) +include(cmake/modules/ClangFlags.cmake) -set_target_properties(tvm_tirx PROPERTIES - ARCHIVE_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" - LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" - RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" -) -if(MSVC) - foreach(config_type Release RELEASE Debug DEBUG RelWithDebInfo RELWITHDEBINFO MinSizeRel MINSIZEREL) - set_target_properties(tvm_tirx PROPERTIES - ARCHIVE_OUTPUT_DIRECTORY_${config_type} "${CMAKE_BINARY_DIR}/lib" - LIBRARY_OUTPUT_DIRECTORY_${config_type} "${CMAKE_BINARY_DIR}/lib" - RUNTIME_OUTPUT_DIRECTORY_${config_type} "${CMAKE_BINARY_DIR}/lib" - ) - endforeach() -endif() +# ######### Library configuration ########## -if(USE_THREADS) - target_link_libraries(tvm_tirx PUBLIC Threads::Threads) -endif() -target_link_libraries(tvm_tirx PRIVATE ${TVM_LINKER_LIBS} ${TVM_RUNTIME_LINKER_LIBS}) +foreach(_target tvm_runtime_objs tvm_tirx_objs tvm_runtime tvm_runtime_static tvm_tirx) + tvm_ffi_add_prefix_map(${_target} "${CMAKE_CURRENT_SOURCE_DIR}") +endforeach() +foreach(_target IN LISTS TVM_RUNTIME_BACKEND_LIBS) + tvm_ffi_add_prefix_map(${_target}_objs "${CMAKE_CURRENT_SOURCE_DIR}") + tvm_ffi_add_prefix_map(${_target} "${CMAKE_CURRENT_SOURCE_DIR}") +endforeach() +foreach(_target tvm_runtime tvm_tirx ${TVM_RUNTIME_BACKEND_LIBS}) + tvm_tirx_configure_target(${_target}) + if(APPLE) + target_link_options(${_target} PRIVATE "-Wl,-undefined,error") + elseif(UNIX) + target_link_options(${_target} PRIVATE "-Wl,--no-undefined") + endif() + tvm_ffi_hide_static_linked_lib_symbols(${_target}) + tvm_ffi_add_apple_dsymutil(${_target}) +endforeach() -if(TVM_IS_DEBUG_BUILD) - target_compile_options(tvm_tirx PRIVATE "${FILE_PREFIX_MAP_FLAG}") - if(FILE_PREFIX_MAP_SUPPORTED) - target_compile_options(tvm_tirx PRIVATE $<$:${FILE_PREFIX_MAP_FLAG}>) - target_compile_options(tvm_tirx_objs PRIVATE $<$:${FILE_PREFIX_MAP_FLAG}>) +# ######### Installation ########## + +if(PROJECT_IS_TOP_LEVEL) + file(GLOB_RECURSE _headers CONFIGURE_DEPENDS RELATIVE + "${CMAKE_CURRENT_SOURCE_DIR}/include" "${CMAKE_CURRENT_SOURCE_DIR}/include/*.h") + foreach(_header IN LISTS _headers) + get_filename_component(_directory "${_header}" DIRECTORY) + install(FILES "${CMAKE_CURRENT_SOURCE_DIR}/include/${_header}" DESTINATION "include/${_directory}") + endforeach() + install(TARGETS tvm_tirx_header tvm_tirx tvm_runtime tvm_runtime_static + ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} + LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} + RUNTIME DESTINATION ${CMAKE_INSTALL_LIBDIR}) + install(FILES "${CMAKE_CURRENT_SOURCE_DIR}/cmake/tvm_tirx-config.cmake" + "${CMAKE_CURRENT_SOURCE_DIR}/python/tvm_tirx/config.py" DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/tvm_tirx") + install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/" + DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/tvm_tirx/utils" FILES_MATCHING PATTERN "Find*.cmake") + if(TVM_TIRX_BUILD_PYTHON_MODULE) + foreach(_extra_lib IN LISTS TVM_TIRX_PACKAGE_EXTRA_LIBS) + if(NOT EXISTS "${_extra_lib}") + message(FATAL_ERROR "TVM_TIRX_PACKAGE_EXTRA_LIBS entry does not exist: ${_extra_lib}") + endif() + install(FILES "${_extra_lib}" DESTINATION lib) + endforeach() + install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/licenses/" DESTINATION licenses) + install(FILES "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE" "${CMAKE_CURRENT_SOURCE_DIR}/NOTICE" DESTINATION .) endif() endif() -tvm_ffi_add_apple_dsymutil(tvm_tirx) -tvm_configure_target_library(tvm_tirx) -install(TARGETS tvm_tirx DESTINATION lib${LIB_SUFFIX}) - -# logging option for libbacktrace -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/Logging.cmake") +# ######### C++ tests ########## -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/ClangFlags.cmake") -if(USE_CUDA AND USE_NVTX) - add_compile_definitions(TVM_NVTX_ENABLED=1) - target_compile_definitions(tvm_runtime INTERFACE TVM_NVTX_ENABLED=1) -endif() -if(INDEX_DEFAULT_I64) - target_compile_definitions(tvm_runtime INTERFACE TVM_INDEX_DEFAULT_I64=1) -endif() -if(USE_KALLOC_ALIGNMENT) - target_compile_definitions(tvm_runtime INTERFACE TVM_KALLOC_ALIGNMENT=${USE_KALLOC_ALIGNMENT}) -endif() -if(TVM_DEBUG_WITH_ABI_CHANGE) - target_compile_definitions(tvm_runtime PUBLIC TVM_DEBUG_WITH_ABI_CHANGE) - target_compile_definitions(tvm_runtime_objs PRIVATE TVM_DEBUG_WITH_ABI_CHANGE) - target_compile_definitions(tvm_tirx_objs PRIVATE TVM_DEBUG_WITH_ABI_CHANGE) +if(TVM_TIRX_BUILD_TESTS) + enable_testing() + add_subdirectory(tests/cpp) endif() -get_target_property(_runtime_compile_definitions tvm_runtime INTERFACE_COMPILE_DEFINITIONS) -if(_runtime_compile_definitions) - target_compile_definitions(tvm_runtime_static INTERFACE ${_runtime_compile_definitions}) -endif() -add_library(tvm_tirx_private_headers INTERFACE) -set_target_properties(tvm_tirx_private_headers PROPERTIES EXPORT_NAME private_headers) -target_include_directories(tvm_tirx_private_headers INTERFACE - "$" - "$" - "$" - "$") -set(_tirx_targets tvm_tirx ${_runtime_targets} tvm_tirx_private_headers ${TVM_RUNTIME_BACKEND_LIBS}) -foreach(_target IN LISTS _tirx_targets) - get_target_property(_export_name ${_target} EXPORT_NAME) - if(NOT _export_name) - set(_export_name ${_target}) - endif() - add_library(tvm_tirx::${_export_name} ALIAS ${_target}) -endforeach() -install(TARGETS ${_tirx_targets} EXPORT tvm_tirxTargets - ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} - LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} - RUNTIME DESTINATION ${CMAKE_INSTALL_LIBDIR}) -install(DIRECTORY include/ DESTINATION ${CMAKE_INSTALL_INCLUDEDIR} FILES_MATCHING PATTERN "*.h") -# CLML's retained OpenCL interface needs this bounded private include closure. -foreach(_header backend/opencl/memory_scope.h - backend/opencl/runtime/opencl_common.h backend/opencl/runtime/texture.h - runtime/metadata.h runtime/pack_args.h runtime/workspace_pool.h) - get_filename_component(_dir "${_header}" DIRECTORY) - install(FILES "src/${_header}" DESTINATION "${CMAKE_INSTALL_INCLUDEDIR}/tvm_tirx/private/${_dir}") -endforeach() -install(FILES 3rdparty/compiler-rt/builtin_fp16.h DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/tvm_tirx/compiler-rt) -include(CMakePackageConfigHelpers) -configure_package_config_file(cmake/tvm_tirxConfig.cmake.in - "${CMAKE_CURRENT_BINARY_DIR}/tvm_tirxConfig.cmake" - INSTALL_DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/tvm_tirx") -install(EXPORT tvm_tirxTargets NAMESPACE tvm_tirx:: DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/tvm_tirx") -install(FILES "${CMAKE_CURRENT_BINARY_DIR}/tvm_tirxConfig.cmake" DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/tvm_tirx") -install(DIRECTORY cmake/utils cmake/modules DESTINATION "${CMAKE_INSTALL_LIBDIR}/cmake/tvm_tirx" FILES_MATCHING PATTERN "*.cmake") set(TVM_TIRX_CMAKE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/cmake") -if(IS_SUBPROJECT) - set(TVM_TIRX_CMAKE_DIR "${TVM_TIRX_CMAKE_DIR}" PARENT_SCOPE) - foreach(_option USE_CUDA USE_ROCM BUILD_FOR_ANDROID) - set(TVM_TIRX_${_option} "${${_option}}" PARENT_SCOPE) - endforeach() -endif() -add_subdirectory(tests/cpp) -if(GTEST_FOUND) - enable_testing() - include(CTest) -endif() -find_and_set_linker(${USE_ALTERNATIVE_LINKER}) -if(SUMMARIZE) - print_summary() -endif() -dump_options_to_file("${TVM_ALL_OPTIONS}") -if(TVM_BUILD_PYTHON_MODULE AND NOT IS_SUBPROJECT) - foreach(_extra_lib IN LISTS TVM_PACKAGE_EXTRA_LIBS) - if(NOT EXISTS "${_extra_lib}") - message(FATAL_ERROR "TVM_PACKAGE_EXTRA_LIBS entry does not exist: ${_extra_lib}") - endif() - install(FILES "${_extra_lib}" DESTINATION lib) +if(NOT PROJECT_IS_TOP_LEVEL) + foreach(_variable TVM_TIRX_VERSION TVM_TIRX_CMAKE_DIR TVM_TIRX_USE_CUDA TVM_TIRX_USE_ROCM) + set(${_variable} "${${_variable}}" PARENT_SCOPE) endforeach() - install(DIRECTORY licenses/ DESTINATION licenses) - install(FILES LICENSE NOTICE DESTINATION .) + set(TVM_TIRX_BUILD_FOR_ANDROID "${BUILD_FOR_ANDROID}" PARENT_SCOPE) + return() endif() + +# ######### Standalone build targets ########## + +add_custom_target(runtime DEPENDS tvm_runtime tvm_runtime_static ${TVM_RUNTIME_BACKEND_LIBS}) diff --git a/tirx/cmake/modules/CUDA.cmake b/tirx/cmake/modules/CUDA.cmake index 29217b891cff..9093c3fb33ca 100644 --- a/tirx/cmake/modules/CUDA.cmake +++ b/tirx/cmake/modules/CUDA.cmake @@ -16,7 +16,7 @@ # under the License. # CUDA Module -find_cuda(${USE_CUDA} OFF) +find_cuda(${TVM_TIRX_USE_CUDA} OFF) if(CUDA_FOUND) # always set the includedir when CUDA is available @@ -24,9 +24,9 @@ if(CUDA_FOUND) include_directories(SYSTEM ${CUDA_INCLUDE_DIRS}) endif(CUDA_FOUND) -if(USE_CUDA) +if(TVM_TIRX_USE_CUDA) if(NOT CUDA_FOUND) - message(FATAL_ERROR "Cannot find CUDA, USE_CUDA=" ${USE_CUDA}) + message(FATAL_ERROR "Cannot find CUDA, TVM_TIRX_USE_CUDA=" ${TVM_TIRX_USE_CUDA}) endif() message(STATUS "Build with CUDA ${CUDA_VERSION} support") enable_language(CUDA) @@ -43,9 +43,7 @@ if(USE_CUDA) # the `compile_commands.json`, the clangd LSP cannot find the # include files. # - # Furthermore, this override cannot be specified in a user's - # `config.cmake` for TVM, because it must be set after CMake's - # built-in CUDA support. + # Set this after enabling CMake's built-in CUDA support. # # [0] https://github.com/Kitware/CMake/commit/6377a438 set(CMAKE_CUDA_USE_RESPONSE_FILE_FOR_INCLUDES 0) @@ -57,37 +55,34 @@ if(USE_CUDA) message(STATUS "CMAKE_CUDA_ARCHITECTURES not set, using native") set(CMAKE_CUDA_ARCHITECTURES native) endif() -endif(USE_CUDA) +endif(TVM_TIRX_USE_CUDA) -if(USE_CUDA) +if(TVM_TIRX_USE_CUDA) message(STATUS "Build cuda device runtime") # tvm_runtime_cuda links libcuda; without it the .so silently drops its # libcuda.so.1 dependency, so fail configure instead. if(NOT CUDA_CUDA_LIBRARY) - message(FATAL_ERROR "USE_CUDA is on but libcuda was not found. " + message(FATAL_ERROR "TVM_TIRX_USE_CUDA is on but libcuda was not found. " "Set -DCUDA_CUDA_LIBRARY= or make the driver stub discoverable.") endif() - tvm_file_glob(GLOB RUNTIME_CUDA_SRCS src/backend/cuda/runtime/*.cc) + file(GLOB RUNTIME_CUDA_SRCS CONFIGURE_DEPENDS src/backend/cuda/runtime/*.cc) add_library(tvm_runtime_cuda_objs OBJECT ${RUNTIME_CUDA_SRCS}) - target_link_libraries(tvm_runtime_cuda_objs PUBLIC tvm_ffi_header) + target_link_libraries(tvm_runtime_cuda_objs PUBLIC tvm_tirx_header) # These sources compile into tvm_runtime_cuda.dll, so their TVM_RUNTIME_DLL / # TVM_FFI_DLL symbols must be dllexport on MSVC (e.g. GetCudaDeviceCount in # cuda_device_api.cc). Mirror tvm_runtime_objs; a no-op on non-MSVC platforms. target_compile_definitions(tvm_runtime_cuda_objs PRIVATE TVM_RUNTIME_EXPORTS TVM_FFI_EXPORTS) set_target_properties(tvm_runtime_cuda_objs PROPERTIES POSITION_INDEPENDENT_CODE ON) - if(TVM_VISIBILITY_FLAG) - target_compile_options(tvm_runtime_cuda_objs PRIVATE "${TVM_VISIBILITY_FLAG}") - endif() add_library(tvm_runtime_cuda SHARED $) + if(PROJECT_IS_TOP_LEVEL) + install(TARGETS tvm_runtime_cuda DESTINATION ${CMAKE_INSTALL_LIBDIR}) + endif() list(APPEND TVM_RUNTIME_BACKEND_LIBS tvm_runtime_cuda) target_link_libraries(tvm_runtime_cuda PUBLIC tvm_runtime PRIVATE ${CUDA_CUDART_LIBRARY} ${CUDA_CUDA_LIBRARY}) - tvm_configure_target_library(tvm_runtime_cuda RUNTIME_MODULE) + tvm_tirx_configure_target(tvm_runtime_cuda_objs) + tvm_tirx_configure_target(tvm_runtime_cuda) - if(USE_NVTX) - message(STATUS "Build with NVTX support") - target_link_libraries(tvm_runtime_cuda PRIVATE ${CUDA_NVTX_LIBRARY}) - endif() -endif(USE_CUDA) +endif(TVM_TIRX_USE_CUDA) diff --git a/tirx/cmake/modules/LLVM.cmake b/tirx/cmake/modules/LLVM.cmake index 2b8d77391bb2..3d660f864706 100644 --- a/tirx/cmake/modules/LLVM.cmake +++ b/tirx/cmake/modules/LLVM.cmake @@ -25,10 +25,10 @@ add_definitions(-DNDEBUG=1) # TODO(@jroesch, @tkonolige): if we actually use targets we can do this. # target_compile_definitions(tvm_compiler PRIVATE NDEBUG=1) -# Test if ${USE_LLVM} is not an explicit boolean false +# Test if ${TVM_TIRX_USE_LLVM} is not an explicit boolean false # It may be a boolean or a string -if(NOT ${USE_LLVM} MATCHES ${IS_FALSE_PATTERN}) - find_llvm(${USE_LLVM}) +if(NOT ${TVM_TIRX_USE_LLVM} MATCHES ${IS_FALSE_PATTERN}) + find_llvm(${TVM_TIRX_USE_LLVM}) if (${TVM_LLVM_VERSION} LESS 150) message(FATAL_ERROR "LLVM version 15.0 or greater is required.") endif() @@ -52,17 +52,12 @@ if(NOT ${USE_LLVM} MATCHES ${IS_FALSE_PATTERN}) # Set flags that are only needed for LLVM target add_definitions(-DTVM_LLVM_VERSION=${TVM_LLVM_VERSION}) add_definitions(-DTVM_LLVM_HAS_AARCH64_TARGET=${TVM_LLVM_HAS_AARCH64_TARGET}) - tvm_file_glob(GLOB COMPILER_LLVM_SRCS - src/codegen/llvm/*.cc + file(GLOB COMPILER_LLVM_SRCS CONFIGURE_DEPENDS src/codegen/llvm/*.cc src/backend/cuda/codegen/llvm/*.cc src/backend/rocm/codegen/llvm/*.cc src/backend/hexagon/codegen/llvm/*.cc ) list(REMOVE_ITEM COMPILER_LLVM_SRCS "${CMAKE_CURRENT_SOURCE_DIR}/src/codegen/llvm/target_kind.cc") - if(USE_HEXAGON_QHL) - set_property(SOURCE src/backend/hexagon/codegen/llvm/intrin_rule_hexagon.cc - APPEND PROPERTY COMPILE_DEFINITIONS ENABLE_QHL=1) - endif() list(APPEND TVM_LINKER_LIBS ${LLVM_LIBS}) list(APPEND COMPILER_SRCS ${COMPILER_LLVM_SRCS}) if(NOT MSVC) diff --git a/tirx/cmake/modules/Logging.cmake b/tirx/cmake/modules/Logging.cmake index 434e284da31f..df7f2a7b7f34 100644 --- a/tirx/cmake/modules/Logging.cmake +++ b/tirx/cmake/modules/Logging.cmake @@ -19,10 +19,11 @@ include(FindPackageHandleStandardArgs) -if(USE_CUSTOM_LOGGING) +if(TVM_TIRX_USE_CUSTOM_LOGGING) # Set and propagate TVM_LOG_CUSTOMIZE flag is custom logging has been requested target_compile_definitions(tvm_tirx_objs PUBLIC TVM_LOG_CUSTOMIZE=1) target_compile_definitions(tvm_runtime_objs PUBLIC TVM_LOG_CUSTOMIZE=1) target_compile_definitions(tvm_tirx PUBLIC TVM_LOG_CUSTOMIZE=1) target_compile_definitions(tvm_runtime PUBLIC TVM_LOG_CUSTOMIZE=1) + target_compile_definitions(tvm_runtime_static INTERFACE TVM_LOG_CUSTOMIZE=1) endif() diff --git a/tirx/cmake/modules/Metal.cmake b/tirx/cmake/modules/Metal.cmake index 247f4fc5d1bb..c733344c256f 100644 --- a/tirx/cmake/modules/Metal.cmake +++ b/tirx/cmake/modules/Metal.cmake @@ -15,23 +15,24 @@ # specific language governing permissions and limitations # under the License. -if(USE_METAL) +if(TVM_TIRX_USE_METAL) message(STATUS "Build metal device runtime") find_library(METAL_LIB Metal) find_library(FOUNDATION_LIB Foundation) - tvm_file_glob(GLOB RUNTIME_METAL_SRCS src/backend/metal/runtime/*.mm) + file(GLOB RUNTIME_METAL_SRCS CONFIGURE_DEPENDS src/backend/metal/runtime/*.mm) add_library(tvm_runtime_metal_objs OBJECT ${RUNTIME_METAL_SRCS}) - target_link_libraries(tvm_runtime_metal_objs PUBLIC tvm_ffi_header) + target_link_libraries(tvm_runtime_metal_objs PUBLIC tvm_tirx_header) set_target_properties(tvm_runtime_metal_objs PROPERTIES POSITION_INDEPENDENT_CODE ON) - if(TVM_VISIBILITY_FLAG) - target_compile_options(tvm_runtime_metal_objs PRIVATE "${TVM_VISIBILITY_FLAG}") - endif() add_library(tvm_runtime_metal SHARED $) + if(PROJECT_IS_TOP_LEVEL) + install(TARGETS tvm_runtime_metal DESTINATION ${CMAKE_INSTALL_LIBDIR}) + endif() list(APPEND TVM_RUNTIME_BACKEND_LIBS tvm_runtime_metal) target_link_libraries(tvm_runtime_metal PUBLIC tvm_runtime PRIVATE ${METAL_LIB} ${FOUNDATION_LIB}) - tvm_configure_target_library(tvm_runtime_metal RUNTIME_MODULE) -endif(USE_METAL) -# When USE_METAL=OFF the codegen-side fallback in + tvm_tirx_configure_target(tvm_runtime_metal_objs) + tvm_tirx_configure_target(tvm_runtime_metal) +endif(TVM_TIRX_USE_METAL) +# When TVM_TIRX_USE_METAL=OFF the codegen-side fallback in # src/backend/metal/codegen/metal_fallback_module.cc handles construction; no opt # stub is needed (it is always compiled via CODEGEN_SRCS in CMakeLists.txt). diff --git a/tirx/cmake/modules/OpenCL.cmake b/tirx/cmake/modules/OpenCL.cmake index e24ec2a24ff2..168977b508d9 100644 --- a/tirx/cmake/modules/OpenCL.cmake +++ b/tirx/cmake/modules/OpenCL.cmake @@ -15,14 +15,13 @@ # specific language governing permissions and limitations # under the License. -if(USE_OPENCL) - tvm_file_glob(GLOB RUNTIME_OPENCL_SRCS src/backend/opencl/runtime/*.cc) +if(TVM_TIRX_USE_OPENCL) + file(GLOB RUNTIME_OPENCL_SRCS CONFIGURE_DEPENDS src/backend/opencl/runtime/*.cc) set(_opencl_libs "") - if(${USE_OPENCL} MATCHES ${IS_TRUE_PATTERN}) + if(${TVM_TIRX_USE_OPENCL} MATCHES ${IS_TRUE_PATTERN}) message(STATUS "Enabled runtime search for OpenCL library location") - tvm_file_glob(GLOB OPENCL_WRAPPER_SRCS - "src/backend/opencl/runtime/opencl_wrapper/opencl_wrapper.cc" + file(GLOB OPENCL_WRAPPER_SRCS CONFIGURE_DEPENDS "src/backend/opencl/runtime/opencl_wrapper/opencl_wrapper.cc" ) list(APPEND RUNTIME_OPENCL_SRCS ${OPENCL_WRAPPER_SRCS}) find_path(TVM_OPENCL_HEADERS_DIR CL/cl.h @@ -30,7 +29,7 @@ if(USE_OPENCL) PATH_SUFFIXES include REQUIRED) set(_opencl_include_dirs "${TVM_OPENCL_HEADERS_DIR}") else() - find_opencl(${USE_OPENCL}) + find_opencl(${TVM_TIRX_USE_OPENCL}) if(NOT OpenCL_FOUND) message(FATAL_ERROR "Error! Cannot find specified OpenCL library") endif() @@ -42,7 +41,7 @@ if(USE_OPENCL) message(STATUS "Build opencl device runtime") add_library(tvm_runtime_opencl_objs OBJECT ${RUNTIME_OPENCL_SRCS}) - target_link_libraries(tvm_runtime_opencl_objs PUBLIC tvm_ffi_header) + target_link_libraries(tvm_runtime_opencl_objs PUBLIC tvm_tirx_header) target_include_directories(tvm_runtime_opencl_objs SYSTEM PRIVATE ${_opencl_include_dirs}) if(WIN32) set(_opencl_export "__declspec(dllexport)") @@ -53,23 +52,24 @@ if(USE_OPENCL) endif() target_compile_definitions(tvm_runtime_opencl_objs PRIVATE "TVM_OPENCL_DLL=${_opencl_export}") set_target_properties(tvm_runtime_opencl_objs PROPERTIES POSITION_INDEPENDENT_CODE ON) - if(TVM_VISIBILITY_FLAG) - target_compile_options(tvm_runtime_opencl_objs PRIVATE "${TVM_VISIBILITY_FLAG}") - endif() add_library(tvm_runtime_opencl SHARED $) + if(PROJECT_IS_TOP_LEVEL) + install(TARGETS tvm_runtime_opencl DESTINATION ${CMAKE_INSTALL_LIBDIR}) + endif() list(APPEND TVM_RUNTIME_BACKEND_LIBS tvm_runtime_opencl) target_link_libraries(tvm_runtime_opencl PUBLIC tvm_runtime PRIVATE ${_opencl_libs}) target_compile_definitions(tvm_runtime_opencl INTERFACE "TVM_OPENCL_DLL=${_opencl_import}") target_include_directories(tvm_runtime_opencl SYSTEM INTERFACE "$") - tvm_configure_target_library(tvm_runtime_opencl RUNTIME_MODULE) + tvm_tirx_configure_target(tvm_runtime_opencl_objs) + tvm_tirx_configure_target(tvm_runtime_opencl) - if(USE_OPENCL_ENABLE_HOST_PTR) + if(TVM_TIRX_USE_OPENCL_ENABLE_HOST_PTR) target_compile_definitions(tvm_runtime_opencl_objs PRIVATE OPENCL_ENABLE_HOST_PTR) - endif(USE_OPENCL_ENABLE_HOST_PTR) - if(USE_OPENCL_EXTN_QCOM) + endif(TVM_TIRX_USE_OPENCL_ENABLE_HOST_PTR) + if(TVM_TIRX_USE_OPENCL_EXTN_QCOM) add_definitions(-DUSE_OPENCL_EXTN_QCOM) find_path(ocl_header cl.h HINTS ${OpenCL_INCLUDE_DIRS} PATH_SUFFIXES CL) set(OCL_VERSION_HEADER "${ocl_header}/cl.h") @@ -79,5 +79,5 @@ if(USE_OPENCL) add_definitions(-DCL_TARGET_OPENCL_VERSION=${CMAKE_MATCH_1}) message(STATUS "Set OpenCL Target version to " ${CMAKE_MATCH_1}) endif() - endif(USE_OPENCL_EXTN_QCOM) -endif(USE_OPENCL) + endif(TVM_TIRX_USE_OPENCL_EXTN_QCOM) +endif(TVM_TIRX_USE_OPENCL) diff --git a/tirx/cmake/modules/OpenMP.cmake b/tirx/cmake/modules/OpenMP.cmake index 5dd9be508342..3eb4d315908b 100644 --- a/tirx/cmake/modules/OpenMP.cmake +++ b/tirx/cmake/modules/OpenMP.cmake @@ -16,7 +16,7 @@ # under the License. # OpenMP Module -if(USE_OPENMP STREQUAL "gnu") +if(TVM_TIRX_USE_OPENMP STREQUAL "gnu") find_package(OpenMP) if(OPENMP_FOUND) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}") @@ -27,7 +27,7 @@ if(USE_OPENMP STREQUAL "gnu") add_definitions(-DTVM_THREADPOOL_USE_OPENMP=0) message(WARNING "OpenMP cannot be found, use TVM threadpool instead.") endif() -elseif(USE_OPENMP STREQUAL "intel") +elseif(TVM_TIRX_USE_OPENMP STREQUAL "intel") find_package(OpenMP) if(OPENMP_FOUND) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}") diff --git a/tirx/cmake/modules/ROCM.cmake b/tirx/cmake/modules/ROCM.cmake index d910674dc2ce..a1be778dd81a 100644 --- a/tirx/cmake/modules/ROCM.cmake +++ b/tirx/cmake/modules/ROCM.cmake @@ -16,7 +16,11 @@ # under the License. # ROCM Module -find_rocm(${USE_ROCM}) +if("${TVM_TIRX_USE_ROCM}" MATCHES ${IS_TRUE_PATTERN} AND IS_DIRECTORY "${TVM_TIRX_ROCM_PATH}") + find_rocm("${TVM_TIRX_ROCM_PATH}") +else() + find_rocm(${TVM_TIRX_USE_ROCM}) +endif() if(ROCM_FOUND) # always set the includedir @@ -26,13 +30,13 @@ if(ROCM_FOUND) add_definitions(-D__HIP_PLATFORM_AMD__=1) endif(ROCM_FOUND) -if(USE_ROCM) +if(TVM_TIRX_USE_ROCM) if(NOT ROCM_FOUND) - message(FATAL_ERROR "Cannot find ROCM, USE_ROCM=" ${USE_ROCM}) + message(FATAL_ERROR "Cannot find ROCM, TVM_TIRX_USE_ROCM=" ${TVM_TIRX_USE_ROCM}) endif() message(STATUS "Build rocm device runtime") - tvm_file_glob(GLOB RUNTIME_ROCM_SRCS src/backend/rocm/runtime/*.cc) + file(GLOB RUNTIME_ROCM_SRCS CONFIGURE_DEPENDS src/backend/rocm/runtime/*.cc) set(_rocm_libs ${ROCM_HIPHCC_LIBRARY}) if(ROCM_HSA_LIBRARY) @@ -40,13 +44,14 @@ if(USE_ROCM) endif() add_library(tvm_runtime_rocm_objs OBJECT ${RUNTIME_ROCM_SRCS}) - target_link_libraries(tvm_runtime_rocm_objs PUBLIC tvm_ffi_header) + target_link_libraries(tvm_runtime_rocm_objs PUBLIC tvm_tirx_header) set_target_properties(tvm_runtime_rocm_objs PROPERTIES POSITION_INDEPENDENT_CODE ON) - if(TVM_VISIBILITY_FLAG) - target_compile_options(tvm_runtime_rocm_objs PRIVATE "${TVM_VISIBILITY_FLAG}") - endif() add_library(tvm_runtime_rocm SHARED $) + if(PROJECT_IS_TOP_LEVEL) + install(TARGETS tvm_runtime_rocm DESTINATION ${CMAKE_INSTALL_LIBDIR}) + endif() list(APPEND TVM_RUNTIME_BACKEND_LIBS tvm_runtime_rocm) target_link_libraries(tvm_runtime_rocm PUBLIC tvm_runtime PRIVATE ${_rocm_libs}) - tvm_configure_target_library(tvm_runtime_rocm RUNTIME_MODULE) -endif(USE_ROCM) + tvm_tirx_configure_target(tvm_runtime_rocm_objs) + tvm_tirx_configure_target(tvm_runtime_rocm) +endif(TVM_TIRX_USE_ROCM) diff --git a/tirx/cmake/modules/Vulkan.cmake b/tirx/cmake/modules/Vulkan.cmake index 34d932b6907d..de6bbfec333c 100644 --- a/tirx/cmake/modules/Vulkan.cmake +++ b/tirx/cmake/modules/Vulkan.cmake @@ -16,20 +16,19 @@ # under the License. # Be compatible with older version of CMake -find_vulkan(${USE_VULKAN} ${USE_KHRONOS_SPIRV}) +find_vulkan(${TVM_TIRX_USE_VULKAN} ${TVM_TIRX_USE_KHRONOS_SPIRV}) -if(USE_VULKAN) +if(TVM_TIRX_USE_VULKAN) if(NOT Vulkan_FOUND) - message(FATAL_ERROR "Cannot find Vulkan, USE_VULKAN=" ${USE_VULKAN}) + message(FATAL_ERROR "Cannot find Vulkan, TVM_TIRX_USE_VULKAN=" ${TVM_TIRX_USE_VULKAN}) endif() - if(USE_SPIRV_KHR_INTEGER_DOT_PRODUCT) + if(TVM_TIRX_USE_SPIRV_KHR_INTEGER_DOT_PRODUCT) add_definitions(-DTVM_SPIRV_KHR_INTEGER_DOT_PRODUCT=1) message(STATUS "Enable SPIRV_KHR_INTEGER_DOT_PRODUCT") endif() include_directories(SYSTEM ${Vulkan_INCLUDE_DIRS}) message(STATUS "Build with Vulkan support") - tvm_file_glob(GLOB COMPILER_VULKAN_SRCS - src/backend/vulkan/codegen/build_vulkan.cc + file(GLOB COMPILER_VULKAN_SRCS CONFIGURE_DEPENDS src/backend/vulkan/codegen/build_vulkan.cc src/backend/vulkan/codegen/codegen_spirv.cc src/backend/vulkan/codegen/intrin_rule_spirv.cc src/backend/vulkan/codegen/ir_builder.cc @@ -39,21 +38,22 @@ if(USE_VULKAN) list(APPEND COMPILER_SRCS ${COMPILER_VULKAN_SRCS}) list(APPEND TVM_LINKER_LIBS ${Vulkan_SPIRV_TOOLS_LIBRARY}) add_definitions(-DTVM_ENABLE_SPIRV=1) -endif(USE_VULKAN) +endif(TVM_TIRX_USE_VULKAN) -if(USE_VULKAN) +if(TVM_TIRX_USE_VULKAN) message(STATUS "Build vulkan device runtime") - tvm_file_glob(GLOB RUNTIME_VULKAN_SRCS src/backend/vulkan/runtime/*.cc) + file(GLOB RUNTIME_VULKAN_SRCS CONFIGURE_DEPENDS src/backend/vulkan/runtime/*.cc) add_library(tvm_runtime_vulkan_objs OBJECT ${RUNTIME_VULKAN_SRCS}) - target_link_libraries(tvm_runtime_vulkan_objs PUBLIC tvm_ffi_header) + target_link_libraries(tvm_runtime_vulkan_objs PUBLIC tvm_tirx_header) set_target_properties(tvm_runtime_vulkan_objs PROPERTIES POSITION_INDEPENDENT_CODE ON) - if(TVM_VISIBILITY_FLAG) - target_compile_options(tvm_runtime_vulkan_objs PRIVATE "${TVM_VISIBILITY_FLAG}") - endif() add_library(tvm_runtime_vulkan SHARED $) + if(PROJECT_IS_TOP_LEVEL) + install(TARGETS tvm_runtime_vulkan DESTINATION ${CMAKE_INSTALL_LIBDIR}) + endif() list(APPEND TVM_RUNTIME_BACKEND_LIBS tvm_runtime_vulkan) target_link_libraries(tvm_runtime_vulkan PUBLIC tvm_runtime PRIVATE ${Vulkan_LIBRARY}) - tvm_configure_target_library(tvm_runtime_vulkan RUNTIME_MODULE) -endif(USE_VULKAN) + tvm_tirx_configure_target(tvm_runtime_vulkan_objs) + tvm_tirx_configure_target(tvm_runtime_vulkan) +endif(TVM_TIRX_USE_VULKAN) diff --git a/tirx/cmake/modules/contrib/Z3.cmake b/tirx/cmake/modules/contrib/Z3.cmake index 6a33b4172393..4543e46909ec 100644 --- a/tirx/cmake/modules/contrib/Z3.cmake +++ b/tirx/cmake/modules/contrib/Z3.cmake @@ -18,12 +18,12 @@ # src/sym/z3_prover.cc is always part of COMPILER_SRCS (picked up by the # src/sym/*.cc glob). It compiles a conservative stub by default and switches # to the real Z3 implementation only when the TVM_USE_Z3 macro is defined below. -if(${USE_Z3} MATCHES ${IS_FALSE_PATTERN}) +if(${TVM_TIRX_USE_Z3} MATCHES ${IS_FALSE_PATTERN}) return() endif() set(TVM_Z3_REQUIRED TRUE) -if("${USE_Z3}" MATCHES "^[Aa][Uu][Tt][Oo]$") +if("${TVM_TIRX_USE_Z3}" MATCHES "^[Aa][Uu][Tt][Oo]$") set(TVM_Z3_REQUIRED FALSE) endif() @@ -77,14 +77,14 @@ elseif(Z3_FOUND OR (Z3_INCLUDE_DIR AND Z3_LIBRARY)) set(Z3_LIBRARY ${Z3_LIBRARIES}) endif() if(NOT Z3_INCLUDE_DIR OR NOT Z3_LIBRARY) - message(FATAL_ERROR "USE_Z3 is ON, but Z3 include directory or library was not found.") + message(FATAL_ERROR "TVM_TIRX_USE_Z3 is ON, but Z3 include directory or library was not found.") endif() include_directories(SYSTEM ${Z3_INCLUDE_DIR}) list(APPEND TVM_LINKER_LIBS ${Z3_LIBRARY}) else() if(TVM_Z3_REQUIRED) message(FATAL_ERROR - "USE_Z3 is ON, but Z3 was not found. Install the static Z3 development " + "TVM_TIRX_USE_Z3 is ON, but Z3 was not found. Install the static Z3 development " "package with `pip install 'mlc-z3-static>=4.16.0'`, or point " "Z3_DIR/CMAKE_PREFIX_PATH at a Z3 installation.") endif() @@ -96,7 +96,7 @@ endif() add_compile_definitions(TVM_USE_Z3) message(STATUS "Build with Z3 SMT solver support") -if(TVM_BUILD_PYTHON_MODULE AND Z3_TARGET) +if(TVM_TIRX_BUILD_PYTHON_MODULE AND Z3_TARGET) get_target_property(_z3_target_type ${Z3_TARGET} TYPE) if(_z3_target_type STREQUAL "STATIC_LIBRARY" AND Python3_EXECUTABLE) execute_process(COMMAND "${Python3_EXECUTABLE}" -c diff --git a/tirx/cmake/tvm_tirx-config.cmake b/tirx/cmake/tvm_tirx-config.cmake new file mode 100644 index 000000000000..4ef057cd43f1 --- /dev/null +++ b/tirx/cmake/tvm_tirx-config.cmake @@ -0,0 +1,56 @@ +# 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(CMakeFindDependencyMacro) +find_dependency(Python 3.10 COMPONENTS Interpreter) + +# Query this installation's helper, independently of any other Python wheel. +# Passing the exact library directory also supports lib64 and multiarch layouts. +get_filename_component(_tvm_tirx_libdir "${CMAKE_CURRENT_LIST_DIR}/../.." ABSOLUTE) +execute_process( + COMMAND "${Python_EXECUTABLE}" "${CMAKE_CURRENT_LIST_DIR}/config.py" + --library-dir "${_tvm_tirx_libdir}" --platform "${CMAKE_SYSTEM_NAME}" --json + OUTPUT_VARIABLE _tvm_tirx_config OUTPUT_STRIP_TRAILING_WHITESPACE + COMMAND_ERROR_IS_FATAL ANY) +string(JSON TVM_TIRX_VERSION GET "${_tvm_tirx_config}" version) +string(JSON tvm_tirx_INCLUDE_DIR GET "${_tvm_tirx_config}" include_dir) +set(tvm_tirx_VERSION "${TVM_TIRX_VERSION}") +set(TVM_TIRX_CMAKE_DIR "${CMAKE_CURRENT_LIST_DIR}") + +if(NOT TARGET tvm_tirx::header) + add_library(tvm_tirx::header INTERFACE IMPORTED) + target_compile_features(tvm_tirx::header INTERFACE cxx_std_20) + target_compile_definitions(tvm_tirx::header INTERFACE TVM_INDEX_DEFAULT_I64=1) + target_include_directories(tvm_tirx::header INTERFACE "${tvm_tirx_INCLUDE_DIR}") +endif() + +foreach(_library tvm_runtime tvm_tirx) + if(TARGET ${_library}::shared) + continue() + endif() + string(JSON _location GET "${_tvm_tirx_config}" libraries ${_library} location) + add_library(${_library}::shared SHARED IMPORTED) + set_target_properties(${_library}::shared PROPERTIES IMPORTED_LOCATION "${_location}") + if(WIN32) + string(JSON _implib GET "${_tvm_tirx_config}" libraries ${_library} implib) + set_target_properties(${_library}::shared PROPERTIES IMPORTED_IMPLIB "${_implib}") + endif() + target_link_libraries(${_library}::shared INTERFACE tvm_tirx::header) + if(_library STREQUAL "tvm_tirx") + target_link_libraries(tvm_tirx::shared INTERFACE tvm_runtime::shared) + endif() +endforeach() diff --git a/tirx/cmake/tvm_tirxConfig.cmake.in b/tirx/cmake/tvm_tirxConfig.cmake.in deleted file mode 100644 index abce76e80a72..000000000000 --- a/tirx/cmake/tvm_tirxConfig.cmake.in +++ /dev/null @@ -1,37 +0,0 @@ -# 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. - -@PACKAGE_INIT@ -include(CMakeFindDependencyMacro) -find_dependency(Python COMPONENTS Interpreter) -if(NOT TARGET tvm_ffi::shared) - if(NOT tvm_ffi_DIR) - execute_process(COMMAND "${Python_EXECUTABLE}" -m tvm_ffi.config --cmakedir - OUTPUT_VARIABLE tvm_ffi_DIR OUTPUT_STRIP_TRAILING_WHITESPACE - COMMAND_ERROR_IS_FATAL ANY) - endif() - find_dependency(tvm_ffi CONFIG) -endif() -if("@USE_THREADS@") - find_dependency(Threads) -endif() -include("${CMAKE_CURRENT_LIST_DIR}/tvm_tirxTargets.cmake") -set(TVM_TIRX_VERSION "@TVM_TIRX_VERSION@") -set(TVM_TIRX_CMAKE_DIR "${CMAKE_CURRENT_LIST_DIR}") -set(TVM_TIRX_USE_CUDA "@USE_CUDA@") -set(TVM_TIRX_USE_ROCM "@USE_ROCM@") -set(TVM_TIRX_BUILD_FOR_ANDROID "@BUILD_FOR_ANDROID@") diff --git a/tirx/cmake/utils/Library.cmake b/tirx/cmake/utils/Library.cmake index 4b4d9f660590..ee26447a800e 100644 --- a/tirx/cmake/utils/Library.cmake +++ b/tirx/cmake/utils/Library.cmake @@ -15,72 +15,11 @@ # specific language governing permissions and limitations # under the License. -# Helpers for configuring library targets. - -# Hide symbols from static archives linked into a shared library. This prevents -# downstream libraries from accidentally binding to private archive symbols such -# as those from libstdc++_nonshared.a. -function(tvm_hide_static_linked_lib_symbols target_name) - if(HIDE_PRIVATE_SYMBOLS - AND CMAKE_SYSTEM_NAME MATCHES "Linux|Android|FreeBSD|NetBSD|OpenBSD" - AND CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang" - ) - target_link_options(${target_name} PRIVATE "-Wl,--exclude-libs,ALL") - endif() -endfunction() - -####################################################### -# tvm_configure_target_library(target_name [RUNTIME_MODULE]) -# -# Configure a TVM library target. The target always gets a relative rpath -# ($ORIGIN / @loader_path) so that sibling shared libraries in the same -# directory resolve each other regardless of the install location (e.g. inside a -# Python wheel). -# -# For shared libraries and runtime modules, hide symbols that come from linked -# static archives so downstream libraries cannot accidentally bind to private -# archive symbols such as those from libstdc++_nonshared.a. -# -# With the RUNTIME_MODULE option -- used for the optional runtime backend -# libraries (tvm_runtime_cuda, tvm_runtime_vulkan, ...) -- the target is also -# emitted into the build "lib" directory and installed; when building the Python -# module it is additionally installed into the package "lib" directory. Targets -# that manage their own output directory / install rules (tvm_compiler, -# tvm_runtime, ...) omit the option and take only the rpath. -# -# No-op if the target does not exist. -function(tvm_configure_target_library target_name) - if(NOT TARGET ${target_name}) - return() - endif() - cmake_parse_arguments(ARG "RUNTIME_MODULE" "" "" ${ARGN}) - +# Keep shared libraries relocatable within their installation directory. +function(tvm_tirx_configure_target target) if(APPLE) - set_target_properties(${target_name} PROPERTIES - BUILD_RPATH "@loader_path" - INSTALL_RPATH "@loader_path" - ) + set_target_properties(${target} PROPERTIES BUILD_RPATH "@loader_path" INSTALL_RPATH "@loader_path") elseif(UNIX) - set_target_properties(${target_name} PROPERTIES - BUILD_RPATH "\$ORIGIN" - INSTALL_RPATH "\$ORIGIN" - ) - endif() - - get_target_property(TVM_TARGET_TYPE ${target_name} TYPE) - if(TVM_TARGET_TYPE STREQUAL "SHARED_LIBRARY" OR TVM_TARGET_TYPE STREQUAL "MODULE_LIBRARY") - tvm_hide_static_linked_lib_symbols(${target_name}) - endif() - - if(ARG_RUNTIME_MODULE) - set_target_properties(${target_name} PROPERTIES - LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" - RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" - ARCHIVE_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib" - ) - install(TARGETS ${target_name} DESTINATION lib${LIB_SUFFIX}) - if(TVM_BUILD_PYTHON_MODULE) - install(TARGETS ${target_name} DESTINATION "lib") - endif() + set_target_properties(${target} PROPERTIES BUILD_RPATH "\$ORIGIN" INSTALL_RPATH "\$ORIGIN") endif() endfunction() diff --git a/tirx/pyproject.toml b/tirx/pyproject.toml index db8d5c071c3e..e0fba2e008bc 100644 --- a/tirx/pyproject.toml +++ b/tirx/pyproject.toml @@ -16,12 +16,8 @@ # under the License. [build-system] -# mlc-z3-static ships the PIC static libz3 + headers consumed by USE_Z3=ON. -requires = [ - "scikit-build-core>=0.11", - "mlc-z3-static>=4.16.0", - "apache-tvm-ffi>=0.1.15rc0", -] +# mlc-z3-static ships the PIC static libz3 + headers consumed by TVM_TIRX_USE_Z3=ON. +requires = ["scikit-build-core>=0.11", "mlc-z3-static>=4.16.0"] build-backend = "scikit_build_core.build" [project] @@ -109,6 +105,7 @@ sdist.include = [ "/NOTICE", "/tests/**", "/3rdparty/compiler-rt/builtin_fp16.h", + "/3rdparty/tvm-ffi/**", "!**/.git", "!**/.git/**", "!**/__pycache__/**", @@ -123,11 +120,11 @@ sdist.exclude = ["**"] logging.level = "INFO" [tool.scikit-build.cmake.define] -TVM_BUILD_PYTHON_MODULE = "ON" -USE_CUDA = "OFF" +TVM_TIRX_BUILD_PYTHON_MODULE = "ON" +TVM_TIRX_USE_CUDA = "OFF" # Statically link Z3 from the mlc-z3-static build dependency by default. -USE_Z3 = "AUTO" -BUILD_TESTING = "OFF" +TVM_TIRX_USE_Z3 = "AUTO" +TVM_TIRX_BUILD_TESTS = "OFF" [tool.pytest.ini_options] testpaths = ["tests"] diff --git a/tirx/python/tvm_tirx/__init__.py b/tirx/python/tvm_tirx/__init__.py index efbdb05a5ece..a66f56d844ad 100644 --- a/tirx/python/tvm_tirx/__init__.py +++ b/tirx/python/tvm_tirx/__init__.py @@ -18,108 +18,122 @@ # pylint: disable=redefined-builtin, wildcard-import """TVM: Open Deep Learning Compiler Stack.""" -import multiprocessing -import os import sys +from typing import TYPE_CHECKING -# ffi module must load first -from tvm_ffi import register_object, register_global_func, get_global_func -# top-level alias -from .libinfo import __version__ -from .base import _RUNTIME_ONLY +def _is_config_mode(): + """Keep the configuration CLI independent of native library loading.""" + argv = getattr(sys, "orig_argv", ()) + return any( + arg == "-m" and argv[index + 1 : index + 2] == ["tvm_tirx.config"] + for index, arg in enumerate(argv) + ) -# tvm.runtime -from .runtime import Object -from .runtime._tensor import device, cpu, cuda, opencl, vulkan, metal -from .runtime._tensor import vpi, rocm, ext_dev, hexagon -from .runtime import DataType, DataTypeCode -# tvm.error -from . import error +if TYPE_CHECKING or not _is_config_mode(): + import multiprocessing + import os -# tvm.ir -from .ir import IRModule, transform -from . import ir + # ffi module must load first + from tvm_ffi import register_object, register_global_func, get_global_func -# tvm.script — must be imported before any dialect package so that -# tvm.script.register_dialect is reachable when dialect __init__.py files run. -from . import script + # top-level alias + from .libinfo import __version__ + from .base import _RUNTIME_ONLY -# tvm.tirx — registers itself via tvm.script.register_dialect in its __init__ -from . import tirx -from .ir import prim + # tvm.runtime + from .runtime import Object + from .runtime._tensor import device, cpu, cuda, opencl, vulkan, metal + from .runtime._tensor import vpi, rocm, ext_dev, hexagon + from .runtime import DataType, DataTypeCode -# tvm.backend — owns backend Python load hooks -from . import backend + # tvm.error + from . import error -# tvm.target -from . import target -from .target import device_from_target + # tvm.ir + from .ir import IRModule, transform + from . import ir -# Shared code generation helpers -from . import codegen + # tvm.script — must be imported before any dialect package so that + # tvm.script.register_dialect is reachable when dialect __init__.py files run. + from . import script -# others -from . import sym + # tvm.tirx — registers itself via tvm.script.register_dialect in its __init__ + from . import tirx + from .ir import prim -# support infra -from . import support + # tvm.backend — owns backend Python load hooks + from . import backend -# Side-effect imports: register CUDA/ROCm FFI callbacks at TVM startup -from .support import rocm as _rocm, nvcc as _nvcc + # tvm.target + from . import target + from .target import device_from_target -# NOTE: This file should be python2 compatible so we can -# raise proper error message when user run the package using -# an older version of the python + # Shared code generation helpers + from . import codegen + # others + from . import sym -def _should_print_backtrace(): - in_pytest = "PYTEST_CURRENT_TEST" in os.environ - tvm_backtrace = os.environ.get("TVM_BACKTRACE", "0") + # support infra + from . import support - try: - tvm_backtrace = bool(int(tvm_backtrace)) - except ValueError: - raise ValueError(f"invalid value for TVM_BACKTRACE {tvm_backtrace}, please set to 0 or 1.") + # Side-effect imports: register CUDA/ROCm FFI callbacks at TVM startup + from .support import rocm as _rocm, nvcc as _nvcc + + # NOTE: This file should be python2 compatible so we can + # raise proper error message when user run the package using + # an older version of the python - return in_pytest or tvm_backtrace + def _should_print_backtrace(): + in_pytest = "PYTEST_CURRENT_TEST" in os.environ + tvm_backtrace = os.environ.get("TVM_BACKTRACE", "0") + try: + tvm_backtrace = bool(int(tvm_backtrace)) + except ValueError: + raise ValueError( + f"invalid value for TVM_BACKTRACE {tvm_backtrace}, please set to 0 or 1." + ) -def tvm_wrap_excepthook(exception_hook): - """Wrap given excepthook with TVM additional work.""" + return in_pytest or tvm_backtrace - def wrapper(exctype, value, trbk): - """Clean subprocesses when TVM is interrupted.""" - if exctype is error.DiagnosticError and not _should_print_backtrace(): - # TODO(@jroesch): consider moving to C++? - print("note: run with `TVM_BACKTRACE=1` environment variable to display a backtrace.") - else: - exception_hook(exctype, value, trbk) + def tvm_wrap_excepthook(exception_hook): + """Wrap given excepthook with TVM additional work.""" - if hasattr(multiprocessing, "active_children"): - # pylint: disable=not-callable - for p in multiprocessing.active_children(): - p.terminate() + def wrapper(exctype, value, trbk): + """Clean subprocesses when TVM is interrupted.""" + if exctype is error.DiagnosticError and not _should_print_backtrace(): + # TODO(@jroesch): consider moving to C++? + print( + "note: run with `TVM_BACKTRACE=1` environment variable to display a backtrace." + ) + else: + exception_hook(exctype, value, trbk) - return wrapper + if hasattr(multiprocessing, "active_children"): + # pylint: disable=not-callable + for p in multiprocessing.active_children(): + p.terminate() + return wrapper -sys.excepthook = tvm_wrap_excepthook(sys.excepthook) + sys.excepthook = tvm_wrap_excepthook(sys.excepthook) -# Autoload loads built-in and out-of-tree backends. Out-of-tree extensions opt -# into being loaded automatically at ``import tvm`` time by declaring an entry -# point in the ``tvm.backends`` group: -# [project.entry-points."tvm.backends"] tvm_foo = "tvm_foo:_autoload". -# Autoload can be disabled via ``TVM_DEVICE_BACKEND_AUTOLOAD=0``. -from .backend._autoload_backends import _autoload_backends + # Autoload loads built-in and out-of-tree backends. Out-of-tree extensions opt + # into being loaded automatically at ``import tvm`` time by declaring an entry + # point in the ``tvm.backends`` group: + # [project.entry-points."tvm.backends"] tvm_foo = "tvm_foo:_autoload". + # Autoload can be disabled via ``TVM_DEVICE_BACKEND_AUTOLOAD=0``. + from .backend._autoload_backends import _autoload_backends -_autoload_backends() + _autoload_backends() -# Load the optional ORC JIT extension for its global FFI registrations. LLVM -# function lookup reports a targeted installation error if this import fails; -# importing TVM itself remains silent for AOT-only workflows. -try: - import tvm_ffi_orcjit as _tvm_ffi_orcjit -except Exception: # pylint: disable=broad-exception-caught - pass + # Load the optional ORC JIT extension for its global FFI registrations. LLVM + # function lookup reports a targeted installation error if this import fails; + # importing TVM itself remains silent for AOT-only workflows. + try: + import tvm_ffi_orcjit as _tvm_ffi_orcjit + except Exception: # pylint: disable=broad-exception-caught + pass diff --git a/tirx/python/tvm_tirx/config.py b/tirx/python/tvm_tirx/config.py new file mode 100644 index 000000000000..de55764a3755 --- /dev/null +++ b/tirx/python/tvm_tirx/config.py @@ -0,0 +1,150 @@ +# 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. +"""Discover TIRx-owned headers and libraries without loading native code. + +The same file is installed beside the CMake package config. CMake invokes it +there with an exact library directory, independently of Python package search. +""" + +import argparse +import json +import sys +from pathlib import Path + +__version__ = "0.1.0a0" + +_LIBRARY_NAMES = ("tvm_tirx", "tvm_runtime") + + +def _find_prefix(start: Path) -> Path: + for candidate in (start, *start.parents): + if (candidate / "include/tvm/runtime/base.h").is_file(): + return candidate + raise RuntimeError(f"Cannot find TIRx public headers above {start}") + + +def _paths(prefix: Path | None, library_dir: Path | None) -> tuple[Path, Path]: + if library_dir is not None: + library_dir = library_dir.resolve() + prefix = prefix.resolve() if prefix is not None else _find_prefix(library_dir) + else: + prefix = prefix.resolve() if prefix is not None else _find_prefix(Path(__file__).parent) + library_dir = prefix / "lib" + for candidate in (prefix / "lib", prefix / "lib64", prefix / "build/lib"): + if all( + any( + (candidate / filename).is_file() + for filename in (f"lib{name}.so", f"lib{name}.dylib", f"{name}.dll") + ) + for name in _LIBRARY_NAMES + ): + library_dir = candidate + break + if not (prefix / "include/tvm/runtime/base.h").is_file(): + raise RuntimeError(f"Cannot find TIRx public headers in {prefix / 'include'}") + return prefix, library_dir + + +def _libraries(library_dir: Path, platform: str) -> dict[str, dict[str, str]]: + libraries = {} + windows = platform in ("Windows", "win32") + apple = platform in ("Darwin", "iOS", "darwin") + for name in _LIBRARY_NAMES: + if windows: + filenames = (f"{name}.dll",) + elif apple: + filenames = (f"lib{name}.dylib", f"lib{name}.so") + else: + filenames = (f"lib{name}.so",) + location = next((library_dir / f for f in filenames if (library_dir / f).is_file()), None) + if location is None: + raise RuntimeError(f"Cannot find TIRx library {name} in {library_dir}") + implib = "" + if windows: + import_names = (f"{name}.lib", f"lib{name}.dll.a") + import_library = next( + (library_dir / f for f in import_names if (library_dir / f).is_file()), None + ) + if import_library is None: + raise RuntimeError(f"Missing import library for {location}") + implib = str(import_library) + libraries[name] = {"location": str(location), "implib": implib} + return libraries + + +def configuration( + prefix: Path | None = None, + library_dir: Path | None = None, + platform: str = sys.platform, +) -> dict: + """Return the public include and core shared-library paths for one installation.""" + prefix, library_dir = _paths(prefix, library_dir) + return { + "version": __version__, + "include_dir": str(prefix / "include"), + "library_dir": str(library_dir), + "cmake_dir": str(library_dir / "cmake/tvm_tirx"), + "libraries": _libraries(library_dir, platform), + } + + +def __main__() -> None: + """Print installed include, library, or CMake package paths.""" + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--prefix", type=Path, help="Use this exact installation prefix") + parser.add_argument("--library-dir", type=Path, help="Use this exact native library directory") + parser.add_argument("--platform", default=sys.platform, help=argparse.SUPPRESS) + outputs = parser.add_mutually_exclusive_group() + outputs.add_argument( + "--includedir", action="store_true", help="Print the public include directory" + ) + outputs.add_argument("--libdir", action="store_true", help="Print the native library directory") + outputs.add_argument( + "--libfiles", action="store_true", help="Print library files, one per line" + ) + outputs.add_argument( + "--cmakedir", action="store_true", help="Print the CMake package directory" + ) + outputs.add_argument( + "--json", action="store_true", help="Print all consumer configuration as JSON" + ) + args = parser.parse_args() + if not any((args.includedir, args.libdir, args.libfiles, args.cmakedir, args.json)): + parser.print_help() + return + try: + if args.json or args.libfiles: + result = configuration(args.prefix, args.library_dir, args.platform) + if args.json: + print(json.dumps(result)) + else: + for library in result["libraries"].values(): + print(library["implib"] or library["location"]) + else: + prefix, library_dir = _paths(args.prefix, args.library_dir) + if args.includedir: + print(prefix / "include") + elif args.libdir: + print(library_dir) + else: + print(library_dir / "cmake/tvm_tirx") + except RuntimeError as error: + parser.exit(1, f"{error}\n") + + +if __name__ == "__main__": + __main__() diff --git a/tirx/tests/cpp/CMakeLists.txt b/tirx/tests/cpp/CMakeLists.txt index 255ea70ed08e..f1fae95c491c 100644 --- a/tirx/tests/cpp/CMakeLists.txt +++ b/tirx/tests/cpp/CMakeLists.txt @@ -15,95 +15,49 @@ # specific language governing permissions and limitations # under the License. -# Enable ctest if gtest is available -if(USE_GTEST AND NOT BUILD_TESTING STREQUAL "OFF") - # Check env var for backward compatibility. A better way to specify package - # locations is to use CMAKE_PREFIX_PATH or other standard CMake mechanism - # (see CMake documentation for `find_package`). - set(GTEST_ROOT "$ENV{GTEST_LIB}") - if("${USE_GTEST}" STREQUAL "AUTO") - # If USE_GTEST is AUTO, treat GTest as optional: enable if found. - find_package(GTest) - elseif("${USE_GTEST}" MATCHES ${IS_TRUE_PATTERN}) - # USE_GTEST is set to ON, TRUE, etc. Treat GTest as a required package. - find_package(GTest REQUIRED) - endif() - if(GTEST_FOUND) - if(NOT TARGET GTest::gmock) - # GMock is formally supported in CMake 3.20; for now, expect libgmock.a in the same directory, - # and require that folks compiling against GTest::gmock also link against GTest::GTest - # (for the includes dir). - add_library(GTest::gmock STATIC IMPORTED GLOBAL) - get_target_property(GTEST_LIB_PATH GTest::GTest IMPORTED_LOCATION) - if("${GTEST_LIB_PATH}" STREQUAL "GTEST_LIB_PATH-NOTFOUND") - # CMake >= 3.20 makes GTest::GTest into a compatibility target. The real import location is in - # GTest::gtest. - get_target_property(GTEST_LIB_PATH GTest::gtest IMPORTED_LOCATION) - if("${GTEST_LIB_PATH}" STREQUAL "GTEST_LIB_PATH-NOTFOUND") - message(FATAL_ERROR "Neither GTest::GTest nor GTest::gtest targets defined IMPORTED_LOCATION") - endif() - endif() - get_filename_component(GTEST_LIB_DIR "${GTEST_LIB_PATH}" DIRECTORY) - set_target_properties(GTest::gmock PROPERTIES - IMPORTED_LOCATION "${GTEST_LIB_DIR}/libgmock.a") - endif() - - enable_testing() - include(GoogleTest) - set(GTEST_FOUND TRUE PARENT_SCOPE) - endif() +# C++ tests always use GTest. Prefer an existing package, otherwise build it. +set(TVM_TIRX_USE_OPENCL_GTEST "" CACHE PATH "Optional OpenCL-specific GTest installation") +if(TVM_TIRX_USE_OPENCL_GTEST) + set(GTest_ROOT "${TVM_TIRX_USE_OPENCL_GTEST}") endif() - -# Create the `tirx_cpptest` target when GTest is available. -if(GTEST_FOUND) - tvm_file_glob(GLOB_RECURSE TEST_SRCS "${CMAKE_CURRENT_SOURCE_DIR}/*.cc") - add_executable(tirx_cpptest ${TEST_SRCS}) - # Preserve the executable location used before this became a subdirectory. - set_target_properties(tirx_cpptest PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}") - # include runtime files for unit testing - target_link_libraries(tirx_cpptest PRIVATE tvm_tirx GTest::GTest GTest::Main GTest::gmock pthread dl) - if(DEFINED LLVM_LIBS) - # The TVM library is linked with LLVM libraries. If the LLVM libraries are - # static and the symbols are not hidden, then don't link them again into - # tirx_cpptest since tirx_cpptest is itself linked against the TVM library. If static - # LLVM libraries are linked in twice, it can cause issues with global - # variable initialization (cl::opt). - # If the LLVM libraries are dynamic, we have to link them again, since the - # TVM library will not contain any LLVM definitions. - unset(LLVM_SO) - foreach(L IN LISTS LLVM_LIBS) - if(L MATCHES "libLLVM.*\.so") - set(LLVM_SO TRUE) - break() - endif() - endforeach() - if(DEFINED LLVM_SO OR HIDE_PRIVATE_SYMBOLS) - target_link_libraries(tirx_cpptest PRIVATE ${LLVM_LIBS}) - endif() - endif() - set_target_properties(tirx_cpptest PROPERTIES EXCLUDE_FROM_ALL 1) - set_target_properties(tirx_cpptest PROPERTIES EXCLUDE_FROM_DEFAULT_BUILD 1) - # tirx_cpptest is built into ${CMAKE_BINARY_DIR} but the TVM shared libs it - # links against (libtvm_tirx.so / libtvm_runtime.so) now live in - # ${CMAKE_BINARY_DIR}/lib. CMake's default behaviour bakes the - # absolute path of build/lib/ into the binary's RUNPATH, which works in tree - # but is not relocatable. Set $ORIGIN/lib (or @loader_path/lib on macOS) - # explicitly so the test binary finds its TVM libs relative to its own - # location, mirroring the install-time RPATH setup. - if(APPLE) - set_target_properties(tirx_cpptest PROPERTIES - BUILD_RPATH "@loader_path/lib" - INSTALL_RPATH "@loader_path/lib" - ) - elseif(UNIX) - set_target_properties(tirx_cpptest PROPERTIES - BUILD_RPATH "\$ORIGIN/lib" - INSTALL_RPATH "\$ORIGIN/lib" - ) - endif() - target_compile_definitions(tirx_cpptest PRIVATE "NDEBUG") - - # For some reason, compile definitions are not propagated correctly, so we manually add them here - target_compile_definitions(tirx_cpptest PUBLIC $) - gtest_discover_tests(tirx_cpptest WORKING_DIRECTORY "${CMAKE_BINARY_DIR}") +find_package(GTest CONFIG QUIET) +if(TARGET GTest::gtest AND TARGET GTest::gtest_main AND TARGET GTest::gmock) + set(_gtest_targets GTest::gtest GTest::gtest_main GTest::gmock) +else() + include(FetchContent) + set(INSTALL_GTEST OFF) + set(BUILD_GMOCK ON) + set(gtest_force_shared_crt ON) + FetchContent_Declare(googletest + GIT_REPOSITORY https://github.com/google/googletest.git + GIT_TAG v1.14.0) + FetchContent_MakeAvailable(googletest) + # Use one coherent fetched build even when a partial installed package exists. + set(_gtest_targets gtest gtest_main gmock) +endif() +include(GoogleTest) +file(GLOB_RECURSE _test_sources CONFIGURE_DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/*.cc") +add_executable(tirx_cpptest ${_test_sources}) +target_link_libraries(tirx_cpptest PRIVATE tvm_tirx ${_gtest_targets} ${CMAKE_DL_LIBS}) +if(TVM_TIRX_USE_THREADS) + target_link_libraries(tirx_cpptest PRIVATE Threads::Threads) +endif() +# LLVM symbols are hidden in the compiler library, including static LLVM builds. +if(DEFINED LLVM_LIBS) + target_link_libraries(tirx_cpptest PRIVATE ${LLVM_LIBS}) +endif() +set_target_properties(tirx_cpptest PROPERTIES + CXX_STANDARD 20 CXX_STANDARD_REQUIRED ON CXX_EXTENSIONS OFF + RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}") +foreach(_config IN LISTS CMAKE_CONFIGURATION_TYPES) + string(TOUPPER "${_config}" _config_upper) + set_target_properties(tirx_cpptest PROPERTIES + RUNTIME_OUTPUT_DIRECTORY_${_config_upper} "${PROJECT_BINARY_DIR}") +endforeach() +target_compile_definitions(tirx_cpptest PRIVATE NDEBUG) +if(APPLE) + set_property(TARGET tirx_cpptest PROPERTY BUILD_RPATH "@loader_path/lib") +elseif(UNIX) + set_property(TARGET tirx_cpptest PROPERTY BUILD_RPATH "\$ORIGIN/lib") endif() +gtest_discover_tests(tirx_cpptest WORKING_DIRECTORY "${PROJECT_BINARY_DIR}" DISCOVERY_MODE PRE_TEST) diff --git a/tirx/tests/python/support/test_config.py b/tirx/tests/python/support/test_config.py new file mode 100644 index 000000000000..76788bc50f32 --- /dev/null +++ b/tirx/tests/python/support/test_config.py @@ -0,0 +1,172 @@ +# 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. +"""Installed discovery stays within its selected prefix and does not load TVM.""" + +import importlib.util +import json +import os +import shutil +import subprocess +import sys +from pathlib import Path + +import pytest + +_CONFIG_PATH = Path(__file__).resolve().parents[3] / "python/tvm_tirx/config.py" +_SPEC = importlib.util.spec_from_file_location("tirx_config", _CONFIG_PATH) +config = importlib.util.module_from_spec(_SPEC) +_SPEC.loader.exec_module(config) + + +def _installation(prefix, libdir="lib", windows=False): + header = prefix / "include/tvm/runtime/base.h" + header.parent.mkdir(parents=True) + header.touch() + library_dir = prefix / libdir + cmake_dir = library_dir / "cmake/tvm_tirx" + cmake_dir.mkdir(parents=True) + for name in ("tvm_tirx", "tvm_runtime"): + filenames = (f"{name}.dll", f"{name}.lib") if windows else (f"lib{name}.so",) + for filename in filenames: + (library_dir / filename).touch() + return library_dir + + +def test_exact_library_directory_and_relocation(tmp_path): + prefix = tmp_path / "native install" + libdir = _installation(prefix, "lib/x86_64-linux-gnu") + (libdir / "libtvm_runtime_static.a").touch() + (libdir / "libtvm_runtime_cuda.so").touch() + (libdir / "libtvm_ffi.so").touch() + result = config.configuration(library_dir=libdir, platform="Linux") + assert result["include_dir"] == str(prefix / "include") + assert set(result["libraries"]) == {"tvm_tirx", "tvm_runtime"} + assert set(result) == {"version", "include_dir", "library_dir", "cmake_dir", "libraries"} + moved = tmp_path / "relocated" + shutil.copytree(prefix, moved) + result = config.configuration(library_dir=moved / "lib/x86_64-linux-gnu", platform="Linux") + assert result["include_dir"] == str(moved / "include") + assert all(str(moved) in library["location"] for library in result["libraries"].values()) + + +def test_prefix_selects_installed_lib64(tmp_path): + prefix = tmp_path / "install" + libdir = _installation(prefix, "lib64") + (prefix / "lib").mkdir() + result = config.configuration(prefix=prefix, platform="Linux") + assert result["library_dir"] == str(libdir) + result = subprocess.run( + [sys.executable, str(_CONFIG_PATH), "--prefix", str(prefix), "--cmakedir"], + check=True, + text=True, + capture_output=True, + ) + assert result.stdout.strip() == str(libdir / "cmake/tvm_tirx") + with pytest.raises(RuntimeError, match="Cannot find TIRx library"): + config.configuration(library_dir=prefix / "lib", platform="Linux") + + +@pytest.mark.parametrize("name", ["tvm_tirx", "tvm_runtime"]) +def test_windows_library_pairing(tmp_path, name): + libdir = _installation(tmp_path / "windows", windows=True) + (libdir / "tvm_runtime_static.lib").touch() + (libdir / "tvm_runtime_cuda.dll").touch() + (libdir / "tvm_runtime_cuda.lib").touch() + result = config.configuration(library_dir=libdir, platform="Windows") + assert set(result["libraries"]) == {"tvm_tirx", "tvm_runtime"} + assert result["libraries"][name] == { + "location": str(libdir / f"{name}.dll"), + "implib": str(libdir / f"{name}.lib"), + } + (libdir / f"{name}.lib").unlink() + with pytest.raises(RuntimeError, match="Missing import library"): + config.configuration(library_dir=libdir, platform="Windows") + (libdir / f"lib{name}.dll.a").touch() + result = config.configuration(library_dir=libdir, platform="Windows") + assert result["libraries"][name]["implib"] == str(libdir / f"lib{name}.dll.a") + + +def test_explicit_prefix_does_not_fall_back(tmp_path, monkeypatch): + good = tmp_path / "good" + _installation(good) + monkeypatch.setenv("PYTHONPATH", str(good)) + monkeypatch.setenv("TVM_TIRX_LIBRARY_PATH", str(good / "lib")) + with pytest.raises(RuntimeError, match="Cannot find TIRx public headers"): + config.configuration(prefix=tmp_path / "missing", platform="Linux") + + +def test_module_cli_without_native_dependencies(tmp_path): + prefix = tmp_path / "install" + _installation(prefix) + env = os.environ.copy() + env["PYTHONPATH"] = str(_CONFIG_PATH.parents[1]) + result = subprocess.run( + [sys.executable, "-S", "-m", "tvm_tirx.config", "--prefix", str(prefix), "--json"], + cwd=tmp_path, + env=env, + check=True, + text=True, + capture_output=True, + ) + assert json.loads(result.stdout)["include_dir"] == str(prefix / "include") + + +@pytest.mark.skipif(shutil.which("cmake") is None, reason="CMake is not installed") +def test_cmake_targets_without_ffi(tmp_path): + libdir = _installation(tmp_path / "install") + cmake_dir = libdir / "cmake/tvm_tirx" + shutil.copy2(_CONFIG_PATH, cmake_dir / "config.py") + shutil.copy2( + _CONFIG_PATH.parents[2] / "cmake/tvm_tirx-config.cmake", + cmake_dir / "tvm_tirx-config.cmake", + ) + source = tmp_path / "consumer" + source.mkdir() + (source / "CMakeLists.txt").write_text( + """cmake_minimum_required(VERSION 3.26) +project(config_contract LANGUAGES NONE) +find_package(tvm_tirx CONFIG REQUIRED) +get_property(imported DIRECTORY PROPERTY IMPORTED_TARGETS) +list(FILTER imported EXCLUDE REGEX "^Python::") +list(SORT imported) +if(NOT "${imported}" STREQUAL "tvm_runtime::shared;tvm_tirx::header;tvm_tirx::shared") + message(FATAL_ERROR "Unexpected imported targets: ${imported}") +endif() +foreach(target tvm_tirx::header tvm_tirx::shared tvm_runtime::shared) + get_target_property(links ${target} INTERFACE_LINK_LIBRARIES) + if(links MATCHES "ffi|OpenMP|Threads") + message(FATAL_ERROR "Unexpected implicit dependency: ${links}") + endif() +endforeach() +""", + encoding="utf-8", + ) + subprocess.run( + [ + "cmake", + "-S", + str(source), + "-B", + str(tmp_path / "build"), + f"-Dtvm_tirx_DIR={cmake_dir}", + f"-DPython_EXECUTABLE={sys.executable}", + "-DCMAKE_DISABLE_FIND_PACKAGE_tvm_ffi=TRUE", + ], + check=True, + text=True, + capture_output=True, + ) diff --git a/web/Makefile b/web/Makefile index 12ab22f471e0..c810feff78fd 100644 --- a/web/Makefile +++ b/web/Makefile @@ -18,8 +18,8 @@ TVM_ROOT=$(realpath $(shell dirname $(firstword $(MAKEFILE_LIST))))/../ INCLUDE_FLAGS = -I$(TVM_ROOT) -I$(TVM_ROOT)/include -I$(TVM_ROOT)/tirx/include\ - -I$(TVM_ROOT)/3rdparty/tvm-ffi/include\ - -I$(TVM_ROOT)/3rdparty/tvm-ffi/3rdparty/dlpack/include\ + -I$(TVM_ROOT)/tirx/3rdparty/tvm-ffi/include\ + -I$(TVM_ROOT)/tirx/3rdparty/tvm-ffi/3rdparty/dlpack/include\ -I$(TVM_ROOT)/tirx/3rdparty/compiler-rt .PHONY: clean all rmtypedep preparetest diff --git a/web/emcc/wasm_runtime.cc b/web/emcc/wasm_runtime.cc index a9f07eb0c0ef..352c63622c96 100644 --- a/web/emcc/wasm_runtime.cc +++ b/web/emcc/wasm_runtime.cc @@ -51,23 +51,23 @@ #include "tirx/src/runtime/timer.cc" #include "tirx/src/runtime/workspace_pool.cc" // relax setup -#include "3rdparty/tvm-ffi/src/ffi/backtrace.cc" -#include "3rdparty/tvm-ffi/src/ffi/container.cc" -#include "3rdparty/tvm-ffi/src/ffi/custom_allocator.cc" -#include "3rdparty/tvm-ffi/src/ffi/dtype.cc" -#include "3rdparty/tvm-ffi/src/ffi/error.cc" -#include "3rdparty/tvm-ffi/src/ffi/extra/env_c_api.cc" -#include "3rdparty/tvm-ffi/src/ffi/extra/env_context.cc" -#include "3rdparty/tvm-ffi/src/ffi/extra/json_parser.cc" -#include "3rdparty/tvm-ffi/src/ffi/extra/json_writer.cc" -#include "3rdparty/tvm-ffi/src/ffi/extra/library_module.cc" -#include "3rdparty/tvm-ffi/src/ffi/extra/library_module_system_lib.cc" -#include "3rdparty/tvm-ffi/src/ffi/extra/module.cc" -#include "3rdparty/tvm-ffi/src/ffi/function.cc" -#include "3rdparty/tvm-ffi/src/ffi/init_once.cc" -#include "3rdparty/tvm-ffi/src/ffi/object.cc" -#include "3rdparty/tvm-ffi/src/ffi/tensor.cc" -#include "3rdparty/tvm-ffi/src/ffi/testing/testing.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/backtrace.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/container.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/custom_allocator.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/dtype.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/error.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/extra/env_c_api.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/extra/env_context.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/extra/json_parser.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/extra/json_writer.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/extra/library_module.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/extra/library_module_system_lib.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/extra/module.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/function.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/init_once.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/object.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/tensor.cc" +#include "tirx/3rdparty/tvm-ffi/src/ffi/testing/testing.cc" // Enum registration requires the FFI type-attribute registry above. #include "src/runtime/vm/attn_backend.cc" #include "src/runtime/vm/builtin.cc"