diff --git a/.ci/docker/ci_commit_pins/pytorch.txt b/.ci/docker/ci_commit_pins/pytorch.txt index 401a0594d98..6c3fe42ddf3 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 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 diff --git a/.ci/docker/common/install_pytorch.sh b/.ci/docker/common/install_pytorch.sh index 0ac5e79cf4a..e51a8886afd 100755 --- a/.ci/docker/common/install_pytorch.sh +++ b/.ci/docker/common/install_pytorch.sh @@ -76,21 +76,31 @@ 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 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}" + # 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" \ - "setuptools>=77.0.0,<82" "cmake>=3.27,<4" 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 + # 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)" + # 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' +torch.linalg.qr(torch.randn(4, 4)) +") + # 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/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/.ci/scripts/utils.sh b/.ci/scripts/utils.sh index 234e162e48e..3d573daa395 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}" @@ -127,18 +127,31 @@ 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 - USE_DISTRIBUTED=1 python setup.py bdist_wheel + # 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)" - - # Invariant: the basename setup.py just produced must match the cache + # 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 # 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 @@ -178,7 +191,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/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 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/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..0b85b09ee51 100644 --- a/backends/aoti/common_shims.h +++ b/backends/aoti/common_shims.h @@ -62,6 +62,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); + // 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..d9e255bf220 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 && tensor->defined(); + 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..0d60f578605 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); +// 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); + AOTI_SHIM_EXPORT int32_t aoti_torch_layout_strided(); // ============================================================ diff --git a/backends/apple/metal/metal_backend.py b/backends/apple/metal/metal_backend.py index 57ca0ddf83e..0b9a852343b 100644 --- a/backends/apple/metal/metal_backend.py +++ b/backends/apple/metal/metal_backend.py @@ -32,16 +32,24 @@ 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, "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, + "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, + "aoti_torch_mps_gated_delta_rule": None, } @classmethod 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/misc/test_portable_int8.py b/backends/cortex_m/test/misc/test_portable_int8.py index 6efeec9e5b0..41f7f254863 100644 --- a/backends/cortex_m/test/misc/test_portable_int8.py +++ b/backends/cortex_m/test/misc/test_portable_int8.py @@ -716,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 206af19a61e..ba5b2da6b78 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, @@ -23,14 +22,13 @@ 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, "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 = { @@ -43,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 67f0937a006..9bc99e4bf2c 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, @@ -28,14 +27,13 @@ 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, "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 @@ -54,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() @@ -63,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/cuda/cuda_backend.py b/backends/cuda/cuda_backend.py index 854ebf8f952..89e1b5ad79a 100644 --- a/backends/cuda/cuda_backend.py +++ b/backends/cuda/cuda_backend.py @@ -669,7 +669,10 @@ def get_supported_fallback_kernels(cls) -> Dict[str, Any]: return {} return { "at::_ops::_weight_int4pack_mm::call": None, + # 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 b/backends/cuda/runtime/aoti_cuda_shims.lib index 8bb03cc1c1e..c0d61c61100 100644 Binary files a/backends/cuda/runtime/aoti_cuda_shims.lib and b/backends/cuda/runtime/aoti_cuda_shims.lib differ 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..af2fb65ca13 --- /dev/null +++ b/backends/cuda/runtime/aoti_cuda_shims.lib.md @@ -0,0 +1,36 @@ +# 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' \ + | 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 \ + && 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. + +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/runtime/shims/memory.cpp b/backends/cuda/runtime/shims/memory.cpp index 8a81916ab6c..976b282a29c 100644 --- a/backends/cuda/runtime/shims/memory.cpp +++ b/backends/cuda/runtime/shims/memory.cpp @@ -195,6 +195,31 @@ AOTITorchError aoti_torch_empty_strided( return Error::Ok; } +AOTITorchError aoti_torch_empty_strided_pinned( + int64_t ndim, + const int64_t* sizes_ptr, + const int64_t* strides_ptr, + int32_t dtype, + 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, + strides_ptr, + dtype, + device_type, + device_index, + ret_new_tensor); +} + AOTITorchError aoti_torch_delete_tensor_object(SlimTensor* tensor) { ET_CHECK_OR_RETURN_ERROR( tensor != nullptr, diff --git a/backends/cuda/runtime/shims/memory.h b/backends/cuda/runtime/shims/memory.h index ca464a9acf5..03e2b3ed18d 100644 --- a/backends/cuda/runtime/shims/memory.h +++ b/backends/cuda/runtime/shims/memory.h @@ -95,6 +95,31 @@ AOTI_SHIM_EXPORT AOTITorchError aoti_torch_empty_strided( int32_t device_index, SlimTensor** ret_new_tensor); +/** + * 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, + const int64_t* sizes_ptr, + const int64_t* strides_ptr, + int32_t dtype, + int32_t device_type, + int32_t device_index, + SlimTensor** ret_new_tensor); + /** * Deletes a tensor object and frees associated resources. * diff --git a/backends/cuda/tests/test_sort_shim.py b/backends/cuda/tests/test_sort_shim.py index fc5f870fc42..da1ee0c3e2c 100644 --- a/backends/cuda/tests/test_sort_shim.py +++ b/backends/cuda/tests/test_sort_shim.py @@ -37,7 +37,9 @@ _CUDA_FALLBACK_KERNELS = frozenset( { "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/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( 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} ) diff --git a/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py b/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py index fca0f1b14c6..751388d9221 100644 --- a/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py +++ b/backends/xnnpack/quantizer/xnnpack_quantizer_utils.py @@ -1105,6 +1105,9 @@ 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 not training, which is how a quantized model is deployed. + torch.ops.aten.dropout.default, + torch.ops.aten.dropout_.default, ] diff --git a/exir/lowered_backend_module.py b/exir/lowered_backend_module.py index a4c2f2cfe79..358bcb25e8c 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,30 @@ 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 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 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: 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..6567ad71564 100644 --- a/kernels/test/ScalarOverflowTestMacros.h +++ b/kernels/test/ScalarOverflowTestMacros.h @@ -11,28 +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); \ +#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); \ } 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/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/shim_et/xplat/executorch/build/runtime_wrapper.bzl b/shim_et/xplat/executorch/build/runtime_wrapper.bzl index e84af76f3d9..445c1aa3b98 100644 --- a/shim_et/xplat/executorch/build/runtime_wrapper.bzl +++ b/shim_et/xplat/executorch/build/runtime_wrapper.bzl @@ -145,30 +145,45 @@ def _has_pytorch_dep(dep_list): return True return False -def _patch_test_compiler_flags(kwargs): +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", + "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: + 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", []) - deps = kwargs.get("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 - "libtorch" in external_deps or - "gtest_aten" in external_deps or - "gmock_aten" in external_deps or - _has_pytorch_dep(deps) or - _has_pytorch_dep(xplat_deps) or - _has_pytorch_dep(fbcode_deps) + "aten_" in name ) - - if not is_aten_test: + if is_aten_test: + kwargs["compiler_flags"] += [ + "-std=c++20", + ] + else: kwargs["compiler_flags"] += [ "-std=c++17", ] @@ -267,9 +282,21 @@ def _patch_kwargs_cxx(kwargs): env.remove_platform_specific_args(kwargs) return _patch_kwargs_common(kwargs) +def _patch_aten_mode_std(kwargs, aten_mode): + """Raises an ATen-mode target to C++20, which PyTorch's headers require. + + A plain compiler flag, which the prelude places after the toolchain's. + """ + 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) @@ -294,8 +321,11 @@ 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) @@ -315,20 +345,25 @@ 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) 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): 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