From 8002a42a3726370f26234e39a90461e6ab2e6401 Mon Sep 17 00:00:00 2001 From: Anthony Shoumikhin Date: Sun, 26 Jul 2026 15:35:51 -0700 Subject: [PATCH 1/2] Pin ExecuTorch to the release/1.4 branch head main pins the ExecuTorch dependency to the release/1.3 branch head (6118688a), which predates the shared extension_cuda caller-stream work. Bump the pin to the release/1.4 branch head (cd380e7aefd18c171271cc228d3a155455095219), the first ExecuTorch release line that contains both pytorch/executorch#20158 (shared extension_cuda caller-stream library) and pytorch/executorch#20498 (caller CUDA stream for H2D/D2H copies). This follows the existing convention of pinning a release-branch head commit; the new_git_repository rule takes a commit, not a branch name. The four Bazel files that pin ExecuTorch are updated identically: MODULE.bazel, docker/MODULE.bazel.docker, docker/MODULE.bazel.ngc, and toolchains/ci_workspaces/MODULE.bazel.tmpl. Swap the branch-head commit for the immutable v1.4.0 tag once ExecuTorch publishes it. --- MODULE.bazel | 6 +++--- docker/MODULE.bazel.docker | 6 +++--- docker/MODULE.bazel.ngc | 6 +++--- toolchains/ci_workspaces/MODULE.bazel.tmpl | 6 +++--- 4 files changed, 12 insertions(+), 12 deletions(-) diff --git a/MODULE.bazel b/MODULE.bazel index 93aca46848..bfaa276b0b 100644 --- a/MODULE.bazel +++ b/MODULE.bazel @@ -42,12 +42,12 @@ new_git_repository = use_repo_rule("@bazel_tools//tools/build_defs/repo:git.bzl" local_torch = use_repo_rule("//toolchains:local_torch.bzl", "local_torch") -# Pinned to the ExecuTorch release/1.3 branch head. +# Pinned to the ExecuTorch release/1.4 branch head. new_git_repository( name = "executorch", build_file = "@//third_party/executorch:BUILD", - # latest commit in release/1.3 branch - commit = "6118688a095fd9697224f5cad72ce42db641c9cd", + # latest commit in release/1.4 branch + commit = "cd380e7aefd18c171271cc228d3a155455095219", patch_cmds = [ "find . -mindepth 2 \\( -name BUILD -o -name BUILD.bazel \\) -delete", "mkdir executorch && find . -mindepth 1 -maxdepth 1 ! -name executorch ! -name BUILD ! -name BUILD.bazel ! -name REPO.bazel -exec cp -a {} executorch/ \\;", diff --git a/docker/MODULE.bazel.docker b/docker/MODULE.bazel.docker index 316fbb6fd7..a60d7db576 100644 --- a/docker/MODULE.bazel.docker +++ b/docker/MODULE.bazel.docker @@ -62,12 +62,12 @@ new_local_repository( build_file = "third_party/libtorch/BUILD" ) -# Pinned to the ExecuTorch release/1.3 branch head. +# Pinned to the ExecuTorch release/1.4 branch head. new_git_repository( name = "executorch", build_file = "@//third_party/executorch:BUILD", - # latest commit in release/1.3 branch - commit = "6118688a095fd9697224f5cad72ce42db641c9cd", + # latest commit in release/1.4 branch + commit = "cd380e7aefd18c171271cc228d3a155455095219", patch_cmds = [ "find . -mindepth 2 \\( -name BUILD -o -name BUILD.bazel \\) -delete", "mkdir executorch && find . -mindepth 1 -maxdepth 1 ! -name executorch ! -name BUILD ! -name BUILD.bazel ! -name REPO.bazel -exec cp -a {} executorch/ \\;", diff --git a/docker/MODULE.bazel.ngc b/docker/MODULE.bazel.ngc index 2079b96b08..38c62b7649 100644 --- a/docker/MODULE.bazel.ngc +++ b/docker/MODULE.bazel.ngc @@ -71,12 +71,12 @@ new_local_repository( build_file = "third_party/libtorch/BUILD" ) -# Pinned to the ExecuTorch release/1.3 branch head. +# Pinned to the ExecuTorch release/1.4 branch head. new_git_repository( name = "executorch", build_file = "@//third_party/executorch:BUILD", - # latest commit in release/1.3 branch - commit = "6118688a095fd9697224f5cad72ce42db641c9cd", + # latest commit in release/1.4 branch + commit = "cd380e7aefd18c171271cc228d3a155455095219", recursive_init_submodules = True, patch_cmds = [ "find . -mindepth 2 \\( -name BUILD -o -name BUILD.bazel \\) -delete", diff --git a/toolchains/ci_workspaces/MODULE.bazel.tmpl b/toolchains/ci_workspaces/MODULE.bazel.tmpl index fface1d10d..ad7e7f6dc6 100644 --- a/toolchains/ci_workspaces/MODULE.bazel.tmpl +++ b/toolchains/ci_workspaces/MODULE.bazel.tmpl @@ -182,12 +182,12 @@ new_local_repository( build_file = "third_party/libtorch/BUILD" ) -# Pinned to the ExecuTorch release/1.3 branch head. +# Pinned to the ExecuTorch release/1.4 branch head. new_git_repository( name = "executorch", build_file = "@//third_party/executorch:BUILD", - # latest commit in release/1.3 branch - commit = "6118688a095fd9697224f5cad72ce42db641c9cd", + # latest commit in release/1.4 branch + commit = "cd380e7aefd18c171271cc228d3a155455095219", recursive_init_submodules = True, patch_cmds = [ "find . -mindepth 2 \\( -name BUILD -o -name BUILD.bazel \\) -delete", From 42bf7d4d2640eb6ed85bd676446d913e9b9e0606 Mon Sep 17 00:00:00 2001 From: shoumikhin Date: Thu, 30 Jul 2026 23:51:44 -0700 Subject: [PATCH 2/2] refactor(executorch): use shared caller stream Replace the TensorRT ExecuTorch backend private caller-stream TLS with ExecuTorch CallerStreamGuard/getCallerStream so CUDA-capable delegates share one process-wide selection. Link and package one shared extension_cuda instance, add ordinary caller-stream inference coverage in the reference runner, and verify both CMake-built and packaged runners consume the shared TLS without libtorch. The ExecuTorch release/1.4 pin is intentionally owned by the preceding version-bump commit. --- .../verify-executorch-reference-runner.sh | 147 ++++++++++++++++-- .github/workflows/_linux-x86_64-core.yml | 1 + .github/workflows/executorch-static-linux.yml | 21 +++ BUILD.bazel | 20 ++- cpp/BUILD | 2 + .../executorch/TensorRTBackend.h | 29 ---- .../torch_tensorrt/executorch/CMakeLists.txt | 77 ++++++++- cpp/src/torch_tensorrt/executorch/README.md | 72 ++++++++- .../executorch/TensorRTBackend.cpp | 27 +--- examples/executorch_reference_runner/BUILD | 2 + .../CMakeLists.txt | 5 +- .../executorch_reference_runner/README.md | 34 ++-- examples/executorch_reference_runner/main.cpp | 32 ++-- tests/cpp/BUILD | 1 + tests/cpp/executorch/BUILD | 10 ++ tests/cpp/executorch/test_caller_stream.cpp | 105 +++++++++++++ third_party/executorch/BUILD | 90 ++++++++++- 17 files changed, 576 insertions(+), 99 deletions(-) create mode 100644 tests/cpp/executorch/test_caller_stream.cpp diff --git a/.github/scripts/verify-executorch-reference-runner.sh b/.github/scripts/verify-executorch-reference-runner.sh index 19e00b4843..a37f8fb363 100755 --- a/.github/scripts/verify-executorch-reference-runner.sh +++ b/.github/scripts/verify-executorch-reference-runner.sh @@ -153,7 +153,18 @@ download_tensorrt_root() { if [[ ! -f "${tensorrt_archive}" ]]; then curl -fL "${tensorrt_url}" -o "${tensorrt_archive}" || return 1 fi - tar -xzf "${tensorrt_archive}" -C "${tensorrt_extract_dir}" || return 1 + case "${tensorrt_archive}" in + *.tar.zst) + tar --zstd -xf "${tensorrt_archive}" -C "${tensorrt_extract_dir}" || return 1 + ;; + *.tar.gz | *.tgz) + tar -xzf "${tensorrt_archive}" -C "${tensorrt_extract_dir}" || return 1 + ;; + *) + echo "Unsupported TensorRT archive format: ${tensorrt_archive}" >&2 + return 1 + ;; + esac if [[ -n "${tensorrt_strip_prefix}" ]]; then tensorrt_root="${tensorrt_extract_dir}/${tensorrt_strip_prefix}" @@ -290,6 +301,8 @@ require_tar_entry() { require_tar_entry "torch_tensorrt/src/torch_tensorrt/executorch/CMakeLists.txt" require_tar_entry "torch_tensorrt/examples/executorch_reference_runner/CMakeLists.txt" +require_tar_entry "torch_tensorrt/bin/example_executorch_runner" +require_tar_entry "torch_tensorrt/lib/libextension_cuda.so" require_tar_entry "torch_tensorrt/BUILD" export TORCH_TENSORRT_ROOT="${verify_root}/torch_tensorrt" @@ -317,19 +330,135 @@ cmake --build "${verify_root}/build-executorch-reference-runner" \ runner_log="${verify_root}/my_runner.log" runner_path="${verify_root}/build-executorch-reference-runner/example_executorch_runner" -if command -v ldd >/dev/null 2>&1 && - ldd "${runner_path}" | + +# Symbol/linkage inspection tools are mandatory on this Linux gate: silently +# skipping them would let a broken single-TLS layout pass unnoticed. +for _tool in ldd readelf nm; do + if ! command -v "${_tool}" >/dev/null 2>&1; then + echo "Required tool '${_tool}' not found; cannot verify caller-stream linkage" >&2 + exit 1 + fi +done + +# The runner must not pull in libtorch: this native path is libtorch-free. +if ldd "${runner_path}" | grep -E "libtorch|libtorch_cpu|libtorch_cuda|libc10" >&2; then echo "example_executorch_runner links PyTorch/libtorch shared libraries" >&2 exit 1 fi +# The runner must declare a real DT_NEEDED dependency on libextension_cuda.so +# (an ldd filename match alone would also accept a "=> not found" line). +if ! readelf -d "${runner_path}" | + grep -E "\(NEEDED\).*libextension_cuda\.so" >&2; then + echo "example_executorch_runner has no DT_NEEDED entry for libextension_cuda.so" >&2 + exit 1 +fi + +# ...and that dependency must actually resolve at load time. +if ldd "${runner_path}" | grep -E "libextension_cuda\.so.*=>.*not found" >&2; then + echo "example_executorch_runner cannot resolve libextension_cuda.so at runtime" >&2 + exit 1 +fi + +# The runner must import the caller-stream API from the shared library rather +# than define it privately. A private definition means a second copy of the +# thread-local, which silently breaks the cross-backend handshake. Assert the +# import (in .dynsym) rather than the absence of a definition: absence-of-symbol +# checks read .symtab, which is stripped from release binaries and would make +# the assertion pass vacuously. A private definition would satisfy the reference +# at link time and leave no import here. +for _symbol in getCallerStream CallerStreamGuard; do + if ! nm -D --undefined-only "${runner_path}" 2>/dev/null | grep -q "${_symbol}"; then + echo "example_executorch_runner does not import ${_symbol} from libextension_cuda.so" >&2 + exit 1 + fi +done + +# Validate the .so the runner ACTUALLY loads (resolved via ldd), not just a +# packaged copy. The runner is CMake-built and may link the CMake-built +# extension_cuda; whichever .so the loader binds to must export the accessor and +# must be the sole definer the runner sees. +loaded_extension_cuda="$( + ldd "${runner_path}" 2>/dev/null | + sed -n 's/.*libextension_cuda\.so[^ ]* => \([^ ]*\).*/\1/p' | + head -n1 +)" +if [[ -z "${loaded_extension_cuda}" || ! -f "${loaded_extension_cuda}" ]]; then + echo "Could not resolve the libextension_cuda.so the runner loads" >&2 + exit 1 +fi +if ! nm --defined-only --dynamic "${loaded_extension_cuda}" 2>/dev/null | + grep -q "getCallerStream"; then + echo "Loaded ${loaded_extension_cuda} does not export getCallerStream" >&2 + exit 1 +fi + +# Packaging integrity (independent of the CMake runner): the Bazel-packaged .so +# must exist and export the accessor, and no other packaged ELF may define the +# caller-stream symbols -- a second definition would reintroduce a duplicate +# thread-local in the shipped artifact. +packaged_runner="${TORCH_TENSORRT_ROOT}/bin/example_executorch_runner" +packaged_extension_cuda="${TORCH_TENSORRT_ROOT}/lib/libextension_cuda.so" +if [[ ! -x "${packaged_runner}" ]]; then + echo "Packaged example_executorch_runner missing or not executable: ${packaged_runner}" >&2 + exit 1 +fi +if [[ ! -f "${packaged_extension_cuda}" ]]; then + echo "Packaged libextension_cuda.so missing at ${packaged_extension_cuda}" >&2 + exit 1 +fi +if ! nm --defined-only --dynamic "${packaged_extension_cuda}" 2>/dev/null | + grep -q "getCallerStream"; then + echo "Packaged libextension_cuda.so does not export getCallerStream" >&2 + exit 1 +fi +if ! readelf -d "${packaged_runner}" | + grep -E "\(NEEDED\).*libextension_cuda\.so" >&2; then + echo "Packaged runner has no DT_NEEDED entry for libextension_cuda.so" >&2 + exit 1 +fi +if ldd "${packaged_runner}" | grep -E "libextension_cuda\.so.*=>.*not found" >&2; then + echo "Packaged runner cannot resolve libextension_cuda.so" >&2 + exit 1 +fi +for _symbol in getCallerStream CallerStreamGuard; do + if ! nm -D --undefined-only "${packaged_runner}" 2>/dev/null | grep -q "${_symbol}"; then + echo "Packaged runner does not import ${_symbol} from libextension_cuda.so" >&2 + exit 1 + fi +done + +# No other packaged ELF may define the caller-stream symbols: a second +# definition would reintroduce a duplicate thread-local. +extra_defs="$( + find "${TORCH_TENSORRT_ROOT}/lib" -maxdepth 1 -type f -name '*.so*' \ + ! -name 'libextension_cuda.so' -print0 2>/dev/null | + while IFS= read -r -d '' _so; do + if nm --defined-only --dynamic "${_so}" 2>/dev/null | + grep -qE "getCallerStream|CallerStreamGuard"; then + echo "${_so}" + fi + done +)" +if [[ -n "${extra_defs}" ]]; then + echo "Unexpected caller-stream definitions outside libextension_cuda.so:" >&2 + echo "${extra_defs}" >&2 + exit 1 +fi + "${runner_path}" \ --model_path="${model_path}" \ --num_runs=1 2>&1 | tee "${runner_log}" - -# The sample model is x + 1, and the reference runner fills inputs with 1.0f, -# so the output sample should contain 2.0000. -grep -q "Inference completed" "${runner_log}" -grep -q "output\\[0\\] shape=" "${runner_log}" -grep -Eq "first [0-9]+ values:.* 2\\.0000" "${runner_log}" +packaged_runner_log="${verify_root}/packaged_runner.log" +"${packaged_runner}" \ + --model_path="${model_path}" \ + --num_runs=1 2>&1 | tee "${packaged_runner_log}" + +# The sample model is x + 1, and both runners fill inputs with 1.0f, so each +# output sample must report the expected shape and values. ET_LOG output is not +# part of the packaged runner contract and may be compiled out. +for _log in "${runner_log}" "${packaged_runner_log}"; do + grep -q "output\\[0\\] shape=" "${_log}" + grep -Eq "first [0-9]+ values:.* 2\\.0000" "${_log}" +done diff --git a/.github/workflows/_linux-x86_64-core.yml b/.github/workflows/_linux-x86_64-core.yml index c0011470b2..9195bcc0bc 100644 --- a/.github/workflows/_linux-x86_64-core.yml +++ b/.github/workflows/_linux-x86_64-core.yml @@ -639,6 +639,7 @@ jobs: needs: [ build, + executorch-static-build, L0-dynamo-converter-tests, L0-dynamo-core-tests, L0-py-core-tests, diff --git a/.github/workflows/executorch-static-linux.yml b/.github/workflows/executorch-static-linux.yml index fe53fdb12c..28d3a56548 100644 --- a/.github/workflows/executorch-static-linux.yml +++ b/.github/workflows/executorch-static-linux.yml @@ -92,6 +92,27 @@ jobs: # this is to build the libtorchtrt.tar.gz bazel build //:libtorchtrt --compilation_mode opt --config=linux + # Run the ExecuTorch backend C++ unit tests, which are otherwise only + # built, never executed. These tests link TensorRT and CUDA, but Bazel's + # cc_import ships the unversioned libnvinfer.so and libcudart.so while the + # loader asks for the versioned sonames, so add the directories holding + # those to LD_LIBRARY_PATH. Append rather than replace: the toolchain + # entries already present are still needed. + bazel_external="$(bazel info output_base)/external" + for _soname in libnvinfer.so libcudart.so; do + _dir="$( + find -L "${bazel_external}" -path "*/lib*/${_soname}.*" -printf '%h\n' 2>/dev/null | + sort -u | head -n1 + )" + if [[ -z "${_dir}" ]]; then + echo "Could not locate a versioned ${_soname} under ${bazel_external}" >&2 + exit 1 + fi + export LD_LIBRARY_PATH="${_dir}${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}" + done + bazel test //tests/cpp/executorch:executorch_backend_tests \ + --compilation_mode opt --config=linux --test_output=errors \ + --test_env=LD_LIBRARY_PATH executorch_cmake_location="$(bazel query @executorch//:executorch/CMakeLists.txt --output=location)" export EXECUTORCH_SOURCE_DIR="$(dirname "${executorch_cmake_location%%:*}")" export EXECUTORCH_ROOT="${EXECUTORCH_SOURCE_DIR}" diff --git a/BUILD.bazel b/BUILD.bazel index 4dbc88f847..93141f93e7 100644 --- a/BUILD.bazel +++ b/BUILD.bazel @@ -163,20 +163,23 @@ alias( actual = ":executorch_trt_backend_archive", ) +# Ship only the shared caller-stream library. The packaged reference runner +# builds the ExecuTorch core and TensorRT backend from source (its CMake uses +# add_subdirectory), so prebuilt .a archives are not consumed by that flow -- but +# libextension_cuda.so must ship so the runner has a single shared caller-stream +# TLS at runtime. Package the .so target directly so the tarball never contains a +# static caller-stream archive. pkg_files( - name = "executorch_lib_pkg_files", - srcs = [ - ":executorch_core_archive", - ":executorch_trt_backend_archive", - ], + name = "executorch_extension_cuda_pkg_files", + srcs = ["@executorch//:libextension_cuda.so"], prefix = "lib/", visibility = ["//visibility:public"], ) pkg_tar( - name = "executorch_lib", - srcs = [":executorch_lib_pkg_files"], - mode = "0644", + name = "executorch_extension_cuda_lib", + srcs = [":executorch_extension_cuda_pkg_files"], + mode = "0755", package_dir = "", ) @@ -275,6 +278,7 @@ pkg_tar( ":rtx_sbsa": [], ":windows": [], "//conditions:default": [ + ":executorch_extension_cuda_lib", ":executorch_source_package", ":include_executorch", ], diff --git a/cpp/BUILD b/cpp/BUILD index b54b2d24e5..5870d52957 100644 --- a/cpp/BUILD +++ b/cpp/BUILD @@ -141,11 +141,13 @@ cc_library( ] + select({ ":linux_x86_64": [ "@executorch//:executorch_headers", + "@executorch//:extension_cuda", "@cuda//:cudart", "@tensorrt//:nvinfer", ], ":sbsa": [ "@executorch//:executorch_headers", + "@executorch//:extension_cuda", "@cuda//:cudart", "@tensorrt_sbsa//:nvinfer", ], diff --git a/cpp/include/torch_tensorrt/executorch/TensorRTBackend.h b/cpp/include/torch_tensorrt/executorch/TensorRTBackend.h index 31383f9e17..2eebb7cb67 100644 --- a/cpp/include/torch_tensorrt/executorch/TensorRTBackend.h +++ b/cpp/include/torch_tensorrt/executorch/TensorRTBackend.h @@ -90,34 +90,5 @@ class TensorRTBackend final : public ::executorch::runtime::BackendInterface { void destroy(::executorch::runtime::DelegateHandle* handle) const override; }; -// Selects, for the calling thread, the CUDA stream the delegate runs TensorRT on; -// scope it around execution. -// -// Confines inference to a CUDA green context's SM partition when the caller -// passes a cuGreenCtxStreamCreate stream: confinement rides the stream (the green -// context need not be current), and cudaStreamPerThread — the no-guard default — -// is rejected while a green context is current. While active, device-resident -// outputs are left enqueued on the stream (no end sync) to compose with later GPU -// work. -// -// Contract: the stream is on the engine's device and outlives the guard; a handle -// is executed by one thread at a time. On the no-end-sync path (guard active, all -// I/O device-resident) execute() returns with the TensorRT enqueue still in flight -// on the stream; the delegate itself orders the next execute() on, and the -// destruction of, that handle after the work completes (via an internal completion -// event), so the caller need only synchronize the stream before reading -// device-resident outputs. -class CudaStreamGuard { - public: - explicit CudaStreamGuard(cudaStream_t stream); - ~CudaStreamGuard(); - CudaStreamGuard(const CudaStreamGuard&) = delete; - CudaStreamGuard& operator=(const CudaStreamGuard&) = delete; - - private: - cudaStream_t prev_stream_; - bool prev_set_; -}; - } // namespace executorch_backend } // namespace torch_tensorrt diff --git a/cpp/src/torch_tensorrt/executorch/CMakeLists.txt b/cpp/src/torch_tensorrt/executorch/CMakeLists.txt index 1b0546b545..6f72a155c8 100644 --- a/cpp/src/torch_tensorrt/executorch/CMakeLists.txt +++ b/cpp/src/torch_tensorrt/executorch/CMakeLists.txt @@ -1,4 +1,4 @@ -cmake_minimum_required(VERSION 3.17) +cmake_minimum_required(VERSION 3.19) project(torch_tensorrt_executorch LANGUAGES CXX) include(GNUInstallDirs) @@ -50,8 +50,10 @@ target_compile_definitions(executorch_trt_backend C10_USING_CUSTOM_GENERATED_MACROS ) +set(_torchtrt_executorch_source_root "") foreach(_executorch_root_candidate IN ITEMS "${EXECUTORCH_SOURCE_DIR}" "${EXECUTORCH_ROOT}") if(_executorch_root_candidate AND EXISTS "${_executorch_root_candidate}/runtime/core/portable_type/c10") + set(_torchtrt_executorch_source_root "${_executorch_root_candidate}") get_filename_component(_executorch_include_parent "${_executorch_root_candidate}" DIRECTORY) target_include_directories(executorch_trt_backend PUBLIC @@ -61,10 +63,83 @@ foreach(_executorch_root_candidate IN ITEMS "${EXECUTORCH_SOURCE_DIR}" "${EXECUT endif() endforeach() +# Select ExecuTorch's shared extension_cuda so every CUDA-capable delegate shares +# one caller-stream thread-local. A static copy linked into a second shared object +# would create a second thread-local and silently break the handshake, so every +# branch below must resolve to one shared library. Precedence: an +# ExecuTorch-provided target, then an explicit prebuilt shared library, then a +# source build. +if(TARGET extension_cuda) + # Provided by ExecuTorch, e.g. add_subdirectory() with EXECUTORCH_BUILD_CUDA=ON. + get_target_property(_extension_cuda_type extension_cuda TYPE) + if(NOT _extension_cuda_type STREQUAL "SHARED_LIBRARY") + message(FATAL_ERROR + "ExecuTorch's extension_cuda target must be shared; got " + "${_extension_cuda_type}. Static copies break caller-stream sharing." + ) + endif() +elseif(EXECUTORCH_EXTENSION_CUDA_LIBRARY) + # Use a prebuilt shared library, e.g. the one shipped in libtorchtrt.tar.gz. + # A caller-supplied override is authoritative: never silently fall back to a + # different source-built library when its path is invalid. + if(NOT EXISTS "${EXECUTORCH_EXTENSION_CUDA_LIBRARY}") + message(FATAL_ERROR + "EXECUTORCH_EXTENSION_CUDA_LIBRARY does not exist: " + "${EXECUTORCH_EXTENSION_CUDA_LIBRARY}" + ) + endif() + # Read the ELF header rather than trusting the file name: the linker accepts + # a static archive or a linker script named ".so", and either one embeds a + # private copy of the caller-stream thread-local. e_ident is always 16 bytes, + # so e_type sits at byte 16 for both ELF32 and ELF64; EI_DATA at byte 5 says + # whether to read it little- or big-endian. + file(READ "${EXECUTORCH_EXTENSION_CUDA_LIBRARY}" _extension_cuda_header LIMIT 18 HEX) + if(_extension_cuda_header MATCHES "^7f454c46..02") + set(_extension_cuda_et_dyn "0003") + else() + set(_extension_cuda_et_dyn "0300") + endif() + if(NOT _extension_cuda_header MATCHES "^7f454c46" OR + NOT _extension_cuda_header MATCHES "${_extension_cuda_et_dyn}$") + message(FATAL_ERROR + "EXECUTORCH_EXTENSION_CUDA_LIBRARY is not an ELF shared object: " + "${EXECUTORCH_EXTENSION_CUDA_LIBRARY}. Static archives, linker scripts " + "and other platform formats are rejected because each would duplicate " + "the caller-stream thread-local." + ) + endif() + add_library(extension_cuda SHARED IMPORTED) + set_target_properties(extension_cuda PROPERTIES + IMPORTED_LOCATION "${EXECUTORCH_EXTENSION_CUDA_LIBRARY}" + ) +elseif(_torchtrt_executorch_source_root AND + EXISTS "${_torchtrt_executorch_source_root}/extension/cuda/CMakeLists.txt") + # Build the single shared library from ExecuTorch's own definition rather + # than a copy of it, so this stays correct if ExecuTorch changes the target + # and so a duplicate definition is a hard error instead of two libraries. + if(NOT _common_include_directories) + set(_common_include_directories "${_executorch_include_parent}") + endif() + add_subdirectory( + "${_torchtrt_executorch_source_root}/extension/cuda" + "${CMAKE_CURRENT_BINARY_DIR}/executorch_extension_cuda" + ) +else() + message(FATAL_ERROR + "Torch-TensorRT's ExecuTorch backend requires ExecuTorch's shared " + "extension_cuda library. Add ExecuTorch first with EXECUTORCH_BUILD_CUDA=ON, " + "or set EXECUTORCH_EXTENSION_CUDA_LIBRARY to a prebuilt shared " + "libextension_cuda, or set EXECUTORCH_SOURCE_DIR/EXECUTORCH_ROOT to an " + "ExecuTorch checkout containing extension/cuda. Do not link a static copy: " + "all delegates must share one caller-stream thread-local." + ) +endif() + set(_torchtrt_executorch_link_libraries CUDA::cudart TensorRT::nvinfer Threads::Threads + extension_cuda ) if(TARGET executorch_core) diff --git a/cpp/src/torch_tensorrt/executorch/README.md b/cpp/src/torch_tensorrt/executorch/README.md index 1f56507ac0..719a8b75ac 100644 --- a/cpp/src/torch_tensorrt/executorch/README.md +++ b/cpp/src/torch_tensorrt/executorch/README.md @@ -10,10 +10,19 @@ user_runner_project/ torch_tensorrt/ ``` -The normal integration path is to add both ExecuTorch and this package from -your runner CMake. Linking `torchtrt::executorch_backend` makes the backend -archive a dependency of your runner target, so you do not need a separate -backend build step. +This backend requires ExecuTorch 1.4 or a source commit containing +`pytorch/executorch#20158` and `pytorch/executorch#20498`. The normal integration +path for a runner that already enables ExecuTorch's CUDA backend is to add both +ExecuTorch (with `EXECUTORCH_BUILD_CUDA=ON`) and this package, so the TensorRT +backend links ExecuTorch's shared `extension_cuda` target directly. A libtorch-free +runner should leave the full CUDA/AOTI backend disabled; this package then builds +only the minimal shared `extension_cuda` caller-stream library from the ExecuTorch +source checkout. Consumers on ELF platforms may instead set +`EXECUTORCH_EXTENSION_CUDA_LIBRARY` to a prebuilt shared library; the value is +checked to be a shared object, because a static copy would give each delegate its +own caller-stream state. Linking `torchtrt::executorch_backend` +makes the backend archive a dependency of your runner target, so you do not need a +separate backend build step. ```cmake add_subdirectory("executorch") @@ -31,6 +40,58 @@ target_link_libraries( The backend archive is available as the `executorch_trt_backend` CMake target and is written to `${CMAKE_BINARY_DIR}/lib/libexecutorch_trt_backend.a`. +`libextension_cuda` remains a shared runtime dependency so every CUDA-capable +delegate in the process observes the same caller-stream TLS instance. + +## Caller Stream API Migration + +`torch_tensorrt::executorch_backend::CudaStreamGuard` has been removed. Use +ExecuTorch's backend-neutral guard instead: + +```cpp +#include + +executorch::extension::cuda::CallerStreamGuard guard(stream); +module.forward(inputs); +``` + +The old class is intentionally not kept as a deprecated alias: the goal is one +backend-neutral primitive and one shared TLS definition, so all CUDA-capable +delegates read the same caller-stream selection. (A deprecated `using` alias to +`executorch::extension::cuda::CallerStreamGuard` would have shared that same TLS, +so this removal is an API-simplification choice, not a correctness requirement.) +This is a source-breaking C++ change; downstream callers must switch to the new +type. + +### Caller-stream contract for the TensorRT backend + +The upstream `CallerStreamGuard` documents the generic contract (per-thread, +nested scoping; the caller owns the stream for the guard's lifetime; the caller +manages host-data lifetime for async work). The TensorRT backend adds these +requirements, which previously lived on +the removed `CudaStreamGuard`: + +- The selected stream must be on the TensorRT engine's device. +- Calls using one delegate handle must not overlap, and must not overlap with + its destruction; the backend serializes `execute()` calls with an internal + mutex, but destruction is not mutex-guarded. +- With a guard active and when no host staging is required (all inputs and + outputs are directly bindable — device, managed, or unified memory), + `execute()` may return with the TensorRT enqueue still in flight on the + stream (no end-of-execute sync). The backend orders the next `execute()` and + the handle's destruction after that work via an internal completion event, but + that event only protects backend-owned state. The caller must therefore keep + all directly bound input/output storage alive and unmodified until the work is + complete, order any cross-stream producers/consumers with their own events, + and synchronize the stream before reading outputs on the host. +- With no guard active, the backend falls back to `cudaStreamPerThread`. +- The reference-runner smoke test runs inference inside a caller-stream guard on + the discrete-GPU CI configuration, where all inputs and outputs are host-backed + and therefore take the synchronized staging path. CI separately asserts that the + runner resolves one shared `libextension_cuda.so`. Device-resident asynchronous + return is not covered end to end. +- CUDA green-context streams require context-aware completion-event handling and + are not yet part of this integration's validated support matrix. ## Standalone Backend Archive @@ -62,5 +123,6 @@ cmake -S torch_tensorrt/src/torch_tensorrt/executorch -B build-torchtrt-executor -DEXECUTORCH_ROOT="${EXECUTORCH_ROOT}" \ -DTensorRT_ROOT="${TensorRT_ROOT}" -cmake --build build-torchtrt-executorch --target executorch_trt_backend -j +cmake --build build-torchtrt-executorch \ + --target executorch_trt_backend -j ``` diff --git a/cpp/src/torch_tensorrt/executorch/TensorRTBackend.cpp b/cpp/src/torch_tensorrt/executorch/TensorRTBackend.cpp index b2e3b08232..e4411e8aa5 100644 --- a/cpp/src/torch_tensorrt/executorch/TensorRTBackend.cpp +++ b/cpp/src/torch_tensorrt/executorch/TensorRTBackend.cpp @@ -20,6 +20,7 @@ #include #include +#include #include #include #include @@ -48,21 +49,6 @@ using ::executorch::runtime::Span; } \ } while (false) -namespace { -thread_local cudaStream_t g_user_stream = nullptr; -thread_local bool g_user_stream_set = false; -} // namespace - -CudaStreamGuard::CudaStreamGuard(cudaStream_t stream) : prev_stream_(g_user_stream), prev_set_(g_user_stream_set) { - g_user_stream = stream; - g_user_stream_set = true; -} - -CudaStreamGuard::~CudaStreamGuard() { - g_user_stream = prev_stream_; - g_user_stream_set = prev_set_; -} - void TRTLogger::log(Severity severity, const char* msg) noexcept { if (severity <= Severity::kERROR) { ET_LOG(Error, "TensorRT: %s", msg); @@ -379,7 +365,9 @@ Error TensorRTBackend::execute(BackendExecutionContext& context, DelegateHandle* return Error::InvalidProgram; } } - cudaStream_t stream = g_user_stream_set ? g_user_stream : cudaStreamPerThread; + const auto caller_stream = ::executorch::extension::cuda::getCallerStream(); + const bool caller_stream_set = caller_stream.has_value(); + cudaStream_t stream = caller_stream.value_or(cudaStreamPerThread); bool output_staged_to_host = false; bool input_staged_from_host = false; @@ -571,9 +559,8 @@ Error TensorRTBackend::execute(BackendExecutionContext& context, DelegateHandle* if (!ctx->enqueueV3(stream)) { ET_LOG( Error, - "TensorRTBackend::execute: enqueueV3 failed. If a CUDA green context is " - "current, scope a CudaStreamGuard with a green-context stream: " - "cudaStreamPerThread is invalid while a green context is current."); + "TensorRTBackend::execute: enqueueV3 failed. Verify that the selected " + "CallerStreamGuard stream belongs to the TensorRT engine device."); return Error::InvalidState; } @@ -587,7 +574,7 @@ Error TensorRTBackend::execute(BackendExecutionContext& context, DelegateHandle* // next execute() and the destructor wait before reusing/freeing exec_ctx. The D2H // copies live in the must_sync branch: an output staged to host always sets // output_staged_to_host, so outputs_needing_copy is empty on the skip path. - const bool must_sync = output_staged_to_host || input_staged_from_host || !g_user_stream_set; + const bool must_sync = output_staged_to_host || input_staged_from_host || !caller_stream_set; if (must_sync) { for (auto& output : outputs_needing_copy) { exec_aten::Tensor et_out = args[num_inputs + output.first]->toTensor(); diff --git a/examples/executorch_reference_runner/BUILD b/examples/executorch_reference_runner/BUILD index 4f53bc6b91..9326a3b664 100644 --- a/examples/executorch_reference_runner/BUILD +++ b/examples/executorch_reference_runner/BUILD @@ -14,9 +14,11 @@ filegroup( cc_binary( name = "example_executorch_runner", srcs = ["main.cpp"], + linkopts = ["-Wl,-rpath,$$ORIGIN/../lib"], deps = [ "//cpp:tensorrt_executorch_backend", "@executorch//:executorch_core", "@executorch//:executorch_file_data_loader", + "@executorch//:extension_cuda", ], ) diff --git a/examples/executorch_reference_runner/CMakeLists.txt b/examples/executorch_reference_runner/CMakeLists.txt index 2bd3544d67..97301cb03a 100644 --- a/examples/executorch_reference_runner/CMakeLists.txt +++ b/examples/executorch_reference_runner/CMakeLists.txt @@ -1,4 +1,4 @@ -cmake_minimum_required(VERSION 3.17) +cmake_minimum_required(VERSION 3.24) project(my_runner LANGUAGES CXX) set(CMAKE_CXX_STANDARD 17) @@ -28,6 +28,9 @@ set(TORCHTRT_EXECUTORCH_SOURCE_DIR set(BUILD_TESTING OFF CACHE BOOL "" FORCE) set(EXECUTORCH_BUILD_PYBIND OFF CACHE BOOL "" FORCE) +# Keep this reference runner libtorch-free. The Torch-TensorRT backend builds +# only ExecuTorch's minimal shared caller-stream extension from source below. +set(EXECUTORCH_BUILD_CUDA OFF CACHE BOOL "" FORCE) set(EXECUTORCH_BUILD_EXTENSION_DATA_LOADER ON CACHE BOOL "" FORCE) set(EXECUTORCH_BUILD_EXTENSION_FLAT_TENSOR ON CACHE BOOL "" FORCE) set(EXECUTORCH_BUILD_EXTENSION_MODULE ON CACHE BOOL "" FORCE) diff --git a/examples/executorch_reference_runner/README.md b/examples/executorch_reference_runner/README.md index a518353bb6..24fcf12743 100644 --- a/examples/executorch_reference_runner/README.md +++ b/examples/executorch_reference_runner/README.md @@ -34,16 +34,15 @@ torch_tensorrt/bin/example_executorch_runner ``` ```bash -# Get the ExecuTorch source code. Set EXECUTORCH_REF to a branch or tag; -# leave it unset for the latest main. -EXECUTORCH_REF="${EXECUTORCH_REF:-main}" -case "${EXECUTORCH_REF}" in - latest|latest-main|latest_main|"latest main") - EXECUTORCH_REF="main" - ;; -esac -git clone --depth 1 --branch "${EXECUTORCH_REF}" --recurse-submodules --shallow-submodules \ +# Get the ExecuTorch release/1.4 source snapshot used by this package. +EXECUTORCH_REF="${EXECUTORCH_REF:-cd380e7aefd18c171271cc228d3a155455095219}" +git clone --filter=blob:none --no-checkout \ https://github.com/pytorch/executorch.git executorch +pushd executorch +git fetch --depth 1 origin "${EXECUTORCH_REF}" +git checkout FETCH_HEAD +git submodule update --init --recursive --depth 1 +popd # download the libtorchtrt.tar.gz tar xvf libtorchtrt.tar.gz @@ -64,16 +63,22 @@ cmake -S "${TORCH_TENSORRT_ROOT}/examples/executorch_reference_runner" \ cmake --build build-executorch-reference-runner --target example_executorch_runner -j ``` +This packaged reference-runner flow currently targets Linux x86_64 and SBSA, +matching the native backend's Bazel compatibility constraints. Windows and +JetPack packaging are outside this flow. + Expected artifact: ```text build-executorch-reference-runner/example_executorch_runner ``` -The build also creates the executorch core and tensorrt backend archive as a dependency: +The build also creates the ExecuTorch core, the shared caller-stream extension, +and the TensorRT backend archive as dependencies: ```text build-executorch-reference-runner/executorch/libexecutorch_core.a +build-executorch-reference-runner/torch_tensorrt_executorch/executorch_extension_cuda/libextension_cuda.so build-executorch-reference-runner/lib/libexecutorch_trt_backend.a ``` @@ -100,6 +105,13 @@ method.get_outputs(...) ``` Loading the method initializes the TensorRT ExecuTorch backend for any -Torch-TensorRT delegate subgraphs embedded in the `.pte`. The Python +Torch-TensorRT delegate subgraphs embedded in the `.pte`. Applications can +scope `executorch::extension::cuda::CallerStreamGuard` around execution to run +TensorRT and CUDA/AOTI delegates on one ordinary caller-owned CUDA stream. On the +discrete-GPU CI configuration, this runner's host-backed inputs and outputs take +the synchronized staging path; that test exercises guarded inference and checks the +output values, not the device-resident asynchronous fast path. Integrated GPUs may bind +host-backed storage directly and can follow the asynchronous contract below. +Green-context streams are not yet in the validated support matrix. The Python `torch_tensorrt` package is needed when exporting the `.pte`; it is not needed by this native runner at inference time. diff --git a/examples/executorch_reference_runner/main.cpp b/examples/executorch_reference_runner/main.cpp index beb39c3925..b42a6af7b5 100644 --- a/examples/executorch_reference_runner/main.cpp +++ b/examples/executorch_reference_runner/main.cpp @@ -21,6 +21,8 @@ #include #include +#include +#include #include #include #include @@ -159,16 +161,28 @@ int main(int argc, char** argv) { input_strides[i].data()); } - for (int run = 0; run < num_runs; ++run) { - for (size_t i = 0; i < num_inputs; ++i) { - exec_aten::Tensor input_tensor(&input_impls[i]); - EValue input_evalue(input_tensor); - Error err = method->set_input(input_evalue, i); - ET_CHECK_MSG(err == Error::Ok, "set_input(%zu) failed: 0x%" PRIx32, i, static_cast(err)); - } + cudaStream_t caller_stream = nullptr; + cudaError_t cuda_status = cudaStreamCreate(&caller_stream); + ET_CHECK_MSG(cuda_status == cudaSuccess, "cudaStreamCreate failed: %s", cudaGetErrorString(cuda_status)); + { + executorch::extension::cuda::CallerStreamGuard caller_stream_guard(caller_stream); + for (int run = 0; run < num_runs; ++run) { + for (size_t i = 0; i < num_inputs; ++i) { + exec_aten::Tensor input_tensor(&input_impls[i]); + EValue input_evalue(input_tensor); + Error err = method->set_input(input_evalue, i); + ET_CHECK_MSG(err == Error::Ok, "set_input(%zu) failed: 0x%" PRIx32, i, static_cast(err)); + } - Error status = method->execute(); - ET_CHECK_MSG(status == Error::Ok, "execute() failed on run %d: 0x%" PRIx32, run, static_cast(status)); + Error status = method->execute(); + ET_CHECK_MSG(status == Error::Ok, "execute() failed on run %d: 0x%" PRIx32, run, static_cast(status)); + } + } + cuda_status = cudaStreamSynchronize(caller_stream); + ET_CHECK_MSG(cuda_status == cudaSuccess, "cudaStreamSynchronize failed: %s", cudaGetErrorString(cuda_status)); + cuda_status = cudaStreamDestroy(caller_stream); + if (cuda_status != cudaSuccess) { + ET_LOG(Error, "cudaStreamDestroy failed: %s", cudaGetErrorString(cuda_status)); } ET_LOG(Info, "Inference completed (%d run(s)).", num_runs); diff --git a/tests/cpp/BUILD b/tests/cpp/BUILD index 7487586b02..ba0ed17596 100644 --- a/tests/cpp/BUILD +++ b/tests/cpp/BUILD @@ -65,6 +65,7 @@ test_suite( test_suite( name = "executorch_backend_tests", tests = [ + "//tests/cpp/executorch:test_caller_stream", "//tests/cpp/executorch:test_executorch_binding_names", "//tests/cpp/executorch:test_executorch_blob_header", ], diff --git a/tests/cpp/executorch/BUILD b/tests/cpp/executorch/BUILD index aee13cbcfd..8043991d45 100644 --- a/tests/cpp/executorch/BUILD +++ b/tests/cpp/executorch/BUILD @@ -5,11 +5,21 @@ package(default_visibility = ["//visibility:public"]) test_suite( name = "executorch_backend_tests", tests = [ + ":test_caller_stream", ":test_executorch_binding_names", ":test_executorch_blob_header", ], ) +cc_test( + name = "test_caller_stream", + srcs = ["test_caller_stream.cpp"], + deps = [ + "@executorch//:extension_cuda", + "@googletest//:gtest_main", + ], +) + cc_test( name = "test_executorch_binding_names", srcs = ["test_executorch_binding_names.cpp"], diff --git a/tests/cpp/executorch/test_caller_stream.cpp b/tests/cpp/executorch/test_caller_stream.cpp new file mode 100644 index 0000000000..983ae4d601 --- /dev/null +++ b/tests/cpp/executorch/test_caller_stream.cpp @@ -0,0 +1,105 @@ +/* + * Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + */ + +// Pins the caller-stream properties the TensorRT backend relies on. The backend +// derives both the stream it enqueues on and whether it may return with work +// still in flight from one getCallerStream() read, so a change to any property +// below would silently change delegate behavior. + +#include +#include +#include + +#include +#include + +namespace { + +namespace cuda = executorch::extension::cuda; + +cudaStream_t sentinel(int index) { + alignas(alignof(std::max_align_t)) static char storage[2][alignof(std::max_align_t)]; + return reinterpret_cast(&storage[index]); +} + +TEST(CallerStreamTest, NoGuardLeavesSelectionEmpty) { + EXPECT_FALSE(cuda::getCallerStream().has_value()); +} + +TEST(CallerStreamTest, GuardSelectsStreamAndRestoresOnExit) { + const cudaStream_t stream = sentinel(0); + { + cuda::CallerStreamGuard guard(stream); + EXPECT_EQ(cuda::getCallerStream(), std::optional(stream)); + } + EXPECT_FALSE(cuda::getCallerStream().has_value()); +} + +// The backend treats "a stream was selected" and "which stream" as separate +// answers, so an explicit null must stay engaged rather than read as no +// selection. This preserves the pre-refactor (stream, is_set) encoding. +TEST(CallerStreamTest, ExplicitNullStreamStaysEngaged) { + { + cuda::CallerStreamGuard guard(nullptr); + const auto selected = cuda::getCallerStream(); + ASSERT_TRUE(selected.has_value()); + EXPECT_EQ(*selected, nullptr); + } + EXPECT_FALSE(cuda::getCallerStream().has_value()); +} + +// cudaStreamPerThread names the same stream the backend falls back to, yet +// selecting it explicitly must still read as a caller selection. +TEST(CallerStreamTest, ExplicitPerThreadStreamStaysEngaged) { + cuda::CallerStreamGuard guard(cudaStreamPerThread); + const auto selected = cuda::getCallerStream(); + ASSERT_TRUE(selected.has_value()); + EXPECT_EQ(*selected, cudaStreamPerThread); +} + +TEST(CallerStreamTest, NestedGuardsRestoreOuterSelection) { + const cudaStream_t outer = sentinel(0); + const cudaStream_t inner = sentinel(1); + + { + cuda::CallerStreamGuard outer_guard(outer); + EXPECT_EQ(cuda::getCallerStream(), std::optional(outer)); + { + cuda::CallerStreamGuard inner_guard(inner); + EXPECT_EQ(cuda::getCallerStream(), std::optional(inner)); + { + cuda::CallerStreamGuard null_guard(nullptr); + EXPECT_EQ(cuda::getCallerStream(), std::optional(nullptr)); + } + EXPECT_EQ(cuda::getCallerStream(), std::optional(inner)); + } + EXPECT_EQ(cuda::getCallerStream(), std::optional(outer)); + } + EXPECT_FALSE(cuda::getCallerStream().has_value()); +} + +// The backend runs one handle per thread, so one thread's guard must not be +// visible to another. +TEST(CallerStreamTest, SelectionIsPerThread) { + const cudaStream_t stream = sentinel(0); + cuda::CallerStreamGuard guard(stream); + ASSERT_EQ(cuda::getCallerStream(), std::optional(stream)); + + std::optional observed_in_worker; + bool worker_saw_selection = true; + std::thread worker([&] { + const auto selected = cuda::getCallerStream(); + worker_saw_selection = selected.has_value(); + observed_in_worker = selected; + }); + worker.join(); + + EXPECT_FALSE(worker_saw_selection); + EXPECT_FALSE(observed_in_worker.has_value()); +} + +} // namespace diff --git a/third_party/executorch/BUILD b/third_party/executorch/BUILD index 119d8dd256..c11f51cf19 100644 --- a/third_party/executorch/BUILD +++ b/third_party/executorch/BUILD @@ -1,4 +1,4 @@ -load("@rules_cc//cc:defs.bzl", "cc_import", "cc_library") +load("@rules_cc//cc:defs.bzl", "cc_binary", "cc_import", "cc_library") load("@rules_foreign_cc//foreign_cc:defs.bzl", "cmake") package(default_visibility = ["//visibility:public"]) @@ -39,13 +39,90 @@ cmake( targets = ["executorch_core"], ) +# extension_cuda MUST resolve to exactly ONE shared library in the process so +# its `thread_local caller_stream_` has a single definition across every DSO +# (the TensorRT backend, the CUDA/AOTI backend, and the runner). A plain +# cc_library with srcs does not guarantee this: cc_binary/cc_test consumers +# default to static-linking their cc_library deps, which would embed a private +# copy of the thread-local and silently break the cross-backend caller-stream +# handshake. +# +# Following this repo's convention for shared libraries (see cpp/lib/BUILD), +# build the implementation once as a cc_binary(linkshared=True), re-import the +# .so via cc_import, and expose it through a srcs-less cc_library that also +# carries the headers (with include remap) and the cudart dep. Consumers depend +# on ":extension_cuda"; caller_stream.cpp is only ever linked into the one .so. + +# Private implementation. It carries hdrs + include_prefix/strip_include_prefix +# so caller_stream.cpp can find its own header +# (the physical tree has an extra executorch/ dir; a cc_binary's srcs headers +# get no such remap). alwayslink=True guarantees the object -- and thus +# getCallerStream/CallerStreamGuard -- is pulled into the shared object below, +# even though that source-less cc_binary references none of its symbols. +cc_library( + name = "extension_cuda_impl", + srcs = ["executorch/extension/cuda/caller_stream.cpp"], + hdrs = [ + "executorch/extension/cuda/caller_stream.h", + "executorch/extension/cuda/export.h", + ], + local_defines = [ + "C10_USING_CUSTOM_GENERATED_MACROS", + "EXECUTORCH_EXTENSION_CUDA_BUILDING", + ], + include_prefix = "executorch", + strip_include_prefix = "executorch", + alwayslink = True, + linkstatic = True, + visibility = ["//visibility:private"], + deps = ["@cuda//:cudart"], +) + +# The single shared object. The explicit -soname pins DT_SONAME to the bare +# filename so consumers get a clean DT_NEEDED=libextension_cuda.so that the +# packaged $ORIGIN/../lib RPATH resolves. alwayslink on the impl above ensures +# the caller-stream symbols are actually present despite the empty srcs. +cc_binary( + name = "libextension_cuda.so", + linkopts = ["-Wl,-soname,libextension_cuda.so"], + linkshared = True, + linkstatic = True, + visibility = ["//visibility:public"], + deps = [":extension_cuda_impl"], +) + +# Import the shared object as a linkable target (native cc_import takes no deps). +cc_import( + name = "extension_cuda_shared", + shared_library = ":libextension_cuda.so", + visibility = ["//visibility:private"], +) + +# Public consumer-facing target: srcs-less, so it never static-links +# caller_stream.cpp. It provides the caller-stream headers (with the executorch/ +# include remap), the imported single .so, and the transitive cudart dep, so +# every consumer dynamically links one shared thread-local. +cc_library( + name = "extension_cuda", + hdrs = [ + "executorch/extension/cuda/caller_stream.h", + "executorch/extension/cuda/export.h", + ], + include_prefix = "executorch", + strip_include_prefix = "executorch", + deps = [ + ":extension_cuda_shared", + "@cuda//:cudart", + ], +) + cc_library( name = "executorch_headers", hdrs = glob( [ - "runtime/**/*.h", - "runtime/backend/**/*.h", - "extension/**/*.h", + "executorch/runtime/**/*.h", + "executorch/runtime/backend/**/*.h", + "executorch/extension/**/*.h", ], allow_empty = True, ), @@ -53,8 +130,9 @@ cc_library( "C10_USING_CUSTOM_GENERATED_MACROS", ], include_prefix = "executorch", + strip_include_prefix = "executorch", includes = [ - "runtime/core/portable_type/c10", + "executorch/runtime/core/portable_type/c10", ], ) @@ -70,7 +148,7 @@ cc_library( name = "executorch_file_data_loader", # buildifier: disable=constant-glob srcs = glob( - ["extension/data_loader/file_data_loader.cpp"], + ["executorch/extension/data_loader/file_data_loader.cpp"], allow_empty = True, ), deps = [