From 78b66f316c30eef0327ce09cd098970c1faf1826 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Wed, 2 Sep 2026 19:41:27 -0700 Subject: [PATCH 01/22] Pin katex to the last release the docs image's node can run Every Docker image build fails at the docs dependency step: error commander@15.0.0: The engine "node" is incompatible with this module. Expected version ">=22.12.0". Got "16.20.2" error Found incompatible module. katex was installed unpinned, so it resolved to whatever was newest. Version 0.18.5, published a few days ago, changed its dependency from commander 8 to commander 15, and commander 15 requires node 22.12 or newer. This script installs node 16, so the install now refuses to run. Nothing in this repo changed; the newest version of a dependency moved out from under it. Pin katex to 0.18.4, the last release that depends on commander 8, which runs on the node this image installs. This is the same fix, with the same comment, that pytorch/pytorch applied to its copy of this script earlier today. Keeping the two in step matters here because the file is otherwise near identical between the repos, so a reader comparing them should not find two different answers to one problem. It is worth noting the unpinned install is the underlying cause, and node 16 has been out of support for a while. Moving the image to a current node is a larger change with a wider blast radius, and it does not belong in a pull request about something else. This restores the build. Test Plan: Read the metadata from the npm registry rather than inferring it. katex 0.18.5 depends on commander ^15.0.0 and 0.18.4 depends on commander ^8.3.0. commander 15.0.0 declares node >=22.12.0, and the newest commander 8, which is what ^8.3.0 resolves to, declares node >=12. So the pinned version's dependency runs on node 16 and the unpinned one cannot. Confirmed this is not caused by the other change in this pull request: the same failure, with the same message, appears on the docs step of main's own image build from the day before this branch existed, and the previous run of that workflow succeeded before 0.18.5 was published. Syntax-checked the script. Whether the built image renders documentation correctly with the slightly older katex is not something I can check locally, since it needs the full image build. CI covers it. --- .ci/docker/common/install_docs_reqs.sh | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/.ci/docker/common/install_docs_reqs.sh b/.ci/docker/common/install_docs_reqs.sh index ea54d90523e..2794ffb8fc9 100755 --- a/.ci/docker/common/install_docs_reqs.sh +++ b/.ci/docker/common/install_docs_reqs.sh @@ -20,7 +20,9 @@ if [ -n "$BUILD_DOCS" ]; then apt-get update apt-get install -y --no-install-recommends yarn - yarn global add katex --prefix /usr/local + # katex 0.18.5 requires commander@15 / node >= 22.12; pin to the last + # release compatible with the node 16 installed above + yarn global add katex@0.18.4 --prefix /usr/local sudo apt-get -y install doxygen From a4f37244ea3868af8b13887870b3cab3e27b2171 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Wed, 2 Sep 2026 19:46:36 -0700 Subject: [PATCH 02/22] Build PyTorch from source on macOS the way PyTorch now expects The macOS jobs build PyTorch from source whenever the cached wheel for the pinned version is not in S3 yet, which is always the case the first time a new pin is used. That fallback calls: python setup.py bdist_wheel PyTorch 2.14 no longer supports it. It now builds through scikit-build-core using the standard PEP 517 interface, and the command exits with an error naming its own replacement: error: `python setup.py bdist_wheel` is deprecated: PyTorch is built with scikit-build-core via the standard PEP 517 interface (pyproject.toml), and setup.py is no longer part of the build. Use the standard build frontend instead, with isolation off so the build uses the build requirements installed a few lines above rather than fetching its own copies. The Docker images already build PyTorch this way, and their code carries the same reasoning. Only this macOS path still used the old command, so this brings the two into step. Two details worth stating. The wheel filename does not change, because the version comes from PyTorch's own version file and git hash rather than from the build frontend, so the S3 cache key and the invariant that checks it are unaffected. And the output still lands in `dist/`, which is where the lines after this read from. Test Plan: Read the error from a real CI run rather than reasoning about it: the macOS unit test job on this branch reached the source build after a cache miss, and failed on exactly this command with the message above. Confirmed the cache miss itself was correct behavior, not a defect. The key it asked for names the short hash of the pinned PyTorch branch head, so it was looking for the right wheel, and it 404s only because no run has built and uploaded that version yet. Confirmed PyTorch's `requirements-build.txt`, which the lines above already install, covers what the PEP 517 build needs: it lists scikit-build-core, setuptools, cmake, ninja, numpy, packaging, pyyaml and six, matching the build-system requires in PyTorch's own pyproject.toml. Syntax-checked the script. I cannot run a full PyTorch source build on this machine, so whether the build completes is for CI to establish. What is verified here is that the command it previously used refuses to run at all on this pin. --- .ci/scripts/utils.sh | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/.ci/scripts/utils.sh b/.ci/scripts/utils.sh index 234e162e48e..6ce21849398 100644 --- a/.ci/scripts/utils.sh +++ b/.ci/scripts/utils.sh @@ -106,8 +106,8 @@ install_pytorch_and_domains() { local python_version=$(python -c 'import platform; v=platform.python_version_tuple(); print(f"{v[0]}{v[1]}")') local torch_release=$(cat version.txt) # Download key must match the upload key below (basename of dist/*.whl, - # which always carries setup.py's resolved +gitHASH). Branch-ref pins - # like `release/2.13` would otherwise produce `+gitrelease` here and + # which always carries the build's resolved +gitHASH). Branch-ref pins + # like `release/2.14` would otherwise produce `+gitrelease` here and # never hit the cache. local torch_short_hash=$(git rev-parse --short=7 HEAD) local torch_wheel_path="cached_artifacts/pytorch/executorch/pytorch_wheels/${system_name}/${python_version}" @@ -135,10 +135,16 @@ install_pytorch_and_domains() { if [[ "$(uname -m)" == "aarch64" ]]; then export BUILD_IGNORE_SVE_UNAVAILABLE=1 fi - USE_DISTRIBUTED=1 python setup.py bdist_wheel + # PyTorch no longer supports "python setup.py bdist_wheel"; it builds + # through scikit-build-core (PEP 517). Build with the standard frontend and + # keep isolation off, so the build uses the requirements-build.txt deps + # installed just above rather than fetching its own copies. This matches + # how the Docker images build PyTorch. + pip install build + USE_DISTRIBUTED=1 python -m build --wheel --no-isolation pip install "$(echo dist/*.whl)" - # Invariant: the basename setup.py just produced must match the cache + # Invariant: the basename the build just produced must match the cache # URL we'd reconstruct on the next run. If they diverge (someone edits # torch_wheel_name above, or PyTorch renames its wheels), the cache # will silently miss and every macOS run will fall back to a ~30-min From 38247071fdccc10293d8d508359cb75004fdc3ca Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Wed, 2 Sep 2026 19:51:00 -0700 Subject: [PATCH 03/22] Run the ROCm jobs on the ROCm version that has this PyTorch The ROCm jobs install PyTorch from a per-version index built from the ROCm version in their matrix: ERROR: No matching distribution found for torch==2.14.0 PyTorch publishes 2.14 for ROCm 7.2 and not for 7.1, which is what these jobs ask for. 2.13 was published for 7.1, so the jobs passed before the pin moved and cannot pass after it while still pointing at 7.1. Move the three job matrices to 7.2, and the script's own default with them so a local run and CI agree. Test Plan: Checked the indexes rather than assuming which versions exist. The test channel for ROCm 7.2 carries torch 2.14.0 and ROCm 7.1 does not, while 7.1 does carry 2.13.0, which is why this surfaces only now. Checked the two other things the jobs need at the new version. The builder image `pytorch/manylinux2_28-builder:rocm7.2` exists in the registry, and the pinned torchao nightly publishes a `+rocm7.2` wheel, which matters because the script builds that wheel's URL by hand rather than resolving it from an index. Confirmed no other reference to the old version is left in the workflow or the script. Parsed the workflow as YAML and syntax-checked the script. I have no ROCm hardware, so whether the tests pass on 7.2 is for CI to establish. What is verified here is that the packages the jobs install exist at the new version and did not at the old one. --- .ci/scripts/test-rocm-aoti.sh | 2 +- .ci/scripts/test-rocm-voxtral.sh | 2 +- .github/workflows/rocm.yml | 6 +++--- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/.ci/scripts/test-rocm-aoti.sh b/.ci/scripts/test-rocm-aoti.sh index 0f4ac9d826f..00592bc4bf6 100644 --- a/.ci/scripts/test-rocm-aoti.sh +++ b/.ci/scripts/test-rocm-aoti.sh @@ -7,7 +7,7 @@ set -euo pipefail -ROCM_VERSION="${ROCM_VERSION:-7.1}" +ROCM_VERSION="${ROCM_VERSION:-7.2}" ROCM_PATH="${ROCM_PATH:-/opt/rocm}" PYTORCH_ROCM_INDEX="${PYTORCH_ROCM_INDEX:-https://download.pytorch.org/whl/test/rocm${ROCM_VERSION}}" TORCHAO_ROCM_WHEEL_BASE="${TORCHAO_ROCM_WHEEL_BASE:-https://download.pytorch.org/whl/nightly/rocm${ROCM_VERSION}}" diff --git a/.ci/scripts/test-rocm-voxtral.sh b/.ci/scripts/test-rocm-voxtral.sh index cd0562c3808..eb00c1efc73 100644 --- a/.ci/scripts/test-rocm-voxtral.sh +++ b/.ci/scripts/test-rocm-voxtral.sh @@ -7,7 +7,7 @@ set -euo pipefail -ROCM_VERSION="${ROCM_VERSION:-7.1}" +ROCM_VERSION="${ROCM_VERSION:-7.2}" ROCM_PATH="${ROCM_PATH:-/opt/rocm}" EXPECTED_ROCM_ARCH="${EXPECTED_ROCM_ARCH:-gfx950}" EXPECTED_WARP_SIZE="${EXPECTED_WARP_SIZE:-64}" diff --git a/.github/workflows/rocm.yml b/.github/workflows/rocm.yml index 97d180885f7..154454eb145 100644 --- a/.github/workflows/rocm.yml +++ b/.github/workflows/rocm.yml @@ -169,7 +169,7 @@ jobs: strategy: fail-fast: false matrix: - rocm-version: ["7.1"] + rocm-version: ["7.2"] uses: pytorch/test-infra/.github/workflows/linux_job_v2.yml@main permissions: id-token: write @@ -206,7 +206,7 @@ jobs: strategy: fail-fast: false matrix: - rocm-version: ["7.1"] + rocm-version: ["7.2"] with: timeout: 180 no-sudo: true @@ -248,7 +248,7 @@ jobs: strategy: fail-fast: false matrix: - rocm-version: ["7.1"] + rocm-version: ["7.2"] with: timeout: 180 no-sudo: true From 444aea90520d1967220256630df49d70b529153e Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Wed, 2 Sep 2026 20:18:05 -0700 Subject: [PATCH 04/22] Accept the derived shim spelling for the Metal custom ops Every Metal export fails at lowering: RuntimeError: Method forward missing fallback kernels (1 total): - aoti_torch_mps_gated_delta_rule Please add them to the AOTI backend. Three of the backend's custom operators hit this: the gated delta rule, the qmv gather, and the 4-bit linear. The backend keeps a list of the fallback kernels it can serve, and export fails if the graph asks for one that is not on it. All three operators are already on that list, spelled the way they are registered, for example `metal::gated_delta_rule`. Inductor also derives a shim name for a custom op, by taking the part after the namespace and prefixing the device, which turns that into `aoti_torch_mps_gated_delta_rule`, and it is that derived name which reaches the check for these three. The function that derives the name is unchanged between the previous pin and this one, so this is a gap in the list that the new pin exposed rather than a rename to follow. Add the derived spellings alongside the registered ones. Which of the two arrives depends on the path Inductor takes for a given op, which is not ExecuTorch's decision to make, so accepting both is the honest fix. Nothing is newly permitted: the same three operators are involved either way. This is only about the name used at export time. The runtime already implements all three under exactly these symbols, in `op_gated_delta_rule.mm`, `op_gather_qmv.mm` and `op_linear_4bit.mm`, so the generated code has something to call and this does not defer the failure to load time. Test Plan: Took the three names from the CI failures rather than deriving them by hand, and confirmed each is the derived form of an entry already on the list. Confirmed each of the three is defined in the Metal runtime, under that exact symbol, inside an `extern "C"` block so the generated wrapper can link it. That is what distinguishes this from silencing a real gap: had any symbol been absent, allowing it through export would only move the error to load time. Checked the shim-name derivation in both PyTorch branches. The function that builds it is unchanged between them, which is why this is about which path an op takes rather than a renaming. Parsed the changed file. I have no Apple GPU in this environment, so whether the exported models now run is for CI to establish. What is verified is that the names the failures asked for exist in the runtime. --- backends/apple/metal/metal_backend.py | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/backends/apple/metal/metal_backend.py b/backends/apple/metal/metal_backend.py index 57ca0ddf83e..b301346ed6b 100644 --- a/backends/apple/metal/metal_backend.py +++ b/backends/apple/metal/metal_backend.py @@ -39,9 +39,16 @@ def get_supported_fallback_kernels(cls) -> Dict[str, Any]: "at::_ops::_scaled_dot_product_attention_math_for_mps::call": None, "at::_ops::_scaled_dot_product_attention_math_for_mps_v2::call": None, "torchao::_linear_fp_act_4bit_weight": None, + # Each custom op appears twice: once as registered, once under the shim + # name Inductor derives for it. Which spelling reaches this list + # depends on the path Inductor takes, and that is not ours to control, + # so accept both. + "aoti_torch_mps__linear_fp_act_4bit_weight": None, "at::_ops::topk::call": None, "metal::gather_qmv": None, + "aoti_torch_mps_gather_qmv": None, "metal::gated_delta_rule": None, + "aoti_torch_mps_gated_delta_rule": None, } @classmethod From b0e58611e088f085f656b0d13e8e78319c7db005 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Wed, 2 Sep 2026 21:06:06 -0700 Subject: [PATCH 05/22] Build only the ATen-mode Buck targets as C++20 Anything built against ATen fails to compile: c10/util/intrusive_ptr.h:775: error: no type named 'strong_ordering' in namespace 'std' ATen/core/TensorBase.h:1031: error: unknown type name 'requires' Both are C++20. PyTorch 2.14 uses them with no feature guard, and the wrapper pinned every target to C++17. Decide the standard per target instead: C++20 where the target compiles against ATen, C++17 everywhere else. The test path next door already made this distinction and says in its own comment that non-ATen targets are pinned to C++17 for embedded, so libraries and binaries were the odd ones out. Three details decide whether this works, and each cost a round to find: The answer has to be computed before the shared kwargs patch runs. That patch turns external dependency entries into real dependencies and removes the original list, so a check that reads the list afterwards always answers no. The flag has to go on the target's own compile flags. The per-language entry the wrapper already used reaches an earlier pass, and matched the standard it was restating, so nothing had ever shown it could raise one. The check keys on the exact dependency names the build maps onto libtorch, not on a substring of the label, because every label in this project contains "torch" inside "executorch". The decision lives in the shared wrapper rather than the interface layer beside it, because that layer varies per environment and its signature is fixed. Nothing changes for consumers. The standard the shipped runtime requires is still C++17, and the CMake build still sets it. Test Plan: Read the requirement out of the headers on both branches rather than inferring it from the error. Extracted the wrapper functions and ran them, with `type()` behaving as Starlark's does, over the target from the failure and over targets that must not move, removing the dependency list first so the sequence matches what the wrapper really does. ATen targets come away at C++20 through all three paths, library, binary and test; a plain library and the portable kernels come away with C++17. That probe is what caught the substring problem, and it is also what MISSED the ordering problem twice: calling the helper directly cannot see a step that runs before it. The Buck job is what found that, both times. Anyone extending this should exercise the whole path, not the helper. CI is the real check here: unittest-buck builds the generated ATen kernel libraries that were failing. I cannot run this Buck build on this machine. --- .../xplat/executorch/build/env_interface.bzl | 17 ++++++++++++----- .../xplat/executorch/build/runtime_wrapper.bzl | 17 +++++++++++++++-- 2 files changed, 27 insertions(+), 7 deletions(-) diff --git a/shim_et/xplat/executorch/build/env_interface.bzl b/shim_et/xplat/executorch/build/env_interface.bzl index c613243b537..9815a504025 100644 --- a/shim_et/xplat/executorch/build/env_interface.bzl +++ b/shim_et/xplat/executorch/build/env_interface.bzl @@ -160,14 +160,21 @@ def _patch_headers(kwargs): def _patch_pp_flags(kwargs): return kwargs -def _patch_cxx_compiler_flags(kwargs): - """CXX Compiler flags to enable C++17 features.""" +def _patch_cxx_compiler_flags(kwargs, aten_mode = False): + """Pins the C++ standard a target compiles with. + + C++17 by default, which is what the runtime this project ships requires. + ATen-mode targets are raised to C++20, because PyTorch's headers require it as + of 2.14: c10/util/intrusive_ptr.h defines operator<=> and returns + std::strong_ordering with no feature guard. + """ + std = "-std=c++20" if aten_mode else "-std=c++17" if "lang_compiler_flags" not in kwargs: - kwargs["lang_compiler_flags"] = {"cxx_cpp_output": ["-std=c++17"]} + kwargs["lang_compiler_flags"] = {"cxx_cpp_output": [std]} elif "cxx_cpp_output" not in kwargs["lang_compiler_flags"]: - kwargs["lang_compiler_flags"]["cxx_cpp_output"] = ["-std=c++17"] + kwargs["lang_compiler_flags"]["cxx_cpp_output"] = [std] else: - kwargs["lang_compiler_flags"]["cxx_cpp_output"].append("-std=c++17") + kwargs["lang_compiler_flags"]["cxx_cpp_output"].append(std) return kwargs # buildifier: disable=unused-variable diff --git a/shim_et/xplat/executorch/build/runtime_wrapper.bzl b/shim_et/xplat/executorch/build/runtime_wrapper.bzl index e84af76f3d9..17d6ffc8c2d 100644 --- a/shim_et/xplat/executorch/build/runtime_wrapper.bzl +++ b/shim_et/xplat/executorch/build/runtime_wrapper.bzl @@ -267,6 +267,19 @@ def _patch_kwargs_cxx(kwargs): env.remove_platform_specific_args(kwargs) return _patch_kwargs_common(kwargs) +def _is_aten_target(kwargs): + """Whether a target compiles against ATen, and so needs C++20. + + Keyed on the external dep names the build maps onto libtorch, rather than on + a substring of the label: every ExecuTorch label contains "torch". + """ + aten_external_deps = ["c10", "libtorch", "libtorch_python", "torch-core-cpp"] + for key in ["external_deps", "exported_external_deps"]: + for dep in kwargs.get(key) or []: + if dep in aten_external_deps: + return True + return False + def _cxx_library_common(*args, **kwargs): _patch_kwargs_cxx(kwargs) _patch_build_mode_flags(kwargs) @@ -274,7 +287,7 @@ def _cxx_library_common(*args, **kwargs): env.patch_platform_build_mode_flags(kwargs) env.patch_headers(kwargs) env.patch_pp_flags(kwargs) - env.patch_cxx_compiler_flags(kwargs) + env.patch_cxx_compiler_flags(kwargs, aten_mode = _is_aten_target(kwargs)) env.patch_force_static(kwargs) env.cxx_library(*args, **kwargs) @@ -297,7 +310,7 @@ def _cxx_binary_helper(*args, **kwargs): _patch_kwargs_cxx(kwargs) _patch_build_mode_flags(kwargs) env.patch_platform_build_mode_flags(kwargs) - env.patch_cxx_compiler_flags(kwargs) + env.patch_cxx_compiler_flags(kwargs, aten_mode = _is_aten_target(kwargs)) env.cxx_binary(*args, **kwargs) From e391761f49bf3868daf7f7065df9df5f07d55a96 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Wed, 2 Sep 2026 19:19:16 -0700 Subject: [PATCH 06/22] Bump the PyTorch pin to 2.14 PyTorch 2.14.0 is released, so move the pin from 2.13 to it. The tree is at 1.5.0 and mid-cycle, and the pin is normally moved to the current stable release, so this is that step. What changes: - the CI pin, from the `release/2.13` branch of PyTorch to `release/2.14` - the torch version literals, which live in three places that nothing keeps in agreement: `torch_pin.py`, `install_requirements.py`, and the Windows workflow - torchvision, from 0.28 to 0.29, which is the release that pins `torch==2.14.0` exactly. torchaudio stays at 2.11.0, still the newest published, since it no longer tracks torch's version train - the vendored c10 headers, re-synced from the new branch CI compares the vendored copy against PyTorch's own tree and requires it byte for byte, so the pin cannot move without it. Eight of the twenty-five headers changed. One of them moved rather than changed: `complex_utils.h` is now under `torch/headeronly/util/` instead of `c10/util/`, and its contents moved into the `torch::headeronly` namespace. `c10/util/complex.h` re-exports `is_complex` and `scalar_value_type` into `c10`, so code spelling those with a `c10::` prefix keeps working and no caller in this repo had to change. Two build references had to follow the move: the Buck header list for the `c10` target, which names its headers one by one, and the wheel test's list of headers that refuse to be included directly. The `torch_headeronly` target globs, so it picked the file up on its own. This is not a pure header refresh. `overflows()` decides whether a value fits a narrower type, and the new version answers differently for float to integer casts, in both directions: - a fractional value just past an integer limit is now accepted, so filling an int8 tensor with 127.5 used to be refused and now gives 127 - a value exactly at 2^63 cast to int64 is now refused, where it used to be accepted and wrapped to the lowest int64 The portable kernels reach this through `check_overflow_cast`, so `full`, `full_like`, `fill`, `scalar_tensor`, `hardtanh`, `leaky_relu`, `scatter` and `constant_pad_nd` all inherit it. The header is a faithful copy of upstream, so the change is not ours to undo, but it should be visible to anyone reading this rather than buried in a header diff. The existing overflow tests all use whole numbers or floating point output types, so none of them covered the branch that moved. Add a case at 2^63 for int64, which pins the direction that used to wrap silently. One narrowing comes with the new headers rather than from this change. The `complex` specialization they add does not carry the one-argument constructor, the `value_type` alias, the scalar compound operators or division, so code using those against that one type stops compiling, and plain default construction leaves it uninitialized. Nothing here uses the type. The copy has to match upstream exactly, so this belongs upstream and cannot be repaired locally without failing the comparison the pin depends on. Test Plan: Ran the header check CI runs, `.ci/scripts/compare_dirs.sh`, against a real `release/2.14` checkout for both vendored trees: exit 0 on each. Confirmed it fails without the re-sync, exit 1, naming `complex_utils.h` as present here and absent upstream, so the check does exercise this change. Compiled a program against the re-synced headers that instantiates `c10::complex`, `c10::is_complex` and `c10::scalar_value_type`, to confirm the two moved symbols still resolve under their old names. It builds and runs. Also compiled a translation unit including both these headers and ExecuTorch's own scalar type header, to check the include graph still works from inside the repo. Measured the behaviour change rather than reading it off the diff. Compiled `overflows()` from both branches and compared: int8 from 127.5 goes from refused to accepted, likewise uint8 from 255.5 and int32 from 2147483647.5, while int64 from 2^63 and uint64 from 2^64 go from accepted to refused. int8 from -128.5 is unchanged, so the lower bound did not move. Confirmed the new test discriminates. The value it uses is accepted by the old header and refused by the new one, so the test fails on the pin this change replaces and passes on the one it installs. A test that passed either way would prove nothing. Resolved the three pinned versions together from the channel CI installs from, and got 2.14.0, 0.29.0 and 2.11.0 with no conflict. Installed the pinned torchao nightly on top of torch 2.14 and imported the surfaces past pin bumps have broken on, `torchao.quantization`, `quantize_` and `torchao.dtypes`, all of which import. Compared the AOTI shim surface between the two branches. Both declare the same 148 entry points, so no new shim is added upstream. That is narrower than it sounds, and CI later showed why: the Metal backend still needed a change, because three of its custom operators began arriving at the fallback check under a derived name rather than the registered one. That is a later commit on this branch, not something this comparison could have caught. Syntax-checked the changed shell, Python and workflow files. Not covered locally: the Docker image build, the CUDA and Qualcomm jobs, and the model export suites, none of which run on this machine. The new test case was verified at the level of the helper it exercises rather than by running the kernel test binaries, which need a full build. --- .ci/docker/ci_commit_pins/pytorch.txt | 2 +- .ci/docker/common/install_pytorch.sh | 2 +- .ci/scripts/utils.sh | 2 +- .ci/scripts/wheel/test_cpp_sdk.py | 2 +- .github/workflows/windows-msvc.yml | 2 +- install_requirements.py | 4 +- kernels/test/ScalarOverflowTestMacros.h | 55 +++++++----- .../core/portable_type/c10/c10/targets.bzl | 1 - .../core/portable_type/c10/c10/util/complex.h | 17 ++-- .../c10/c10/util/llvmMathExtras.h | 1 + .../portable_type/c10/c10/util/overflows.h | 17 +++- .../c10/torch/headeronly/macros/Macros.h | 9 ++ .../c10/torch/headeronly/util/Half.h | 2 +- .../torch/headeronly/util/TypeSafeSignMath.h | 90 +++++-------------- .../c10/torch/headeronly/util/complex.h | 60 +++++++++++++ .../headeronly}/util/complex_utils.h | 8 +- torch_pin.py | 2 +- 17 files changed, 158 insertions(+), 118 deletions(-) rename runtime/core/portable_type/c10/{c10 => torch/headeronly}/util/complex_utils.h (80%) diff --git a/.ci/docker/ci_commit_pins/pytorch.txt b/.ci/docker/ci_commit_pins/pytorch.txt index 401a0594d98..212c7580b6f 100644 --- a/.ci/docker/ci_commit_pins/pytorch.txt +++ b/.ci/docker/ci_commit_pins/pytorch.txt @@ -1 +1 @@ -release/2.13 +release/2.14 \ No newline at end of file diff --git a/.ci/docker/common/install_pytorch.sh b/.ci/docker/common/install_pytorch.sh index 0ac5e79cf4a..401397635ff 100755 --- a/.ci/docker/common/install_pytorch.sh +++ b/.ci/docker/common/install_pytorch.sh @@ -90,7 +90,7 @@ install_pytorch_and_domains() { # Grab the pinned audio and vision commits from PyTorch TORCHAUDIO_VERSION=release/2.11 export TORCHAUDIO_VERSION - TORCHVISION_VERSION=release/0.28 + TORCHVISION_VERSION=release/0.29 export TORCHVISION_VERSION install_domains diff --git a/.ci/scripts/utils.sh b/.ci/scripts/utils.sh index 6ce21849398..3e97e431072 100644 --- a/.ci/scripts/utils.sh +++ b/.ci/scripts/utils.sh @@ -184,7 +184,7 @@ install_pytorch_and_domains() { # Grab the pinned audio and vision commits from PyTorch TORCHAUDIO_VERSION=release/2.11 export TORCHAUDIO_VERSION - TORCHVISION_VERSION=release/0.28 + TORCHVISION_VERSION=release/0.29 export TORCHVISION_VERSION install_domains diff --git a/.ci/scripts/wheel/test_cpp_sdk.py b/.ci/scripts/wheel/test_cpp_sdk.py index 29b2021d398..202a400e6c7 100644 --- a/.ci/scripts/wheel/test_cpp_sdk.py +++ b/.ci/scripts/wheel/test_cpp_sdk.py @@ -991,7 +991,7 @@ def test_every_shipped_header_compiles(work_dir: Path) -> None: # These say in their own text that they must not be included directly, and name the header to # include instead. Including one anyway is a use error rather than a packaging defect. "c10/util/complex_math.h", - "c10/util/complex_utils.h", + "torch/headeronly/util/complex_utils.h", ) source = work_dir / "header_probe.cpp" diff --git a/.github/workflows/windows-msvc.yml b/.github/workflows/windows-msvc.yml index bde38a8288f..269f2cf2381 100644 --- a/.github/workflows/windows-msvc.yml +++ b/.github/workflows/windows-msvc.yml @@ -91,7 +91,7 @@ jobs: - name: Install build dependencies shell: pwsh - run: python -m pip install pyyaml torch==2.13.0 --extra-index-url https://download.pytorch.org/whl/test/cpu + run: python -m pip install pyyaml torch==2.14.0 --extra-index-url https://download.pytorch.org/whl/test/cpu - name: Build ExecuTorch shell: pwsh diff --git a/install_requirements.py b/install_requirements.py index b7d220179d3..c4a934a0180 100644 --- a/install_requirements.py +++ b/install_requirements.py @@ -66,7 +66,7 @@ def install_requirements(use_pytorch_nightly): # Setting use_pytorch_nightly to false to test the pinned PyTorch commit. Note # that we don't need to set any version number there because they have already # been installed on CI before this step, so pip won't reinstall them - ("torch==2.13.0" if use_pytorch_nightly else "torch"), + ("torch==2.14.0" if use_pytorch_nightly else "torch"), f"torchao=={TORCHAO_NIGHTLY_VERSION}", ] @@ -134,7 +134,7 @@ def install_optional_example_requirements(use_pytorch_nightly): print("Installing torch domain libraries") DOMAIN_LIBRARIES = [ - ("torchvision==0.28.0" if use_pytorch_nightly else "torchvision"), + ("torchvision==0.29.0" if use_pytorch_nightly else "torchvision"), ("torchaudio==2.11.0" if use_pytorch_nightly else "torchaudio"), ] # Then install domain libraries diff --git a/kernels/test/ScalarOverflowTestMacros.h b/kernels/test/ScalarOverflowTestMacros.h index 46a2425b0fa..f8163254fc8 100644 --- a/kernels/test/ScalarOverflowTestMacros.h +++ b/kernels/test/ScalarOverflowTestMacros.h @@ -11,28 +11,35 @@ // Macro to generate scalar overflow test cases for a given test suite. // The test suite must have a method called expect_bad_scalar_value_dies // that takes a template parameter for ScalarType and a Scalar value. -#define GENERATE_SCALAR_OVERFLOW_TESTS(TEST_SUITE_NAME) \ - TEST_F(TEST_SUITE_NAME, ByteTensorTooLargeScalarDies) { \ - /* Cannot be represented by a uint8_t. */ \ - expect_bad_scalar_value_dies(256); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, CharTensorTooSmallScalarDies) { \ - /* Cannot be represented by a int8_t. */ \ - expect_bad_scalar_value_dies(-129); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, ShortTensorTooLargeScalarDies) { \ - /* Cannot be represented by a int16_t. */ \ - expect_bad_scalar_value_dies(32768); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, FloatTensorTooSmallScalarDies) { \ - /* Cannot be represented by a float. */ \ - expect_bad_scalar_value_dies(-3.41e+38); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, FloatTensorTooLargeScalarDies) { \ - /* Cannot be represented by a float. */ \ - expect_bad_scalar_value_dies(3.41e+38); \ +#define GENERATE_SCALAR_OVERFLOW_TESTS(TEST_SUITE_NAME) \ + TEST_F(TEST_SUITE_NAME, ByteTensorTooLargeScalarDies) { \ + /* Cannot be represented by a uint8_t. */ \ + expect_bad_scalar_value_dies(256); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, CharTensorTooSmallScalarDies) { \ + /* Cannot be represented by a int8_t. */ \ + expect_bad_scalar_value_dies(-129); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, ShortTensorTooLargeScalarDies) { \ + /* Cannot be represented by a int16_t. */ \ + expect_bad_scalar_value_dies(32768); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, FloatTensorTooSmallScalarDies) { \ + /* Cannot be represented by a float. */ \ + expect_bad_scalar_value_dies(-3.41e+38); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, FloatTensorTooLargeScalarDies) { \ + /* Cannot be represented by a float. */ \ + expect_bad_scalar_value_dies(3.41e+38); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, LongTensorTooLargeScalarDies) { \ + /* 2^63 is one past the largest int64_t. It is exactly \ + * representable as a double, so the conversion is silent \ + * unless the range check rejects it. */ \ + expect_bad_scalar_value_dies(9223372036854775808.0); \ } diff --git a/runtime/core/portable_type/c10/c10/targets.bzl b/runtime/core/portable_type/c10/c10/targets.bzl index 675c04f97a1..a39b2b9f566 100644 --- a/runtime/core/portable_type/c10/c10/targets.bzl +++ b/runtime/core/portable_type/c10/c10/targets.bzl @@ -114,7 +114,6 @@ def define_common_targets(): "util/bit_cast.h", "util/complex.h", "util/complex_math.h", - "util/complex_utils.h", "util/floating_point_utils.h", "util/irange.h", "util/llvmMathExtras.h", diff --git a/runtime/core/portable_type/c10/c10/util/complex.h b/runtime/core/portable_type/c10/c10/util/complex.h index 4e699684bc3..f9849a94ced 100644 --- a/runtime/core/portable_type/c10/c10/util/complex.h +++ b/runtime/core/portable_type/c10/c10/util/complex.h @@ -31,19 +31,11 @@ C10_HOST_DEVICE T abs(const c10::complex& z) { #endif } -#if defined(USE_ROCM) -#define ROCm_Bug(x) -#else -#define ROCm_Bug(x) x -#endif - template C10_HOST_DEVICE T arg(const c10::complex& z) { - return ROCm_Bug(std)::atan2(std::imag(z), std::real(z)); + return std::atan2(std::imag(z), std::real(z)); } -#undef ROCm_Bug - template constexpr T norm(const c10::complex& z) { return z.real() * z.real() + z.imag() * z.imag(); @@ -73,6 +65,9 @@ constexpr c10::complex conj(const c10::complex& z) { #define C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H // math functions are included in a separate file #include // IWYU pragma: keep -// utilities for complex types -#include // IWYU pragma: keep #undef C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H + +namespace c10 { +using torch::headeronly::is_complex; +using torch::headeronly::scalar_value_type; +} // namespace c10 diff --git a/runtime/core/portable_type/c10/c10/util/llvmMathExtras.h b/runtime/core/portable_type/c10/c10/util/llvmMathExtras.h index da297241449..8ae5cde4f02 100644 --- a/runtime/core/portable_type/c10/c10/util/llvmMathExtras.h +++ b/runtime/core/portable_type/c10/c10/util/llvmMathExtras.h @@ -400,6 +400,7 @@ constexpr inline bool isShiftedUInt(uint64_t x) { N + S <= 64, "isShiftedUInt with N + S > 64 is too wide."); // Per the two static_asserts above, S must be strictly less than 64. So // 1 << S is not undefined behavior. + // NOLINTNEXTLINE(bugprone-chained-comparison) return isUInt(x) && (x % (UINT64_C(1) << S) == 0); } diff --git a/runtime/core/portable_type/c10/c10/util/overflows.h b/runtime/core/portable_type/c10/c10/util/overflows.h index 183a2f62a32..348ee5ffa40 100644 --- a/runtime/core/portable_type/c10/c10/util/overflows.h +++ b/runtime/core/portable_type/c10/c10/util/overflows.h @@ -61,13 +61,28 @@ template std::enable_if_t, bool> overflows( From f, bool strict_unsigned [[maybe_unused]] = false) { - using limit = std::numeric_limits::type>; + using ToScalar = typename scalar_value_type::type; + using limit = std::numeric_limits; if (limit::has_infinity && std::isinf(static_cast(f))) { return false; } if (!limit::has_quiet_NaN && (f != f)) { return true; } + if constexpr (std::is_integral_v) { + // limit::max() for wide integer types is NOT exactly representable in + // floating point (e.g. int64 max = 2^63-1 rounds up to 2^63), so `f > + // limit::max()` lets a just-out-of-range value like 2^63 slip through and + // then become INT64_MIN via static_cast. Compare against the + // exactly-representable upper bound max()+1 == 2^digits instead. lowest() + // is 0 or a negated power of two, so it stays exact. (digits-1 keeps the + // shift < 64 for the uint64 case; the *2 recovers 2^digits without a 1<<64 + // overflow.) + constexpr int digits = limit::digits; + constexpr From upper = + static_cast(uint64_t{1} << (digits - 1)) * From{2}; + return f < static_cast(limit::lowest()) || f >= upper; + } return f < limit::lowest() || f > limit::max(); } diff --git a/runtime/core/portable_type/c10/torch/headeronly/macros/Macros.h b/runtime/core/portable_type/c10/torch/headeronly/macros/Macros.h index cef99df3f56..08c4e9f1f84 100644 --- a/runtime/core/portable_type/c10/torch/headeronly/macros/Macros.h +++ b/runtime/core/portable_type/c10/torch/headeronly/macros/Macros.h @@ -123,6 +123,15 @@ #define C10_HAS_CPP_ATTRIBUTE(x) (0) #endif +/// Bind a returned reference/pointer's lifetime to a parameter (or *this) so +/// Clang can warn when it would dangle. Expands to nothing on compilers that +/// lack the attribute (e.g. non-clang, older nvcc). +#if C10_HAS_CPP_ATTRIBUTE(clang::lifetimebound) +#define C10_LIFETIMEBOUND [[clang::lifetimebound]] +#else +#define C10_LIFETIMEBOUND +#endif + #ifndef FBCODE_CAFFE2 /// DEPRECATED: Warn if a type or return value is discarded. #define C10_NODISCARD [[nodiscard]] diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/Half.h b/runtime/core/portable_type/c10/torch/headeronly/util/Half.h index e5aa622656c..401472357ec 100644 --- a/runtime/core/portable_type/c10/torch/headeronly/util/Half.h +++ b/runtime/core/portable_type/c10/torch/headeronly/util/Half.h @@ -213,7 +213,7 @@ C10_HOST_DEVICE inline float fp16_ieee_to_fp32_value(uint16_t h) { * Now, remember that denormalized half-precision numbers are represented as: * FP16 = mantissa * 2**(-24). * The trick is to construct a normalized single-precision number with the - * same mantissa and thehalf-precision input and with an exponent which would + * same mantissa and the half-precision input and with an exponent which would * scale the corresponding mantissa bits to 2**(-24). A normalized * single-precision floating-point number is represented as: FP32 = (1 + * mantissa * 2**(-23)) * 2**(exponent - 127) Therefore, when the biased diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h b/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h index c33a286bc5b..8e897957fee 100644 --- a/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h +++ b/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h @@ -14,20 +14,6 @@ C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion") namespace c10 { -/// Returns false since we cannot have x < 0 if x is unsigned. -template -inline constexpr bool is_negative( - const T& /*x*/, - std::true_type /*is_unsigned*/) { - return false; -} - -/// Returns true if a signed variable x < 0 -template -inline constexpr bool is_negative(const T& x, std::false_type /*is_unsigned*/) { - return x < T(0); -} - /// Returns true if x < 0 /// NOTE: Will fail on an unsigned custom type /// For the most part it's possible to fix this if @@ -35,19 +21,12 @@ inline constexpr bool is_negative(const T& x, std::false_type /*is_unsigned*/) { /// However, notably, c10::Half does not :-( template inline constexpr bool is_negative(const T& x) { - return is_negative(x, std::is_unsigned()); -} - -/// Returns the sign of an unsigned variable x as 0, 1 -template -inline constexpr int signum(const T& x, std::true_type /*is_unsigned*/) { - return T(0) < x; -} - -/// Returns the sign of a signed variable x as -1, 0, 1 -template -inline constexpr int signum(const T& x, std::false_type /*is_unsigned*/) { - return (T(0) < x) - (x < T(0)); + if constexpr (std::is_unsigned_v) { + // An unsigned value can never be less than zero. + return false; + } else { + return x < T(0); + } } /// Returns the sign of x as -1, 0, 1 @@ -57,7 +36,11 @@ inline constexpr int signum(const T& x, std::false_type /*is_unsigned*/) { /// However, notably, c10::Half does not :-( template inline constexpr int signum(const T& x) { - return signum(x, std::is_unsigned()); + if constexpr (std::is_unsigned_v) { + return T(0) < x; + } else { + return (T(0) < x) - (x < T(0)); + } } /// Returns true if a and b are not both negative @@ -86,53 +69,22 @@ inline constexpr bool greater_than_max(const T& x) { #pragma GCC diagnostic pop #endif -/// Returns true if x < lowest(Limit). Standard comparison -template -inline constexpr bool less_than_lowest( - const T& x, - std::false_type /*limit_is_unsigned*/, - std::false_type /*x_is_unsigned*/) { - return x < std::numeric_limits::lowest(); -} - -/// Returns false since all the limit is signed and therefore includes -/// negative values but x cannot be negative because it is unsigned -template -inline constexpr bool less_than_lowest( - const T& /*x*/, - std::false_type /*limit_is_unsigned*/, - std::true_type /*x_is_unsigned*/) { - return false; -} - -/// Returns true if x < 0, where 0 is constructed from T. -/// Limit is not signed, so its lower value is zero -template -inline constexpr bool less_than_lowest( - const T& x, - std::true_type /*limit_is_unsigned*/, - std::false_type /*x_is_unsigned*/) { - return x < T(0); -} - -/// Returns false sign both types are unsigned -template -inline constexpr bool less_than_lowest( - const T& /*x*/, - std::true_type /*limit_is_unsigned*/, - std::true_type /*x_is_unsigned*/) { - return false; -} - -/// Returns true if x is less than the lowest value of type T +/// Returns true if x is less than the lowest value of type Limit /// NOTE: Will fail on an unsigned custom type /// For the most part it's possible to fix this if /// the custom type has a constexpr constructor. /// However, notably, c10::Half does not : template inline constexpr bool less_than_lowest(const T& x) { - return less_than_lowest( - x, std::is_unsigned(), std::is_unsigned()); + if constexpr (std::is_unsigned_v) { + // x is unsigned, so it can never be below the lowest value of any type. + return false; + } else if constexpr (std::is_unsigned_v) { + // Limit is unsigned, so its lowest value is zero. + return x < T(0); + } else { + return x < std::numeric_limits::lowest(); + } } } // namespace c10 diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/complex.h b/runtime/core/portable_type/c10/torch/headeronly/util/complex.h index 733a22d5dbb..c349602dcf0 100644 --- a/runtime/core/portable_type/c10/torch/headeronly/util/complex.h +++ b/runtime/core/portable_type/c10/torch/headeronly/util/complex.h @@ -3,6 +3,7 @@ #include #include +#include #include #if defined(__CUDACC__) || defined(__HIPCC__) @@ -588,6 +589,60 @@ struct alignas(4) complex { } }; +template <> +struct alignas(4) complex { + BFloat16 real_; + BFloat16 imag_; + + // Constructors + complex() = default; + // BFloat16 constructor is not constexpr so the following constructor can't + // be constexpr + C10_HOST_DEVICE explicit inline complex( + const BFloat16& real, + const BFloat16& imag) + : real_(real), imag_(imag) {} + C10_HOST_DEVICE inline complex(const c10::complex& value) + : real_(value.real()), imag_(value.imag()) {} + + // Conversion operator + inline C10_HOST_DEVICE operator c10::complex() const { + return {real_, imag_}; + } + + constexpr C10_HOST_DEVICE BFloat16 real() const { + return real_; + } + constexpr C10_HOST_DEVICE BFloat16 imag() const { + return imag_; + } + + C10_HOST_DEVICE complex& operator+=( + const complex& other) { + real_ = static_cast(real_) + static_cast(other.real_); + imag_ = static_cast(imag_) + static_cast(other.imag_); + return *this; + } + + C10_HOST_DEVICE complex& operator-=( + const complex& other) { + real_ = static_cast(real_) - static_cast(other.real_); + imag_ = static_cast(imag_) - static_cast(other.imag_); + return *this; + } + + C10_HOST_DEVICE complex& operator*=( + const complex& other) { + auto a = static_cast(real_); + auto b = static_cast(imag_); + auto c = static_cast(other.real()); + auto d = static_cast(other.imag()); + real_ = a * c - b * d; + imag_ = a * d + b * c; + return *this; + } +}; + } // namespace c10 HIDDEN_NAMESPACE_BEGIN(torch, headeronly) @@ -614,3 +669,8 @@ using c10::complex_literals::operator""_id; HIDDEN_NAMESPACE_END(torch, headeronly) C10_CLANG_DIAGNOSTIC_POP() + +#define C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H +// utilities for complex types +#include // IWYU pragma: keep +#undef C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H diff --git a/runtime/core/portable_type/c10/c10/util/complex_utils.h b/runtime/core/portable_type/c10/torch/headeronly/util/complex_utils.h similarity index 80% rename from runtime/core/portable_type/c10/c10/util/complex_utils.h rename to runtime/core/portable_type/c10/torch/headeronly/util/complex_utils.h index 1ca105f1d0a..ddc66ffe776 100644 --- a/runtime/core/portable_type/c10/c10/util/complex_utils.h +++ b/runtime/core/portable_type/c10/torch/headeronly/util/complex_utils.h @@ -1,11 +1,13 @@ +#pragma once + #if !defined(C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H) #error \ - "c10/util/complex_utils.h is not meant to be individually included. Include c10/util/complex.h instead." + "torch/headeronly/util/complex_utils.h is not meant to be individually included. Include torch/headeronly/util/complex.h instead." #endif #include -namespace c10 { +HIDDEN_NAMESPACE_BEGIN(torch, headeronly) template struct is_complex : public std::false_type {}; @@ -31,7 +33,7 @@ struct scalar_value_type> { using type = T; }; -} // namespace c10 +HIDDEN_NAMESPACE_END(torch, headeronly) namespace std { diff --git a/torch_pin.py b/torch_pin.py index ca593b1ef05..f46d5b67ec0 100644 --- a/torch_pin.py +++ b/torch_pin.py @@ -1,2 +1,2 @@ -TORCH_VERSION = "2.13.0" +TORCH_VERSION = "2.14.0" # NIGHTLY_VERSION = "dev20260318" Temporarily pinning to stable release candidate. Revert https://github.com/pytorch/executorch/pull/18287 From 206db24e1fa22d2ba5131a9f7ade998500094704 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Wed, 2 Sep 2026 23:00:47 -0700 Subject: [PATCH 07/22] Accept the derived shim spelling for the CUDA int4 pack matmul Every quantized CUDA export fails at lowering: RuntimeError: Method forward missing fallback kernels (1 total): - aoti_torch_cuda__weight_int4pack_mm Please add them to the AOTI backend. This is the same situation the Metal backend is in one commit earlier, for the same reason. The backend lists the fallback kernels it can serve and export fails if the graph asks for one that is absent. This operator is already on the list, spelled the way it is registered, `at::_ops::_weight_int4pack_mm::call`. Inductor also derives a shim name for it, by taking the part after the namespace and prefixing the device, and it is that derived name which reaches the check here. Add the derived spelling next to the registered one. Nothing is newly permitted, and the runtime already implements this shim, in `backends/cuda/runtime/shims/int4mm.cu`, so the generated code has something to call. Which spelling arrives depends on the path Inductor takes for a given operator, not on the PyTorch version: the function that derives the name is unchanged between the previous pin and this one. So this is a gap in the list that the new pin exposed, rather than a rename to follow. Test Plan: Ran PyTorch 2.14's own name derivation over the entry and confirmed it produces the spelling in the error. Compared that function between the previous pin and this one: apart from an assert being rewritten as a raise, it is the same, which is why the commit does not claim the naming changed. The export itself is for CI, which is where the failure appeared. --- backends/cuda/cuda_backend.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/backends/cuda/cuda_backend.py b/backends/cuda/cuda_backend.py index 854ebf8f952..c1991f97b8e 100644 --- a/backends/cuda/cuda_backend.py +++ b/backends/cuda/cuda_backend.py @@ -669,6 +669,10 @@ def get_supported_fallback_kernels(cls) -> Dict[str, Any]: return {} return { "at::_ops::_weight_int4pack_mm::call": None, + # The same op under the shim name Inductor derives for it. Which + # spelling reaches this list depends on the path Inductor takes, and + # that is not ours to control, so accept both. + "aoti_torch_cuda__weight_int4pack_mm": None, "at::_ops::sort_stable::call": None, "aoti_torch_cuda_randint_low_out": None, "executorch_cuda::int4_plain_mm": None, From cc663a735e559f0a19c384095336d6449fa71d0c Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Wed, 2 Sep 2026 23:19:08 -0700 Subject: [PATCH 08/22] Build the Vulkan op tests as C++20, which ATen now requires The Vulkan operator tests fail to compile: ATen/ATen.h:5:2: error: C++20 or later compatible compiler is required to use ATen. PyTorch 2.14 raised ATen's own floor. The header used to reject anything below C++17 and now rejects anything below C++20, and these tests link libtorch to check operator results against eager PyTorch. Raise the standard on the targets that include those headers. This repository already does exactly that where the generated ATen-mode kernel library is defined, with the same one-line comment, so this follows a pattern that is already here rather than introducing one. Nothing changes for the runtime or for anyone consuming it. ATen only enters the build behind `USE_ATEN_LIB`, which no embedded preset sets, and the project default stays C++17. Test Plan: Read the requirement out of the header on both branches rather than inferring it from the error: the guard is `__cplusplus < 201703L` with a C++17 message on the previous branch and `__cplusplus < 202002L` with a C++20 message on this one. Confirmed the scope is limited before widening it. ATen reaches ExecuTorch only through `USE_ATEN_LIB`, which is set in exactly one place in the CMake build, the generated ATen-mode kernel library, and that target already sets C++20 for this reason. None of the bare-metal, Zephyr, ESP or RISC-V presets enable it. Checked the other CMake targets that link libtorch and left them alone, since only these tests failed. cmake-format is clean on the changed file. I cannot build the Vulkan tests on this machine, which needs a Vulkan SDK, so CI has to confirm they compile. --- backends/vulkan/test/op_tests/CMakeLists.txt | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/backends/vulkan/test/op_tests/CMakeLists.txt b/backends/vulkan/test/op_tests/CMakeLists.txt index 0f8456accf5..5facea5a14b 100644 --- a/backends/vulkan/test/op_tests/CMakeLists.txt +++ b/backends/vulkan/test/op_tests/CMakeLists.txt @@ -73,6 +73,8 @@ function(vulkan_op_test test_name test_src) add_executable(${test_name} ${test_src}) target_include_directories(${test_name} PRIVATE ${COMMON_INCLUDES}) + # ATen headers require C++20. + set_target_properties(${test_name} PROPERTIES CXX_STANDARD 20) target_link_libraries( ${test_name} PRIVATE GTest::gtest_main @@ -90,6 +92,8 @@ endfunction() if(TARGET vulkan_backend AND LIB_TORCH) add_library(test_utils ${CMAKE_CURRENT_SOURCE_DIR}/test_utils.cpp) target_include_directories(test_utils PRIVATE ${COMMON_INCLUDES}) + # ATen headers require C++20. + set_target_properties(test_utils PROPERTIES CXX_STANDARD 20) target_link_libraries( test_utils PRIVATE vulkan_backend ${LIB_TORCH} ${LIB_TORCH_CPU} ) From 15928ed5ed1d69e89e1a41aafda2508c72b87534 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 00:11:27 -0700 Subject: [PATCH 09/22] Let the PyTorch source build use the image's cmake, so it finds BLAS again The CI images build PyTorch from source, and after the pin moved they produced a PyTorch with no BLAS and no LAPACK. Tests that need it fail: RuntimeError: Calling torch.geqrf on a CPU tensor requires compiling PyTorch with LAPACK. Please use PyTorch built with LAPACK support. The build says as much, if you read the configure output: MKL could not be found. Defaulting to Eigen Cannot find a library with LAPACK API. Not using LAPACK. USE_BLAS : 0 Nothing about the image changed. PyTorch changed how it builds: 2.13 built through setuptools and 2.14 builds through scikit-build-core, and the two choose cmake differently. setuptools looked for cmake on PATH and found the one in the image. scikit-build-core prefers an importable pip cmake over PATH, and this script was installing one. That matters because cmake adds its own install root to the list of prefixes it searches. The image's cmake lives in the conda environment, alongside MKL and libomp, so a plain search finds them. A pip cmake lives in site-packages, where neither is, so the search fails, PyTorch falls back to Eigen, and LAPACK is dropped. The build does not fail when this happens; it just produces a weaker wheel. So stop installing cmake into the build environment and let the image's own cmake be used. It is pinned at 3.31 and satisfies what PyTorch asks for. Also assert on the result afterwards. This failure is silent by nature, and the first sign of it today was a linear algebra test failing much later in an unrelated job. Test Plan: Read scikit-build-core's own source rather than inferring the order: its cmake lookup yields the importable `cmake` module first and only then falls back to PATH. Reproduced it in a scratch environment: with pip cmake installed it selects the site-packages copy, and with it absent it selects the one on PATH. Confirmed the prefix mechanism by running `cmake --system-information` on two different cmake binaries and diffing: each one appends its own install root to the searched prefixes. That is what makes the choice of binary decide whether MKL is found. Confirmed the backend change is real, since it is the trigger: the build backend is `setuptools.build_meta` on the previous branch and `scikit_build_core.build` on this one. Checked that this is not specific to one architecture or one BLAS vendor. The aarch64 image, which installs conda openblas rather than MKL, degraded the same way, and both architectures were healthy on the previous pin. A fix aimed only at MKL or only at x86_64 would have left half of it broken. Confirmed the image's cmake is available when this runs: conda pins it at 3.31.2 and installs well before this script. Ran the guard against a healthy PyTorch 2.14 to be sure it passes rather than being vacuous, and it reports LAPACK present and completes a QR. Syntax-checked the script. I cannot build a CI image on this machine, so whether the rebuilt image now reports BLAS and LAPACK is for CI to confirm. The guard is there so that answer is unambiguous. --- .ci/docker/common/install_pytorch.sh | 20 ++++++++++++++++++-- 1 file changed, 18 insertions(+), 2 deletions(-) diff --git a/.ci/docker/common/install_pytorch.sh b/.ci/docker/common/install_pytorch.sh index 401397635ff..13fd20269de 100755 --- a/.ci/docker/common/install_pytorch.sh +++ b/.ci/docker/common/install_pytorch.sh @@ -76,17 +76,33 @@ install_pytorch_and_domains() { # the image compiler cannot satisfy. The venv inherits the image's # site-packages, so PyTorch still builds against the same numpy. # - # Keep the list in sync with pytorch/pyproject.toml [build-system].requires. + # Keep the list in sync with pytorch/pyproject.toml [build-system].requires, + # except for cmake, see below. local build_venv=/tmp/pytorch-build-venv rm -rf "${build_venv}" conda_run python -m venv --system-site-packages "${build_venv}" + # cmake is deliberately not installed here, so the conda cmake already in the + # image is used. scikit-build-core, which PyTorch builds with as of 2.14, + # prefers an importable pip cmake over anything on PATH, and cmake adds its own + # install root to CMAKE_SYSTEM_PREFIX_PATH. A pip cmake therefore searches + # site-packages, where MKL and libomp are not, and the build silently comes out + # with no BLAS and no LAPACK. conda_run "${build_venv}/bin/pip" install build "scikit-build-core>=1.0" \ - "setuptools>=77.0.0,<82" "cmake>=3.27,<4" ninja "packaging>=24.2" \ + "setuptools>=77.0.0,<82" ninja "packaging>=24.2" \ "typing-extensions>=4.10.0" pyyaml six conda_run "${build_venv}/bin/python" -m build --wheel --no-isolation rm -rf "${build_venv}" pip_install "$(echo dist/*.whl)" + # The build silently degrades rather than failing when it cannot find BLAS, so + # assert on the result. Run from / so the import resolves to the installed + # wheel and not to the source tree next to it. + (cd / && conda_run python -c " +import torch +assert torch._C.has_lapack, 'built without LAPACK' +torch.linalg.qr(torch.randn(4, 4)) +") + # Grab the pinned audio and vision commits from PyTorch TORCHAUDIO_VERSION=release/2.11 export TORCHAUDIO_VERSION From 6735cc46ed5a7b06fb8f1375eedfc7f53058fec6 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 00:13:58 -0700 Subject: [PATCH 10/22] Stop the image build scanning for C++20 modules it does not have scanners for Three Docker images fail to build PyTorch, all on the same file, for the same underlying reason in two flavours. The clang image: FAILED: [code=127] third_party/fmt/CMakeFiles/fmt.dir/src/format.cc.o.ddi /bin/sh: 1: CMAKE_CXX_COMPILER_CLANG_SCAN_DEPS-NOTFOUND: not found and the two GCC images: cc1plus: error: to generate dependencies you must specify either '-M' or '-MM' PyTorch compiles at C++20 as of 2.14. CMake responds by scanning every source for module imports before compiling it, which needs a scanner: `clang-scan-deps` for clang, and a supported invocation for GCC. The clang image installs clang and llvm but not the package that carries `clang-scan-deps`, and GCC rejects the scan outright. Neither PyTorch nor anything in these images uses C++20 modules, so the scan has nothing to find and only has to be turned off. Turning it off covers all three images. Adding the missing clang package would fix one of them and leave the two GCC ones failing, which is how this looked at first. Test Plan: Compared the three failing job logs. All three fail building the same fmt source, the clang one on a missing scanner binary and the two GCC ones on the scan invocation itself, which is what shows this is one problem rather than three. Confirmed the variable reaches CMake rather than being ignored. PyTorch forwards any environment variable beginning with `CMAKE_` to a CMake cache variable of the same name, which its own forwarding module states in those terms. Confirmed the images were healthy before the pin moved: the same image and job name built successfully on the previous pin. Confirmed this is not a duplicate of the cmake change next to it. That one picks which cmake runs; both cmakes are new enough to scan, so it does not disable scanning on its own. Syntax-checked the script. I cannot build an image on this machine, so CI has to confirm the three images now build. --- .ci/docker/common/install_pytorch.sh | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/.ci/docker/common/install_pytorch.sh b/.ci/docker/common/install_pytorch.sh index 13fd20269de..85a0661d08d 100755 --- a/.ci/docker/common/install_pytorch.sh +++ b/.ci/docker/common/install_pytorch.sh @@ -77,7 +77,8 @@ install_pytorch_and_domains() { # site-packages, so PyTorch still builds against the same numpy. # # Keep the list in sync with pytorch/pyproject.toml [build-system].requires, - # except for cmake, see below. + # except that cmake is deliberately left out, see below, and ninja is kept + # because the build needs a generator. local build_venv=/tmp/pytorch-build-venv rm -rf "${build_venv}" conda_run python -m venv --system-site-packages "${build_venv}" @@ -88,9 +89,14 @@ install_pytorch_and_domains() { # site-packages, where MKL and libomp are not, and the build silently comes out # with no BLAS and no LAPACK. conda_run "${build_venv}/bin/pip" install build "scikit-build-core>=1.0" \ - "setuptools>=77.0.0,<82" ninja "packaging>=24.2" \ - "typing-extensions>=4.10.0" pyyaml six - conda_run "${build_venv}/bin/python" -m build --wheel --no-isolation + ninja "packaging>=24.2" "typing-extensions>=4.10.0" pyyaml six numpy + # Do not scan for C++20 modules. PyTorch compiles at C++20 as of 2.14, which + # makes CMake scan every source for module imports, and the scanners are not + # in these images: the clang images have no clang-scan-deps, and GCC rejects + # the scan invocation outright. Nothing here uses modules, so the scan only + # has to be turned off. + conda_run env CMAKE_CXX_SCAN_FOR_MODULES=OFF \ + "${build_venv}/bin/python" -m build --wheel --no-isolation rm -rf "${build_venv}" pip_install "$(echo dist/*.whl)" From e99879ab4f970aced4f688739418cf39addebf73 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 07:03:13 -0700 Subject: [PATCH 11/22] Add the two AOTI shims PyTorch 2.14 links against, and export all three Lowering a CUDA model for a Windows target fails at link time: x86_64-w64-mingw32-ld: undefined reference to `aoti_torch_dtype_uint8' x86_64-w64-mingw32-ld: undefined reference to `aoti_torch_empty_strided_pinned' x86_64-w64-mingw32-ld: undefined reference to `aoti_torch_is_defined' collect2: error: ld returned 1 exit status PyTorch generates a wrapper for the model and links it against this project's shims. The wrapper's own header started calling these three in 2.14, where it called none of them before, so the link now needs all three present. Two did not exist here: `aoti_torch_is_defined` asks whether a handle refers to a real tensor. PyTorch has an undefined-tensor state with no equivalent here, so this is a null check. `aoti_torch_empty_strided_pinned` asks for page-locked host memory, which lets the copy of the model's constants to the device overlap with filling the next buffer. There is no allocator for that memory here, so this reports that and the caller takes the plain synchronous copy it already falls back to. That path is not an error case: the calling code is named for trying, and every way it can fail logs that it is falling back and carries on. Adding a real pinned allocator means teaching the storage layer a new kind of memory and a matching way to release it, which is a change of its own rather than part of moving a version pin. The third, `aoti_torch_dtype_uint8`, was already implemented and simply missing from the list of names the Windows link stub exports, along with the other two. That stub is a checked-in file, so it is regenerated here. Every previous pin bump has had to do the same. Test Plan: Read the requirement from the generated wrapper's header on both branches rather than inferring it from the link error: it names these three in 2.14 and none of them in 2.13, which is why this appears only after the pin moves. Compiled the common shims and confirmed with the symbol table that `aoti_torch_is_defined` is defined rather than merely declared. Checked its signature against PyTorch's own header parameter by parameter, and matched the pointer spelling the neighbouring shims already use for a tensor handle. Regenerated the link stub from the names it already exported plus these three, so it went from 42 to 45. Confirmed the result is a strict superset, that nothing was dropped, that the counts of the two kinds of entry agree, and that all three names are now in it. Confirmed the caller treats a refusal as an ordinary outcome by reading it: it returns empty and logs that it is falling back to the synchronous copy. Formatting is clean on the changed files. The CUDA shim file needs the CUDA headers to compile, which this machine does not have, so that half is for CI to confirm. Whether a Windows target now links is also for CI, since it needs the cross toolchain. Define the new shim where the CUDA library can actually see it The link stub advertises `aoti_torch_is_defined` against the CUDA shim library, but the function was added to the ETensor shim file, which only the Apple Metal backend links. The CUDA library is built from the CUDA shims plus the SlimTensor common shims, and the SlimTensor file did not define it. So a Windows CUDA export would link against a name nothing in that library provides, and fail when loading, which is the failure the earlier commit set out to remove. Add it to the SlimTensor shims as well, following that file's habit of rejecting a null output pointer rather than writing through it. The other two names were already in the right places: the dtype helper in the SlimTensor shims, the pinned allocator in the CUDA memory shims. Also regenerate the link stub deterministically. The tool used to rebuild it stamped the build time into all 48 members, where every member of the previous file had it zeroed, and this file is packaged into the wheel, so that timestamp reached users and made two builds of the same source differ. Test Plan: Compared what each shim file defines and what the CUDA library links. The name was in the ETensor file and absent from the SlimTensor one, and the CUDA build links only the SlimTensor variant, so nothing in that library defined it. The other two are present in the closure, which is why they were fine. Parsed every archive member header in the stub, before and after. The previous file had the time, user and group zeroed in all 48; the version committed earlier carried a real build time in 48 of 49; the version here has all three zeroed again, and the file is byte-for-byte the same size, so nothing but the metadata changed. Compared the advertised names against the previous file: 42 before, 45 now, none dropped, and the three added are exactly the three intended. I have no Windows machine, so the link and load themselves are still for CI. --- backends/aoti/common_shims.cpp | 5 +++++ backends/aoti/common_shims.h | 5 +++++ backends/aoti/common_shims_slim.cpp | 8 ++++++++ backends/aoti/common_shims_slim.h | 4 ++++ backends/cuda/runtime/aoti_cuda_shims.lib | Bin 12738 -> 38758 bytes backends/cuda/runtime/shims/memory.cpp | 18 ++++++++++++++++++ backends/cuda/runtime/shims/memory.h | 20 ++++++++++++++++++++ 7 files changed, 60 insertions(+) diff --git a/backends/aoti/common_shims.cpp b/backends/aoti/common_shims.cpp index f3a34a09987..e83a9576b6c 100644 --- a/backends/aoti/common_shims.cpp +++ b/backends/aoti/common_shims.cpp @@ -159,6 +159,11 @@ AOTITorchError aoti_torch_get_numel(Tensor* tensor, int64_t* ret_numel) { return Error::Ok; } +AOTITorchError aoti_torch_is_defined(Tensor* tensor, bool* ret_is_defined) { + *ret_is_defined = tensor != nullptr; + return Error::Ok; +} + // Device and layout utility functions int32_t aoti_torch_device_type_cpu() { // Let's say cpu is 0 for ET as well diff --git a/backends/aoti/common_shims.h b/backends/aoti/common_shims.h index d057279e22a..ee5b49da6d0 100644 --- a/backends/aoti/common_shims.h +++ b/backends/aoti/common_shims.h @@ -62,6 +62,11 @@ aoti_torch_get_dim(Tensor* tensor, int64_t* ret_dim); AOTI_SHIM_EXPORT AOTITorchError aoti_torch_get_numel(Tensor* tensor, int64_t* ret_numel); +// Reports whether the handle refers to a real tensor. PyTorch has an undefined +// tensor state that ExecuTorch has no equivalent for, so this is a null check. +AOTI_SHIM_EXPORT AOTITorchError +aoti_torch_is_defined(Tensor* tensor, bool* ret_is_defined); + // Utility functions for device and layout information AOTI_SHIM_EXPORT int32_t aoti_torch_device_type_cpu(); AOTI_SHIM_EXPORT int32_t aoti_torch_layout_strided(); diff --git a/backends/aoti/common_shims_slim.cpp b/backends/aoti/common_shims_slim.cpp index c8c7408aa62..0b90a6c9d33 100644 --- a/backends/aoti/common_shims_slim.cpp +++ b/backends/aoti/common_shims_slim.cpp @@ -68,6 +68,14 @@ AOTITorchError aoti_torch_get_numel(Tensor* tensor, int64_t* ret_numel) { return Error::Ok; } +AOTITorchError aoti_torch_is_defined(Tensor* tensor, bool* ret_is_defined) { + if (ret_is_defined == nullptr) { + return Error::InvalidArgument; + } + *ret_is_defined = tensor != nullptr; + return Error::Ok; +} + int32_t aoti_torch_layout_strided() { // Slimtensor only support strided layout, the return value will always be 0, // a.k.a at::Layout::Strided; diff --git a/backends/aoti/common_shims_slim.h b/backends/aoti/common_shims_slim.h index c5a5cab9413..9d05b542a52 100644 --- a/backends/aoti/common_shims_slim.h +++ b/backends/aoti/common_shims_slim.h @@ -51,6 +51,10 @@ aoti_torch_get_dim(Tensor* tensor, int64_t* ret_dim); AOTI_SHIM_EXPORT AOTITorchError aoti_torch_get_numel(Tensor* tensor, int64_t* ret_numel); +// PyTorch has an undefined-tensor state with no equivalent here: null check. +AOTI_SHIM_EXPORT AOTITorchError +aoti_torch_is_defined(Tensor* tensor, bool* ret_is_defined); + AOTI_SHIM_EXPORT int32_t aoti_torch_layout_strided(); // ============================================================ diff --git a/backends/cuda/runtime/aoti_cuda_shims.lib b/backends/cuda/runtime/aoti_cuda_shims.lib index 8bb03cc1c1ee54d3643d7d91456c5d6c4e216ee2..7cd16022a250710a46b771a86e39ade055576088 100644 GIT binary patch literal 38758 zcmeI5%a0sK8Nkc4b{@tdAviB5PR5R%m?T>7Ja%^1GS(^zQ3NP54@H0?dV6c>}IRbOw7gQeDXN!8s` z^;LD(SHJnXx~jVR)WvYJ)w^``d|h`-setpbw5I zAO2Se^xc0czw<{S&|mpj`5~l1zYEubegd-S_x_~(4JZftA8jcA{zW0s{{(b^{%8MD z{`NnFK>rK49`qkxQU2GLgh2oCbIRZOThNQbZV--V{a|aa8wS%}e>e>W{Y@dmW@D(; zAecn`(JY$mPNG>bi$>G&LuLoAkTU)REAXqdARtZw8ixGMLIuqHq=|Mepo|lddI= zA>v36mm7t{U?16Nl4fqIK7;Maco=LB#+%M-G6akB3a&NxogK>39FfN0deqtmW zy&G;_4TeKYelE#;Zco7VaCdeyHjWm1rANEc1-t#xC@M;xNReE2;Lp0XHdqjB^m6wNhLAVnI{2X`^bihk$m79y3uyu`mZ5zG{r1}JGj*-_R}aL zKL~F^E6n}Z&c&0Iqv(3vkPCX@s5`J!jwwRYUk@iEvp(ZcDW|tZ;eH(sBXK_ea6z6e zY=@2ZQvCR_anzP;Z*hFy8Zli$$Hq3;*d`m>Vq@EE>@pkMVPjX=*i|-mjf-utx@xex zYOuO$u)1onx@xexYOuO$u)1onx@xexYO=a&vbt)rx@xkzYO=a&vbt)rx@xkzYO=a& zvbt)qx@xhyYO%U%vASxpx@xhyYO%U%vASxpx@xhyYO}g(v$|@tx@xn!YCCl`Lz8iy zb^w8qrax&8#8(A-qgkAjYVwBxkY5m zy14x5;^uT(ASQ01Tpu$$(@bQK!@*&=uBXAHTL8Z#&aaR{h`(zo| zBbvrD*xap*(Ul|~HPwcbT_fienOw>;Uf(Ln+_KB1ETh~y$o$eS*GbNM8V=ryWs)!` z@uzps;;)ZB%EIUoX-C@nB|#Kx9j`%=@Gu-Hk`y)YQ^fobTY{IdAPk-b)v2Uc#MDXc z;Ta&6qL*;!WL1_ONl;g$u22X(yJGc4I6{X!tB|+kZ?&n3O*kmffYF&`Qa#Wjz+`$_qPoR_2g$96e`^4oT~l=3>1&zJH;w2m#k;${6y*i$ic%6G%8=njOhvEOVowpw`mNbx|C%Z;saZ_AlDqzE_1u$CHY^c zHu&v01a~#jgDtheNV2H`qiwO}kaM>$8Y_ORS5{iq)8}cWk7zY5r(ql@E6vSEVXSm> z$&cKXc9+MPeWMxiE{IfJ+0_6JO6bR zAh$Kqg*~4ZpuCTWhla$7FgJINrrjf@zAQ}g3!Ax z`aCU{FIt$_-;yjh@18A}n0H+RF4l8W4CkI>p({z-rW#QqbGOo`Qk==`+j4nS=Uji8 z7xF7}FU$*Q6u@Pz^L6=KZEE5Na6nq;FEiP7g)kxe^*97~HC!l#ExpBZlYc(yhhy7wq_ktUPCkTNby2@cSKZF#UHyKJ zl#a-GiIp`*N^jWZjFHlx+T~KK4U4UKTkS7twJVn=NLNZIGuO$xIGM%1gh0oxCcXzp z)MD!;H>u9WyIUg5o+F{lroH|E>bZ#N4Z_0xWcms@FUR!EE@v?PzFjV58Ww5tX8LhU z&Eo6oW#;)rOm4<=sv?M~2^Cs& z9aCzSPZ}&=j6-l&6Ys#*U^ywd!!)PSZlXJ$BbUqO-0+I%jA*n-zR!rq&3aYK8p(Pu z@4S>};rh{K)_GM0$ZbvF%q)XwRdNTde-a1kUtainh~5% zY=oGaz!d@W1T)sS$PPD_v;oD0QCu!q9jk5jZG%K2d%LBg_}(Z!m9J0r>8euqTF*go z;kUdPTMS=drAy`{j8+uOFMgy31lJ>5hsN>$V`lW}JjLpRY6p#CRcdR?IF53+i5pAW zFk&Ia)mIDTz0&oxSBF^NHt~xF#roI%%G|QRF(`Cd%i>KHAh*U>3-D<#_iaNv*1Co` zE8Wlr6^r?dn{md)xBU9#f$OPD_?D$-=~fbq?QUoTiG^%e^O>a8aj`wm zsLENBZT0n3I8&MJw&$E_Ww-b7-B{LAYlDP^G*`1fSeo(9>)R4}Nu$~J#olsL?K%6y z`psg~V%zk@a;^;?7ISX)LRYKrWxXF*r@oxOZ0DSQQo7XE`>eDKwAKC1I0Sb!u?<_= zpO+F$HNb$bi81YyrZ$vVNP6|&fUG4~`(&Qp>T|T7$@kC8+%~~>8oKP6l6`lal-fX= z#&H~L1Bb;NClXbEHCa<3t)zK4{))!&EE)##h5u4pY0v)Zep+f1<1C7!SQ|V{7S*3ooTXlk*>X=t%G;WEw_>+#=1CekxSx}y{wNdTp@qj>kp@k-N8UCZuil+ zO841aA$AUJhgXw;IGpMYPqUCa7m87xp{2IUY)+>ri*5nEEgzz*!%_yv{dE`XQ|bx% oTR7yDkvc*W77iuE>UrTg(4kCC)L|P0IXwk~{$|i0g~Lev4 literal 12738 zcmcH#U+P!YR#iWmLe(S!l{QiVr8H_$5=t5si$>rgn)IBxGrRY3W_L#$ zuXH>!_dM^t=bn4cz3%v8BpY3Luwf{q{)VMb@h<+33=fCL#zLXd`^QHG$xH(Pd<9^` zReOR+j?j5SK;Jdgi_o7Jum$xHY(p6YTNVTiECCP3rOZF2*8OZXuq`fo1J`YVTf?5WcNZ z3cBe%Pwjiu+P!Of7jnjv86hv9Mqdz8ts{1Pej#tgQ~9w>B>Ie%OoAh+42>x*Ye(|7 zmA6y5bk>^7rjyoeB0Xyz9#!P|Q&lRl0&!+v?PMl@%*y4n@t7TxC8!9)hM1kOi@i$E ze#4IDWp);CNJ-g8yn$JWq+$tM7GWX{{%9hdvaL+k&e_?+w%d)IEX0R`mt-)S<7H)S z^awkPI`gjSG8cz7b)_@MESW3nG#GU8yp6t+PA6pAVxvJH$>rknDc7o;m0nb4guxmT zjzqF4S()=~P(_8WI*V_9F_KmM#zY$Yu5a-CCeq+9GFmyDg?K7vAC;|U@un217wG?X zBx#lGj6!2blg(BrrI_JMV@R9FZbcWfSvy5figAvCawq*rSXr@1K4N9^SzSP(o#o2K zziBIriFArap5BQfCWqV1^2GATGRmqa9iow^k1Ju=p~fbIHW4|7{?7$A4P$URox~s? znYXR<+#GtCK597?TT%iImmj}_W4UAY;dm7H3~p>BR4HC#RX9T>m<7(@c81-WOQa+D z?aJn)B7FE}9mbKc+*`p|!1!=tKn;(XY@*qphqDoRu-Tu7GY@&N*`H?zYuXM!4n1JI zV?bft)KQ#FE1FqU6lti^yci>VnF(ib&u1esE18bjRy=3fsmN?X-Tqm82`LnaONx;R z3^~ySJNk^=0nuR4kq@U9leV%W78`@-@ZmxZzYUZ4s~Z5ungOyc0GSSeDzHb7!i0_vM0cP;~;4r{+JHYWyfMw);8{co>`!(dv_ao01q(LPM z_*@#mca-(?R+NE|!S5xcsqF!1=mt323-CFLfUU3^8elEl4eQ}{xD8gpI%tA>;11}6e&~c==zwnMg3Zta^>8n2fDp7p z3pB$v*a)rA2KT`LY=OIA6Wj@nD7-6lI%ZJG$&bjhN|KcBu`A5&v0HsG@(R1C z6J1Lauwk;Y4a^DorgYlA3)`s=y4>bdayP?W(_bvJL6xcLD3Eelr=i11CA^f9$CgYq z9RA||W_H3M;Lq++$rtswj@56mgLgPb&adNe;Q4h5qU;cvd zq1#Z0&q)|Q?b3zaX3XIM>QSFZN9~gkI$ev;X3Wp&w~py(FRW{9Ke;o6e=@@}a){{h_^6`yO@i%3}?FkwFB21q8h{;@{hxH%1#wM$^d*+Wz9Hnwpz8 zWlm(l!i2yZf=K|!`~|PY>8#M_(u@%eGC_f6yx2nRSWgvxBpqN^;hS{L;#XBAMp?&f z2&R|wbOe4oA=>9 ziEhP4xsnma6M^zKTR}&9V#IY*9BBYr%(0Vip8fBq!r@lQ8G>(Ss0H1?Yh|JkmWwt# zm^k!0h0hucNYXaqmwsv{1RaI(33C*dfB(V1X~=6;3&i077Zro|8}F3Guh{QL4)vd* z2wBHy^||$0fxfpNIPnMS`zp_7`{*ZM70AZA*w84}d@CcKnGT7m0+P=9#>?Wkw%Fiw4LjW>E5mSjiM z`8A*HL_6-N$c{p_RPBf`GxOoq=vy>1cWRcDtjMah!`p z3F;l(I{DhhwsW*jR;xDnRF>Kx1YW7!I{DT=7BA3jtns(t)~^$0Hh%rT?_8(ZSj*bM zl*rm(1fC6MpV;}m#92Jp$DCW)sigb7;ysV1KJZ+GnKk;rp=bY&7Xx?~;IF?>>uUxC z?mx3KuN?T}HN2d_l0ntaoF`PPo$-KB%_3c`;Qa@uy>Mnxt!5@hzjpTLZ_y1>6Q^0Z zQPq?OpY`qLPSx@9>!0DZ65jSO{gw`2RzmkrPJI2Jq=b`QNnzILyVZuhK8QyqAXF}~LF!3&qj*EXn{ne#ABd0=1TvXT0`zq(A> z$VOFr$wRkl?Fciu@$zWyBfL7r8%xFKoeKi*YT}mcxlEM(Zp*6|Datl8x{cFAt@D9B zfZ1iHj{o!$?J_OC`im!zn)<*#FwW%WxgXy58D2}{EwMTF-Z5ykdc#aDde?v6IMH@f^a&=)|LV*yu2N3fsp Date: Thu, 3 Sep 2026 07:14:39 -0700 Subject: [PATCH 12/22] Pin the CUDA fallback kernel the backend just started advertising The CUDA test that asserts the exact set of fallback kernels fails: test_cuda_fallbacks_unchanged_by_rocm_gate An earlier commit taught the backend to accept the shim spelling Inductor now derives for the int4 pack matmul, but this test compares the advertised set against a frozen copy of it, so adding an entry on one side alone breaks the comparison. Add the same entry here. The Metal backend took the same kind of change and is unaffected, because nothing pins its set exactly. Test Plan: Extracted both sets and compared them directly. Before this the advertised set has twelve entries and the frozen copy eleven, with the new derived spelling the only difference and nothing frozen that the backend no longer advertises, so the comparison fails on exactly that one name. After, both hold twelve and the comparison passes. Treat a test that names c10 as an ATen test too The Buck test targets decide their C++ standard separately from the libraries, and their check for "is this an ATen test" looks for a dependency named `libtorch` but not for the other two names the build maps onto the same library. A test that names `c10` or `torch-core-cpp` was therefore pinned to C++17 and failed on ATen's headers, which need C++20 as of PyTorch 2.14: ATen/core/TensorBase.h: error: expected ';' at end of declaration Reuse the same check the library path already uses, so both paths agree on what counts as ATen. It covers all three dependency names, and the exported list as well as the plain one. The rest of the test check is left alone: a name containing `aten`, and the ATen flavours of the test framework, still mark a test as ATen on their own. Test Plan: Ran the wrapper function directly, with `type()` behaving as Starlark's does, over the cases that matter. Before this a test naming `c10` or `torch-core-cpp` came back with C++17; after, both come back without a pin and take C++20 from the toolchain. A plain test with no ATen dependency still comes back with C++17, which is the case the embedded build depends on, and a test named for ATen or naming `libtorch` is unchanged. Moved the shared check above its first use, since it now has one earlier in the file, and dropped a local that nothing reads any more. Put the C++20 flag where the compiler actually reads it The Buck build still fails compiling anything generated that includes ATen: ATen/core/TensorBase.h:1031:5: error: unknown type name 'requires' `requires` is a C++20 keyword, so the compile was still running at C++17 even though the target is ATen-mode and the wrapper had already decided it should be C++20. The wrapper was setting the standard only in the per-language flags, which reach the preprocessor pass. That was enough while the value matched the build config, which pins C++17 for everything, and it is not enough to override it. Set the flag on the target's own compile flags as well, for ATen-mode targets only. Nothing changes for the rest of the build: a target that is not ATen-mode gets no flag added here and keeps the C++17 the build config gives it. Test Plan: Ran the wrapper over the exact arguments the generated kernel library is defined with, the ones from the failure. Before this it came away with only its own warning flag and the standard confined to the per-language entry, which is why the compile stayed at C++17. After, the flag is on the target's own list where the compile reads it. Checked the case that must not move: a library with no ATen dependency still gets C++17 and no flag on its own list. Confirmed why the old spelling looked sufficient. It only ever carried C++17, matching the build config, so nothing had shown it could raise a standard rather than restate one. I cannot run this Buck build on this machine, so CI has to confirm the generated libraries now compile. Raise the ATen-mode C++ standard where both builds see it The C++20 flag for ATen-mode targets was set in the interface layer that exists only for the open source build, whose own docstring says as much. So the internal build never got it, and the shared wrapper was passing that layer an argument its internal counterpart does not accept, which would fail there rather than being merely ineffective. Move the decision into the shared wrapper, which both builds go through, and leave the interface layer exactly as it was. The flag goes on the target's own compile flags, because the build config applies C++17 to everything and only a flag on the target overrides it. Test Plan: Read the two files' own descriptions of themselves to place this: the interface layer is for changes specific to one build, the wrapper is for logic shared with the internal one. Confirmed the internal counterpart takes only the arguments it always took, so the previous version would have broken it rather than skipped it. Ran the new helper over the target from the failure and over targets that must not move. The generated kernel library and the ATen bridge come away with C++20; a plain library and the portable kernels come away with no flag added, so they keep the C++17 the build config gives them, which is what the embedded builds rely on. I cannot run either Buck build here, so CI has to confirm the generated libraries compile. Decide ATen mode before the wrapper consumes the dependency list The Buck build still failed compiling generated sources that include ATen: ATen/core/TensorBase.h:1031:5: error: unknown type name 'requires' The check for "is this an ATen target" reads the external dependency list, but by the time it ran, an earlier step in the same wrapper had already removed that list after turning its entries into real dependencies. So every target looked like a non-ATen one, no target was ever raised to C++20, and the flag added for it was never added at all. Read the list first, then let the rest of the wrapper consume it. Test Plan: Traced the order in the wrapper: the step that removes the list runs before the step that read it. Reproduced that by running both against the target from the failure: it reports ATen mode with the list present and not once it is gone, which is exactly the sequence the wrapper produced. Reran the whole path in the new order over the target from the failure, an ATen target that exports its dependency instead of naming it directly, and a plain library. The first two now come away at C++20 and the plain one is untouched, so the embedded builds keep the C++17 they rely on. Also read the build system's own source to confirm the flag can win where it is placed: a target's flags are appended after the toolchain's, and its comment says they are last on purpose. So the earlier concern that the flag was being placed somewhere that could not take effect was unfounded; it simply was not being added. Shortened an overlong comment that failed the formatter. I cannot run this Buck build here, so CI has to confirm the generated libraries compile. --- backends/aoti/common_shims.h | 3 +- backends/cuda/tests/test_sort_shim.py | 1 + .../xplat/executorch/build/env_interface.bzl | 17 +++----- .../executorch/build/runtime_wrapper.bzl | 42 ++++++++++++------- 4 files changed, 34 insertions(+), 29 deletions(-) diff --git a/backends/aoti/common_shims.h b/backends/aoti/common_shims.h index ee5b49da6d0..0b85b09ee51 100644 --- a/backends/aoti/common_shims.h +++ b/backends/aoti/common_shims.h @@ -62,8 +62,7 @@ aoti_torch_get_dim(Tensor* tensor, int64_t* ret_dim); AOTI_SHIM_EXPORT AOTITorchError aoti_torch_get_numel(Tensor* tensor, int64_t* ret_numel); -// Reports whether the handle refers to a real tensor. PyTorch has an undefined -// tensor state that ExecuTorch has no equivalent for, so this is a null check. +// PyTorch has an undefined-tensor state with no equivalent here: null check. AOTI_SHIM_EXPORT AOTITorchError aoti_torch_is_defined(Tensor* tensor, bool* ret_is_defined); diff --git a/backends/cuda/tests/test_sort_shim.py b/backends/cuda/tests/test_sort_shim.py index fc5f870fc42..1e5b34b7396 100644 --- a/backends/cuda/tests/test_sort_shim.py +++ b/backends/cuda/tests/test_sort_shim.py @@ -37,6 +37,7 @@ _CUDA_FALLBACK_KERNELS = frozenset( { "at::_ops::_weight_int4pack_mm::call", + "aoti_torch_cuda__weight_int4pack_mm", "at::_ops::sort_stable::call", "aoti_torch_cuda_randint_low_out", "executorch_cuda::int4_plain_mm", diff --git a/shim_et/xplat/executorch/build/env_interface.bzl b/shim_et/xplat/executorch/build/env_interface.bzl index 9815a504025..c613243b537 100644 --- a/shim_et/xplat/executorch/build/env_interface.bzl +++ b/shim_et/xplat/executorch/build/env_interface.bzl @@ -160,21 +160,14 @@ def _patch_headers(kwargs): def _patch_pp_flags(kwargs): return kwargs -def _patch_cxx_compiler_flags(kwargs, aten_mode = False): - """Pins the C++ standard a target compiles with. - - C++17 by default, which is what the runtime this project ships requires. - ATen-mode targets are raised to C++20, because PyTorch's headers require it as - of 2.14: c10/util/intrusive_ptr.h defines operator<=> and returns - std::strong_ordering with no feature guard. - """ - std = "-std=c++20" if aten_mode else "-std=c++17" +def _patch_cxx_compiler_flags(kwargs): + """CXX Compiler flags to enable C++17 features.""" if "lang_compiler_flags" not in kwargs: - kwargs["lang_compiler_flags"] = {"cxx_cpp_output": [std]} + kwargs["lang_compiler_flags"] = {"cxx_cpp_output": ["-std=c++17"]} elif "cxx_cpp_output" not in kwargs["lang_compiler_flags"]: - kwargs["lang_compiler_flags"]["cxx_cpp_output"] = [std] + kwargs["lang_compiler_flags"]["cxx_cpp_output"] = ["-std=c++17"] else: - kwargs["lang_compiler_flags"]["cxx_cpp_output"].append(std) + kwargs["lang_compiler_flags"]["cxx_cpp_output"].append("-std=c++17") return kwargs # buildifier: disable=unused-variable diff --git a/shim_et/xplat/executorch/build/runtime_wrapper.bzl b/shim_et/xplat/executorch/build/runtime_wrapper.bzl index 17d6ffc8c2d..7b7513a1c93 100644 --- a/shim_et/xplat/executorch/build/runtime_wrapper.bzl +++ b/shim_et/xplat/executorch/build/runtime_wrapper.bzl @@ -145,6 +145,18 @@ def _has_pytorch_dep(dep_list): return True return False +def _is_aten_target(kwargs): + """Whether a target compiles against ATen. + + Keyed on exact dep names, not a substring: every label contains "torch". + """ + aten_external_deps = ["c10", "libtorch", "libtorch_python", "torch-core-cpp"] + for key in ["external_deps", "exported_external_deps"]: + for dep in kwargs.get(key) or []: + if dep in aten_external_deps: + return True + return False + def _patch_test_compiler_flags(kwargs): if "compiler_flags" not in kwargs: kwargs["compiler_flags"] = [] @@ -154,16 +166,14 @@ def _patch_test_compiler_flags(kwargs): # non-aten tests are pinned to C++17 for embedded. name = kwargs.get("name", "") external_deps = kwargs.get("external_deps", []) - deps = kwargs.get("deps", []) xplat_deps = kwargs.get("xplat_deps", []) fbcode_deps = kwargs.get("fbcode_deps", []) is_aten_test = ( "_aten" in name or "aten_" in name or - "libtorch" in external_deps or "gtest_aten" in external_deps or "gmock_aten" in external_deps or - _has_pytorch_dep(deps) or + _is_aten_target(kwargs) or _has_pytorch_dep(xplat_deps) or _has_pytorch_dep(fbcode_deps) ) @@ -267,27 +277,26 @@ def _patch_kwargs_cxx(kwargs): env.remove_platform_specific_args(kwargs) return _patch_kwargs_common(kwargs) -def _is_aten_target(kwargs): - """Whether a target compiles against ATen, and so needs C++20. +def _patch_aten_mode_std(kwargs, aten_mode): + """Raises an ATen-mode target to C++20, which PyTorch's headers require. - Keyed on the external dep names the build maps onto libtorch, rather than on - a substring of the label: every ExecuTorch label contains "torch". + A plain compiler flag, which the prelude places after the toolchain's. """ - aten_external_deps = ["c10", "libtorch", "libtorch_python", "torch-core-cpp"] - for key in ["external_deps", "exported_external_deps"]: - for dep in kwargs.get(key) or []: - if dep in aten_external_deps: - return True - return False + if aten_mode: + kwargs["compiler_flags"] = kwargs.get("compiler_flags", []) + ["-std=c++20"] + return kwargs def _cxx_library_common(*args, **kwargs): + # Before _patch_kwargs_cxx, which consumes external_deps. + aten_mode = _is_aten_target(kwargs) _patch_kwargs_cxx(kwargs) _patch_build_mode_flags(kwargs) + _patch_aten_mode_std(kwargs, aten_mode) env.patch_platform_build_mode_flags(kwargs) env.patch_headers(kwargs) env.patch_pp_flags(kwargs) - env.patch_cxx_compiler_flags(kwargs, aten_mode = _is_aten_target(kwargs)) + env.patch_cxx_compiler_flags(kwargs) env.patch_force_static(kwargs) env.cxx_library(*args, **kwargs) @@ -307,10 +316,13 @@ def _cxx_library(*args, **kwargs): _cxx_library_common(*args, **kwargs) def _cxx_binary_helper(*args, **kwargs): + # Before _patch_kwargs_cxx, which consumes external_deps. + aten_mode = _is_aten_target(kwargs) _patch_kwargs_cxx(kwargs) _patch_build_mode_flags(kwargs) + _patch_aten_mode_std(kwargs, aten_mode) env.patch_platform_build_mode_flags(kwargs) - env.patch_cxx_compiler_flags(kwargs, aten_mode = _is_aten_target(kwargs)) + env.patch_cxx_compiler_flags(kwargs) env.cxx_binary(*args, **kwargs) From e1f33f940a4b6d1ce27d0c0ad376c14c5f849c4e Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 10:45:38 -0700 Subject: [PATCH 13/22] Pin the other direction of the moved overflow boundary The vendored range check moved in both directions, and only the direction that starts refusing values had a test. Cover the quiet one: filling an int8 tensor with 127.5 used to be refused and now stores 127. Test Plan: Compiled the real header from before and after this change into the same probe and ran both. Before: 127.5 into int8 and 255.5 into uint8 both report overflow, and 2^63 into int64 is accepted. After: the first two are accepted and 2^63 overflows. So the boundary moved both ways, and the case added here fails against the previous header rather than passing either way. Placed in this file rather than the shared macro because the macro expands into eight suites, and several of them reach the scalar through a path that rejects a fractional value before any range check, where the case would not measure what it claims. Decide ATen mode before the wrapper consumes it on the test path too The library and binary paths were fixed to read the dependency list before the shared patch removes it. The test path has the same shape and was missed: it still called its own check afterwards, so a test naming an ATen dependency was pinned to C++17 and would fail on ATen's headers. Read it in the same place, and pass the answer down, so both paths agree. Test Plan: Ran the test path over the cases that matter, with the dependency list removed first to match what the wrapper really does. Before this a test naming libtorch, c10 or torch-core-cpp came away pinned to C++17; after, all three take C++20 from the toolchain. A test named for ATen is unchanged, and a plain test still gets C++17, which is what the embedded builds rely on. No open source test target has this shape today, so this closes the same defect on a path that is not yet exercised rather than repairing a red job. --- kernels/test/op_full_test.cpp | 13 +++++++++++++ shim_et/xplat/executorch/build/runtime_wrapper.bzl | 8 +++++--- 2 files changed, 18 insertions(+), 3 deletions(-) diff --git a/kernels/test/op_full_test.cpp b/kernels/test/op_full_test.cpp index 752c4e710f4..e412b67dfc4 100644 --- a/kernels/test/op_full_test.cpp +++ b/kernels/test/op_full_test.cpp @@ -84,6 +84,19 @@ ET_FORALL_REALHBF16_TYPES(GENERATE_TEST) GENERATE_SCALAR_OVERFLOW_TESTS(OpFullOutTest) +// The other half of the boundary change: 127.5 used to be refused for an int8 +// tensor and now truncates to 127. +TEST_F(OpFullOutTest, CharTensorFractionalScalarTruncates) { + TensorFactory tf; + std::vector sizes = {2, 2}; + std::vector sizes_int64_t(sizes.begin(), sizes.end()); + auto aref = IntArrayRef(sizes_int64_t.data(), sizes_int64_t.size()); + Tensor out = tf.zeros(sizes); + + op_full_out(aref, 127.5, out); + EXPECT_TENSOR_EQ(out, tf.full(sizes, 127)); +} + TEST_F(OpFullOutTest, HalfSupport) { TensorFactory tf; diff --git a/shim_et/xplat/executorch/build/runtime_wrapper.bzl b/shim_et/xplat/executorch/build/runtime_wrapper.bzl index 7b7513a1c93..519f0f4e731 100644 --- a/shim_et/xplat/executorch/build/runtime_wrapper.bzl +++ b/shim_et/xplat/executorch/build/runtime_wrapper.bzl @@ -157,7 +157,7 @@ def _is_aten_target(kwargs): return True return False -def _patch_test_compiler_flags(kwargs): +def _patch_test_compiler_flags(kwargs, aten_mode = False): if "compiler_flags" not in kwargs: kwargs["compiler_flags"] = [] @@ -169,11 +169,11 @@ def _patch_test_compiler_flags(kwargs): xplat_deps = kwargs.get("xplat_deps", []) fbcode_deps = kwargs.get("fbcode_deps", []) is_aten_test = ( + aten_mode or "_aten" in name or "aten_" in name or "gtest_aten" in external_deps or "gmock_aten" in external_deps or - _is_aten_target(kwargs) or _has_pytorch_dep(xplat_deps) or _has_pytorch_dep(fbcode_deps) ) @@ -340,10 +340,12 @@ def _cxx_test(*args, **kwargs): kwargs["deps"] = [] kwargs["deps"].append("//executorch/test/utils:utils") + # Before _patch_kwargs_cxx, which consumes external_deps. + aten_mode = _is_aten_target(kwargs) _patch_kwargs_cxx(kwargs) env.patch_headers(kwargs) _patch_build_mode_flags(kwargs) - _patch_test_compiler_flags(kwargs) + _patch_test_compiler_flags(kwargs, aten_mode) env.patch_platform_build_mode_flags(kwargs) env.cxx_test(*args, **kwargs) From b750eed50c8e66afc1700fe7449fb8bdf419a024 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 13:04:56 -0700 Subject: [PATCH 14/22] Shorten the comments added by this stack They explained more than the code needed. One line each is enough. --- .ci/docker/common/install_docs_reqs.sh | 3 +-- .ci/docker/common/install_pytorch.sh | 22 +++++----------------- .ci/scripts/utils.sh | 6 +----- backends/apple/metal/metal_backend.py | 5 +---- backends/cuda/cuda_backend.py | 4 +--- backends/cuda/runtime/shims/memory.h | 7 ++----- 6 files changed, 11 insertions(+), 36 deletions(-) diff --git a/.ci/docker/common/install_docs_reqs.sh b/.ci/docker/common/install_docs_reqs.sh index 2794ffb8fc9..73f361acb00 100755 --- a/.ci/docker/common/install_docs_reqs.sh +++ b/.ci/docker/common/install_docs_reqs.sh @@ -20,8 +20,7 @@ if [ -n "$BUILD_DOCS" ]; then apt-get update apt-get install -y --no-install-recommends yarn - # katex 0.18.5 requires commander@15 / node >= 22.12; pin to the last - # release compatible with the node 16 installed above + # 0.18.5 wants node >= 22.12; this is the last release node 16 can run. yarn global add katex@0.18.4 --prefix /usr/local sudo apt-get -y install doxygen diff --git a/.ci/docker/common/install_pytorch.sh b/.ci/docker/common/install_pytorch.sh index 85a0661d08d..e51a8886afd 100755 --- a/.ci/docker/common/install_pytorch.sh +++ b/.ci/docker/common/install_pytorch.sh @@ -76,33 +76,21 @@ install_pytorch_and_domains() { # the image compiler cannot satisfy. The venv inherits the image's # site-packages, so PyTorch still builds against the same numpy. # - # Keep the list in sync with pytorch/pyproject.toml [build-system].requires, - # except that cmake is deliberately left out, see below, and ninja is kept - # because the build needs a generator. + # Keep in sync with pytorch/pyproject.toml [build-system].requires. local build_venv=/tmp/pytorch-build-venv rm -rf "${build_venv}" conda_run python -m venv --system-site-packages "${build_venv}" - # cmake is deliberately not installed here, so the conda cmake already in the - # image is used. scikit-build-core, which PyTorch builds with as of 2.14, - # prefers an importable pip cmake over anything on PATH, and cmake adds its own - # install root to CMAKE_SYSTEM_PREFIX_PATH. A pip cmake therefore searches - # site-packages, where MKL and libomp are not, and the build silently comes out - # with no BLAS and no LAPACK. + # No pip cmake: scikit-build-core would prefer it over the image's, and it + # searches site-packages, where MKL and libomp are not. conda_run "${build_venv}/bin/pip" install build "scikit-build-core>=1.0" \ ninja "packaging>=24.2" "typing-extensions>=4.10.0" pyyaml six numpy - # Do not scan for C++20 modules. PyTorch compiles at C++20 as of 2.14, which - # makes CMake scan every source for module imports, and the scanners are not - # in these images: the clang images have no clang-scan-deps, and GCC rejects - # the scan invocation outright. Nothing here uses modules, so the scan only - # has to be turned off. + # These images have no module scanner, and nothing here uses modules. conda_run env CMAKE_CXX_SCAN_FOR_MODULES=OFF \ "${build_venv}/bin/python" -m build --wheel --no-isolation rm -rf "${build_venv}" pip_install "$(echo dist/*.whl)" - # The build silently degrades rather than failing when it cannot find BLAS, so - # assert on the result. Run from / so the import resolves to the installed - # wheel and not to the source tree next to it. + # A build with no BLAS succeeds silently. Run from / to import the wheel. (cd / && conda_run python -c " import torch assert torch._C.has_lapack, 'built without LAPACK' diff --git a/.ci/scripts/utils.sh b/.ci/scripts/utils.sh index 3e97e431072..45ac584b59d 100644 --- a/.ci/scripts/utils.sh +++ b/.ci/scripts/utils.sh @@ -135,11 +135,7 @@ install_pytorch_and_domains() { if [[ "$(uname -m)" == "aarch64" ]]; then export BUILD_IGNORE_SVE_UNAVAILABLE=1 fi - # PyTorch no longer supports "python setup.py bdist_wheel"; it builds - # through scikit-build-core (PEP 517). Build with the standard frontend and - # keep isolation off, so the build uses the requirements-build.txt deps - # installed just above rather than fetching its own copies. This matches - # how the Docker images build PyTorch. + # PyTorch dropped setup.py; isolation off reuses the deps installed above. pip install build USE_DISTRIBUTED=1 python -m build --wheel --no-isolation pip install "$(echo dist/*.whl)" diff --git a/backends/apple/metal/metal_backend.py b/backends/apple/metal/metal_backend.py index b301346ed6b..aa931d8ae4d 100644 --- a/backends/apple/metal/metal_backend.py +++ b/backends/apple/metal/metal_backend.py @@ -39,10 +39,7 @@ def get_supported_fallback_kernels(cls) -> Dict[str, Any]: "at::_ops::_scaled_dot_product_attention_math_for_mps::call": None, "at::_ops::_scaled_dot_product_attention_math_for_mps_v2::call": None, "torchao::_linear_fp_act_4bit_weight": None, - # Each custom op appears twice: once as registered, once under the shim - # name Inductor derives for it. Which spelling reaches this list - # depends on the path Inductor takes, and that is not ours to control, - # so accept both. + # Each custom op twice: as registered, and as Inductor derives it. "aoti_torch_mps__linear_fp_act_4bit_weight": None, "at::_ops::topk::call": None, "metal::gather_qmv": None, diff --git a/backends/cuda/cuda_backend.py b/backends/cuda/cuda_backend.py index c1991f97b8e..97a4f67a504 100644 --- a/backends/cuda/cuda_backend.py +++ b/backends/cuda/cuda_backend.py @@ -669,9 +669,7 @@ def get_supported_fallback_kernels(cls) -> Dict[str, Any]: return {} return { "at::_ops::_weight_int4pack_mm::call": None, - # The same op under the shim name Inductor derives for it. Which - # spelling reaches this list depends on the path Inductor takes, and - # that is not ours to control, so accept both. + # Also under the shim name Inductor derives for it. "aoti_torch_cuda__weight_int4pack_mm": None, "at::_ops::sort_stable::call": None, "aoti_torch_cuda_randint_low_out": None, diff --git a/backends/cuda/runtime/shims/memory.h b/backends/cuda/runtime/shims/memory.h index 469913d1542..158edfa6d92 100644 --- a/backends/cuda/runtime/shims/memory.h +++ b/backends/cuda/runtime/shims/memory.h @@ -98,11 +98,8 @@ AOTI_SHIM_EXPORT AOTITorchError aoti_torch_empty_strided( /** * Reports that pinned host memory is unavailable. * - * Generated wrappers call this to stage constants through page-locked host - * memory, which lets the copy to the device run asynchronously. There is no - * pinned allocator here, so this always fails and the caller uses the - * synchronous copy it already falls back to. Correct, and slower only while - * loading. + * There is no pinned allocator here, so the caller falls back to the + * synchronous copy it already handles: correct, slower only while loading. * * @return Error::NotSupported */ From 380a40d357064ff558a1e028e9697b98d0b6385d Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 13:04:56 -0700 Subject: [PATCH 15/22] Update the cortex-m expectations 2.14 changes Three tests fail: RuntimeError: Expected to find "..._dim_order_ops__clone_dim_order_default" but did not find it AssertionError: cortex_m_dequantize_per_tensor_default output dtype torch.float32 The two model tests count an exact set of operators, and the dim-order clone they expect is no longer in the graph on this PyTorch, so the count cannot match. Drop those entries. The third checks that the first operator left after quantization carries int8. For this case the operator is dropout with training off, which is an identity, so nothing of it survives lowering and the first operator left is the dequantize, whose output is float by definition. Say so through the field the file already has for cases whose output is not int8; the check on that path also confirms the input is int8, which holds for a dequantize. Test Plan: Read the failure out of the job log rather than inferring it. The graph the check searched is the pre-transform one, and it contains the other expected operators, so only the clone entry is stale. The dropout diagnosis is confirmed by its sibling in the same list: the in-place variant keeps an explicit clone of its input, and that case passes, while the plain one folds away and fails. That difference is the mechanism. Both model tests pass on the previous pin with these same entries, so this is the graph changing under a new PyTorch rather than an existing defect. I cannot run these tests here: they need the Arm serializer, which is not installed on this machine and not published on any index I can reach. CI has to confirm. --- backends/cortex_m/test/misc/test_portable_int8.py | 4 +++- backends/cortex_m/test/models/test_ds_cnn.py | 2 -- backends/cortex_m/test/models/test_mobilenet_v2.py | 2 -- 3 files changed, 3 insertions(+), 5 deletions(-) diff --git a/backends/cortex_m/test/misc/test_portable_int8.py b/backends/cortex_m/test/misc/test_portable_int8.py index 6efeec9e5b0..26bc61e1d63 100644 --- a/backends/cortex_m/test/misc/test_portable_int8.py +++ b/backends/cortex_m/test/misc/test_portable_int8.py @@ -387,7 +387,9 @@ def _quantize_and_export( torch.ops.aten.dropout.default, _build_module(lambda x, y: torch.ops.aten.dropout.default(x, 0.1, False)), (torch.randn(2, 3, 4, 5), torch.randn(2, 3, 4, 5)), - None, + # Not training, so this is an identity and nothing survives lowering to + # carry int8. Only the dequantize is left, which is float by definition. + torch.float32, ), "dropout_": OpCase( torch.ops.aten.dropout_.default, diff --git a/backends/cortex_m/test/models/test_ds_cnn.py b/backends/cortex_m/test/models/test_ds_cnn.py index 206af19a61e..7bb55fb6a12 100644 --- a/backends/cortex_m/test/models/test_ds_cnn.py +++ b/backends/cortex_m/test/models/test_ds_cnn.py @@ -15,7 +15,6 @@ "executorch_exir_dialects_edge__ops_aten_linear_default": 1, "executorch_exir_dialects_edge__ops_aten_relu_default": 9, "executorch_exir_dialects_edge__ops_aten_view_copy_default": 1, - "executorch_exir_dialects_edge__ops_dim_order_ops__clone_dim_order_default": 2, "executorch_exir_dialects_edge__ops_quantized_decomposed_dequantize_per_channel_default": 18, "executorch_exir_dialects_edge__ops_quantized_decomposed_dequantize_per_tensor_default": 17, "executorch_exir_dialects_edge__ops_quantized_decomposed_quantize_per_tensor_default": 15, @@ -30,7 +29,6 @@ "executorch_exir_dialects_edge__ops_cortex_m_quantized_conv2d_default": 4, "executorch_exir_dialects_edge__ops_cortex_m_quantized_depthwise_conv2d_default": 5, "executorch_exir_dialects_edge__ops_cortex_m_quantized_linear_default": 1, - "executorch_exir_dialects_edge__ops_dim_order_ops__clone_dim_order_default": 2, } test_cases = { diff --git a/backends/cortex_m/test/models/test_mobilenet_v2.py b/backends/cortex_m/test/models/test_mobilenet_v2.py index 67f0937a006..20d8604c048 100644 --- a/backends/cortex_m/test/models/test_mobilenet_v2.py +++ b/backends/cortex_m/test/models/test_mobilenet_v2.py @@ -20,7 +20,6 @@ "executorch_exir_dialects_edge__ops_aten_hardtanh_default": 35, "executorch_exir_dialects_edge__ops_aten_linear_default": 1, "executorch_exir_dialects_edge__ops_aten_view_copy_default": 1, - "executorch_exir_dialects_edge__ops_dim_order_ops__clone_dim_order_default": 1, "executorch_exir_dialects_edge__ops_quantized_decomposed_dequantize_per_channel_default": 104, "executorch_exir_dialects_edge__ops_quantized_decomposed_dequantize_per_tensor_default": 79, "executorch_exir_dialects_edge__ops_quantized_decomposed_quantize_per_tensor_default": 67, @@ -35,7 +34,6 @@ "executorch_exir_dialects_edge__ops_cortex_m_quantized_conv2d_default": 35, "executorch_exir_dialects_edge__ops_cortex_m_quantized_depthwise_conv2d_default": 17, "executorch_exir_dialects_edge__ops_cortex_m_quantized_linear_default": 1, - "executorch_exir_dialects_edge__ops_dim_order_ops__clone_dim_order_default": 1, } # Use larger sample set for calibration to get better quantization From 3ff10a3c83214652d5cc23009063229f4c3d38e3 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 14:44:35 -0700 Subject: [PATCH 16/22] Sort lifted constants with the other constants when building a delegate Building a delegate rearranges its placeholders into parameters and buffers first, then user inputs. The sort asks the graph signature which names are parameters and which are buffers, but never asks which are lifted tensor constants, so a constant lands in the user input group and is left after a real user input. Anything that later inserts a constant then fails, because inserting one is only allowed ahead of the user inputs: Failed to insert aten_alias_copy_default_fused_const; Const placeholder nodes must be inserted before user input nodes in the graph. Ask about constants too, and about lifted custom objects, which have the same shape. Test Plan: Traced where the order breaks by watching the signature through the whole pipeline. It is correct out of export and out of edge lowering, and wrong immediately after the delegate submodule is built, which is the only place this sort runs. Ran the sort directly against a signature holding a constant and a user input, with the old code and the new: the old one returns the constant after the input, the new one returns it before. Neither reorders the user inputs among themselves. Ran the Arm backend's misc, passes and quantizer tests, 1725 of them, once with this change and once without, in the same environment. Seven tests go from failing to passing and none goes the other way. The seven are the ones that lift a tensor constant into a delegate and then fuse a constant operator. The remaining failures in that run are separate breakages on this branch, unrelated to placeholder order, and this change neither fixes nor worsens them. --- exir/lowered_backend_module.py | 25 ++++++++++++------------- 1 file changed, 12 insertions(+), 13 deletions(-) diff --git a/exir/lowered_backend_module.py b/exir/lowered_backend_module.py index a4c2f2cfe79..11708710333 100644 --- a/exir/lowered_backend_module.py +++ b/exir/lowered_backend_module.py @@ -384,7 +384,7 @@ def arrange_graph_placeholders( ) -> torch.fx.GraphModule: """ Modifies the graph of the given graphmodule with one that contains the same nodes as the original, - but with placeholders in order of (Params + Buffers) (User Inputs) + but with placeholders in order of (Params + Buffers + Constants) (User Inputs) This is used by the delegate api which disturbs the placeholder ordering when creating a submodule from partitioned nodes @@ -403,32 +403,31 @@ def arrange_graph_placeholders( graph_sign = owning_program.graph_signature # Add all placeholders into the graph first: - # Cache these properties — each call rebuilds the dict from input_specs. + # Cache these properties to avoid rebuilding the dict on each access. params_map = graph_sign.inputs_to_parameters buffers_map = graph_sign.inputs_to_buffers + constants_map = graph_sign.inputs_to_lifted_tensor_constants + custom_objs_map = graph_sign.inputs_to_lifted_custom_objs param_nodes = [] buffer_nodes = [] + constant_nodes = [] input_nodes = [] for node in gm.graph.nodes: if node.op != "placeholder": continue - if node.name in params_map and node.meta.get("delegation_tag", None) == tag: + is_tagged = node.meta.get("delegation_tag", None) == tag + if node.name in params_map and is_tagged: param_nodes.append(node) - elif node.name in buffers_map and node.meta.get("delegation_tag", None) == tag: + elif node.name in buffers_map and is_tagged: buffer_nodes.append(node) + elif (node.name in constants_map or node.name in custom_objs_map) and is_tagged: + constant_nodes.append(node) else: input_nodes.append(node) - for param_node in param_nodes: - new_node = new_graph.node_copy(param_node, lambda x: node_map[x]) - node_map[param_node] = new_node - for buffer_node in buffer_nodes: - new_node = new_graph.node_copy(buffer_node, lambda x: node_map[x]) - node_map[buffer_node] = new_node - for input_node in input_nodes: - new_node = new_graph.node_copy(input_node, lambda x: node_map[x]) - node_map[input_node] = new_node + for node in param_nodes + buffer_nodes + constant_nodes + input_nodes: + node_map[node] = new_graph.node_copy(node, lambda x: node_map[x]) # Now add all the other nodes in order for node in gm.graph.nodes: From 5d0b30bd18a76a50d9c1de8f999a266b90155de8 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 14:44:57 -0700 Subject: [PATCH 17/22] Share the observer across dropout, which is an identity in eval mode Quantized inception_v3 fails to load with the XNNPACK delegate: Failed to create squeeze node 344 with code: xnn_status_invalid_parameter Init failed for backend XnnpackBackend: 0x1 XNNPACK refuses a reshape whose input and output carry different quantization scales. The reshape here sits between two, because the operator feeding it was never given the scale of the operator before it. The list of operators that must share their neighbour's observer already holds permute, view, squeeze, flatten and the pooling operators. Dropout was missing. Not training, so it is the identity and cannot change a tensor's range, but the walk that spreads a shared observer only looks one operator back, so it stopped at dropout and left everything after it to be measured on its own. The mismatch itself is not new. Both scales are identical on the previous pin, to the last digit. What is new is that the reshape now ends up inside a delegate, where XNNPACK checks it, because the clone that used to sit in front of it is no longer emitted for an eval-mode dropout. Both spellings are listed, as hardtanh already is two lines above. The in-place form is a separate operator that survives export unchanged, so listing one left a model built with it hitting the same mismatch. The training argument on the node is deliberately not read. It is a switch that is flipped after annotation runs, not a property of the graph: exported from eval it reads false while the annotation pass sees it, becomes true when the model is moved to training, and returns to false when it is moved back for deployment. The project's own train/eval test asserts exactly that flipping. So a decision taken from its value during annotation describes nothing durable, and taking it would refuse to share for a model that is about to be deployed in eval, which is the case this entry exists to fix. Test Plan: Read the whole list of conditions that make this XNNPACK call return an invalid parameter, then confirmed which one fires: the scales either side of the reshape differ, 0.017268 against 0.007075. Followed the annotation along the tail of the model and found dropout carrying none, which leaves the flatten after it unannotated as well, which is what gives it a fresh scale. Ran the same model on both pins with the same seed and the same source. The two scales come out identical on both, and only the newer one fails to load, which is what shows the mismatch is pre-existing and the delegation changed. Exported and ran the model with this change on the new pin: it loads and runs. On the previous pin it still loads and runs, unchanged. All 57 tests under the XNNPACK quantizer directory pass with this change. Confirmed the argument cannot be read at annotation time by following it through the whole recipe: false when the annotation pass sees it, false after preparing, true once the model is moved to training, false again once it is moved back. Only the last of those describes how the model runs. --- backends/xnnpack/quantizer/xnnpack_quantizer_utils.py | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py b/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py index fca0f1b14c6..43659d2a061 100644 --- a/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py +++ b/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py @@ -1105,6 +1105,11 @@ def _is_share_obs_or_fq_op(op: Callable) -> bool: torch.ops.aten.slice.Tensor, torch.ops.aten.slice_copy.Tensor, torch.ops.aten.flatten.using_ints, + # Identity once the model is not training, which is how it is deployed. + # The training argument on the node cannot be read here: it is a switch + # that is flipped after this runs, which the train/eval test asserts. + torch.ops.aten.dropout.default, + torch.ops.aten.dropout_.default, ] From 286cd1a54973dd434dca81f0af7ec53a3b2c82b3 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 16:26:23 -0700 Subject: [PATCH 18/22] Update the operator counts 2.14 changes, measured not guessed Three tests pin an exact number of operators and now disagree with the graph: Expected to not find "..._cortex_m_dequantize_per_tensor_default" but found it Expected 2 occurrences of TOSA op TRANSPOSE but found 1 The two cortex-m tests count quantize and dequantize pairs. Sharing the observer across dropout, one commit earlier, means fewer places need a change of scale, which changes how many of those pairs survive. The Arm test counts layout conversions, and this PyTorch needs one fewer for a group norm on a channels-last input. Test Plan: Took the numbers from the graph rather than from the error message. Ran each model through its own pipeline, stopped after the passes, and counted the operators in the graph: three of each pair for the smaller model, two for mobilenet, one transpose for the group norm. Confirmed the graphs are right and not merely different, which is the part that decides whether to change a count at all. Each of these tests compares its outputs against eager after the counting stage, and that stage is what fails first, so it had never run. Ran all three with the counts above and the comparison passes: the smaller model reports counts matching and outputs compared, and the group norm case runs to completion once the count no longer stops it. The Arm file's other ten cases pass untouched, so the change is specific to the one that moved. --- backends/arm/test/misc/test_transpose_counts.py | 2 +- backends/cortex_m/test/models/test_ds_cnn.py | 4 ++-- backends/cortex_m/test/models/test_mobilenet_v2.py | 4 ++-- .../ir/converter/node_converter/test_clone_converter.py | 5 +++-- 4 files changed, 8 insertions(+), 7 deletions(-) diff --git a/backends/arm/test/misc/test_transpose_counts.py b/backends/arm/test/misc/test_transpose_counts.py index bd73ddfe0cb..643a14b3301 100644 --- a/backends/arm/test/misc/test_transpose_counts.py +++ b/backends/arm/test/misc/test_transpose_counts.py @@ -543,7 +543,7 @@ def forward(self, x: torch.Tensor): "groupnorm_channels_last": TransposeCountCase( GroupNormModule(), (torch.randn(1, 4, 4, 4).to(memory_format=torch.channels_last),), - 2, + 1, ), "cumsum_rank4_dim3_channels_last": TransposeCountCase( CumsumModule(), diff --git a/backends/cortex_m/test/models/test_ds_cnn.py b/backends/cortex_m/test/models/test_ds_cnn.py index 7bb55fb6a12..898047980d9 100644 --- a/backends/cortex_m/test/models/test_ds_cnn.py +++ b/backends/cortex_m/test/models/test_ds_cnn.py @@ -22,9 +22,9 @@ ops_after_transforms: dict[str, int] = { "executorch_exir_dialects_edge__ops_aten_view_copy_default": 1, - "executorch_exir_dialects_edge__ops_cortex_m_dequantize_per_tensor_default": 1, + "executorch_exir_dialects_edge__ops_cortex_m_dequantize_per_tensor_default": 3, "executorch_exir_dialects_edge__ops_cortex_m_pad_default": 1, - "executorch_exir_dialects_edge__ops_cortex_m_quantize_per_tensor_default": 1, + "executorch_exir_dialects_edge__ops_cortex_m_quantize_per_tensor_default": 3, "executorch_exir_dialects_edge__ops_cortex_m_quantized_avg_pool2d_default": 1, "executorch_exir_dialects_edge__ops_cortex_m_quantized_conv2d_default": 4, "executorch_exir_dialects_edge__ops_cortex_m_quantized_depthwise_conv2d_default": 5, diff --git a/backends/cortex_m/test/models/test_mobilenet_v2.py b/backends/cortex_m/test/models/test_mobilenet_v2.py index 20d8604c048..7230657e548 100644 --- a/backends/cortex_m/test/models/test_mobilenet_v2.py +++ b/backends/cortex_m/test/models/test_mobilenet_v2.py @@ -27,8 +27,8 @@ ops_after_transforms: dict[str, int] = { "executorch_exir_dialects_edge__ops_aten_view_copy_default": 1, - "executorch_exir_dialects_edge__ops_cortex_m_dequantize_per_tensor_default": 1, - "executorch_exir_dialects_edge__ops_cortex_m_quantize_per_tensor_default": 1, + "executorch_exir_dialects_edge__ops_cortex_m_dequantize_per_tensor_default": 2, + "executorch_exir_dialects_edge__ops_cortex_m_quantize_per_tensor_default": 2, "executorch_exir_dialects_edge__ops_cortex_m_quantized_add_default": 10, "executorch_exir_dialects_edge__ops_cortex_m_quantized_avg_pool2d_default": 1, "executorch_exir_dialects_edge__ops_cortex_m_quantized_conv2d_default": 35, diff --git a/backends/nxp/tests/ir/converter/node_converter/test_clone_converter.py b/backends/nxp/tests/ir/converter/node_converter/test_clone_converter.py index 5ee3db6752f..1238e31e246 100644 --- a/backends/nxp/tests/ir/converter/node_converter/test_clone_converter.py +++ b/backends/nxp/tests/ir/converter/node_converter/test_clone_converter.py @@ -220,8 +220,9 @@ def test_conv_dropout_no_quant( ], ) - # Clone with inplace=True should not produce clone edge op and vice versa - assert inplace_dropout ^ has_clone + # Neither spelling leaves a clone behind on this PyTorch: the out-of-place + # one used to and no longer does. + assert not has_clone @parameterized.expand([("QAT", True), ("PTQ", False)]) def test_clone_pool_view_copy_quant( From 6c6ad01e4c36db93777f5a9b9451ec8d38009491 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 17:40:09 -0700 Subject: [PATCH 19/22] Make the two new shims answer correctly, and pin the ATen test standard Three defects in the work this stack already added. The pinned allocation refused every request. Only one caller in the generated code tolerates that: the constants staging pool, which logs that it is falling back. The ordinary buffer allocator asks for the same thing whenever a buffer is pinned, and it wraps the call in a check that throws on any failure, so a model with such a buffer would have failed to initialize. Allocate ordinary memory instead. Pinned memory only lets a copy to the device overlap other work, so this is correct and slower, which is the outcome the other caller settles for anyway. The defined check asked whether the handle was null. This tensor type has an undefined state of its own: default construction and reset both leave a real object with no storage, and both were reported as defined. Ask the tensor. ATen-mode tests were left without a standard rather than given one, which relied on the toolchain default being C++20 while the commit that added the libraries argued a flag on the target is needed because the config pins C++17. Both cannot be true. Set it explicitly, the same way the library path does. Also fold the two ATen flavours of the test framework into the shared check. They were read from the dependency list after the wrapper had already consumed it, so they could never match. Test Plan: Read both callers of the pinned allocation in the installed PyTorch. The staging pool tests the result and logs a fallback; the ordinary allocator emits the call inside the macro that throws, so a refusal there ends initialization rather than degrading. Compared the defined check against the tensor's own: it returns whether storage is present, while the shim returned whether the pointer is non-null, and both the default constructor and reset leave the first false with the second true. Ran the test rule over the cases that matter, with the dependency list consumed first so the sequence matches the real one. Every ATen signal now comes away at C++20: naming libtorch, naming c10, naming either flavour of the test framework, or having ATen in the target name. A plain test still comes away at C++17, which is what the embedded builds rely on. Formatting is clean on the changed files. --- backends/aoti/common_shims_slim.cpp | 2 +- backends/cuda/runtime/shims/memory.cpp | 16 ++++++++-------- backends/cuda/runtime/shims/memory.h | 8 ++++---- .../xplat/executorch/build/runtime_wrapper.bzl | 17 +++++++++++++---- 4 files changed, 26 insertions(+), 17 deletions(-) diff --git a/backends/aoti/common_shims_slim.cpp b/backends/aoti/common_shims_slim.cpp index 0b90a6c9d33..d9e255bf220 100644 --- a/backends/aoti/common_shims_slim.cpp +++ b/backends/aoti/common_shims_slim.cpp @@ -72,7 +72,7 @@ AOTITorchError aoti_torch_is_defined(Tensor* tensor, bool* ret_is_defined) { if (ret_is_defined == nullptr) { return Error::InvalidArgument; } - *ret_is_defined = tensor != nullptr; + *ret_is_defined = tensor != nullptr && tensor->defined(); return Error::Ok; } diff --git a/backends/cuda/runtime/shims/memory.cpp b/backends/cuda/runtime/shims/memory.cpp index 199da37557d..ecc740ac8fc 100644 --- a/backends/cuda/runtime/shims/memory.cpp +++ b/backends/cuda/runtime/shims/memory.cpp @@ -203,14 +203,14 @@ AOTITorchError aoti_torch_empty_strided_pinned( int32_t device_type, int32_t device_index, SlimTensor** ret_new_tensor) { - (void)ndim; - (void)sizes_ptr; - (void)strides_ptr; - (void)dtype; - (void)device_type; - (void)device_index; - (void)ret_new_tensor; - return Error::NotSupported; + return aoti_torch_empty_strided( + ndim, + sizes_ptr, + strides_ptr, + dtype, + device_type, + device_index, + ret_new_tensor); } AOTITorchError aoti_torch_delete_tensor_object(SlimTensor* tensor) { diff --git a/backends/cuda/runtime/shims/memory.h b/backends/cuda/runtime/shims/memory.h index 158edfa6d92..1ed8277d7ef 100644 --- a/backends/cuda/runtime/shims/memory.h +++ b/backends/cuda/runtime/shims/memory.h @@ -96,12 +96,12 @@ AOTI_SHIM_EXPORT AOTITorchError aoti_torch_empty_strided( SlimTensor** ret_new_tensor); /** - * Reports that pinned host memory is unavailable. + * Allocates ordinary memory where pinned host memory was asked for. * - * There is no pinned allocator here, so the caller falls back to the - * synchronous copy it already handles: correct, slower only while loading. + * There is no pinned allocator here. Pinned memory only lets the copy to the + * device overlap other work, so ordinary memory is correct and slower. * - * @return Error::NotSupported + * @return Error::Ok on success */ AOTI_SHIM_EXPORT AOTITorchError aoti_torch_empty_strided_pinned( int64_t ndim, diff --git a/shim_et/xplat/executorch/build/runtime_wrapper.bzl b/shim_et/xplat/executorch/build/runtime_wrapper.bzl index 519f0f4e731..9f690d4c8f4 100644 --- a/shim_et/xplat/executorch/build/runtime_wrapper.bzl +++ b/shim_et/xplat/executorch/build/runtime_wrapper.bzl @@ -150,7 +150,14 @@ def _is_aten_target(kwargs): Keyed on exact dep names, not a substring: every label contains "torch". """ - aten_external_deps = ["c10", "libtorch", "libtorch_python", "torch-core-cpp"] + aten_external_deps = [ + "c10", + "gmock_aten", + "gtest_aten", + "libtorch", + "libtorch_python", + "torch-core-cpp", + ] for key in ["external_deps", "exported_external_deps"]: for dep in kwargs.get(key) or []: if dep in aten_external_deps: @@ -172,13 +179,15 @@ def _patch_test_compiler_flags(kwargs, aten_mode = False): aten_mode or "_aten" in name or "aten_" in name or - "gtest_aten" in external_deps or - "gmock_aten" in external_deps or _has_pytorch_dep(xplat_deps) or _has_pytorch_dep(fbcode_deps) ) - if not is_aten_test: + if is_aten_test: + kwargs["compiler_flags"] += [ + "-std=c++20", + ] + else: kwargs["compiler_flags"] += [ "-std=c++17", ] From b4134d2702e702f15769d06fa5567cd73237d40e Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 18:18:49 -0700 Subject: [PATCH 20/22] Close the rest of the review findings on this stack Four more derived shim spellings. The same gap this stack already closed for one CUDA operator and three Metal ones is open for four others, and the runtime already defines every one of the derived names, so each is a single entry. Added, and the set the CUDA test pins is updated with them. The macOS source build. It installed a pip cmake immediately before the same build the Docker change fixed, and it had no check on the result, so the same silent loss of the maths libraries could happen there unseen. Drop the pip cmake and assert on the built wheel, the way the image build now does. The Windows link stub. Nothing in the tree said how to produce it, which is how the last regeneration ended up in a different archive format with a build timestamp in it. Written down beside the file, including that a name only resolves if the library that becomes the DLL defines it. The custom object case in the placeholder sort. The helper partitioners use to tag constants consults parameters, buffers and lifted tensor constants and never custom objects, so that branch cannot be reached through the normal path, and the signature builder does not copy their values either. Removed; supporting them is its own change. The new overflow case in the add fixtures. Add rejects a floating point alpha before it reaches the range check, so there this case repeats coverage those fixtures already have. Said so beside it rather than leaving the impression it tests the new behaviour everywhere the macro expands. Test Plan: Derived the missing shim names with PyTorch's own function rather than by eye, over both backends' lists, and confirmed each of the four has a definition in the runtime. Compared the CUDA list against the set its test pins: thirteen against thirteen. Confirmed the overflow case dies early in the add fixtures by building the real scalar and asking it: two to the sixty three reports itself as floating point, which is what the alpha check rejects. Confirmed the custom object branch is unreachable by reading the tagging helper: it looks at three maps and custom objects are not among them. Read the macOS path against the Docker one to confirm the exposure is the same shape, and matched the repair. I have no macOS runner, so whether that path ever loses the maths libraries in practice is still unproven; the check is there so it cannot happen quietly. Formatting is clean on the changed files, checked with the versions CI pins. --- .ci/scripts/utils.sh | 9 +++ backends/apple/metal/metal_backend.py | 3 + backends/cuda/cuda_backend.py | 1 + backends/cuda/runtime/aoti_cuda_shims.lib.md | 31 ++++++++++ backends/cuda/tests/test_sort_shim.py | 1 + exir/lowered_backend_module.py | 3 +- kernels/test/ScalarOverflowTestMacros.h | 64 ++++++++++---------- 7 files changed, 79 insertions(+), 33 deletions(-) create mode 100644 backends/cuda/runtime/aoti_cuda_shims.lib.md diff --git a/.ci/scripts/utils.sh b/.ci/scripts/utils.sh index 45ac584b59d..b7f4fee6b4c 100644 --- a/.ci/scripts/utils.sh +++ b/.ci/scripts/utils.sh @@ -136,9 +136,18 @@ install_pytorch_and_domains() { export BUILD_IGNORE_SVE_UNAVAILABLE=1 fi # PyTorch dropped setup.py; isolation off reuses the deps installed above. + # Drop the pip cmake first: scikit-build-core prefers it over the one on + # PATH, and it searches site-packages, where the maths libraries are not. + pip uninstall -y cmake || true pip install build USE_DISTRIBUTED=1 python -m build --wheel --no-isolation pip install "$(echo dist/*.whl)" + # A build with no BLAS succeeds silently, so check rather than assume. + (cd / && python -c " +import torch +assert torch._C.has_lapack, 'built without LAPACK' +torch.linalg.qr(torch.randn(4, 4)) +") # Invariant: the basename the build just produced must match the cache # URL we'd reconstruct on the next run. If they diverge (someone edits diff --git a/backends/apple/metal/metal_backend.py b/backends/apple/metal/metal_backend.py index aa931d8ae4d..9914e5b2e83 100644 --- a/backends/apple/metal/metal_backend.py +++ b/backends/apple/metal/metal_backend.py @@ -37,11 +37,14 @@ def get_supported_fallback_kernels(cls) -> Dict[str, Any]: "aoti_torch_mps_convolution": None, "aoti_torch_mps_mm_out": None, "at::_ops::_scaled_dot_product_attention_math_for_mps::call": None, + "aoti_torch_mps__scaled_dot_product_attention_math_for_mps": None, "at::_ops::_scaled_dot_product_attention_math_for_mps_v2::call": None, + "aoti_torch_mps__scaled_dot_product_attention_math_for_mps_v2": None, "torchao::_linear_fp_act_4bit_weight": None, # Each custom op twice: as registered, and as Inductor derives it. "aoti_torch_mps__linear_fp_act_4bit_weight": None, "at::_ops::topk::call": None, + "aoti_torch_mps_topk": None, "metal::gather_qmv": None, "aoti_torch_mps_gather_qmv": None, "metal::gated_delta_rule": None, diff --git a/backends/cuda/cuda_backend.py b/backends/cuda/cuda_backend.py index 97a4f67a504..89e1b5ad79a 100644 --- a/backends/cuda/cuda_backend.py +++ b/backends/cuda/cuda_backend.py @@ -672,6 +672,7 @@ def get_supported_fallback_kernels(cls) -> Dict[str, Any]: # Also under the shim name Inductor derives for it. "aoti_torch_cuda__weight_int4pack_mm": None, "at::_ops::sort_stable::call": None, + "aoti_torch_cuda_sort_stable": None, "aoti_torch_cuda_randint_low_out": None, "executorch_cuda::int4_plain_mm": None, "aoti_torch_cuda_int4_plain_mm": None, diff --git a/backends/cuda/runtime/aoti_cuda_shims.lib.md b/backends/cuda/runtime/aoti_cuda_shims.lib.md new file mode 100644 index 00000000000..a2e03cb3e5e --- /dev/null +++ b/backends/cuda/runtime/aoti_cuda_shims.lib.md @@ -0,0 +1,31 @@ +# aoti_cuda_shims.lib + +Import library for `aoti_cuda_shims.dll`. Lowering a CUDA model for a Windows target +links the generated wrapper against this file, so it has to advertise every shim name +the wrapper can reference. It is checked in because the build that produces the DLL +does not run on the machines that lower for Windows. + +Regenerate it whenever a shim is added, which in practice means whenever the PyTorch +pin moves and the generated wrapper starts calling something new: + +``` +nm --defined-only aoti_cuda_shims.lib \ + | sed -n 's/.* T _\?\(aoti_torch_[a-z_0-9]*\)$/\1/p' \ + | grep -v '^_imp_' | sort -u > exports.txt +# add the new names to exports.txt, then +{ echo 'LIBRARY aoti_cuda_shims.dll'; echo 'EXPORTS'; sed 's/^/ /' exports.txt; } \ + > aoti_cuda_shims.def +x86_64-w64-mingw32-dlltool -d aoti_cuda_shims.def -l aoti_cuda_shims.lib \ + --dllname aoti_cuda_shims.dll +# zero the archive metadata so the file is reproducible +mkdir extract && cd extract && x86_64-w64-mingw32-ar x ../aoti_cuda_shims.lib \ + && x86_64-w64-mingw32-ar rcsD ../aoti_cuda_shims.lib * +``` + +Then check the result is a superset of what it replaced and that no member carries a +timestamp, since this file ships in the wheel and a stamped one makes two builds of +the same source differ. + +A name only resolves if something in the DLL defines it. The DLL is built from the +CUDA shims plus the SlimTensor common shims, so a shim added to the ETensor common +shims will link and then fail to load. diff --git a/backends/cuda/tests/test_sort_shim.py b/backends/cuda/tests/test_sort_shim.py index 1e5b34b7396..da1ee0c3e2c 100644 --- a/backends/cuda/tests/test_sort_shim.py +++ b/backends/cuda/tests/test_sort_shim.py @@ -39,6 +39,7 @@ "at::_ops::_weight_int4pack_mm::call", "aoti_torch_cuda__weight_int4pack_mm", "at::_ops::sort_stable::call", + "aoti_torch_cuda_sort_stable", "aoti_torch_cuda_randint_low_out", "executorch_cuda::int4_plain_mm", "aoti_torch_cuda_int4_plain_mm", diff --git a/exir/lowered_backend_module.py b/exir/lowered_backend_module.py index 11708710333..358bcb25e8c 100644 --- a/exir/lowered_backend_module.py +++ b/exir/lowered_backend_module.py @@ -407,7 +407,6 @@ def arrange_graph_placeholders( params_map = graph_sign.inputs_to_parameters buffers_map = graph_sign.inputs_to_buffers constants_map = graph_sign.inputs_to_lifted_tensor_constants - custom_objs_map = graph_sign.inputs_to_lifted_custom_objs param_nodes = [] buffer_nodes = [] constant_nodes = [] @@ -421,7 +420,7 @@ def arrange_graph_placeholders( param_nodes.append(node) elif node.name in buffers_map and is_tagged: buffer_nodes.append(node) - elif (node.name in constants_map or node.name in custom_objs_map) and is_tagged: + elif node.name in constants_map and is_tagged: constant_nodes.append(node) else: input_nodes.append(node) diff --git a/kernels/test/ScalarOverflowTestMacros.h b/kernels/test/ScalarOverflowTestMacros.h index f8163254fc8..4ab219f8acc 100644 --- a/kernels/test/ScalarOverflowTestMacros.h +++ b/kernels/test/ScalarOverflowTestMacros.h @@ -11,35 +11,37 @@ // Macro to generate scalar overflow test cases for a given test suite. // The test suite must have a method called expect_bad_scalar_value_dies // that takes a template parameter for ScalarType and a Scalar value. -#define GENERATE_SCALAR_OVERFLOW_TESTS(TEST_SUITE_NAME) \ - TEST_F(TEST_SUITE_NAME, ByteTensorTooLargeScalarDies) { \ - /* Cannot be represented by a uint8_t. */ \ - expect_bad_scalar_value_dies(256); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, CharTensorTooSmallScalarDies) { \ - /* Cannot be represented by a int8_t. */ \ - expect_bad_scalar_value_dies(-129); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, ShortTensorTooLargeScalarDies) { \ - /* Cannot be represented by a int16_t. */ \ - expect_bad_scalar_value_dies(32768); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, FloatTensorTooSmallScalarDies) { \ - /* Cannot be represented by a float. */ \ - expect_bad_scalar_value_dies(-3.41e+38); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, FloatTensorTooLargeScalarDies) { \ - /* Cannot be represented by a float. */ \ - expect_bad_scalar_value_dies(3.41e+38); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, LongTensorTooLargeScalarDies) { \ - /* 2^63 is one past the largest int64_t. It is exactly \ - * representable as a double, so the conversion is silent \ - * unless the range check rejects it. */ \ - expect_bad_scalar_value_dies(9223372036854775808.0); \ +#define GENERATE_SCALAR_OVERFLOW_TESTS(TEST_SUITE_NAME) \ + TEST_F(TEST_SUITE_NAME, ByteTensorTooLargeScalarDies) { \ + /* Cannot be represented by a uint8_t. */ \ + expect_bad_scalar_value_dies(256); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, CharTensorTooSmallScalarDies) { \ + /* Cannot be represented by a int8_t. */ \ + expect_bad_scalar_value_dies(-129); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, ShortTensorTooLargeScalarDies) { \ + /* Cannot be represented by a int16_t. */ \ + expect_bad_scalar_value_dies(32768); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, FloatTensorTooSmallScalarDies) { \ + /* Cannot be represented by a float. */ \ + expect_bad_scalar_value_dies(-3.41e+38); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, FloatTensorTooLargeScalarDies) { \ + /* Cannot be represented by a float. */ \ + expect_bad_scalar_value_dies(3.41e+38); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, LongTensorTooLargeScalarDies) { \ + /* 2^63 is one past the largest int64_t. It is exactly \ + * representable as a double, so the conversion is silent \ + * unless the range check rejects it. The add suites reject it earlier, \ + * on the alpha type, so there this case only repeats their existing \ + * floating-point alpha coverage. */ \ + expect_bad_scalar_value_dies(9223372036854775808.0); \ } From d3b37637e56ddb2a190a285b56a1f7591c14d692 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 23:46:13 -0700 Subject: [PATCH 21/22] Correct the build and shim work this change adds The macOS source build installed PyTorch's build requirements into the active environment. Those requirements pin their own CMake, which replaces the one already there with a launcher script, and removing it afterwards deletes the file rather than putting the original back, so a machine with no second CMake could not start the build. Build in a throwaway environment that can see the active one instead, which is what the container path already does. The Windows link stub advertised two of the seven shim names the backend offers, so lowering a model that needs any of the other five failed to link. Regenerated with all of them: fifty one names, none lost. The note beside the file also built its archive in place, which keeps the generator's reverse member order, so following it did not reproduce the file it documents. It now builds at a fresh path. Asking for pinned memory forwarded the device type without looking at it, so asking for a device gave device memory from a function that says it returns host memory. It now refuses anything but CPU, as PyTorch's own does, and documents each argument. The check for whether a tensor is defined asked only whether the handle was null, while the tensor type has an undefined state of its own. The declaration beside it said the opposite of what the body does. The ATen standard decision read two dependency lists after the wrapper had already consumed them, so three tests that declare ATen only through those lists were compiled at C++17 against headers that need C++20. It now reads them before they are consumed. The extension rule was also left out when the other three were wired up, so a module naming the Python library kept the old standard. Test Plan: Reproduced the CMake deletion twice, once in a container-style environment and once in a plain one: installing the pinned CMake replaces the existing binary with a small launcher, and removing it leaves nothing behind. Compared the names the backend offers against the names the stub carries: seven offered, two present before, all seven present now. Ran the documented regeneration against its own output and the result is identical to the file, which it was not before. Ran the standard decision over the cases that matter, with the dependency lists consumed first so the sequence matches the real one. The three tests that were at C++17 now come away at C++20, and a plain test is unchanged from before this change, which I checked by running both versions side by side. --- .ci/scripts/utils.sh | 22 ++++++++++-------- backends/aoti/common_shims_slim.h | 2 +- backends/cuda/runtime/aoti_cuda_shims.lib | Bin 38758 -> 43802 bytes backends/cuda/runtime/aoti_cuda_shims.lib.md | 9 +++++-- backends/cuda/runtime/shims/memory.cpp | 7 ++++++ backends/cuda/runtime/shims/memory.h | 20 +++++++++++----- .../executorch/build/runtime_wrapper.bzl | 19 +++++++-------- 7 files changed, 50 insertions(+), 29 deletions(-) diff --git a/.ci/scripts/utils.sh b/.ci/scripts/utils.sh index b7f4fee6b4c..3d573daa395 100644 --- a/.ci/scripts/utils.sh +++ b/.ci/scripts/utils.sh @@ -127,20 +127,22 @@ install_pytorch_and_domains() { if [[ "${torch_wheel_not_found}" == "1" ]]; then echo "No cached wheel found, continue with building PyTorch at ${TORCH_VERSION}" - # Install PyTorch's own build-time deps so the source build does not - # silently inherit them from whatever else happens to be in the env - # (e.g. executorch's requirements-ci.txt). - pip install -r requirements-build.txt git submodule update --init --recursive if [[ "$(uname -m)" == "aarch64" ]]; then export BUILD_IGNORE_SVE_UNAVAILABLE=1 fi - # PyTorch dropped setup.py; isolation off reuses the deps installed above. - # Drop the pip cmake first: scikit-build-core prefers it over the one on - # PATH, and it searches site-packages, where the maths libraries are not. - pip uninstall -y cmake || true - pip install build - USE_DISTRIBUTED=1 python -m build --wheel --no-isolation + # PyTorch dropped setup.py. Build in a throwaway environment that can see the + # active one, so its build requirements, which pin a cmake that would be + # preferred over the one on PATH, cannot disturb what is installed here. + # + # Keep in sync with pytorch/pyproject.toml [build-system].requires. + local build_venv=/tmp/pytorch-build-venv + rm -rf "${build_venv}" + python -m venv --system-site-packages "${build_venv}" + "${build_venv}/bin/pip" install build "scikit-build-core>=1.0" ninja \ + "packaging>=24.2" "typing-extensions>=4.10.0" pyyaml six numpy + USE_DISTRIBUTED=1 "${build_venv}/bin/python" -m build --wheel --no-isolation + rm -rf "${build_venv}" pip install "$(echo dist/*.whl)" # A build with no BLAS succeeds silently, so check rather than assume. (cd / && python -c " diff --git a/backends/aoti/common_shims_slim.h b/backends/aoti/common_shims_slim.h index 9d05b542a52..0d60f578605 100644 --- a/backends/aoti/common_shims_slim.h +++ b/backends/aoti/common_shims_slim.h @@ -51,7 +51,7 @@ aoti_torch_get_dim(Tensor* tensor, int64_t* ret_dim); AOTI_SHIM_EXPORT AOTITorchError aoti_torch_get_numel(Tensor* tensor, int64_t* ret_numel); -// PyTorch has an undefined-tensor state with no equivalent here: null check. +// Undefined means either a null handle or a tensor whose storage was released. AOTI_SHIM_EXPORT AOTITorchError aoti_torch_is_defined(Tensor* tensor, bool* ret_is_defined); diff --git a/backends/cuda/runtime/aoti_cuda_shims.lib b/backends/cuda/runtime/aoti_cuda_shims.lib index 7cd16022a250710a46b771a86e39ade055576088..c0d61c611008c3897350d7b0950a8e0229607ba3 100644 GIT binary patch delta 1975 zcmZvbZA_C_6vyu^FRsuQOA!{r+uTfu1o{AdpuE3~4kBzShyj9mTB#zer3hsZa8_g@ zPNd**S;j)5WEwx1MRS(;VViNEmM#0REygWbw!mWMK1_yYOSWugJGZx11e*MA`~2@Y z|L5FuPk;GPH-A?*TM|%h_N>sy8OInq!dTK2V=4P2silkw(o-dwcNi0FpONesWlWH5 zljI~aCdiAGD1R{~C|s2k-D6Bp@|C3gT7a?gYm8M*OKLrg2^>w51~L{jZk4>c!kA#s z&ytoe858XPK+=AKF+t}(iI&Egz(qL%_a~C>w;2=k2HI#Lyjc<-1qlXjONOp8CK#EQ zjP)`mIR1)cobm({RFGhjVg+aVB-2#9AebbXqfQ7e+>#W$dxbH<{H$bwiV5G7>C?P$$!$qkr^0M>l7WKMO@qSflc1ta zgTm^QAp534VO0dh54if=^h(X+8wRT`SqYW2M$~*me&+J|do;gqpzDYh+8^`{__aa5 z>!9Zl?9H(xh{LVbmhBnvxa37ljTze#r`nb9*n(o*h94QUnAkx@QEVz)GN<#R&{`Z? z?V+^<9vaegszP7Sqa|@{2mYJV%wfrHOkCG6&Fh7lv`=Aa)Vg_xH%%x`gR)Wu|4C`# zaI5U8X{7BNO(R?4+9IYAl@iF*S|Lp{O&c0U9gf`BM)%0pGwu;px~Ej<%RbE4JS{nL zzZ*>>ThEwAR4?3cq{E|HD{2)j{DrL>N?tbuaKpK>b^UONiAh>7iisifCcNKG4c=34 zlxb_{W3DdjwN~70kVK7(4Oa98TR2a|WqSvQo<;>5}SVGpYC8VYTjiQSQe> z3Z@()wQ|^uV~5K*jF05wuibR2=utEL_nZ;iddfJs#`3YDn$&y8gt`UuJk*zp<5Y>4 zvTp~>(65ii7~f6~@#FO7eo|*nn&5rC9=|x6%i)g+8ZV)~amob8hV&Tplls%7g3f;O zbxzRptIrv$i zkweCLImU1is2da8@cPL*4tFo);r6%oah?lv6Uh*qG~j!u99(BpaB1=c=XvOwa_}fy zp$bRab3-L@D`ZY5W7XMQ(u-95#u4CL!QapArRNeAUB&dA56kfxxIA0Ib!q`-ohPdT TyaXzbQdJdP#W@)L&OzP(>kObv delta 1158 zcmX}qUrdu%6aesZtCT{bB9k)y5T8YTCh=DZA`pP zL3rDiPMx-<1nOTB%q73GPGHR@<3riA)o*aTx(06+OIq&n3X~;kaT)&Xn&Sm4&Y3sh zn>8O{^fO7ZVNO`|xq+kJbc)oK32fg|K>vFAo%3Y|HTb?1SY0k@G{||g6$Uo=41Afm zkSca=61c8X()NqIGQE|zXiOzX!#aWXDoFzYuMmsz$(~G_-@Qg)teTtY|VquAlEt!^a(Fx_fY~pzNk0fsc|68a%{RvAud;k%cqKE69I@r@h-M>Gv+~qttI; zamrGvJ<1cGYLoP!o1?aU2IUlP8_%6 zc54zv@l%1X`rZ1j{U>ZVXHCW_tA$!ezTjR?y6M?Dp4AmmE9q4IL%G1t0ZHEl>V$Bj zFfz$CmM?JqV7$NSsDZnT6#DxD_m^?nt-pcuLk3>IoJzl3S}*YJ88?0H^a@iR;Ea(? d-%s#m{yOX?ag~q4MZRl%A(Q9wX?n87^gnWWaO(g7 diff --git a/backends/cuda/runtime/aoti_cuda_shims.lib.md b/backends/cuda/runtime/aoti_cuda_shims.lib.md index a2e03cb3e5e..af2fb65ca13 100644 --- a/backends/cuda/runtime/aoti_cuda_shims.lib.md +++ b/backends/cuda/runtime/aoti_cuda_shims.lib.md @@ -11,7 +11,7 @@ pin moves and the generated wrapper starts calling something new: ``` nm --defined-only aoti_cuda_shims.lib \ | sed -n 's/.* T _\?\(aoti_torch_[a-z_0-9]*\)$/\1/p' \ - | grep -v '^_imp_' | sort -u > exports.txt + | sort -u > exports.txt # add the new names to exports.txt, then { echo 'LIBRARY aoti_cuda_shims.dll'; echo 'EXPORTS'; sed 's/^/ /' exports.txt; } \ > aoti_cuda_shims.def @@ -19,9 +19,14 @@ x86_64-w64-mingw32-dlltool -d aoti_cuda_shims.def -l aoti_cuda_shims.lib \ --dllname aoti_cuda_shims.dll # zero the archive metadata so the file is reproducible mkdir extract && cd extract && x86_64-w64-mingw32-ar x ../aoti_cuda_shims.lib \ - && x86_64-w64-mingw32-ar rcsD ../aoti_cuda_shims.lib * + && rm -f ../aoti_cuda_shims.lib \ + && x86_64-w64-mingw32-ar rcsD ../aoti_cuda_shims.lib $(ls | sort) ``` +The archive has to be built at a fresh path. Updating it in place keeps the reverse +member order the generator produced, so the same export list would not give the same +bytes twice. + Then check the result is a superset of what it replaced and that no member carries a timestamp, since this file ships in the wheel and a stamped one makes two builds of the same source differ. diff --git a/backends/cuda/runtime/shims/memory.cpp b/backends/cuda/runtime/shims/memory.cpp index ecc740ac8fc..976b282a29c 100644 --- a/backends/cuda/runtime/shims/memory.cpp +++ b/backends/cuda/runtime/shims/memory.cpp @@ -203,6 +203,13 @@ AOTITorchError aoti_torch_empty_strided_pinned( int32_t device_type, int32_t device_index, SlimTensor** ret_new_tensor) { + ET_CHECK_OR_RETURN_ERROR( + static_cast(device_type) == DeviceType::CPU, + InvalidArgument, + "aoti_torch_empty_strided_pinned: pinned memory is host memory, so the " + "device type must be CPU, got %d", + device_type); + return aoti_torch_empty_strided( ndim, sizes_ptr, diff --git a/backends/cuda/runtime/shims/memory.h b/backends/cuda/runtime/shims/memory.h index 1ed8277d7ef..03e2b3ed18d 100644 --- a/backends/cuda/runtime/shims/memory.h +++ b/backends/cuda/runtime/shims/memory.h @@ -96,12 +96,20 @@ AOTI_SHIM_EXPORT AOTITorchError aoti_torch_empty_strided( SlimTensor** ret_new_tensor); /** - * Allocates ordinary memory where pinned host memory was asked for. - * - * There is no pinned allocator here. Pinned memory only lets the copy to the - * device overlap other work, so ordinary memory is correct and slower. - * - * @return Error::Ok on success + * Allocates ordinary host memory where pinned host memory was asked for. + * + * There is no pinned allocator here. Pinning only lets a copy to the device + * overlap other work, so ordinary memory is correct and slower. + * + * @param ndim Number of dimensions + * @param sizes_ptr Pointer to the sizes, ndim of them + * @param strides_ptr Pointer to the strides, ndim of them, or null for + * contiguous + * @param dtype Element type, as a scalar type value + * @param device_type Must be CPU, since pinned memory is host memory + * @param device_index Device index, unused for CPU + * @param ret_new_tensor Receives the new tensor + * @return Error::Ok on success, Error::InvalidArgument if the device is not CPU */ AOTI_SHIM_EXPORT AOTITorchError aoti_torch_empty_strided_pinned( int64_t ndim, diff --git a/shim_et/xplat/executorch/build/runtime_wrapper.bzl b/shim_et/xplat/executorch/build/runtime_wrapper.bzl index 9f690d4c8f4..445c1aa3b98 100644 --- a/shim_et/xplat/executorch/build/runtime_wrapper.bzl +++ b/shim_et/xplat/executorch/build/runtime_wrapper.bzl @@ -162,27 +162,23 @@ def _is_aten_target(kwargs): for dep in kwargs.get(key) or []: if dep in aten_external_deps: return True + for key in ["xplat_deps", "fbcode_deps"]: + if _has_pytorch_dep(kwargs.get(key)): + return True return False def _patch_test_compiler_flags(kwargs, aten_mode = False): if "compiler_flags" not in kwargs: kwargs["compiler_flags"] = [] - # Determine C++ standard based on whether this is an aten test. - # Aten tests require at least C++20 to compile against PyTorch, while - # non-aten tests are pinned to C++17 for embedded. + # A test that compiles against ATen needs C++20, which PyTorch's headers + # require. Every other test stays at C++17, which the embedded builds use. name = kwargs.get("name", "") - external_deps = kwargs.get("external_deps", []) - xplat_deps = kwargs.get("xplat_deps", []) - fbcode_deps = kwargs.get("fbcode_deps", []) is_aten_test = ( aten_mode or "_aten" in name or - "aten_" in name or - _has_pytorch_dep(xplat_deps) or - _has_pytorch_dep(fbcode_deps) + "aten_" in name ) - if is_aten_test: kwargs["compiler_flags"] += [ "-std=c++20", @@ -360,11 +356,14 @@ def _cxx_test(*args, **kwargs): env.cxx_test(*args, **kwargs) def _cxx_python_extension(*args, **kwargs): + # Before _patch_kwargs_common, which consumes external_deps. + aten_mode = _is_aten_target(kwargs) _patch_kwargs_common(kwargs) _remove_caffe2_deps(kwargs) kwargs["srcs"] = _patch_executorch_references(kwargs["srcs"]) if "types" in kwargs: kwargs["types"] = _patch_executorch_references(kwargs["types"]) + _patch_aten_mode_std(kwargs, aten_mode) env.cxx_python_extension(*args, **kwargs) def _export_file(*args, **kwargs): From 5ada9cf73474d54476d2851475276ff48d6ca151 Mon Sep 17 00:00:00 2001 From: PyTorch Bot Date: Thu, 3 Sep 2026 23:46:34 -0700 Subject: [PATCH 22/22] Say what the tests and comments actually check The two model tests counted a dim order clone that this PyTorch no longer emits. Those entries were deleted, but the count check only looks at the keys it is given, so deleting a key stops checking it rather than asserting it is gone. The harness already has a way to say an operator must be absent, so the clone is now asserted absent and the counts stay separate. One case expected a float output where nothing survives lowering to carry the integer type, which sent it into a branch whose checks any dequantize satisfies. It is skipped with a reason now, the way the other unsupported case in that file is. Four comments said things that are not true. Dropout is listed under both spellings, so the note about each operator appearing twice did not match the list and sat against one entry rather than the list. Converting a value past the end of the integer range is undefined rather than silent. The pinned node floor comes from a dependency rather than from the package named. And the pin file was missing the newline both its siblings have. Test Plan: Ran the absent check directly against a graph with the clone and one without: it raises on the first and passes on the second, so it can fail. Read the count check to confirm it iterates only the keys given, which is why deleting one removes the signal. Measured the comment claims rather than reasoning about them: the list has six entries with a second spelling and four with only one, converting two to the sixty third power reports an out of range error under a sanitizer, and the package registry shows the node floor comes from the dependency, not the package. Formatting checked with the versions the lint job pins. --- .ci/docker/ci_commit_pins/pytorch.txt | 2 +- .ci/docker/common/install_docs_reqs.sh | 3 +- backends/apple/metal/metal_backend.py | 3 +- .../cortex_m/test/misc/test_portable_int8.py | 9 +-- backends/cortex_m/test/models/test_ds_cnn.py | 12 +++- .../cortex_m/test/models/test_mobilenet_v2.py | 6 ++ backends/cortex_m/test/tester.py | 4 +- .../quantizer/xnnpack_quantizer_utils.py | 4 +- kernels/test/ScalarOverflowTestMacros.h | 65 +++++++++---------- 9 files changed, 63 insertions(+), 45 deletions(-) diff --git a/.ci/docker/ci_commit_pins/pytorch.txt b/.ci/docker/ci_commit_pins/pytorch.txt index 212c7580b6f..6c3fe42ddf3 100644 --- a/.ci/docker/ci_commit_pins/pytorch.txt +++ b/.ci/docker/ci_commit_pins/pytorch.txt @@ -1 +1 @@ -release/2.14 \ No newline at end of file +release/2.14 diff --git a/.ci/docker/common/install_docs_reqs.sh b/.ci/docker/common/install_docs_reqs.sh index 73f361acb00..2794ffb8fc9 100755 --- a/.ci/docker/common/install_docs_reqs.sh +++ b/.ci/docker/common/install_docs_reqs.sh @@ -20,7 +20,8 @@ if [ -n "$BUILD_DOCS" ]; then apt-get update apt-get install -y --no-install-recommends yarn - # 0.18.5 wants node >= 22.12; this is the last release node 16 can run. + # katex 0.18.5 requires commander@15 / node >= 22.12; pin to the last + # release compatible with the node 16 installed above yarn global add katex@0.18.4 --prefix /usr/local sudo apt-get -y install doxygen diff --git a/backends/apple/metal/metal_backend.py b/backends/apple/metal/metal_backend.py index 9914e5b2e83..0b9a852343b 100644 --- a/backends/apple/metal/metal_backend.py +++ b/backends/apple/metal/metal_backend.py @@ -32,6 +32,8 @@ def get_device_name(cls) -> str: @classmethod def get_supported_fallback_kernels(cls) -> Dict[str, Any]: return { + # An operator named the way it is registered also needs the name + # Inductor derives for it, so several appear under both. "aoti_torch_mps_addmm_out": None, "aoti_torch_mps_bmm_out": None, "aoti_torch_mps_convolution": None, @@ -41,7 +43,6 @@ def get_supported_fallback_kernels(cls) -> Dict[str, Any]: "at::_ops::_scaled_dot_product_attention_math_for_mps_v2::call": None, "aoti_torch_mps__scaled_dot_product_attention_math_for_mps_v2": None, "torchao::_linear_fp_act_4bit_weight": None, - # Each custom op twice: as registered, and as Inductor derives it. "aoti_torch_mps__linear_fp_act_4bit_weight": None, "at::_ops::topk::call": None, "aoti_torch_mps_topk": None, diff --git a/backends/cortex_m/test/misc/test_portable_int8.py b/backends/cortex_m/test/misc/test_portable_int8.py index 26bc61e1d63..41f7f254863 100644 --- a/backends/cortex_m/test/misc/test_portable_int8.py +++ b/backends/cortex_m/test/misc/test_portable_int8.py @@ -387,9 +387,7 @@ def _quantize_and_export( torch.ops.aten.dropout.default, _build_module(lambda x, y: torch.ops.aten.dropout.default(x, 0.1, False)), (torch.randn(2, 3, 4, 5), torch.randn(2, 3, 4, 5)), - # Not training, so this is an identity and nothing survives lowering to - # carry int8. Only the dequantize is left, which is float by definition. - torch.float32, + None, ), "dropout_": OpCase( torch.ops.aten.dropout_.default, @@ -718,7 +716,10 @@ def _quantize_and_export( OP_CASES, xfails=xfails, strict=False, - skips={"while_loop": "Has been observed to hang randomly."}, + skips={ + "while_loop": "Has been observed to hang randomly.", + "dropout": "Not training, so it folds away and no node survives to carry int8.", + }, ) def test_shared_qspec_portable_int8_ops(op_case: OpCase) -> None: tester = CortexMTester(op_case.module, op_case.example_inputs) diff --git a/backends/cortex_m/test/models/test_ds_cnn.py b/backends/cortex_m/test/models/test_ds_cnn.py index 898047980d9..ba5b2da6b78 100644 --- a/backends/cortex_m/test/models/test_ds_cnn.py +++ b/backends/cortex_m/test/models/test_ds_cnn.py @@ -41,11 +41,21 @@ } +ops_absent_after_transforms: list[str] = [ + "executorch_exir_dialects_edge__ops_dim_order_ops__clone_dim_order_default", +] + + @parametrize("test_case", test_cases) def test_dialect_ds_cnn(test_case): inputs = test_case.get_example_inputs() tester = CortexMTester(test_case.model, inputs) - tester.test_dialect(ops_before_transforms, ops_after_transforms, qtol=1) + tester.test_dialect( + ops_before_transforms, + ops_after_transforms, + qtol=1, + ops_absent_after_transforms=ops_absent_after_transforms, + ) @parametrize("test_case", test_cases) diff --git a/backends/cortex_m/test/models/test_mobilenet_v2.py b/backends/cortex_m/test/models/test_mobilenet_v2.py index 7230657e548..9bc99e4bf2c 100644 --- a/backends/cortex_m/test/models/test_mobilenet_v2.py +++ b/backends/cortex_m/test/models/test_mobilenet_v2.py @@ -52,6 +52,11 @@ } +ops_absent_after_transforms: list[str] = [ + "executorch_exir_dialects_edge__ops_dim_order_ops__clone_dim_order_default", +] + + @parametrize("test_case", test_cases) def test_dialect_mv2(test_case): inputs = test_case.get_example_inputs() @@ -61,6 +66,7 @@ def test_dialect_mv2(test_case): ops_after_transforms, qtol=10, calibration_samples=calibration_samples, + ops_absent_after_transforms=ops_absent_after_transforms, ) # assert that top 1 output matches diff --git a/backends/cortex_m/test/tester.py b/backends/cortex_m/test/tester.py index a1b5245b80b..b644db4e6c1 100644 --- a/backends/cortex_m/test/tester.py +++ b/backends/cortex_m/test/tester.py @@ -132,6 +132,7 @@ def test_dialect( qtol=0, atol=1e-03, calibration_samples=None, + ops_absent_after_transforms=None, ): """ Test the python dialect op implementation. @@ -142,13 +143,14 @@ def test_dialect( ) else: quantization_stage = None - self.quantize(quantization_stage) self.export() self.to_edge() self.check_count(ops_before_transforms) self.run_passes() self.check_count(ops_after_transforms) + if ops_absent_after_transforms: + self.check_not(ops_absent_after_transforms) self.run_method_and_compare_outputs( inputs=self.example_inputs, qtol=qtol, atol=atol ) diff --git a/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py b/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py index 43659d2a061..751388d9221 100644 --- a/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py +++ b/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py @@ -1105,9 +1105,7 @@ def _is_share_obs_or_fq_op(op: Callable) -> bool: torch.ops.aten.slice.Tensor, torch.ops.aten.slice_copy.Tensor, torch.ops.aten.flatten.using_ints, - # Identity once the model is not training, which is how it is deployed. - # The training argument on the node cannot be read here: it is a switch - # that is flipped after this runs, which the train/eval test asserts. + # Identity once not training, which is how a quantized model is deployed. torch.ops.aten.dropout.default, torch.ops.aten.dropout_.default, ] diff --git a/kernels/test/ScalarOverflowTestMacros.h b/kernels/test/ScalarOverflowTestMacros.h index 4ab219f8acc..6567ad71564 100644 --- a/kernels/test/ScalarOverflowTestMacros.h +++ b/kernels/test/ScalarOverflowTestMacros.h @@ -11,37 +11,36 @@ // Macro to generate scalar overflow test cases for a given test suite. // The test suite must have a method called expect_bad_scalar_value_dies // that takes a template parameter for ScalarType and a Scalar value. -#define GENERATE_SCALAR_OVERFLOW_TESTS(TEST_SUITE_NAME) \ - TEST_F(TEST_SUITE_NAME, ByteTensorTooLargeScalarDies) { \ - /* Cannot be represented by a uint8_t. */ \ - expect_bad_scalar_value_dies(256); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, CharTensorTooSmallScalarDies) { \ - /* Cannot be represented by a int8_t. */ \ - expect_bad_scalar_value_dies(-129); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, ShortTensorTooLargeScalarDies) { \ - /* Cannot be represented by a int16_t. */ \ - expect_bad_scalar_value_dies(32768); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, FloatTensorTooSmallScalarDies) { \ - /* Cannot be represented by a float. */ \ - expect_bad_scalar_value_dies(-3.41e+38); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, FloatTensorTooLargeScalarDies) { \ - /* Cannot be represented by a float. */ \ - expect_bad_scalar_value_dies(3.41e+38); \ - } \ - \ - TEST_F(TEST_SUITE_NAME, LongTensorTooLargeScalarDies) { \ - /* 2^63 is one past the largest int64_t. It is exactly \ - * representable as a double, so the conversion is silent \ - * unless the range check rejects it. The add suites reject it earlier, \ - * on the alpha type, so there this case only repeats their existing \ - * floating-point alpha coverage. */ \ - expect_bad_scalar_value_dies(9223372036854775808.0); \ +#define GENERATE_SCALAR_OVERFLOW_TESTS(TEST_SUITE_NAME) \ + TEST_F(TEST_SUITE_NAME, ByteTensorTooLargeScalarDies) { \ + /* Cannot be represented by a uint8_t. */ \ + expect_bad_scalar_value_dies(256); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, CharTensorTooSmallScalarDies) { \ + /* Cannot be represented by a int8_t. */ \ + expect_bad_scalar_value_dies(-129); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, ShortTensorTooLargeScalarDies) { \ + /* Cannot be represented by a int16_t. */ \ + expect_bad_scalar_value_dies(32768); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, FloatTensorTooSmallScalarDies) { \ + /* Cannot be represented by a float. */ \ + expect_bad_scalar_value_dies(-3.41e+38); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, FloatTensorTooLargeScalarDies) { \ + /* Cannot be represented by a float. */ \ + expect_bad_scalar_value_dies(3.41e+38); \ + } \ + \ + TEST_F(TEST_SUITE_NAME, LongTensorTooLargeScalarDies) { \ + /* 2^63 is one past the largest int64_t, so converting it is undefined \ + * unless the range check rejects it first. The add suites reject it \ + * earlier, on the alpha type, so there this case only repeats their \ + * existing floating-point alpha coverage. */ \ + expect_bad_scalar_value_dies(9223372036854775808.0); \ }