diff --git a/.bazelrc b/.bazelrc deleted file mode 100644 index fc2995dc838c5..0000000000000 --- a/.bazelrc +++ /dev/null @@ -1,114 +0,0 @@ -build --cxxopt=--std=c++17 -build --copt=-I. -# Bazel does not support including its cc_library targets as system -# headers. We work around this for generated code -# (e.g. torch/headeronly/macros/cmake_macros.h) by making the generated directory a -# system include path. -build --copt=-isystem --copt bazel-out/k8-fastbuild/bin -build --copt=-isystem --copt bazel-out/darwin-fastbuild/bin -build --experimental_ui_max_stdouterr_bytes=2048576 - -# Configuration to disable tty features for environments like CI -build:no-tty --curses no -build:no-tty --progress_report_interval 10 -build:no-tty --show_progress_rate_limit 10 - -# Build with GPU support by default. -build --define=cuda=true -# rules_cuda configuration -build --@rules_cuda//cuda:enable_cuda -build --@rules_cuda//cuda:cuda_targets=sm_52 -build --@rules_cuda//cuda:compiler=nvcc -build --repo_env=CUDA_PATH=/usr/local/cuda - -# Configuration to build without GPU support -build:cpu-only --define=cuda=false -# define a separate build folder for faster switching between configs -build:cpu-only --platform_suffix=-cpu-only -# See the note on the config-less build for details about why we are -# doing this. We must also do it for the "-cpu-only" platform suffix. -build --copt=-isystem --copt=bazel-out/k8-fastbuild-cpu-only/bin -# rules_cuda configuration -build:cpu-only --@rules_cuda//cuda:enable_cuda=False - -# Definition of --config=shell -# interactive shell immediately before execution -build:shell --run_under="//tools/bazel_tools:shellwrap" - -# Disable all warnings for external repositories. We don't care about -# their warnings. -build --per_file_copt=^external/@-w - -# Set additional warnings to error level. -# -# Implementation notes: -# * we use file extensions to determine if we are using the C++ -# compiler or the cuda compiler -# * we use ^// at the start of the regex to only permit matching -# PyTorch files. This excludes external repos. -# -# Note that because this is logically a command-line flag, it is -# considered the word on what warnings are enabled. This has the -# unfortunate consequence of preventing us from disabling an error at -# the target level because those flags will come before these flags in -# the action invocation. Instead we provide per-file exceptions after -# this. -# -# On the bright side, this means we don't have to more broadly apply -# the exceptions to an entire target. -# -# Looking for CUDA flags? We have a cu_library macro that we can edit -# directly. Look in //tools/rules:cu.bzl for details. Editing the -# macro over this has the following advantages: -# * making changes does not require discarding the Bazel analysis -# cache -# * it allows for selective overrides on individual targets since the -# macro-level opts will come earlier than target level overrides - -build --per_file_copt='^//.*\.(cpp|cc)$'@-Werror=all -# The following warnings come from -Wall. We downgrade them from error -# to warnings here. -# -# We intentionally use #pragma unroll, which is compiler specific. -build --per_file_copt='^//.*\.(cpp|cc)$'@-Wno-error=unknown-pragmas - -build --per_file_copt='^//.*\.(cpp|cc)$'@-Werror=extra -# The following warnings come from -Wextra. We downgrade them from error -# to warnings here. -# -# unused-parameter-compare has a tremendous amount of violations in the -# codebase. It will be a lot of work to fix them, just disable it for -# now. -build --per_file_copt='^//.*\.(cpp|cc)$'@-Wno-unused-parameter -# missing-field-parameters has both a large number of violations in -# the codebase, but it also is used pervasively in the Python C -# API. There are a couple of catches though: -# * we use multiple versions of the Python API and hence have -# potentially multiple different versions of each relevant -# struct. They may have different numbers of fields. It will be -# unwieldy to support multiple versions in the same source file. -# * Python itself for many of these structs recommends only -# initializing a subset of the fields. We should respect the API -# usage conventions of our dependencies. -# -# Hence, we just disable this warning altogether. We may want to clean -# up some of the clear-cut cases that could be risky, but we still -# likely want to have this disabled for the most part. -build --per_file_copt='^//.*\.(cpp|cc)$'@-Wno-missing-field-initializers - -build --per_file_copt='^//.*\.(cpp|cc)$'@-Wno-unused-function -build --per_file_copt='^//.*\.(cpp|cc)$'@-Wno-unused-variable - -build --per_file_copt='//:aten/src/ATen/RegisterCompositeExplicitAutograd\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterCompositeImplicitAutograd\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterMkldnnCPU\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterNestedTensorCPU\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterQuantizedCPU\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterSparseCPU\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterSparseCsrCPU\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterNestedTensorMeta\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterSparseMeta\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterQuantizedMeta\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:aten/src/ATen/RegisterZeroTensor\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:torch/csrc/lazy/generated/RegisterAutogradLazy\.cpp$'@-Wno-error=unused-function -build --per_file_copt='//:torch/csrc/lazy/generated/RegisterLazy\.cpp$'@-Wno-error=unused-function diff --git a/.bazelversion b/.bazelversion deleted file mode 100644 index f22d756da39d4..0000000000000 --- a/.bazelversion +++ /dev/null @@ -1 +0,0 @@ -6.5.0 diff --git a/.ci/docker/README.md b/.ci/docker/README.md index 32d3a21a1897b..56d6ff1f2fe27 100644 --- a/.ci/docker/README.md +++ b/.ci/docker/README.md @@ -25,7 +25,6 @@ See `build.sh` for valid build environments (it's the giant switch). * `conda` - Dockerfile and build.sh to build Docker images used in nightly conda builds * `manywheel` - Dockerfile and build.sh to build Docker images used in nightly manywheel builds -* `libtorch` - Dockerfile and build.sh to build Docker images used in nightly libtorch builds ## Usage @@ -108,8 +107,6 @@ If your new Docker image needs a library installed from a specific pinned commit GCC_VERSION=11 VISION=yes KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} TRITON=yes NEW_ARG_1=yes ;; diff --git a/.ci/docker/almalinux/Dockerfile b/.ci/docker/almalinux/Dockerfile index 08a43513a0904..d061353187c55 100644 --- a/.ci/docker/almalinux/Dockerfile +++ b/.ci/docker/almalinux/Dockerfile @@ -19,15 +19,12 @@ RUN git config --global --add safe.directory '*' ENV PATH=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/bin:$PATH # cmake-3.18.4 from pip +# NS: Apr 1 2026 3.18.4 is gone, reported here https://github.com/scikit-build/cmake-python-distributions/issues/693 RUN yum install -y python3-pip && \ - python3 -mpip install cmake==3.18.4 && \ + python3 -mpip install cmake==3.18.4.post1 && \ ln -s /usr/local/bin/cmake /usr/bin/cmake3 RUN rm -rf /usr/local/cuda-* -FROM base as openssl -ADD ./common/install_openssl.sh install_openssl.sh -RUN bash ./install_openssl.sh && rm install_openssl.sh - FROM base as patchelf # Install patchelf ADD ./common/install_patchelf.sh install_patchelf.sh @@ -53,7 +50,6 @@ ENV CUDA_VERSION=${CUDA_VERSION} # Make things in our path by default ENV PATH=/usr/local/cuda-${CUDA_VERSION}/bin:/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/bin:$PATH - FROM cuda as cuda12.6 RUN bash ./install_cuda.sh 12.6 ENV DESIRED_CUDA=12.6 @@ -70,6 +66,10 @@ FROM cuda as cuda13.0 RUN bash ./install_cuda.sh 13.0 ENV DESIRED_CUDA=13.0 +FROM cuda as cuda13.2 +RUN bash ./install_cuda.sh 13.2 +ENV DESIRED_CUDA=13.2 + FROM ${ROCM_IMAGE} as rocm_base ARG DEVTOOLSET_VERSION=13 ENV LC_ALL en_US.UTF-8 @@ -81,6 +81,8 @@ RUN yum -y update && \ yum -y install glibc-langpack-en && \ yum install -y sudo wget curl perl util-linux xz bzip2 git patch which perl zlib-devel openssl-devel yum-utils autoconf automake make gcc-toolset-${DEVTOOLSET_VERSION}-gcc gcc-toolset-${DEVTOOLSET_VERSION}-gcc-c++ gcc-toolset-${DEVTOOLSET_VERSION}-gcc-gfortran gcc-toolset-${DEVTOOLSET_VERSION}-gdb RUN git config --global --add safe.directory '*' +# All rocm clang cfg files load the same rocm.cfg, make sure it points to the right toolchain. +RUN echo "--gcc-toolchain=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr" >> /opt/rocm/llvm/bin/rocm.cfg ENV PATH=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/bin:$PATH FROM rocm_base as rocm @@ -101,11 +103,11 @@ COPY --from=cuda12.6 /usr/local/cuda-12.6 /usr/local/cuda-12.6 COPY --from=cuda12.8 /usr/local/cuda-12.8 /usr/local/cuda-12.8 COPY --from=cuda12.9 /usr/local/cuda-12.9 /usr/local/cuda-12.9 COPY --from=cuda13.0 /usr/local/cuda-13.0 /usr/local/cuda-13.0 +COPY --from=cuda13.2 /usr/local/cuda-13.2 /usr/local/cuda-13.2 # Final step FROM ${BASE_TARGET} as final ARG DEVTOOLSET_VERSION=13 -COPY --from=openssl /opt/openssl /opt/openssl COPY --from=patchelf /patchelf /usr/local/bin/patchelf COPY --from=conda /opt/conda /opt/conda diff --git a/.ci/docker/build.sh b/.ci/docker/build.sh index 37c082e7d378e..1eca06471a110 100755 --- a/.ci/docker/build.sh +++ b/.ci/docker/build.sh @@ -81,13 +81,6 @@ elif [[ "$image" == *riscv* ]]; then DOCKERFILE="ubuntu-cross-riscv/Dockerfile" fi -_UCX_COMMIT=7836b165abdbe468a2f607e7254011c07d788152 -_UCC_COMMIT=430e241bf5d38cbc73fc7a6b89155397232e3f96 -if [[ "$image" == *rocm* ]]; then - _UCX_COMMIT=29831d319e6be55cb8c768ca61de335c934ca39e - _UCC_COMMIT=9f4b242cbbd8b1462cbc732eb29316cdfa124b77 -fi - tag=$(echo $image | awk -F':' '{print $2}') # If no tag (no colon in image name), use the image name itself if [[ -z "$tag" ]]; then @@ -102,20 +95,14 @@ case "$tag" in CUDA_VERSION=12.4 ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=11 - VISION=yes KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} TRITON=yes ;; pytorch-linux-jammy-cuda12.8-cudnn9-py3-gcc11) CUDA_VERSION=12.8.1 ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=11 - VISION=yes KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} TRITON=yes INSTALL_MINGW=yes ;; @@ -123,69 +110,48 @@ case "$tag" in CUDA_VERSION=13.0.2 ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=11 - VISION=yes KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} TRITON=yes - ;; - pytorch-linux-jammy-cuda12.8-cudnn9-py3-gcc11-inductor-benchmarks) - CUDA_VERSION=12.8.1 - ANACONDA_PYTHON_VERSION=3.10 - GCC_VERSION=11 - VISION=yes - KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} - TRITON=yes - INDUCTOR_BENCHMARKS=yes + INSTALL_MINGW=yes ;; pytorch-linux-jammy-cuda13.0-cudnn9-py3-gcc11-inductor-benchmarks) CUDA_VERSION=13.0.2 ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=11 - VISION=yes KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} TRITON=yes INDUCTOR_BENCHMARKS=yes ;; - pytorch-linux-jammy-cuda12.9-cudnn9-py3.12-gcc11-vllm) - CUDA_VERSION=12.9.1 + pytorch-linux-jammy-cuda13.0-cudnn9-py3.12-gcc11-vllm) + CUDA_VERSION=13.0.2 ANACONDA_PYTHON_VERSION=3.12 GCC_VERSION=11 - VISION=yes KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} TRITON=yes ;; - pytorch-linux-jammy-py3-clang15-onnx) + pytorch-linux-jammy-py3.10-clang18) ANACONDA_PYTHON_VERSION=3.10 - CLANG_VERSION=15 - VISION=yes + CLANG_VERSION=18 + GCC_VERSION=11 + KATEX=yes + DOCS=yes ONNX=yes ;; - pytorch-linux-jammy-py3.10-clang15) - ANACONDA_PYTHON_VERSION=3.10 - CLANG_VERSION=15 - ;; - pytorch-linux-jammy-py3.11-clang15) + pytorch-linux-jammy-py3.11-clang18) ANACONDA_PYTHON_VERSION=3.11 - CLANG_VERSION=15 + CLANG_VERSION=18 ;; - pytorch-linux-jammy-py3.12-clang15) + pytorch-linux-jammy-py3.12-clang18) ANACONDA_PYTHON_VERSION=3.12 - CLANG_VERSION=15 + CLANG_VERSION=18 ;; - pytorch-linux-jammy-py3.13-clang15) + pytorch-linux-jammy-py3.13-clang18) ANACONDA_PYTHON_VERSION=3.13 - CLANG_VERSION=15 + CLANG_VERSION=18 ;; - pytorch-linux-jammy-py3.14-clang15) + pytorch-linux-jammy-py3.14-clang18) ANACONDA_PYTHON_VERSION=3.14 - CLANG_VERSION=15 + CLANG_VERSION=18 ;; pytorch-linux-jammy-rocm-n-py3 | pytorch-linux-jammy-rocm-n-py3-benchmarks | pytorch-linux-noble-rocm-n-py3) if [[ $tag =~ "jammy" ]]; then @@ -193,14 +159,10 @@ case "$tag" in else ANACONDA_PYTHON_VERSION=3.12 fi - GCC_VERSION=11 - VISION=yes - ROCM_VERSION=7.1 - NINJA_VERSION=1.9.0 + GCC_VERSION=13 + ROCM_VERSION=7.2 TRITON=yes KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} PYTORCH_ROCM_ARCH="gfx90a;gfx942;gfx950;gfx1100" if [[ $tag =~ "benchmarks" ]]; then INDUCTOR_BENCHMARKS=yes @@ -208,32 +170,28 @@ case "$tag" in ;; pytorch-linux-noble-rocm-nightly-py3) ANACONDA_PYTHON_VERSION=3.12 - GCC_VERSION=11 - VISION=yes + GCC_VERSION=13 ROCM_VERSION=nightly - NINJA_VERSION=1.9.0 TRITON=yes KATEX=yes - UCX_COMMIT=${_UCX_COMMIT} - UCC_COMMIT=${_UCC_COMMIT} PYTORCH_ROCM_ARCH="gfx942" ;; pytorch-linux-jammy-xpu-n-1-py3) ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=11 - VISION=yes XPU_VERSION=2025.2 XPU_DRIVER_TYPE=LTS - NINJA_VERSION=1.9.0 TRITON=yes ;; - pytorch-linux-noble-xpu-n-py3 | pytorch-linux-noble-xpu-n-py3-inductor-benchmarks) + pytorch-linux-noble-xpu-n-py3 | pytorch-linux-noble-xpu-n-py3-client | pytorch-linux-noble-xpu-n-py3-inductor-benchmarks) ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=13 - VISION=yes XPU_VERSION=2025.3 - XPU_DRIVER_TYPE=LTS - NINJA_VERSION=1.9.0 + if [[ $tag =~ "client" ]]; then + XPU_DRIVER_TYPE=CLIENT + else + XPU_DRIVER_TYPE=LTS + fi TRITON=yes if [[ $tag =~ "benchmarks" ]]; then INDUCTOR_BENCHMARKS=yes @@ -242,36 +200,19 @@ case "$tag" in pytorch-linux-jammy-py3-gcc11-inductor-benchmarks) ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=11 - VISION=yes KATEX=yes - TRITON=yes DOCS=yes INDUCTOR_BENCHMARKS=yes ;; - pytorch-linux-jammy-cuda12.8-cudnn9-py3.10-clang15) + pytorch-linux-jammy-cuda12.8-cudnn9-py3.10-clang18) ANACONDA_PYTHON_VERSION=3.10 CUDA_VERSION=12.8.1 - CLANG_VERSION=15 - VISION=yes - TRITON=yes - ;; - pytorch-linux-jammy-py3-clang18-asan) - ANACONDA_PYTHON_VERSION=3.10 CLANG_VERSION=18 - VISION=yes - ;; - pytorch-linux-jammy-py3.10-gcc11) - ANACONDA_PYTHON_VERSION=3.10 - GCC_VERSION=11 - VISION=yes - KATEX=yes TRITON=yes - DOCS=yes - UNINSTALL_DILL=yes ;; - pytorch-linux-jammy-py3-clang15-executorch) + pytorch-linux-jammy-py3-clang18-executorch) ANACONDA_PYTHON_VERSION=3.10 - CLANG_VERSION=15 + CLANG_VERSION=18 EXECUTORCH=yes ;; pytorch-linux-jammy-py3.12-halide) @@ -316,31 +257,13 @@ case "$tag" in ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=13 ACL=yes - VISION=yes OPENBLAS=yes - # snadampal: skipping llvm src build install because the current version - # from pytorch/llvm:9.0.1 is x86 specific - SKIP_LLVM_SRC_BUILD_INSTALL=yes - ;; - pytorch-linux-jammy-aarch64-py3.10-clang21) - ANACONDA_PYTHON_VERSION=3.10 - CLANG_VERSION=21 - ACL=yes - VISION=yes - OPENBLAS=yes - # snadampal: skipping llvm src build install because the current version - # from pytorch/llvm:9.0.1 is x86 specific - SKIP_LLVM_SRC_BUILD_INSTALL=yes ;; pytorch-linux-jammy-aarch64-py3.10-gcc13-inductor-benchmarks) ANACONDA_PYTHON_VERSION=3.10 GCC_VERSION=13 ACL=yes - VISION=yes OPENBLAS=yes - # snadampal: skipping llvm src build install because the current version - # from pytorch/llvm:9.0.1 is x86 specific - SKIP_LLVM_SRC_BUILD_INSTALL=yes INDUCTOR_BENCHMARKS=yes ;; pytorch-linux-noble-riscv64-py3.12-gcc14) @@ -348,7 +271,6 @@ case "$tag" in ;; *) # Catch-all for builds that are not hardcoded. - VISION=yes echo "image '$image' did not match an existing build configuration" if [[ "$image" == *py* ]]; then extract_version_from_image_name py ANACONDA_PYTHON_VERSION @@ -365,14 +287,10 @@ case "$tag" in if [[ -z "$ROCM_VERSION" ]]; then extract_version_from_image_name rocm ROCM_VERSION fi - NINJA_VERSION=1.9.0 TRITON=yes # To ensure that any ROCm config will build using conda cmake # and thus have LAPACK/MKL enabled fi - if [[ "$image" == *centos7* ]]; then - NINJA_VERSION=1.10.2 - fi if [[ "$image" == *gcc* ]]; then extract_version_from_image_name gcc GCC_VERSION fi @@ -404,7 +322,6 @@ docker buildx build \ ${progress_flag} \ --build-arg "BUILD_ENVIRONMENT=${image}" \ --build-arg "LLVMDEV=${LLVMDEV:-}" \ - --build-arg "VISION=${VISION:-}" \ --build-arg "UBUNTU_VERSION=${UBUNTU_VERSION}" \ --build-arg "DEVTOOLSET_VERSION=${DEVTOOLSET_VERSION}" \ --build-arg "GLIBC_VERSION=${GLIBC_VERSION}" \ @@ -414,13 +331,10 @@ docker buildx build \ --build-arg "PYTHON_VERSION=${PYTHON_VERSION}" \ --build-arg "GCC_VERSION=${GCC_VERSION}" \ --build-arg "CUDA_VERSION=${CUDA_VERSION}" \ - --build-arg "NINJA_VERSION=${NINJA_VERSION:-}" \ --build-arg "KATEX=${KATEX:-}" \ --build-arg "ROCM_VERSION=${ROCM_VERSION:-}" \ --build-arg "PYTORCH_ROCM_ARCH=${PYTORCH_ROCM_ARCH}" \ --build-arg "IMAGE_NAME=${IMAGE_NAME}" \ - --build-arg "UCX_COMMIT=${UCX_COMMIT}" \ - --build-arg "UCC_COMMIT=${UCC_COMMIT}" \ --build-arg "TRITON=${TRITON}" \ --build-arg "TRITON_CPU=${TRITON_CPU}" \ --build-arg "ONNX=${ONNX}" \ @@ -432,11 +346,9 @@ docker buildx build \ --build-arg "TPU=${TPU}" \ --build-arg "XPU_VERSION=${XPU_VERSION}" \ --build-arg "XPU_DRIVER_TYPE=${XPU_DRIVER_TYPE}" \ - --build-arg "UNINSTALL_DILL=${UNINSTALL_DILL}" \ --build-arg "ACL=${ACL:-}" \ --build-arg "OPENBLAS=${OPENBLAS:-}" \ --build-arg "SKIP_SCCACHE_INSTALL=${SKIP_SCCACHE_INSTALL:-}" \ - --build-arg "SKIP_LLVM_SRC_BUILD_INSTALL=${SKIP_LLVM_SRC_BUILD_INSTALL:-}" \ --build-arg "INSTALL_MINGW=${INSTALL_MINGW:-}" \ -f $(dirname ${DOCKERFILE})/Dockerfile \ --load \ diff --git a/.ci/docker/ci_commit_pins/huggingface-requirements.txt b/.ci/docker/ci_commit_pins/huggingface-requirements.txt index 408343c9099c8..51a16f10e0632 100644 --- a/.ci/docker/ci_commit_pins/huggingface-requirements.txt +++ b/.ci/docker/ci_commit_pins/huggingface-requirements.txt @@ -1,2 +1,2 @@ -transformers==4.57.5 +transformers==5.5.3 soxr==0.5.0 diff --git a/.ci/docker/ci_commit_pins/nccl-cu126.txt b/.ci/docker/ci_commit_pins/nccl-cu126.txt new file mode 100644 index 0000000000000..1706c910183ce --- /dev/null +++ b/.ci/docker/ci_commit_pins/nccl-cu126.txt @@ -0,0 +1 @@ +v2.29.3-1 diff --git a/.ci/docker/ci_commit_pins/nccl.txt b/.ci/docker/ci_commit_pins/nccl.txt index 7c451d9fad29a..9ad2e5cfc6595 100644 --- a/.ci/docker/ci_commit_pins/nccl.txt +++ b/.ci/docker/ci_commit_pins/nccl.txt @@ -1 +1 @@ -v2.28.9-1 +v2.29.7-1 diff --git a/.ci/docker/ci_commit_pins/torchbench.txt b/.ci/docker/ci_commit_pins/torchbench.txt index fdcdfa96ac643..36a3b62eaec96 100644 --- a/.ci/docker/ci_commit_pins/torchbench.txt +++ b/.ci/docker/ci_commit_pins/torchbench.txt @@ -1 +1 @@ -7be477a510cdbe9543b2f11357598db51c22e94e +0fff8bf35400ec7b733d5080628fc79551205e36 diff --git a/.ci/docker/ci_commit_pins/triton-xpu.txt b/.ci/docker/ci_commit_pins/triton-xpu.txt index 7453190b8a7f1..3bec45a60a706 100644 --- a/.ci/docker/ci_commit_pins/triton-xpu.txt +++ b/.ci/docker/ci_commit_pins/triton-xpu.txt @@ -1 +1 @@ -64bb0de394fe1b9e163758380ce8eed25511bb5e +21033c4e2be9b42c9e6ce7a39a70ead2aba279b4 diff --git a/.ci/docker/ci_commit_pins/triton.txt b/.ci/docker/ci_commit_pins/triton.txt index 23407b4d540c4..cdcf66491de62 100644 --- a/.ci/docker/ci_commit_pins/triton.txt +++ b/.ci/docker/ci_commit_pins/triton.txt @@ -1 +1 @@ -9844da955a9db14ec69c9aac828ee9803085e288 +4cff872ced001ea92d9fcf05b3f6517e2b486d19 diff --git a/.ci/docker/common/install_amdsmi.sh b/.ci/docker/common/install_amdsmi.sh index 8e0ee620da679..759e2bababe25 100644 --- a/.ci/docker/common/install_amdsmi.sh +++ b/.ci/docker/common/install_amdsmi.sh @@ -7,7 +7,7 @@ source /etc/rocm_env.sh # For theRock nightly, amd_smi may already be installed or in a different location if [ -d "${ROCM_PATH}/share/amd_smi" ]; then echo "Installing amdsmi from: ${ROCM_PATH}/share/amd_smi" - cd ${ROCM_PATH}/share/amd_smi && pip install . + cd ${ROCM_PATH}/share/amd_smi && python3 -m pip install . else echo "AMD SMI not found at ${ROCM_PATH}/share/amd_smi - skipping (may already be installed via pip)" fi diff --git a/.ci/docker/common/install_base.sh b/.ci/docker/common/install_base.sh index 7d8ae247d7a0b..27d9ca0601779 100755 --- a/.ci/docker/common/install_base.sh +++ b/.ci/docker/common/install_base.sh @@ -11,36 +11,26 @@ install_ubuntu() { # "$UBUNTU_VERSION" == "18.04" if [[ "$UBUNTU_VERSION" == "20.04"* ]]; then cmake3="cmake=3.16*" - maybe_libiomp_dev="" elif [[ "$UBUNTU_VERSION" == "22.04"* ]]; then cmake3="cmake=3.22*" - maybe_libiomp_dev="" elif [[ "$UBUNTU_VERSION" == "24.04"* ]]; then cmake3="cmake=3.28*" - maybe_libiomp_dev="" else - cmake3="cmake=3.5*" - maybe_libiomp_dev="libiomp-dev" - fi - - if [[ "$CLANG_VERSION" == 15 ]]; then - maybe_libomp_dev="libomp-15-dev" - elif [[ "$CLANG_VERSION" == 12 ]]; then - maybe_libomp_dev="libomp-12-dev" - elif [[ "$CLANG_VERSION" == 10 ]]; then - maybe_libomp_dev="libomp-10-dev" - else - maybe_libomp_dev="" + echo "Unknown Ubuntu version $UBUNTU_VERSION" + exit 1 fi # Install common dependencies apt-get update + # Install prerequisites for add-apt-repository (needs gpg-agent for PPA key import) + apt-get install -y --no-install-recommends software-properties-common gpg-agent + # Add git-core PPA for a newer version of git + add-apt-repository ppa:git-core/ppa -y + apt-get update # TODO: Some of these may not be necessary - ccache_deps="asciidoc docbook-xml docbook-xsl xsltproc" deploy_deps="libffi-dev libbz2-dev libreadline-dev libncurses5-dev libncursesw5-dev libgdbm-dev libsqlite3-dev uuid-dev tk-dev" numpy_deps="gfortran" apt-get install -y --no-install-recommends \ - $ccache_deps \ $numpy_deps \ ${deploy_deps} \ ${cmake3} \ @@ -53,14 +43,13 @@ install_ubuntu() { git \ libatlas-base-dev \ libc6-dbg \ - ${maybe_libiomp_dev} \ libyaml-dev \ libz-dev \ libjemalloc2 \ + libgl1 \ libjpeg-dev \ libasound2-dev \ libsndfile-dev \ - ${maybe_libomp_dev} \ software-properties-common \ wget \ sudo \ @@ -71,7 +60,9 @@ install_ubuntu() { unzip \ gpg-agent \ gdb \ - bc + bc \ + zip \ + valgrind # Should resolve issues related to various apt package repository cert issues # see: https://github.com/pytorch/pytorch/issues/65931 @@ -82,70 +73,14 @@ install_ubuntu() { rm -rf /var/lib/apt/lists/* /tmp/* /var/tmp/* } -install_centos() { - # Need EPEL for many packages we depend on. - # See http://fedoraproject.org/wiki/EPEL - yum --enablerepo=extras install -y epel-release - - ccache_deps="asciidoc docbook-dtds docbook-style-xsl libxslt" - numpy_deps="gcc-gfortran" - yum install -y \ - $ccache_deps \ - $numpy_deps \ - autoconf \ - automake \ - bzip2 \ - cmake \ - cmake3 \ - curl \ - gcc \ - gcc-c++ \ - gflags-devel \ - git \ - glibc-devel \ - glibc-headers \ - glog-devel \ - libstdc++-devel \ - libsndfile-devel \ - make \ - opencv-devel \ - sudo \ - wget \ - vim \ - unzip \ - gdb - - # Cleanup - yum clean all - rm -rf /var/cache/yum - rm -rf /var/lib/yum/yumdb - rm -rf /var/lib/yum/history -} - # Install base packages depending on the base OS ID=$(grep -oP '(?<=^ID=).+' /etc/os-release | tr -d '"') case "$ID" in ubuntu) install_ubuntu ;; - centos) - install_centos - ;; *) echo "Unable to determine OS..." exit 1 ;; esac - -# Install Valgrind separately since the apt-get version is too old. -mkdir valgrind_build && cd valgrind_build -VALGRIND_VERSION=3.20.0 -wget https://ossci-linux.s3.amazonaws.com/valgrind-${VALGRIND_VERSION}.tar.bz2 -tar -xjf valgrind-${VALGRIND_VERSION}.tar.bz2 -cd valgrind-${VALGRIND_VERSION} -./configure --prefix=/usr/local -make -j$[$(nproc) - 2] -sudo make install -cd ../../ -rm -rf valgrind_build -alias valgrind="/usr/local/bin/valgrind" diff --git a/.ci/docker/common/install_cache.sh b/.ci/docker/common/install_cache.sh index 3eab145967854..837ba50098029 100644 --- a/.ci/docker/common/install_cache.sh +++ b/.ci/docker/common/install_cache.sh @@ -18,7 +18,8 @@ install_ubuntu() { cp target/release/sccache-dist /opt/cache/bin echo "Cleaning up" cd .. - rm -rf sccache .cargo + rm -rf sccache + rustup self uninstall -y apt-get remove -y pkg-config libssl-dev apt-get autoclean && apt-get clean @@ -77,10 +78,15 @@ EOF chmod a+x "/opt/cache/bin/$1" } -write_sccache_stub cc -write_sccache_stub c++ -write_sccache_stub gcc -write_sccache_stub g++ +# Skip all sccache wrapping for theRock nightly: sccache PATH wrappers +# intercept assembly (.s) compilation and fail because the assembler does not +# produce the .d dependency file that sccache expects. +if [ "$ROCM_VERSION" != "nightly" ]; then + write_sccache_stub cc + write_sccache_stub c++ + write_sccache_stub gcc + write_sccache_stub g++ +fi # NOTE: See specific ROCM_VERSION case below. if [ "x$ROCM_VERSION" = x ]; then diff --git a/.ci/docker/common/install_clang.sh b/.ci/docker/common/install_clang.sh index 93daeee919b3d..20eea1ee157fe 100755 --- a/.ci/docker/common/install_clang.sh +++ b/.ci/docker/common/install_clang.sh @@ -15,7 +15,7 @@ if [ -n "$CLANG_VERSION" ]; then sudo apt-get update if [[ $CLANG_VERSION -ge 18 ]]; then - apt-get install -y libomp-${CLANG_VERSION}-dev libclang-rt-${CLANG_VERSION}-dev clang-"$CLANG_VERSION" llvm-"$CLANG_VERSION" + apt-get install -y --no-install-recommends libomp-${CLANG_VERSION}-dev libclang-rt-${CLANG_VERSION}-dev clang-"$CLANG_VERSION" llvm-"$CLANG_VERSION" else apt-get install -y --no-install-recommends clang-"$CLANG_VERSION" llvm-"$CLANG_VERSION" fi diff --git a/.ci/docker/common/install_conda.sh b/.ci/docker/common/install_conda.sh index 57c9845d76a0b..547eec0d401f5 100755 --- a/.ci/docker/common/install_conda.sh +++ b/.ci/docker/common/install_conda.sh @@ -67,13 +67,6 @@ if [ -n "$ANACONDA_PYTHON_VERSION" ]; then conda_install sqlite fi - # Install PyTorch conda deps, as per https://github.com/pytorch/pytorch README - if [[ $(uname -m) != "aarch64" ]]; then - pip_install mkl==2024.2.0 - pip_install mkl-static==2024.2.0 - pip_install mkl-include==2024.2.0 - fi - # Install llvm-8 as it is required to compile llvmlite-0.30.0 from source # and libpython-static for torch deploy conda_install llvmdev=8.0.0 "libpython-static=${ANACONDA_PYTHON_VERSION}" @@ -103,5 +96,8 @@ if [ -n "$ANACONDA_PYTHON_VERSION" ]; then pip_install -r /opt/conda/requirements-docs.txt fi + # Clean conda package cache + as_jenkins conda clean -ya + popd fi diff --git a/.ci/docker/common/install_conda_docker.sh b/.ci/docker/common/install_conda_docker.sh index dc377075750ac..9665a799b6ff5 100755 --- a/.ci/docker/common/install_conda_docker.sh +++ b/.ci/docker/common/install_conda_docker.sh @@ -13,8 +13,8 @@ rm $(basename "$MINICONDA_URL") export PATH=/opt/conda/bin:$PATH # See https://github.com/pytorch/builder/issues/1473 # Pin conda to 23.5.2 as it's the last one compatible with openssl-1.1.1 -conda install -y conda=23.5.2 conda-build anaconda-client git ninja +conda install -y conda=23.5.2 conda-build anaconda-client git # The cmake version here needs to match with the minimum version of cmake -# supported by PyTorch (3.18). There is only 3.18.2 on anaconda -/opt/conda/bin/pip3 install cmake==3.18.2 +# supported by PyTorch (3.18). +/opt/conda/bin/pip3 install cmake==3.18.4.post1 ninja conda remove -y --force patchelf diff --git a/.ci/docker/common/install_cuda.sh b/.ci/docker/common/install_cuda.sh index 2d1db795d9cb4..9d7cd7ad78c05 100644 --- a/.ci/docker/common/install_cuda.sh +++ b/.ci/docker/common/install_cuda.sh @@ -82,21 +82,23 @@ function install_nvshmem { function install_124 { CUDNN_VERSION=9.1.0.70 - echo "Installing CUDA 12.4.1 and cuDNN ${CUDNN_VERSION} and NCCL and cuSparseLt-0.6.2" + CUSPARSELT_VERSION=0.6.2.3 + echo "Installing CUDA 12.4.1 and cuDNN ${CUDNN_VERSION} and NCCL and cuSparseLt-${CUSPARSELT_VERSION}" install_cuda 12.4.1 cuda_12.4.1_550.54.15_linux install_cudnn 12 $CUDNN_VERSION CUDA_VERSION=12.4 bash install_nccl.sh - CUDA_VERSION=12.4 bash install_cusparselt.sh + CUDA_VERSION=12.4 bash install_cusparselt.sh $CUSPARSELT_VERSION ldconfig } function install_126 { CUDNN_VERSION=9.10.2.21 - echo "Installing CUDA 12.6.3 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-0.7.1" + CUSPARSELT_VERSION=0.7.1.0 + echo "Installing CUDA 12.6.3 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-${CUSPARSELT_VERSION}" install_cuda 12.6.3 cuda_12.6.3_560.35.05_linux install_cudnn 12 $CUDNN_VERSION @@ -105,14 +107,15 @@ function install_126 { CUDA_VERSION=12.6 bash install_nccl.sh - CUDA_VERSION=12.6 bash install_cusparselt.sh + CUDA_VERSION=12.6 bash install_cusparselt.sh $CUSPARSELT_VERSION ldconfig } function install_129 { - CUDNN_VERSION=9.17.1.4 - echo "Installing CUDA 12.9.1 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-0.7.1" + CUDNN_VERSION=9.20.0.48 + CUSPARSELT_VERSION=0.8.1.1 + echo "Installing CUDA 12.9.1 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-${CUSPARSELT_VERSION}" # install CUDA 12.9.1 in the same container install_cuda 12.9.1 cuda_12.9.1_575.57.08_linux @@ -123,14 +126,15 @@ function install_129 { CUDA_VERSION=12.9 bash install_nccl.sh - CUDA_VERSION=12.9 bash install_cusparselt.sh + CUDA_VERSION=12.9 bash install_cusparselt.sh $CUSPARSELT_VERSION ldconfig } function install_128 { - CUDNN_VERSION=9.17.1.4 - echo "Installing CUDA 12.8.1 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-0.7.1" + CUDNN_VERSION=9.20.0.48 + CUSPARSELT_VERSION=0.7.1.0 + echo "Installing CUDA 12.8.1 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-${CUSPARSELT_VERSION}" # install CUDA 12.8.1 in the same container install_cuda 12.8.1 cuda_12.8.1_570.124.06_linux @@ -141,14 +145,15 @@ function install_128 { CUDA_VERSION=12.8 bash install_nccl.sh - CUDA_VERSION=12.8 bash install_cusparselt.sh + CUDA_VERSION=12.8 bash install_cusparselt.sh $CUSPARSELT_VERSION ldconfig } function install_130 { - CUDNN_VERSION=9.17.1.4 - echo "Installing CUDA 13.0 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-0.7.1" + CUDNN_VERSION=9.20.0.48 + CUSPARSELT_VERSION=0.8.1.1 + echo "Installing CUDA 13.0 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-${CUSPARSELT_VERSION}" # install CUDA 13.0 in the same container install_cuda 13.0.2 cuda_13.0.2_580.95.05_linux @@ -159,7 +164,26 @@ function install_130 { CUDA_VERSION=13.0 bash install_nccl.sh - CUDA_VERSION=13.0 bash install_cusparselt.sh + CUDA_VERSION=13.0 bash install_cusparselt.sh $CUSPARSELT_VERSION + + ldconfig +} + +function install_132 { + CUDNN_VERSION=9.20.0.48 + CUSPARSELT_VERSION=0.8.1.1 + echo "Installing CUDA 13.2 and cuDNN ${CUDNN_VERSION} and NVSHMEM and NCCL and cuSparseLt-${CUSPARSELT_VERSION}" + # install CUDA 13.2 in the same container + install_cuda 13.2.1 cuda_13.2.1_595.58.03_linux + + # cuDNN license: https://developer.nvidia.com/cudnn/license_agreement + install_cudnn 13 $CUDNN_VERSION + + install_nvshmem 13 $NVSHMEM_VERSION + + CUDA_VERSION=13.2 bash install_nccl.sh + + CUDA_VERSION=13.2 bash install_cusparselt.sh $CUSPARSELT_VERSION ldconfig } @@ -178,6 +202,8 @@ do ;; 13.0|13.0.*) install_130; ;; + 13.2|13.2.*) install_132; + ;; *) echo "bad argument $1"; exit 1 ;; esac diff --git a/.ci/docker/common/install_cusparselt.sh b/.ci/docker/common/install_cusparselt.sh index b532c086371f1..0568dd1a18f55 100644 --- a/.ci/docker/common/install_cusparselt.sh +++ b/.ci/docker/common/install_cusparselt.sh @@ -5,34 +5,29 @@ set -ex # cuSPARSELt license: https://docs.nvidia.com/cuda/cusparselt/license.html mkdir tmp_cusparselt && cd tmp_cusparselt -if [[ ${CUDA_VERSION:0:4} =~ "13" ]]; then - arch_path='sbsa' - export TARGETARCH=${TARGETARCH:-$(uname -m)} - if [ ${TARGETARCH} = 'amd64' ] || [ "${TARGETARCH}" = 'x86_64' ]; then - arch_path='x86_64' - fi - CUSPARSELT_NAME="libcusparse_lt-linux-${arch_path}-0.8.0.4_cuda13-archive" - curl --retry 3 -OLs https://developer.download.nvidia.com/compute/cusparselt/redist/libcusparse_lt/linux-${arch_path}/${CUSPARSELT_NAME}.tar.xz -elif [[ ${CUDA_VERSION:0:4} =~ ^12\.[5-9]$ ]]; then - arch_path='sbsa' - export TARGETARCH=${TARGETARCH:-$(uname -m)} - if [ ${TARGETARCH} = 'amd64' ] || [ "${TARGETARCH}" = 'x86_64' ]; then - arch_path='x86_64' - fi - CUSPARSELT_NAME="libcusparse_lt-linux-${arch_path}-0.7.1.0-archive" - curl --retry 3 -OLs https://developer.download.nvidia.com/compute/cusparselt/redist/libcusparse_lt/linux-${arch_path}/${CUSPARSELT_NAME}.tar.xz -elif [[ ${CUDA_VERSION:0:4} == "12.4" ]]; then - arch_path='sbsa' - export TARGETARCH=${TARGETARCH:-$(uname -m)} - if [ ${TARGETARCH} = 'amd64' ] || [ "${TARGETARCH}" = 'x86_64' ]; then - arch_path='x86_64' - fi - CUSPARSELT_NAME="libcusparse_lt-linux-${arch_path}-0.6.2.3-archive" - curl --retry 3 -OLs https://developer.download.nvidia.com/compute/cusparselt/redist/libcusparse_lt/linux-${arch_path}/${CUSPARSELT_NAME}.tar.xz +cusparselt_version=$1 + +arch_path='sbsa' +export TARGETARCH=${TARGETARCH:-$(uname -m)} +if [ ${TARGETARCH} = 'amd64' ] || [ "${TARGETARCH}" = 'x86_64' ]; then + arch_path='x86_64' +fi + +if [[ -z "${cusparselt_version}" ]]; then + echo "Usage: install_cusparselt.sh " + exit 1 +fi + +cuda_major_version=${CUDA_VERSION%%.*} +cusparselt_minor=$(echo "${cusparselt_version}" | cut -d. -f2) +# Starting from 0.8.0, NVIDIA ships separate archives per CUDA major version +if [[ "${cusparselt_minor}" -ge 8 ]]; then + CUSPARSELT_NAME="libcusparse_lt-linux-${arch_path}-${cusparselt_version}_cuda${cuda_major_version}-archive" else - echo "Not sure which libcusparselt version to install for this ${CUDA_VERSION}" + CUSPARSELT_NAME="libcusparse_lt-linux-${arch_path}-${cusparselt_version}-archive" fi +curl --retry 3 -OLs https://developer.download.nvidia.com/compute/cusparselt/redist/libcusparse_lt/linux-${arch_path}/${CUSPARSELT_NAME}.tar.xz tar xf ${CUSPARSELT_NAME}.tar.xz cp -a ${CUSPARSELT_NAME}/include/* /usr/local/cuda/include/ cp -a ${CUSPARSELT_NAME}/lib/* /usr/local/cuda/lib64/ diff --git a/.ci/docker/common/install_devtoolset.sh b/.ci/docker/common/install_devtoolset.sh deleted file mode 100755 index bdae637598138..0000000000000 --- a/.ci/docker/common/install_devtoolset.sh +++ /dev/null @@ -1,10 +0,0 @@ -#!/bin/bash - -set -ex - -[ -n "$DEVTOOLSET_VERSION" ] - -yum install -y centos-release-scl -yum install -y devtoolset-$DEVTOOLSET_VERSION - -echo "source scl_source enable devtoolset-$DEVTOOLSET_VERSION" > "/etc/profile.d/devtoolset-$DEVTOOLSET_VERSION.sh" diff --git a/.ci/docker/common/install_docs_reqs.sh b/.ci/docker/common/install_docs_reqs.sh index c907145f2ec62..c06160373a05e 100644 --- a/.ci/docker/common/install_docs_reqs.sh +++ b/.ci/docker/common/install_docs_reqs.sh @@ -17,7 +17,7 @@ if [ -n "$KATEX" ]; then apt-get install -y --no-install-recommends yarn yarn global add katex --prefix /usr/local - sudo apt-get -y install doxygen + sudo apt-get -y install doxygen lcov apt-get autoclean && apt-get clean rm -rf /var/lib/apt/lists/* /tmp/* /var/tmp/* diff --git a/.ci/docker/common/install_glibc.sh b/.ci/docker/common/install_glibc.sh deleted file mode 100755 index c98791e2bf85b..0000000000000 --- a/.ci/docker/common/install_glibc.sh +++ /dev/null @@ -1,34 +0,0 @@ -#!/bin/bash - -set -ex - -[ -n "$GLIBC_VERSION" ] -if [[ -n "$CENTOS_VERSION" ]]; then - [ -n "$DEVTOOLSET_VERSION" ] -fi - -yum install -y wget sed - -mkdir -p /packages && cd /packages -wget -q http://ftp.gnu.org/gnu/glibc/glibc-$GLIBC_VERSION.tar.gz -tar xzf glibc-$GLIBC_VERSION.tar.gz -if [[ "$GLIBC_VERSION" == "2.26" ]]; then - cd glibc-$GLIBC_VERSION - sed -i 's/$name ne "nss_test1"/$name ne "nss_test1" \&\& $name ne "nss_test2"/' scripts/test-installation.pl - cd .. -fi -mkdir -p glibc-$GLIBC_VERSION-build && cd glibc-$GLIBC_VERSION-build - -if [[ -n "$CENTOS_VERSION" ]]; then - export PATH=/opt/rh/devtoolset-$DEVTOOLSET_VERSION/root/usr/bin:$PATH -fi - -../glibc-$GLIBC_VERSION/configure --prefix=/usr CFLAGS='-Wno-stringop-truncation -Wno-format-overflow -Wno-restrict -Wno-format-truncation -g -O2' -make -j$(nproc) -make install - -# Cleanup -rm -rf /packages -rm -rf /var/cache/yum/* -rm -rf /var/lib/rpm/__db.* -yum clean all diff --git a/.ci/docker/common/install_inductor_benchmark_deps.sh b/.ci/docker/common/install_inductor_benchmark_deps.sh index 674b141efcfb2..c54b8a44f0632 100644 --- a/.ci/docker/common/install_inductor_benchmark_deps.sh +++ b/.ci/docker/common/install_inductor_benchmark_deps.sh @@ -18,18 +18,16 @@ function install_timm() { function install_torchbench() { local commit commit=$(get_pinned_commit torchbench) - git clone https://github.com/pytorch/benchmark torchbench + mkdir torchbench && chown jenkins torchbench + as_jenkins git clone https://github.com/pytorch/benchmark torchbench pushd torchbench - git checkout "$commit" + as_jenkins git checkout "$commit" - python install.py --continue_on_fail + conda_run python install.py --continue_on_fail echo "Print all dependencies after TorchBench is installed" - python -mpip freeze + conda_run python -mpip freeze popd - - chown -R jenkins torchbench - chown -R jenkins /opt/conda } # Pango is needed for weasyprint which is needed for doctr diff --git a/.ci/docker/common/install_mingw.sh b/.ci/docker/common/install_mingw.sh index 6232a0d0245c7..e82a666ff4352 100644 --- a/.ci/docker/common/install_mingw.sh +++ b/.ci/docker/common/install_mingw.sh @@ -4,7 +4,7 @@ set -ex # Install MinGW-w64 for Windows cross-compilation apt-get update -apt-get install -y g++-mingw-w64-x86-64-posix +apt-get install -y g++-mingw-w64-x86-64-posix mingw-w64-tools echo "MinGW-w64 installed successfully" x86_64-w64-mingw32-g++ --version diff --git a/.ci/docker/common/install_miopen.sh b/.ci/docker/common/install_miopen.sh index 3dbc67b90abaf..039458add8406 100644 --- a/.ci/docker/common/install_miopen.sh +++ b/.ci/docker/common/install_miopen.sh @@ -16,7 +16,7 @@ case "$ID" in ubuntu) IS_UBUNTU=1 ;; - centos|almalinux) + almalinux) IS_UBUNTU=0 ;; *) diff --git a/.ci/docker/common/install_nccl.sh b/.ci/docker/common/install_nccl.sh index 486604140a983..f505e4c3f249a 100644 --- a/.ci/docker/common/install_nccl.sh +++ b/.ci/docker/common/install_nccl.sh @@ -18,6 +18,11 @@ NCCL_VERSION=$(cat ci_commit_pins/nccl.txt) # exit 1 # fi +# Use the NCCL version for CUDA 12.6 due to sm50 support +if [[ ${CUDA_VERSION:0:4} == "12.6" ]]; then + NCCL_VERSION=$(cat ci_commit_pins/nccl-cu126.txt) +fi + if [[ -n "${NCCL_VERSION}" ]]; then # NCCL license: https://docs.nvidia.com/deeplearning/nccl/#licenses # Follow build: https://github.com/NVIDIA/nccl/tree/master?tab=readme-ov-file#build diff --git a/.ci/docker/common/install_ninja.sh b/.ci/docker/common/install_ninja.sh deleted file mode 100644 index fa380722bdc2f..0000000000000 --- a/.ci/docker/common/install_ninja.sh +++ /dev/null @@ -1,18 +0,0 @@ -#!/bin/bash - -set -ex - -[ -n "$NINJA_VERSION" ] - -arch=$(uname -m) -if [ "$arch" == "aarch64" ]; then - url="https://github.com/ninja-build/ninja/releases/download/v${NINJA_VERSION}/ninja-linux-aarch64.zip" -else - url="https://github.com/ninja-build/ninja/releases/download/v${NINJA_VERSION}/ninja-linux.zip" -fi - -pushd /tmp -wget --no-verbose --output-document=ninja-linux.zip "$url" -unzip ninja-linux.zip -d /usr/local/bin -rm -f ninja-linux.zip -popd \ No newline at end of file diff --git a/.ci/docker/common/install_onnx.sh b/.ci/docker/common/install_onnx.sh index 36ce5b11d9135..eacc7bbcad157 100755 --- a/.ci/docker/common/install_onnx.sh +++ b/.ci/docker/common/install_onnx.sh @@ -11,15 +11,11 @@ retry () { # ONNXRuntime should be installed before installing # onnx-weekly. Otherwise, onnx-weekly could be # overwritten by onnx. +# Note: parameterized, pytest-subtests, tabulate, packaging are already +# installed via requirements-ci.txt pip_install \ - parameterized==0.8.1 \ - pytest-cov==4.0.0 \ - pytest-subtests==0.10.0 \ - tabulate==0.9.0 \ - transformers==4.36.2 - -pip_install coloredlogs packaging -pip_install onnxruntime==1.23.1 + transformers==4.36.2 \ + onnxruntime==1.23.1 # Cache the transformers model to be used later by ONNX tests. We need to run the transformers # package to download the model. By default, the model is cached at ~/.cache/huggingface/hub/ @@ -34,4 +30,5 @@ conda_run python "${IMPORT_SCRIPT_FILENAME}" # Cleaning up conda_run pip uninstall -y torch +conda_run pip cache purge rm "${IMPORT_SCRIPT_FILENAME}" || true diff --git a/.ci/docker/common/install_openssl.sh b/.ci/docker/common/install_openssl.sh deleted file mode 100644 index c73c9c333c002..0000000000000 --- a/.ci/docker/common/install_openssl.sh +++ /dev/null @@ -1,17 +0,0 @@ -#!/bin/bash - -set -ex - -OPENSSL=openssl-1.1.1k - -wget -q -O "${OPENSSL}.tar.gz" "https://ossci-linux.s3.amazonaws.com/${OPENSSL}.tar.gz" -tar xf "${OPENSSL}.tar.gz" -cd "${OPENSSL}" -./config --prefix=/opt/openssl -d '-Wl,--enable-new-dtags,-rpath,$(LIBRPATH)' -# NOTE: openssl install errors out when built with the -j option -NPROC=$[$(nproc) - 2] -make -j${NPROC}; make install_sw -# Link the ssl libraries to the /usr/lib folder. -sudo ln -s /opt/openssl/lib/lib* /usr/lib -cd .. -rm -rf "${OPENSSL}" diff --git a/.ci/docker/common/install_rocSHMEM.sh b/.ci/docker/common/install_rocSHMEM.sh new file mode 100644 index 0000000000000..ff59951f94875 --- /dev/null +++ b/.ci/docker/common/install_rocSHMEM.sh @@ -0,0 +1,38 @@ +#!/usr/bin/env bash +# Script used only in CD pipeline to build and install rocSHMEM + +set -eou pipefail + +function do_install() { + ROCSHMEM_VERSION=ea5c137103f18a9aadd570d09d72e78ec52f0a3a + rocm_dir="${ROCM_HOME:-}" + if [[ -z "${rocm_dir}" && -f /etc/rocm_env.sh ]]; then + source /etc/rocm_env.sh + rocm_dir="${ROCM_HOME:-}" + fi + rocm_dir="${rocm_dir:-/opt/rocm}" + echo "install_rocSHMEM.sh: using ROCM install prefix ${rocm_dir}" + if [[ -f "${rocm_dir}/lib/librocshmem.a" ]]; then + echo "install_rocSHMEM.sh: librocshmem.a already present in ${rocm_dir}/lib, skipping build" + return + fi + ( + set -x + curr_dir=$(pwd) + tmp_dir=$(mktemp -d) + + git clone --no-checkout --filter=blob:none https://github.com/ROCm/rocm-systems.git ${tmp_dir}/rocm-systems + cd ${tmp_dir}/rocm-systems + git sparse-checkout set --cone projects/rocshmem + git checkout ${ROCSHMEM_VERSION} + + cd ${tmp_dir}/rocm-systems/projects/rocshmem + mkdir build + cd build + INSTALL_PREFIX="${rocm_dir}" ../scripts/build_configs/all_backends + cd ${curr_dir} + + ) +} + +do_install diff --git a/.ci/docker/common/install_rocm.sh b/.ci/docker/common/install_rocm.sh index 21e5968016bd6..b4f6d70af3676 100644 --- a/.ci/docker/common/install_rocm.sh +++ b/.ci/docker/common/install_rocm.sh @@ -26,87 +26,131 @@ install_ubuntu() { apt-get install -y libc++1 apt-get install -y libc++abi1 - # When ROCM_VERSION=nightly, install ROCm from TheRock nightly wheels + # When ROCM_VERSION=nightly, install ROCm from TheRock nightly tarballs + # Mirrors: https://github.com/ROCm/TheRock/blob/main/dockerfiles/install_rocm_tarball.sh if [[ "${ROCM_VERSION}" == "nightly" ]]; then - echo "install_rocm.sh: installing ROCm from TheRock nightly wheels" + apt-get install -y --no-install-recommends pkg-config - # Clean any previous ROCm installation in the base CI image. if [[ -d /opt/rocm ]]; then - echo "Removing existing /opt/rocm from base image" rm -rf /opt/rocm fi - # Determine theRock nightly URL based on GPU architecture - # Check BUILD_ENVIRONMENT or PYTORCH_ROCM_ARCH for the target GPU - if [[ -z "${THEROCK_NIGHTLY_INDEX_URL:-}" ]]; then + # Determine GPU family based on target architecture + AMDGPU_FAMILY="${THEROCK_AMDGPU_FAMILY:-}" + if [[ -z "${AMDGPU_FAMILY}" ]]; then if [[ "${BUILD_ENVIRONMENT}" == *"gfx950"* ]] || [[ "${PYTORCH_ROCM_ARCH}" == *"gfx950"* ]]; then - # MI350 (gfx950) - THEROCK_NIGHTLY_INDEX_URL="https://rocm.nightlies.amd.com/v2/gfx950-dcgpu/" - echo "Detected gfx950 architecture - using MI350 theRock nightly repository" + AMDGPU_FAMILY="gfx950-dcgpu" else - # Default to MI300 (gfx942/gfx94X) - THEROCK_NIGHTLY_INDEX_URL="https://rocm.nightlies.amd.com/v2/gfx94X-dcgpu/" - echo "Using gfx94X (MI300) theRock nightly repository" + AMDGPU_FAMILY="gfx94X-dcgpu" fi fi - export THEROCK_NIGHTLY_INDEX_URL - echo "TheRock Index URL: ${THEROCK_NIGHTLY_INDEX_URL}" + # Auto-detect latest nightly version if not pinned + VERSION="${THEROCK_VERSION:-}" + if [[ -z "${VERSION}" ]]; then + VERSION=$(curl -fsSL "https://rocm.nightlies.amd.com/tarball/" \ + | grep -oP "therock-dist-linux-${AMDGPU_FAMILY}-\K[^\"]+(?=\.tar\.gz)" \ + | grep -v ADHOCBUILD \ + | sort -V \ + | tail -1) + if [[ -z "${VERSION}" ]]; then + echo "Error: Could not find a nightly tarball for ${AMDGPU_FAMILY}" + exit 1 + fi + fi + + # URL-encode '+' as '%2B' in VERSION (required for devreleases) + VERSION_ENCODED="${VERSION//+/%2B}" - python3 -m pip install \ - --index-url "${THEROCK_NIGHTLY_INDEX_URL}" \ - "rocm[libraries,devel]" + TARBALL_URL="https://rocm.nightlies.amd.com/tarball/therock-dist-linux-${AMDGPU_FAMILY}-${VERSION_ENCODED}.tar.gz" - # Use the rocm-sdk CLI helper to populate environment defaults - ROCM_HOME="$(rocm-sdk path --root)" - ROCM_BIN="$(rocm-sdk path --bin)" - ROCM_CMAKE_PREFIX="$(rocm-sdk path --cmake)" + echo "==============================================" + echo "ROCm Tarball Installation" + echo "==============================================" + echo "Version: ${VERSION}" + echo "AMDGPU Family: ${AMDGPU_FAMILY}" + echo "Tarball URL: ${TARBALL_URL}" + echo "==============================================" - echo "ROCM_HOME=${ROCM_HOME}" - echo "ROCM_BIN=${ROCM_BIN}" - echo "ROCM_CMAKE_PREFIX=${ROCM_CMAKE_PREFIX}" + # Download tarball + TARBALL_FILE="/tmp/rocm-tarball.tar.gz" - export ROCM_HOME - export ROCM_PATH="${ROCM_HOME}" - export PATH="${ROCM_BIN}:${PATH}" - export CMAKE_PREFIX_PATH="${ROCM_CMAKE_PREFIX}:${CMAKE_PREFIX_PATH:-}" + echo "Downloading tarball..." + curl -fsSL -o "$TARBALL_FILE" "$TARBALL_URL" || { + echo "Error: Failed to download tarball from $TARBALL_URL" + exit 1 + } + + # Verify download + if [ ! -f "$TARBALL_FILE" ] || [ ! -s "$TARBALL_FILE" ]; then + echo "Error: Downloaded file is empty or does not exist" + exit 1 + fi - # theRock bundles system dependencies like libdrm, liblzma in rocm_sysdeps - ROCM_SYSDEPS="${ROCM_HOME}/lib/rocm_sysdeps" - ROCM_SYSDEPS_INCLUDE="${ROCM_SYSDEPS}/include" - ROCM_SYSDEPS_PKGCONFIG="${ROCM_SYSDEPS}/lib/pkgconfig" + # Install directory is fixed to /opt/rocm-{VERSION} + ROCM_INSTALL_DIR="/opt/rocm-${VERSION}" - # Write environment to file that can be sourced by CI scripts and users + # Extract tarball to versioned directory + echo "Extracting tarball to ${ROCM_INSTALL_DIR}..." + mkdir -p "$ROCM_INSTALL_DIR" + tar -xzf "$TARBALL_FILE" -C "$ROCM_INSTALL_DIR" + + # Clean up downloaded file + rm -f "$TARBALL_FILE" + echo "Tarball extracted and cleaned up" + + # Create symlink /opt/rocm -> /opt/rocm-{VERSION} for compatibility + ln -sfn "$ROCM_INSTALL_DIR" /opt/rocm + echo "Created symlink: /opt/rocm -> $ROCM_INSTALL_DIR" + + # Verify bin and lib folder exists after extraction + echo "Verifying installation..." + for dir in bin clients include lib libexec share; do + if [ ! -d "$ROCM_INSTALL_DIR/$dir" ]; then + echo "Error: ROCm $dir directory not found" + exit 1 + fi + echo "ROCm $dir found in $ROCM_INSTALL_DIR/$dir" + done + + echo "==============================================" + echo "ROCm installed successfully to $ROCM_INSTALL_DIR" + echo "ROCM_PATH=$ROCM_INSTALL_DIR" + echo "PATH should include: $ROCM_INSTALL_DIR/bin" + echo "==============================================" + + # Write environment file (sourced by CI scripts and interactive shells) cat > /etc/rocm_env.sh << ROCM_ENV # ROCm paths -export ROCM_PATH="${ROCM_HOME}" -export ROCM_HOME="${ROCM_HOME}" -export ROCM_SOURCE_DIR="${ROCM_HOME}" -export ROCM_BIN="${ROCM_BIN}" -export ROCM_CMAKE="${ROCM_CMAKE_PREFIX}" -export PATH="${ROCM_BIN}:\${PATH}" -export CMAKE_PREFIX_PATH="${ROCM_CMAKE_PREFIX}:\${CMAKE_PREFIX_PATH:-}" -# Device library paths -export HIP_DEVICE_LIB_PATH="${ROCM_HOME}/lib/llvm/amdgcn/bitcode" -export ROCM_DEVICE_LIB_PATH="${ROCM_HOME}/lib/llvm/amdgcn/bitcode" -# theRock system dependencies -export ROCM_SYSDEPS_INCLUDE="${ROCM_SYSDEPS_INCLUDE}" -export CPLUS_INCLUDE_PATH="${ROCM_SYSDEPS_INCLUDE}:\${CPLUS_INCLUDE_PATH:-}" -export C_INCLUDE_PATH="${ROCM_SYSDEPS_INCLUDE}:\${C_INCLUDE_PATH:-}" -export PKG_CONFIG_PATH="${ROCM_SYSDEPS_PKGCONFIG}:\${PKG_CONFIG_PATH:-}" -export LD_LIBRARY_PATH="${ROCM_SYSDEPS}/lib:\${LD_LIBRARY_PATH:-}" -export LIBRARY_PATH="${ROCM_SYSDEPS}/lib:\${LIBRARY_PATH:-}" -export MAGMA_HOME="${ROCM_HOME}/magma" +export ROCM_PATH=/opt/rocm +export ROCM_HOME=/opt/rocm +export ROCM_SOURCE_DIR=/opt/rocm +export ROCM_BIN=/opt/rocm/bin +export ROCM_CMAKE=/opt/rocm +export PATH=/opt/rocm/bin:/opt/rocm/llvm/bin:\${PATH} +export LD_LIBRARY_PATH=/opt/rocm/lib:\${LD_LIBRARY_PATH:-} +# Sysdeps include paths (libdrm headers, etc.) +export CPLUS_INCLUDE_PATH=/opt/rocm/lib/rocm_sysdeps/include:\${CPLUS_INCLUDE_PATH:-} +export C_INCLUDE_PATH=/opt/rocm/lib/rocm_sysdeps/include:\${C_INCLUDE_PATH:-} +# Device library path +export HIP_DEVICE_LIB_PATH=/opt/rocm/amdgcn/bitcode +export MAGMA_HOME=/opt/rocm/magma +# Tarball bundles sysdeps (libdrm, liblzma, etc.); expose their libs and .pc files +if [ -d /opt/rocm/lib/rocm_sysdeps/lib ]; then + export LD_LIBRARY_PATH=/opt/rocm/lib/rocm_sysdeps/lib:\${LD_LIBRARY_PATH} + export PKG_CONFIG_PATH=/opt/rocm/lib/rocm_sysdeps/lib/pkgconfig:\${PKG_CONFIG_PATH:-} +fi # Disable MSLK for theRock nightly (not yet supported) export USE_MSLK=0 ROCM_ENV - # Append to bash.bashrc so interactive shells get the env vars echo "source /etc/rocm_env.sh" >> /etc/bash.bashrc - echo "install_rocm.sh: TheRock nightly ROCm install complete" - exit 0 - fi + # --- End of theRock nightly tarball installation --- + else + # ========================================================================= + # Non-nightly: install ROCm from repo.radeon.com apt packages + # ========================================================================= # Make sure rocm packages from repo.radeon.com have highest priority cat << EOF > /etc/apt/preferences.d/rocm-pin-600 @@ -120,6 +164,11 @@ EOF ROCM_VERSION="${ROCM_VERSION}.2" fi + # we want the patch version of 7.2 instead + if [[ $(ver $ROCM_VERSION) -eq $(ver 7.2) ]]; then + ROCM_VERSION="${ROCM_VERSION}.1" + fi + # Default url values rocm_baseurl="http://repo.radeon.com/rocm/apt/${ROCM_VERSION}" UBUNTU_VERSION_NAME=`cat /etc/os-release | grep UBUNTU_CODENAME | awk -F= '{print $2}'` @@ -194,7 +243,7 @@ EOF pip_install "git+https://github.com/rocm/composable_kernel@$ROCM_COMPOSABLE_KERNEL_VERSION" - # Write environment to file that can be sourced by CI scripts and users + # Write environment file (sourced by CI scripts and interactive shells) cat > /etc/rocm_env.sh << ROCM_ENV # ROCm paths export ROCM_PATH=/opt/rocm @@ -203,85 +252,18 @@ export ROCM_SOURCE_DIR=/opt/rocm export ROCM_BIN=/opt/rocm/bin export ROCM_CMAKE=/opt/rocm export PATH=/opt/rocm/bin:/opt/rocm/llvm/bin:\${PATH} -# Device library paths -export ROCM_DEVICE_LIB_PATH=/opt/rocm/amdgcn/bitcode +export LD_LIBRARY_PATH=/opt/rocm/lib:\${LD_LIBRARY_PATH:-} +# Device library path export HIP_DEVICE_LIB_PATH=/opt/rocm/amdgcn/bitcode export MAGMA_HOME=/opt/rocm/magma ROCM_ENV - # Append to bash.bashrc so interactive shells get the env vars echo "source /etc/rocm_env.sh" >> /etc/bash.bashrc # Cleanup apt-get autoclean && apt-get clean rm -rf /var/lib/apt/lists/* /tmp/* /var/tmp/* -} - -install_centos() { - - yum update -y - yum install -y kmod - yum install -y wget - yum install -y openblas-devel - - yum install -y epel-release - yum install -y dkms kernel-headers-`uname -r` kernel-devel-`uname -r` - - # Add amdgpu repository - local amdgpu_baseurl - if [[ $OS_VERSION == 9 ]]; then - amdgpu_baseurl="https://repo.radeon.com/amdgpu/${ROCM_VERSION}/rhel/9.0/main/x86_64" - else - amdgpu_baseurl="https://repo.radeon.com/amdgpu/${ROCM_VERSION}/rhel/7.9/main/x86_64" - fi - echo "[AMDGPU]" > /etc/yum.repos.d/amdgpu.repo - echo "name=AMDGPU" >> /etc/yum.repos.d/amdgpu.repo - echo "baseurl=${amdgpu_baseurl}" >> /etc/yum.repos.d/amdgpu.repo - echo "enabled=1" >> /etc/yum.repos.d/amdgpu.repo - echo "gpgcheck=1" >> /etc/yum.repos.d/amdgpu.repo - echo "gpgkey=http://repo.radeon.com/rocm/rocm.gpg.key" >> /etc/yum.repos.d/amdgpu.repo - - local rocm_baseurl="http://repo.radeon.com/rocm/yum/${ROCM_VERSION}" - echo "[ROCm]" > /etc/yum.repos.d/rocm.repo - echo "name=ROCm" >> /etc/yum.repos.d/rocm.repo - echo "baseurl=${rocm_baseurl}" >> /etc/yum.repos.d/rocm.repo - echo "enabled=1" >> /etc/yum.repos.d/rocm.repo - echo "gpgcheck=1" >> /etc/yum.repos.d/rocm.repo - echo "gpgkey=http://repo.radeon.com/rocm/rocm.gpg.key" >> /etc/yum.repos.d/rocm.repo - - yum update -y - - yum install -y \ - rocm-dev \ - rocm-utils \ - rocm-libs \ - rccl \ - rocprofiler-dev \ - roctracer-dev \ - amd-smi-lib - - # precompiled miopen kernels; search for all unversioned packages - # if search fails it will abort this script; use true to avoid case where search fails - MIOPENHIPGFX=$(yum -q search miopen-hip-gfx | grep miopen-hip-gfx | awk '{print $1}'| grep -F kdb. || true) - if [[ "x${MIOPENHIPGFX}" = x ]]; then - echo "miopen-hip-gfx package not available" && exit 1 - else - yum install -y ${MIOPENHIPGFX} - fi - - # ROCm 6.0 had a regression where journal_mode was enabled on the kdb files resulting in permission errors at runtime - for kdb in /opt/rocm/share/miopen/db/*.kdb - do - sqlite3 $kdb "PRAGMA journal_mode=off; PRAGMA VACUUM;" - done - - pip_install "git+https://github.com/rocm/composable_kernel@$ROCM_COMPOSABLE_KERNEL_VERSION" - - # Cleanup - yum clean all - rm -rf /var/cache/yum - rm -rf /var/lib/yum/yumdb - rm -rf /var/lib/yum/history + fi } # Install Python packages depending on the base OS @@ -290,9 +272,6 @@ case "$ID" in ubuntu) install_ubuntu ;; - centos) - install_centos - ;; *) echo "Unable to determine OS..." exit 1 diff --git a/.ci/docker/common/install_rocm_drm.sh b/.ci/docker/common/install_rocm_drm.sh index c70f5880f2c5c..a6b0fe2c03924 100644 --- a/.ci/docker/common/install_rocm_drm.sh +++ b/.ci/docker/common/install_rocm_drm.sh @@ -14,7 +14,7 @@ case "$ID" in apt-get install -y libpciaccess-dev pkg-config apt-get clean ;; - centos|almalinux) + almalinux) yum install -y libpciaccess-devel pkgconfig ;; *) diff --git a/.ci/docker/common/install_torch_tpu.sh b/.ci/docker/common/install_torch_tpu.sh index cfddd4badc741..c4e4104edfe42 100644 --- a/.ci/docker/common/install_torch_tpu.sh +++ b/.ci/docker/common/install_torch_tpu.sh @@ -60,7 +60,7 @@ fetch_secret() { set +x fi - if ! gcloud secrets versions access latest --secret="torchtpu-readonly-key" --project="ml-velocity-actions-testing" > "temp_ssh_key"; then + if ! gcloud secrets versions access latest --secret="torchtpu-read-key" --project="ml-velocity-actions-testing" > "temp_ssh_key"; then echo "Error: Failed to fetch secret. Ensure you are authenticated with gcloud." # Restore xtrace if it was enabled, before exiting @@ -82,7 +82,7 @@ clone_repo() { # Use GIT_SSH_COMMAND to specify the key and disable strict host key checking for automation export GIT_SSH_COMMAND="ssh -i temp_ssh_key -o IdentitiesOnly=yes -o StrictHostKeyChecking=no" - if git clone --recursive "git@github.com:google-ml-infra/torch_tpu.git"; then + if git clone --recursive "git@github.com:google-pytorch/torch_tpu.git"; then echo "Repository cloned successfully." else echo "Error: Failed to clone repository." @@ -110,7 +110,7 @@ pull_torch_tpu() { # sleep 28800 # Debug sleep to connect to runner to streamline debugging, do not submit # 3. Configuration -TORCH_TPU_REPO="${TORCH_TPU_REPO:-https://github.com/google-ml-infra/torch_tpu.git}" +TORCH_TPU_REPO="${TORCH_TPU_REPO:-https://github.com/google-pytorch/torch_tpu.git}" TORCH_TPU_BRANCH="${TORCH_TPU_BRANCH:-main}" # Pin File Configuration @@ -124,7 +124,11 @@ fi if ! command -v bazel &> /dev/null; then echo "Bazel not found. Installing Bazelisk..." temp_dir=$(mktemp -d) + # Download Bazelisk v1.27.0 curl -L https://github.com/bazelbuild/bazelisk/releases/download/v1.27.0/bazelisk-linux-amd64 -o "${temp_dir}/bazel" + # Verify Checksum (SHA256 for v1.27.0 linux-amd64) + # Source: https://github.com/bazelbuild/bazelisk/releases/tag/v1.27.0 + echo "e1508323f347ad1465a887bc5d2bfb91cffc232d11e8e997b623227c6b32fb76 ${temp_dir}/bazel" | sha256sum --check sudo mv "${temp_dir}/bazel" /usr/local/bin/bazel sudo chmod +x /usr/local/bin/bazel rm -rf "${temp_dir}" diff --git a/.ci/docker/common/install_triton.sh b/.ci/docker/common/install_triton.sh index 1b68e3c247839..b2fdebdcc4747 100755 --- a/.ci/docker/common/install_triton.sh +++ b/.ci/docker/common/install_triton.sh @@ -21,7 +21,7 @@ elif [ -n "${TRITON_CPU}" ]; then TRITON_REPO="https://github.com/triton-lang/triton-cpu" TRITON_TEXT_FILE="triton-cpu" else - TRITON_REPO="https://github.com/triton-lang/triton" + TRITON_REPO="https://github.com/ROCm/triton" TRITON_TEXT_FILE="triton" fi diff --git a/.ci/docker/common/install_ucc.sh b/.ci/docker/common/install_ucc.sh deleted file mode 100755 index 6c97edb0e0287..0000000000000 --- a/.ci/docker/common/install_ucc.sh +++ /dev/null @@ -1,94 +0,0 @@ -#!/bin/bash - -set -ex - -if [[ -d "/usr/local/cuda/" ]]; then - with_cuda=/usr/local/cuda/ -else - with_cuda=no -fi - -if [[ -f /etc/rocm_env.sh ]]; then - source /etc/rocm_env.sh -fi - -if [[ -d "${ROCM_PATH}" ]]; then - with_rocm="${ROCM_PATH}" -else - with_rocm=no -fi - -function install_ucx() { - set -ex - git clone --recursive https://github.com/openucx/ucx.git - pushd ucx - git checkout ${UCX_COMMIT} - git submodule update --init --recursive - - ./autogen.sh - ./configure --prefix=$UCX_HOME \ - --enable-mt \ - --with-cuda=$with_cuda \ - --with-rocm=$with_rocm \ - --enable-profiling \ - --enable-stats - time make -j - sudo make install - - popd - rm -rf ucx -} - -function install_ucc() { - set -ex - git clone --recursive https://github.com/openucx/ucc.git - pushd ucc - git checkout ${UCC_COMMIT} - git submodule update --init --recursive - - ./autogen.sh - - if [[ -n "$CUDA_VERSION" && $CUDA_VERSION == 13* ]]; then - NVCC_GENCODE="-gencode=arch=compute_86,code=compute_86" - else - # We only run distributed tests on Tesla M60 and A10G - NVCC_GENCODE="-gencode=arch=compute_52,code=sm_52 -gencode=arch=compute_86,code=compute_86" - fi - - if [[ -n "$ROCM_VERSION" ]]; then - if [[ -n "$PYTORCH_ROCM_ARCH" ]]; then - amdgpu_targets=`echo $PYTORCH_ROCM_ARCH | sed 's/;/ /g'` - else - amdgpu_targets=`rocm_agent_enumerator | grep -v gfx000 | sort -u | xargs` - fi - for arch in $amdgpu_targets; do - HIP_OFFLOAD="$HIP_OFFLOAD --offload-arch=$arch" - done - HIP_OFFLOAD="$HIP_OFFLOAD --rocm-path=${ROCM_PATH}" - - # Set device library path if detected (handles TheRock vs traditional ROCm) - if [ -n "${ROCM_DEVICE_LIB_PATH}" ] && [ -d "${ROCM_DEVICE_LIB_PATH}" ]; then - HIP_OFFLOAD="$HIP_OFFLOAD --rocm-device-lib-path=${ROCM_DEVICE_LIB_PATH}" - fi - else - HIP_OFFLOAD="all-arch-no-native" - fi - - ./configure --prefix=$UCC_HOME \ - --with-ucx=$UCX_HOME \ - --with-cuda=$with_cuda \ - --with-nvcc-gencode="${NVCC_GENCODE}" \ - --with-rocm=$with_rocm \ - --with-rocm-arch="${HIP_OFFLOAD}" - # First observed by ROCm nightly builds, ucc rccl sources fail compile with - # error: #warning "NCCL C++ API is disabled because C compiler is being used. [-Werror=cpp] - # Work-around by adding make CFLAGS=-Wno-error=cpp - time make -j CFLAGS=-Wno-error=cpp - sudo make install - - popd - rm -rf ucc -} - -install_ucx -install_ucc diff --git a/.ci/docker/common/install_vision.sh b/.ci/docker/common/install_vision.sh deleted file mode 100755 index 78c445568ddcd..0000000000000 --- a/.ci/docker/common/install_vision.sh +++ /dev/null @@ -1,46 +0,0 @@ -#!/bin/bash - -set -ex - -install_ubuntu() { - apt-get update - apt-get install -y --no-install-recommends \ - libopencv-dev - - # Cleanup - apt-get autoclean && apt-get clean - rm -rf /var/lib/apt/lists/* /tmp/* /var/tmp/* -} - -install_centos() { - # Need EPEL for many packages we depend on. - # See http://fedoraproject.org/wiki/EPEL - yum --enablerepo=extras install -y epel-release - - yum install -y \ - opencv-devel - - # Cleanup - yum clean all - rm -rf /var/cache/yum - rm -rf /var/lib/yum/yumdb - rm -rf /var/lib/yum/history -} - -# Install base packages depending on the base OS -ID=$(grep -oP '(?<=^ID=).+' /etc/os-release | tr -d '"') -case "$ID" in - ubuntu) - install_ubuntu - ;; - centos) - install_centos - ;; - *) - echo "Unable to determine OS..." - exit 1 - ;; -esac - -# Cache vision models used by the test -source "$(dirname "${BASH_SOURCE[0]}")/cache_vision_models.sh" diff --git a/.ci/docker/libtorch/Dockerfile b/.ci/docker/libtorch/Dockerfile deleted file mode 100644 index 9c4be5abe459d..0000000000000 --- a/.ci/docker/libtorch/Dockerfile +++ /dev/null @@ -1,117 +0,0 @@ -ARG BASE_TARGET=base -ARG GPU_IMAGE=ubuntu:20.04 -FROM ${GPU_IMAGE} as base - -ENV DEBIAN_FRONTEND=noninteractive - -RUN apt-get clean && apt-get update -RUN apt-get install -y curl locales g++ git-all autoconf automake make cmake wget unzip sudo -# Just add everything as a safe.directory for git since these will be used in multiple places with git -RUN git config --global --add safe.directory '*' - -RUN locale-gen en_US.UTF-8 - -ENV LC_ALL en_US.UTF-8 -ENV LANG en_US.UTF-8 -ENV LANGUAGE en_US.UTF-8 - -# Install openssl -FROM base as openssl -ADD ./common/install_openssl.sh install_openssl.sh -RUN bash ./install_openssl.sh && rm install_openssl.sh - -# Install python -FROM base as python -ADD common/install_cpython.sh install_cpython.sh -RUN apt-get update -y && \ - apt-get install build-essential gdb lcov libbz2-dev libffi-dev \ - libgdbm-dev liblzma-dev libncurses5-dev libreadline6-dev \ - libsqlite3-dev libssl-dev lzma lzma-dev tk-dev uuid-dev zlib1g-dev -y && \ - bash ./install_cpython.sh && \ - rm install_cpython.sh && \ - apt-get clean - -FROM base as conda -ADD ./common/install_conda_docker.sh install_conda.sh -RUN bash ./install_conda.sh && rm install_conda.sh - -FROM base as cpu -# Install Anaconda -COPY --from=conda /opt/conda /opt/conda -# Install python -COPY --from=python /opt/python /opt/python -COPY --from=python /opt/_internal /opt/_internal -ENV PATH=/opt/conda/bin:/usr/local/cuda/bin:$PATH -# Install MKL -ADD ./common/install_mkl.sh install_mkl.sh -RUN bash ./install_mkl.sh && rm install_mkl.sh - -FROM cpu as cuda -ADD ./common/install_cuda.sh install_cuda.sh -ADD ./common/install_magma.sh install_magma.sh -COPY ./common/install_nccl.sh install_nccl.sh -COPY ./ci_commit_pins/nccl* /ci_commit_pins/ -COPY ./common/install_cusparselt.sh install_cusparselt.sh -ENV CUDA_HOME /usr/local/cuda - -FROM cuda as cuda12.6 -RUN bash ./install_cuda.sh 12.6 -RUN bash ./install_magma.sh 12.6 -RUN ln -sf /usr/local/cuda-12.6 /usr/local/cuda - -FROM cuda as cuda12.8 -RUN bash ./install_cuda.sh 12.8 -RUN bash ./install_magma.sh 12.8 -RUN ln -sf /usr/local/cuda-12.8 /usr/local/cuda - -FROM cuda as cuda12.9 -RUN bash ./install_cuda.sh 12.9 -RUN bash ./install_magma.sh 12.9 -RUN ln -sf /usr/local/cuda-12.9 /usr/local/cuda - -FROM cuda as cuda13.0 -RUN bash ./install_cuda.sh 13.0 -RUN bash ./install_magma.sh 13.0 -RUN ln -sf /usr/local/cuda-13.0 /usr/local/cuda - -# Install libibverbs for libtorch and copy to CUDA directory -RUN apt-get update -y && \ - apt-get install -y libibverbs-dev librdmacm-dev && \ - cp /usr/lib/x86_64-linux-gnu/libmlx5.so* /usr/local/cuda/lib64/ && \ - cp /usr/lib/x86_64-linux-gnu/librdmacm.so* /usr/local/cuda/lib64/ && \ - cp /usr/lib/x86_64-linux-gnu/libibverbs.so* /usr/local/cuda/lib64/ && \ - cp /usr/lib/x86_64-linux-gnu/libnl* /usr/local/cuda/lib64/ - -FROM cpu as rocm -ARG ROCM_VERSION -ARG PYTORCH_ROCM_ARCH -ENV PYTORCH_ROCM_ARCH ${PYTORCH_ROCM_ARCH} -ENV MKLROOT /opt/intel -# Adding ROCM_PATH env var so that LoadHip.cmake (even with logic updated for ROCm6.0) -# find HIP works for ROCm5.7. Not needed for ROCm6.0 and above. -# Remove below when ROCm5.7 is not in support matrix anymore. -ENV ROCM_PATH /opt/rocm -# No need to install ROCm as base docker image should have full ROCm install -#ADD ./common/install_rocm.sh install_rocm.sh -ADD ./common/install_rocm_drm.sh install_rocm_drm.sh -ADD ./common/install_rocm_magma.sh install_rocm_magma.sh -# gfortran and python needed for building magma from source for ROCm -RUN apt-get update -y && \ - apt-get install gfortran -y && \ - apt-get install python3 python-is-python3 -y && \ - apt-get clean - -RUN bash ./install_rocm_drm.sh /opt/amdgpu && rm install_rocm_drm.sh -RUN bash ./install_rocm_magma.sh ${ROCM_VERSION} && rm install_rocm_magma.sh - -FROM ${BASE_TARGET} as final -COPY --from=openssl /opt/openssl /opt/openssl -# Install patchelf -ADD ./common/install_patchelf.sh install_patchelf.sh -RUN bash ./install_patchelf.sh && rm install_patchelf.sh -# Install Anaconda -COPY --from=conda /opt/conda /opt/conda -# Install python -COPY --from=python /opt/python /opt/python -COPY --from=python /opt/_internal /opt/_internal -ENV PATH=/opt/conda/bin:/usr/local/cuda/bin:$PATH diff --git a/.ci/docker/libtorch/build.sh b/.ci/docker/libtorch/build.sh deleted file mode 100755 index 5bfe70f34347e..0000000000000 --- a/.ci/docker/libtorch/build.sh +++ /dev/null @@ -1,75 +0,0 @@ -#!/usr/bin/env bash -# Script used only in CD pipeline - -set -eoux pipefail - -image="$1" -shift - -if [ -z "${image}" ]; then - echo "Usage: $0 IMAGENAME:ARCHTAG" - exit 1 -fi - -TOPDIR=$(git rev-parse --show-toplevel) - -DOCKER=${DOCKER:-docker} - -# Go from imagename:tag to tag -DOCKER_TAG_PREFIX=$(echo "${image}" | awk -F':' '{print $2}') - -GPU_ARCH_VERSION="" -if [[ "${DOCKER_TAG_PREFIX}" == cuda* ]]; then - # extract cuda version from image name. e.g. manylinux2_28-builder:cuda12.8 returns 12.8 - GPU_ARCH_VERSION=$(echo "${DOCKER_TAG_PREFIX}" | awk -F'cuda' '{print $2}') -elif [[ "${DOCKER_TAG_PREFIX}" == rocm* ]]; then - # extract rocm version from image name. e.g. manylinux2_28-builder:rocm6.2.4 returns 6.2.4 - GPU_ARCH_VERSION=$(echo "${DOCKER_TAG_PREFIX}" | awk -F'rocm' '{print $2}') -fi - -case ${DOCKER_TAG_PREFIX} in - cpu) - BASE_TARGET=cpu - GPU_IMAGE=ubuntu:20.04 - DOCKER_GPU_BUILD_ARG="" - ;; - cuda*) - BASE_TARGET=cuda${GPU_ARCH_VERSION} - GPU_IMAGE=ubuntu:20.04 - DOCKER_GPU_BUILD_ARG="" - ;; - rocm*) - # we want the patch version of 7.1 instead - if [[ "$GPU_ARCH_VERSION" == *"7.1"* ]]; then - GPU_ARCH_VERSION="${GPU_ARCH_VERSION}.1" - fi - # we want the patch version of 7.0 instead - if [[ "$GPU_ARCH_VERSION" == *"7.0"* ]]; then - GPU_ARCH_VERSION="${GPU_ARCH_VERSION}.2" - fi - # we want the patch version of 6.4 instead - if [[ "$GPU_ARCH_VERSION" == *"6.4"* ]]; then - GPU_ARCH_VERSION="${GPU_ARCH_VERSION}.4" - fi - BASE_TARGET=rocm - GPU_IMAGE=rocm/dev-ubuntu-22.04:${GPU_ARCH_VERSION}-complete - PYTORCH_ROCM_ARCH="gfx900;gfx906;gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1200;gfx1201;gfx950;gfx1150;gfx1151" - DOCKER_GPU_BUILD_ARG="--build-arg PYTORCH_ROCM_ARCH=${PYTORCH_ROCM_ARCH} --build-arg ROCM_VERSION=${GPU_ARCH_VERSION}" - ;; - *) - echo "ERROR: Unrecognized DOCKER_TAG_PREFIX: ${DOCKER_TAG_PREFIX}" - exit 1 - ;; -esac - -tmp_tag=$(basename "$(mktemp -u)" | tr '[:upper:]' '[:lower:]') - -DOCKER_BUILDKIT=1 ${DOCKER} build \ - --target final \ - ${DOCKER_GPU_BUILD_ARG} \ - --build-arg "GPU_IMAGE=${GPU_IMAGE}" \ - --build-arg "BASE_TARGET=${BASE_TARGET}" \ - -t "${tmp_tag}" \ - $@ \ - -f "${TOPDIR}/.ci/docker/libtorch/Dockerfile" \ - "${TOPDIR}/.ci/docker/" diff --git a/.ci/docker/manywheel/Dockerfile_2_28 b/.ci/docker/manywheel/Dockerfile_2_28 index 4055e6b872539..2ad5cd6498249 100644 --- a/.ci/docker/manywheel/Dockerfile_2_28 +++ b/.ci/docker/manywheel/Dockerfile_2_28 @@ -12,20 +12,12 @@ RUN yum install -y sudo wget curl perl util-linux xz bzip2 git patch which perl ENV PATH=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/bin:$PATH ENV LD_LIBRARY_PATH=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/lib64:/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/lib:$LD_LIBRARY_PATH -# cmake-3.18.4 from pip +# cmake-3.18.4.post1 from pip +# NS: Apr 1 2026 3.18.4 is gone, reported here https://github.com/scikit-build/cmake-python-distributions/issues/693 RUN yum install -y python3-pip && \ - python3 -mpip install cmake==3.18.4 && \ + python3 -mpip install cmake==3.18.4.post1 && \ ln -s /usr/local/bin/cmake /usr/bin/cmake3 -FROM base as openssl -# Install openssl (this must precede `build python` step) -# (In order to have a proper SSL module, Python is compiled -# against a recent openssl [see env vars above], which is linked -# statically. We delete openssl afterwards.) -ADD ./common/install_openssl.sh install_openssl.sh -RUN bash ./install_openssl.sh && rm install_openssl.sh - - FROM base as cuda ARG BASE_CUDA_VERSION=12.6 # Install CUDA @@ -95,7 +87,6 @@ RUN git config --global --add safe.directory "*" ENV SSL_CERT_FILE=/opt/_internal/certs.pem # Install LLVM version -COPY --from=openssl /opt/openssl /opt/openssl COPY --from=base /opt/python /opt/python COPY --from=base /usr/local/lib/ /usr/local/lib/ COPY --from=base /opt/_internal /opt/_internal @@ -127,9 +118,9 @@ RUN for cpython_version in "cp312-cp312" "cp313-cp313" "cp313-cp313t"; do \ ADD ./common/patch_libstdc.sh patch_libstdc.sh RUN bash ./patch_libstdc.sh && rm patch_libstdc.sh -# cmake-3.18.4 from pip; force in case cmake3 already exists +# cmake-3.18.4.post1 from pip; force in case cmake3 already exists RUN yum install -y python3-pip && \ - python3 -mpip install cmake==3.18.4 && \ + python3 -mpip install cmake==3.18.4.post1 && \ ln -sf /usr/local/bin/cmake /usr/bin/cmake3 FROM cpu_final as cuda_final @@ -144,7 +135,8 @@ ARG ROCM_VERSION=6.0 ARG PYTORCH_ROCM_ARCH ENV PYTORCH_ROCM_ARCH ${PYTORCH_ROCM_ARCH} ARG DEVTOOLSET_VERSION=13 -ENV LDFLAGS="-Wl,-rpath=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/lib64 -Wl,-rpath=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/lib" +# All rocm clang cfg files load the same rocm.cfg, make sure it points to the right toolchain. +RUN echo "--gcc-toolchain=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr" >> /opt/rocm/llvm/bin/rocm.cfg # Somewhere in ROCm stack, we still use non-existing /opt/rocm/hip path, # below workaround helps avoid error ENV ROCM_PATH /opt/rocm @@ -160,6 +152,10 @@ RUN yum install -y libdrm-devel ENV MKLROOT /opt/intel ADD ./common/install_rocm_magma.sh install_rocm_magma.sh RUN bash ./install_rocm_magma.sh ${ROCM_VERSION} && rm install_rocm_magma.sh + +ADD ./common/install_rocSHMEM.sh install_rocSHMEM.sh +RUN bash ./install_rocSHMEM.sh ${ROCM_VERSION} && rm install_rocSHMEM.sh + ADD ./common/install_miopen.sh install_miopen.sh RUN bash ./install_miopen.sh ${ROCM_VERSION} && rm install_miopen.sh diff --git a/.ci/docker/manywheel/Dockerfile_2_28_aarch64 b/.ci/docker/manywheel/Dockerfile_2_28_aarch64 index b5bf2ffc1c081..477e4221cb49f 100644 --- a/.ci/docker/manywheel/Dockerfile_2_28_aarch64 +++ b/.ci/docker/manywheel/Dockerfile_2_28_aarch64 @@ -39,12 +39,7 @@ RUN yum install -y \ gcc-toolset-${GCCTOOLSET_VERSION}-gcc-c++ \ gcc-toolset-${GCCTOOLSET_VERSION}-gcc-gfortran \ gcc-toolset-${GCCTOOLSET_VERSION}-gdb - -# (optional) Install non-default Ninja version -ARG NINJA_VERSION -COPY ./common/install_ninja.sh install_ninja.sh -RUN if [ -n "${NINJA_VERSION}" ]; then bash ./install_ninja.sh; fi -RUN rm install_ninja.sh +RUN yum install -y --enablerepo=powertools ninja-build # Ensure the expected devtoolset is used ENV PATH=/opt/rh/gcc-toolset-${GCCTOOLSET_VERSION}/root/usr/bin:$PATH diff --git a/.ci/docker/manywheel/Dockerfile_cuda_aarch64 b/.ci/docker/manywheel/Dockerfile_cuda_aarch64 index 794a791b2721a..2a3f266fc413f 100644 --- a/.ci/docker/manywheel/Dockerfile_cuda_aarch64 +++ b/.ci/docker/manywheel/Dockerfile_cuda_aarch64 @@ -50,16 +50,7 @@ ENV LD_LIBRARY_PATH=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/lib64:/op RUN git config --global --add safe.directory "*" -FROM base as openssl -# Install openssl (this must precede `build python` step) -# (In order to have a proper SSL module, Python is compiled -# against a recent openssl [see env vars above], which is linked -# statically. We delete openssl afterwards.) -ADD ./common/install_openssl.sh install_openssl.sh -RUN bash ./install_openssl.sh && rm install_openssl.sh -ENV SSL_CERT_FILE=/opt/_internal/certs.pem - -FROM openssl as final +FROM base as final FROM base as cuda ARG BASE_CUDA_VERSION diff --git a/.ci/docker/manywheel/Dockerfile_s390x b/.ci/docker/manywheel/Dockerfile_s390x index 1cf83acb1c736..1367b004ee8a3 100644 --- a/.ci/docker/manywheel/Dockerfile_s390x +++ b/.ci/docker/manywheel/Dockerfile_s390x @@ -84,7 +84,7 @@ RUN cp $(which patchelf) /patchelf FROM patchelf as python # build python -COPY manywheel/build_scripts /build_scripts +COPY manywheel/s390_scripts /build_scripts ADD ./common/install_cpython.sh /build_scripts/install_cpython.sh ENV SSL_CERT_FILE= RUN bash build_scripts/build.sh && rm -r build_scripts diff --git a/.ci/docker/manywheel/build.sh b/.ci/docker/manywheel/build.sh index b0047f98290b3..05e468d72dda9 100755 --- a/.ci/docker/manywheel/build.sh +++ b/.ci/docker/manywheel/build.sh @@ -40,7 +40,7 @@ case ${image} in manylinux2_28_aarch64-builder:cpu-aarch64) TARGET=final GPU_IMAGE=arm64v8/almalinux:8 - DOCKER_GPU_BUILD_ARG=" --build-arg DEVTOOLSET_VERSION=13 --build-arg NINJA_VERSION=1.12.1" + DOCKER_GPU_BUILD_ARG=" --build-arg DEVTOOLSET_VERSION=13" MANY_LINUX_VERSION="2_28_aarch64" ;; manylinuxs390x-builder:cpu-s390x) @@ -75,6 +75,10 @@ case ${image} in DOCKERFILE_SUFFIX="_cuda_aarch64" ;; manylinux2_28-builder:rocm*) + # we want the patch version of 7.2 instead + if [[ "$GPU_ARCH_VERSION" == *"7.2"* ]]; then + GPU_ARCH_VERSION="${GPU_ARCH_VERSION}.1" + fi # we want the patch version of 7.1 instead if [[ "$GPU_ARCH_VERSION" == *"7.1"* ]]; then GPU_ARCH_VERSION="${GPU_ARCH_VERSION}.1" @@ -89,7 +93,7 @@ case ${image} in fi TARGET=rocm_final MANY_LINUX_VERSION="2_28" - DEVTOOLSET_VERSION="11" + DEVTOOLSET_VERSION="13" GPU_IMAGE=rocm/dev-almalinux-8:${GPU_ARCH_VERSION}-complete PYTORCH_ROCM_ARCH="gfx900;gfx906;gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1200;gfx1201;gfx950;gfx1150;gfx1151" DOCKER_GPU_BUILD_ARG="--build-arg ROCM_VERSION=${GPU_ARCH_VERSION} --build-arg PYTORCH_ROCM_ARCH=${PYTORCH_ROCM_ARCH} --build-arg DEVTOOLSET_VERSION=${DEVTOOLSET_VERSION}" diff --git a/.ci/docker/manywheel/build_scripts/manylinux1-check.py b/.ci/docker/manywheel/build_scripts/manylinux1-check.py deleted file mode 100644 index f6b9b9fc2393e..0000000000000 --- a/.ci/docker/manywheel/build_scripts/manylinux1-check.py +++ /dev/null @@ -1,63 +0,0 @@ -# Logic copied from PEP 513 - - -def is_manylinux1_compatible(): - # Only Linux, and only x86-64 / i686 - from distutils.util import get_platform - - if get_platform() not in ["linux-x86_64", "linux-i686", "linux-s390x"]: - return False - - # Check for presence of _manylinux module - try: - import _manylinux - - return bool(_manylinux.manylinux1_compatible) - except (ImportError, AttributeError): - # Fall through to heuristic check below - pass - - # Check glibc version. CentOS 5 uses glibc 2.5. - return have_compatible_glibc(2, 5) - - -def have_compatible_glibc(major, minimum_minor): - import ctypes - - process_namespace = ctypes.CDLL(None) - try: - gnu_get_libc_version = process_namespace.gnu_get_libc_version - except AttributeError: - # Symbol doesn't exist -> therefore, we are not linked to - # glibc. - return False - - # Call gnu_get_libc_version, which returns a string like "2.5". - gnu_get_libc_version.restype = ctypes.c_char_p - version_str = gnu_get_libc_version() - # py2 / py3 compatibility: - if not isinstance(version_str, str): - version_str = version_str.decode("ascii") - - # Parse string and check against requested version. - version = [int(piece) for piece in version_str.split(".")] - if len(version) != 2: - raise AssertionError( - f"Expected version to have 2 components (major.minor), got {len(version)}: {version_str}" - ) - if major != version[0]: - return False - if minimum_minor > version[1]: - return False - return True - - -import sys - - -if is_manylinux1_compatible(): - print(f"{sys.executable} is manylinux1 compatible") - sys.exit(0) -else: - print(f"{sys.executable} is NOT manylinux1 compatible") - sys.exit(1) diff --git a/.ci/docker/manywheel/build_scripts/ssl-check.py b/.ci/docker/manywheel/build_scripts/ssl-check.py deleted file mode 100644 index c4df0eacbb7fd..0000000000000 --- a/.ci/docker/manywheel/build_scripts/ssl-check.py +++ /dev/null @@ -1,26 +0,0 @@ -# cf. https://github.com/pypa/manylinux/issues/53 - -import sys -from urllib.request import urlopen - - -GOOD_SSL = "https://google.com" -BAD_SSL = "https://self-signed.badssl.com" - - -print("Testing SSL certificate checking for Python:", sys.version) - -EXC = OSError - -print(f"Connecting to {GOOD_SSL} should work") -urlopen(GOOD_SSL) -print("...it did, yay.") - -print(f"Connecting to {BAD_SSL} should fail") -try: - urlopen(BAD_SSL) - # If we get here then we failed: - print("...it DIDN'T!!!!!11!!1one!") - sys.exit(1) -except EXC: - print("...it did, yay.") diff --git a/.ci/docker/manywheel/build_scripts/build.sh b/.ci/docker/manywheel/s390_scripts/build.sh similarity index 90% rename from .ci/docker/manywheel/build_scripts/build.sh rename to .ci/docker/manywheel/s390_scripts/build.sh index b6a70f0a72787..13141dfd4ae33 100644 --- a/.ci/docker/manywheel/build_scripts/build.sh +++ b/.ci/docker/manywheel/s390_scripts/build.sh @@ -18,13 +18,7 @@ AUTOCONF_HASH=954bd69b391edc12d6a4a51a2dd1476543da5c6bbf05a95b59dc0dd6fd4c2969 # Dependencies for compiling Python that we want to remove from # the final image after compiling Python -PYTHON_COMPILE_DEPS="zlib-devel bzip2-devel ncurses-devel sqlite-devel readline-devel tk-devel gdbm-devel libpcap-devel xz-devel libffi-devel" - -if [ "$(uname -m)" != "s390x" ] ; then - PYTHON_COMPILE_DEPS="${PYTHON_COMPILE_DEPS} db4-devel" -else - PYTHON_COMPILE_DEPS="${PYTHON_COMPILE_DEPS} libdb-devel" -fi +PYTHON_COMPILE_DEPS="zlib-devel bzip2-devel ncurses-devel sqlite-devel readline-devel tk-devel gdbm-devel libpcap-devel xz-devel libffi-devel libdb-devel" # Libraries that are allowed as part of the manylinux1 profile MANYLINUX1_DEPS="glibc-devel libstdc++-devel glib2-devel libX11-devel libXext-devel libXrender-devel mesa-libGL-devel libICE-devel libSM-devel ncurses-devel" @@ -103,13 +97,6 @@ find /opt/_internal \ -o \( -type f -a -name '*.pyc' -o -name '*.pyo' \) \ -print0 | xargs -0 rm -f -for PYTHON in /opt/python/*/bin/python; do - # Smoke test to make sure that our Pythons work, and do indeed detect as - # being manylinux compatible: - $PYTHON $MY_DIR/manylinux1-check.py - # Make sure that SSL cert checking works - $PYTHON $MY_DIR/ssl-check.py -done # Fix libc headers to remain compatible with C99 compilers. find /usr/include/ -type f -exec sed -i 's/\bextern _*inline_*\b/extern __inline __attribute__ ((__gnu_inline__))/g' {} + diff --git a/.ci/docker/manywheel/build_scripts/build_utils.sh b/.ci/docker/manywheel/s390_scripts/build_utils.sh similarity index 100% rename from .ci/docker/manywheel/build_scripts/build_utils.sh rename to .ci/docker/manywheel/s390_scripts/build_utils.sh diff --git a/.ci/docker/requirements-ci.txt b/.ci/docker/requirements-ci.txt index 9a033b90fcb46..2ed9f431e7ae6 100644 --- a/.ci/docker/requirements-ci.txt +++ b/.ci/docker/requirements-ci.txt @@ -15,13 +15,13 @@ build==1.3.0 #Pinned versions: 1.3.0 #test that import: -click +click==8.3.1 #Description: Command Line Interface Creation Kit -#Pinned versions: +#Pinned versions: 8.3.1 #test that import: coremltools==5.0b5 ; python_version < "3.12" -coremltools==8.3 ; python_version == "3.12" +coremltools==8.3.0 ; python_version == "3.12" #Description: Apple framework for ML integration #Pinned versions: 5.0b5 #test that import: @@ -42,7 +42,7 @@ expecttest==0.3.0 #Pinned versions: 0.3.0 #test that import: -fbscribelogger==0.1.7 +fbscribelogger==0.1.7 ; python_version < "3.14" #Description: write to scribe from authenticated jobs on CI #Pinned versions: 0.1.6 #test that import: @@ -68,23 +68,23 @@ lark==0.12.0 #Pinned versions: 0.12.0 #test that import: -librosa>=0.6.2 ; python_version < "3.11" and platform_machine != "s390x" +librosa==0.10.2 ; python_version < "3.11" and platform_machine != "s390x" librosa==0.10.2 ; python_version == "3.12" and platform_machine != "s390x" #Description: A python package for music and audio analysis #Pinned versions: >=0.6.2 #test that import: test_spectral_ops.py #librosa depends on numba; disable it for s390x while numba is disabled too -#mkl #this breaks linux-bionic-rocm4.5-py3.7 +# Only mkl-static and mkl-include are needed; the mkl package contains +# dynamic libraries that are not discoverable by our build scripts. +mkl-static==2024.2.0 ; platform_machine != "aarch64" and sys_platform != "darwin" +mkl-include==2024.2.0 ; platform_machine != "aarch64" and sys_platform != "darwin" #Description: Intel oneAPI Math Kernel Library -#Pinned versions: +#Pinned versions: 2024.2.0 #test that import: test_profiler.py, test_public_bindings.py, test_testing.py, #test_nn.py, test_mkldnn.py, test_jit.py, test_fx_experimental.py, #test_autograd.py -#mkl-devel -# see mkl - #mock #Description: A testing library that allows you to replace parts of your #system under test with mock objects @@ -105,7 +105,8 @@ mypy==1.16.0 ; platform_system == "Linux" #Pinned versions: 1.16.0 #test that import: test_typing.py, test_type_hints.py -networkx==2.8.8 +networkx==2.8.8 ; python_version < "3.13" +networkx==3.0.0 ; python_version >= "3.13" #Description: creation, manipulation, and study of #the structure, dynamics, and functions of complex networks #Pinned versions: 2.8.8 @@ -117,8 +118,9 @@ ninja==1.11.1.4 #Pinned versions: 1.11.1.4 #test that import: run_test.py, test_cpp_extensions_aot.py,test_determination.py -numba==0.57.1 ; python_version == "3.10" and platform_machine != "s390x" -numba==0.60.0 ; python_version == "3.12" and platform_machine != "s390x" +numba==0.61.2 ; python_version < "3.14" and platform_machine != "s390x" +numba==0.64.0 ; python_version >= "3.14" and platform_machine != "s390x" + #Description: Just-In-Time Compiler for Numerical Functions #Pinned versions: 0.55.2, 0.60.0 #test that import: test_numba_integration.py @@ -136,13 +138,10 @@ numba==0.60.0 ; python_version == "3.12" and platform_machine != "s390x" #test_nn.py, test_namedtensor.py, test_linalg.py, test_jit_cuda_fuser.py, #test_jit.py, test_indexing.py, test_datapipe.py, test_dataloader.py, #test_binary_ufuncs.py -numpy==1.23.2; python_version == "3.10" -numpy==1.26.2; python_version == "3.11" or python_version == "3.12" -numpy==2.1.2; python_version >= "3.13" and python_version < "3.14" +numpy==2.1.2 ; python_version < "3.14" numpy==2.3.4; python_version >= "3.14" -pandas==2.0.3; python_version < "3.12" -pandas==2.2.3; python_version >= "3.12" and python_version < "3.14" +pandas==2.2.3; python_version < "3.14" pandas==2.3.3; python_version >= "3.14" #onnxruntime @@ -150,9 +149,9 @@ pandas==2.3.3; python_version >= "3.14" #Pinned versions: 1.9.0 #test that import: -opt-einsum==3.3 +opt-einsum==3.3.0 #Description: Python library to optimize tensor contraction order, used in einsum -#Pinned versions: 3.3 +#Pinned versions: 3.3.0 #test that import: test_linalg.py optree==0.13.0 ; python_version < "3.14" @@ -169,7 +168,7 @@ optree==0.17.0 ; python_version >= "3.14" #test_pointwise_ops.py, test_dtensor_ops.py, test_torchinductor.py, test_fx.py, #test_fake_tensor.py, test_mps.py -pillow==11.0.0 +pillow==12.2.0 #Description: Python Imaging Library fork #Pinned versions: 11.0.0 #test that import: @@ -179,9 +178,9 @@ protobuf==6.33.5 #Pinned versions: 6.33.2 #test that import: test_tensorboard.py, test/onnx/* -psutil +psutil==7.2.2 #Description: information on running processes and system utilization -#Pinned versions: +#Pinned versions: 7.2.2 #test that import: test_profiler.py, test_openmp.py, test_dataloader.py pytest==7.3.2 @@ -199,9 +198,9 @@ pytest-flakefinder==1.1.0 #Pinned versions: 1.1.0 #test that import: -pytest-rerunfailures>=10.3 +pytest-rerunfailures==14.0 #Description: plugin for rerunning failure tests in pytest -#Pinned versions: +#Pinned versions: 14.0 #test that import: pytest-subtests==0.13.1 @@ -224,7 +223,7 @@ xdoctest==1.3.0 #Pinned versions: 1.1.0 #test that import: -pygments==2.15.0 +pygments==2.20.0 #Description: support doctest highlighting #Pinned versions: 2.12.0 #test that import: the doctests @@ -244,7 +243,8 @@ pygments==2.15.0 #Pinned versions: 14.1.0 #test that import: -scikit-image==0.22.0 +scikit-image==0.22.0 ; python_version < "3.13" +scikit-image==0.26.0 ; python_version >= "3.13" #Description: image processing routines #Pinned versions: 0.22.0 #test that import: test_nn.py @@ -254,9 +254,9 @@ scikit-image==0.22.0 #Pinned versions: 0.20.3 #test that import: -scipy==1.10.1 ; python_version <= "3.11" -scipy==1.14.1 ; python_version > "3.11" and python_version < "3.14" +scipy==1.14.1 ; python_version < "3.14" scipy==1.16.2 ; python_version >= "3.14" + # Pin SciPy because of failing distribution tests (see #60347) #Description: scientific python #Pinned versions: 1.10.1 @@ -270,8 +270,7 @@ scipy==1.16.2 ; python_version >= "3.14" #test that import: # needed by torchgen utils -typing-extensions==4.12.2 ; python_version < "3.14" -typing-extensions==4.15.0 ; python_version >= "3.14" +typing-extensions==4.15.0 #Description: type hints for python #Pinned versions: #test that import: @@ -281,18 +280,22 @@ typing-extensions==4.15.0 ; python_version >= "3.14" #Pinned versions: #test that import: -unittest-xml-reporting<=3.2.0,>=2.0.0 +unittest-xml-reporting==3.2.0 #Description: saves unit test results to xml #Pinned versions: #test that import: -#lintrunner is supported on aarch64-linux only from 0.12.4 version -lintrunner==0.12.11 +lintrunner==0.13.0 #Description: all about linters! -#Pinned versions: 0.12.11 +#Pinned versions: 0.13.0 +#test that import: + +spin==0.17 +#Description: developer CLI for common build/lint tasks +#Pinned versions: 0.17 #test that import: -redis>=4.0.0 +redis==7.4.0 #Description: redis database #test that import: anything that tests OSS caching/mocking (inductor/test_codecache.py, inductor/test_max_autotune.py) @@ -316,14 +319,14 @@ z3-solver==4.15.1.0 ; platform_machine != "s390x" #Pinned versions: #test that import: -tensorboard==2.13.0 ; python_version < "3.13" -tensorboard==2.18.0 ; python_version >= "3.13" +tensorboard==2.18.0 #Description: Also included in .ci/docker/requirements-docs.txt #Pinned versions: #test that import: test_tensorboard pywavelets==1.4.1 ; python_version < "3.12" -pywavelets==1.7.0 ; python_version >= "3.12" +pywavelets==1.7.0 ; python_version >= "3.12" and python_version < "3.14" +pywavelets==1.9.0 ; python_version >= "3.14" #Description: This is a requirement of scikit-image, we need to pin # it here because 1.5.0 conflicts with numpy 1.21.2 used in CI #Pinned versions: 1.4.1 @@ -340,7 +343,7 @@ sympy==1.13.3 #Pinned versions: #test that import: -onnx==1.20.0 +onnx==1.21.0 #Description: Required by the torch.onnx exporter #Pinned versions: #test that import: @@ -370,11 +373,13 @@ pwlf==2.2.1 #test that import: test_sac_estimator.py # To build PyTorch itself +pip==26.0.1 pyyaml==6.0.3 -pyzstd -setuptools==78.1.1 -packaging==24.0 -six +pyzstd==0.16.2 ; python_version < "3.14" +pyzstd==0.18.0 ; python_version >= "3.14" +setuptools==79.0.1 +packaging==25.0 +six==1.17.0 scons==4.5.2 ; platform_machine == "aarch64" @@ -397,7 +402,7 @@ tlparse==0.4.0 filelock==3.20.3 #Description: required for inductor testing -cuda-bindings>=12.0,<13.0 ; platform_machine != "s390x" and platform_system != "Darwin" +cuda-bindings==12.9.6 ; platform_machine != "s390x" and platform_system != "Darwin" #Description: required for testing CUDAGraph::raw_cuda_graph(). See https://nvidia.github.io/cuda-python/cuda-bindings/latest/support.html for how this version was chosen. Note "Any fix in the latest bindings would be backported to the prior major version" means that only the newest version of cuda-bindings will get fixes. Depending on the latest version of 12.x is okay because all 12.y versions will be supported via "CUDA minor version compatibility". Pytorch builds against 13.z versions of cuda toolkit work with 12.x versions of cuda-bindings as well because newer drivers work with old toolkits. #test that import: test_cuda.py @@ -407,9 +412,9 @@ pyre-extensions==0.0.32 tabulate==0.9.0 #Description: These package are needed to build FBGEMM and torchrec on PyTorch CI -tqdm>=4.66.0 +tqdm==4.67.3 #Description: progress bar library required for dynamo benchmarks #test that import: benchmarks/dynamo/* -aiohttp==3.13.3 +aiohttp==3.13.4 #Description: required for torch.distributed.debug diff --git a/.ci/docker/requirements-docs.txt b/.ci/docker/requirements-docs.txt index 7f3e0b5cc9215..484d99ec1152e 100644 --- a/.ci/docker/requirements-docs.txt +++ b/.ci/docker/requirements-docs.txt @@ -2,17 +2,14 @@ sphinx==7.2.6 #Description: This is used to generate PyTorch docs #Pinned versions: 7.2.6 -pytorch_sphinx_theme2==0.4.3 +pytorch_sphinx_theme2==0.4.9 #Description: This is needed to generate PyTorch docs -#Pinned versions: 0.4.3 +#Pinned versions: 0.4.9 -# TODO: sphinxcontrib.katex 0.9.0 adds a local KaTeX server to speed up pre-rendering -# but it doesn't seem to work and hangs around idly. The initial thought that it is probably -# something related to Docker setup. We can investigate this later. - -sphinxcontrib.katex==0.8.6 +sphinxcontrib.katex==0.9.11 #Description: This is used to generate PyTorch docs -#Pinned versions: 0.8.6 +#Pinned versions: 0.9.11 (0.9.0+ uses a persistent KaTeX server instead of +# spawning a subprocess per math expression, ~20% faster writes) sphinxext-opengraph==0.9.1 #Description: This is used to generate PyTorch docs @@ -48,6 +45,10 @@ docutils==0.20 #Description: This is used to generate PyTorch C++ docs #Pinned versions: 0.20 +coverxygen==1.8.1 +#Description: This is used to measure C++ API doc coverage from Doxygen XML +#Pinned versions: 1.8.1 + bs4==0.0.1 #Description: This is used to generate PyTorch C++ docs #Pinned versions: 0.0.1 diff --git a/.ci/docker/triton_version.txt b/.ci/docker/triton_version.txt index 40c341bdcdbe8..19811903a7f75 100644 --- a/.ci/docker/triton_version.txt +++ b/.ci/docker/triton_version.txt @@ -1 +1 @@ -3.6.0 +3.8.0 diff --git a/.ci/docker/triton_xpu_version.txt b/.ci/docker/triton_xpu_version.txt index 7c69a55dbb185..a76ccff2a6e0d 100644 --- a/.ci/docker/triton_xpu_version.txt +++ b/.ci/docker/triton_xpu_version.txt @@ -1 +1 @@ -3.7.0 +3.7.1 diff --git a/.ci/docker/ubuntu-cross-riscv/Dockerfile b/.ci/docker/ubuntu-cross-riscv/Dockerfile index 08201dc83216c..8cf540f8414af 100644 --- a/.ci/docker/ubuntu-cross-riscv/Dockerfile +++ b/.ci/docker/ubuntu-cross-riscv/Dockerfile @@ -38,7 +38,7 @@ COPY ./common/install_user.sh install_user.sh RUN bash ./install_user.sh && rm install_user.sh FROM base as python -ARG ZLIB_VERSION=1.3.1 +ARG ZLIB_VERSION=1.3.2 ARG FFI_VERSION=3.4.6 ARG BZ2_VERSION=1.0.8 ARG XZ_VERSION=5.4.6 diff --git a/.ci/docker/ubuntu-rocm/Dockerfile b/.ci/docker/ubuntu-rocm/Dockerfile index 3f487eb12809e..a3d697f1b27f2 100644 --- a/.ci/docker/ubuntu-rocm/Dockerfile +++ b/.ci/docker/ubuntu-rocm/Dockerfile @@ -13,6 +13,7 @@ ENV PYTORCH_ROCM_ARCH ${PYTORCH_ROCM_ARCH} # Install common dependencies (so that this step can be cached separately) COPY ./common/install_base.sh install_base.sh RUN bash ./install_base.sh && rm install_base.sh +RUN apt-get update && apt-get install -y --no-install-recommends libtbb-dev && rm -rf /var/lib/apt/lists/* # Install user COPY ./common/install_user.sh install_user.sh @@ -43,13 +44,6 @@ ARG CLANG_VERSION COPY ./common/install_clang.sh install_clang.sh RUN bash ./install_clang.sh && rm install_clang.sh -# (optional) Install vision packages like OpenCV -ARG VISION -COPY ./common/install_vision.sh ./common/cache_vision_models.sh ./common/common_utils.sh ./ -RUN if [ -n "${VISION}" ]; then bash ./install_vision.sh; fi -RUN rm install_vision.sh cache_vision_models.sh common_utils.sh -ENV INSTALLED_VISION ${VISION} - # Install rocm ARG ROCM_VERSION ENV ROCM_VERSION=${ROCM_VERSION} @@ -63,14 +57,27 @@ RUN rm -r ci_commit_pins COPY ./common/install_rocm_magma.sh install_rocm_magma.sh RUN if [ "${ROCM_VERSION}" != "nightly" ]; then bash ./install_rocm_magma.sh ${ROCM_VERSION}; fi RUN rm install_rocm_magma.sh +COPY ./common/install_rocSHMEM.sh install_rocSHMEM.sh +RUN bash ./install_rocSHMEM.sh ${ROCM_VERSION} +RUN rm install_rocSHMEM.sh ADD ./common/install_miopen.sh install_miopen.sh RUN if [ "${ROCM_VERSION}" != "nightly" ]; then bash ./install_miopen.sh ${ROCM_VERSION}; fi && rm install_miopen.sh ADD ./common/install_rocm_drm.sh install_rocm_drm.sh RUN if [ "${ROCM_VERSION}" != "nightly" ]; then bash ./install_rocm_drm.sh /usr ; fi && rm install_rocm_drm.sh -# ROCm environment variables are set in /etc/rocm_env.sh by install_rocm.sh -# and sourced via /etc/bash.bashrc for interactive shells. -# CI scripts should source /etc/rocm_env.sh directly. +# Default ROCm environment; /etc/rocm_env.sh (created by install_rocm.sh) may +# override these at runtime for different install methods (tarballs vs wheels). +ENV ROCM_PATH=/opt/rocm \ + ROCM_HOME=/opt/rocm \ + ROCM_SOURCE_DIR=/opt/rocm \ + ROCM_BIN=/opt/rocm/bin \ + ROCM_CMAKE=/opt/rocm \ + ROCM_DEVICE_LIB_PATH=/opt/rocm/amdgcn/bitcode \ + HIP_DEVICE_LIB_PATH=/opt/rocm/amdgcn/bitcode \ + MAGMA_HOME=/opt/rocm/magma +ENV PATH=/opt/rocm/bin:/opt/rocm/llvm/bin:$PATH +ENV LD_LIBRARY_PATH=/opt/rocm/lib:${LD_LIBRARY_PATH:-} + ENV LANG C.UTF-8 ENV LC_ALL C.UTF-8 @@ -79,22 +86,6 @@ COPY ./common/install_amdsmi.sh install_amdsmi.sh RUN bash ./install_amdsmi.sh RUN rm install_amdsmi.sh -# (optional) Install UCC -ARG UCX_COMMIT -ARG UCC_COMMIT -ENV UCX_COMMIT $UCX_COMMIT -ENV UCC_COMMIT $UCC_COMMIT -ENV UCX_HOME /usr -ENV UCC_HOME /usr -ADD ./common/install_ucc.sh install_ucc.sh -RUN if [ -n "${UCX_COMMIT}" ] && [ -n "${UCC_COMMIT}" ]; then bash ./install_ucc.sh; fi -RUN rm install_ucc.sh - -COPY ./common/install_openssl.sh install_openssl.sh -ENV OPENSSL_ROOT_DIR /opt/openssl -RUN bash ./install_openssl.sh -ENV OPENSSL_DIR /opt/openssl - ARG INDUCTOR_BENCHMARKS ARG ANACONDA_PYTHON_VERSION ENV ANACONDA_PYTHON_VERSION=$ANACONDA_PYTHON_VERSION @@ -106,12 +97,6 @@ COPY ci_commit_pins/torchbench.txt torchbench.txt RUN if [ -n "${INDUCTOR_BENCHMARKS}" ]; then bash ./install_inductor_benchmark_deps.sh; fi RUN rm install_inductor_benchmark_deps.sh common_utils.sh timm.txt huggingface-requirements.txt torchbench.txt -# (optional) Install non-default Ninja version -ARG NINJA_VERSION -COPY ./common/install_ninja.sh install_ninja.sh -RUN if [ -n "${NINJA_VERSION}" ]; then bash ./install_ninja.sh; fi -RUN rm install_ninja.sh - ARG TRITON # Install triton, this needs to be done before sccache because the latter will # try to reach out to S3, which docker build runners don't have access @@ -137,8 +122,5 @@ RUN rm install_openmpi.sh ARG BUILD_ENVIRONMENT ENV BUILD_ENVIRONMENT ${BUILD_ENVIRONMENT} -# Install LLVM dev version (Defined in the pytorch/builder github repository) -COPY --from=pytorch/llvm:9.0.1 /opt/llvm /opt/llvm - USER jenkins CMD ["bash"] diff --git a/.ci/docker/ubuntu-xpu/Dockerfile b/.ci/docker/ubuntu-xpu/Dockerfile index c61612882032d..3aba712da78d7 100644 --- a/.ci/docker/ubuntu-xpu/Dockerfile +++ b/.ci/docker/ubuntu-xpu/Dockerfile @@ -47,12 +47,6 @@ RUN bash ./install_gcc.sh && rm install_gcc.sh COPY ./common/install_lcov.sh install_lcov.sh RUN bash ./install_lcov.sh && rm install_lcov.sh -COPY ./common/install_openssl.sh install_openssl.sh -RUN bash ./install_openssl.sh -ENV OPENSSL_ROOT_DIR /opt/openssl -ENV OPENSSL_DIR /opt/openssl -RUN rm install_openssl.sh - ARG INDUCTOR_BENCHMARKS ARG ANACONDA_PYTHON_VERSION ENV ANACONDA_PYTHON_VERSION=$ANACONDA_PYTHON_VERSION @@ -80,19 +74,6 @@ COPY triton_xpu_version.txt triton_version.txt RUN if [ -n "${TRITON}" ]; then bash ./install_triton.sh; fi RUN rm install_triton.sh common_utils.sh triton-xpu.txt triton_version.txt -# (optional) Install vision packages like OpenCV -ARG VISION -COPY ./common/install_vision.sh ./common/cache_vision_models.sh ./common/common_utils.sh ./ -RUN if [ -n "${VISION}" ]; then bash ./install_vision.sh; fi -RUN rm install_vision.sh cache_vision_models.sh common_utils.sh -ENV INSTALLED_VISION ${VISION} - -# (optional) Install non-default Ninja version -ARG NINJA_VERSION -COPY ./common/install_ninja.sh install_ninja.sh -RUN if [ -n "${NINJA_VERSION}" ]; then bash ./install_ninja.sh; fi -RUN rm install_ninja.sh - # Install ccache/sccache (do this last, so we get priority in PATH) COPY ./common/install_cache.sh install_cache.sh ENV PATH /opt/cache/bin:$PATH @@ -102,8 +83,5 @@ RUN bash ./install_cache.sh && rm install_cache.sh ARG BUILD_ENVIRONMENT ENV BUILD_ENVIRONMENT ${BUILD_ENVIRONMENT} -# Install LLVM dev version (Defined in the pytorch/builder github repository) -COPY --from=pytorch/llvm:9.0.1 /opt/llvm /opt/llvm - USER jenkins CMD ["bash"] diff --git a/.ci/docker/ubuntu/Dockerfile b/.ci/docker/ubuntu/Dockerfile index e719df220dbaf..6c98d6639fb16 100644 --- a/.ci/docker/ubuntu/Dockerfile +++ b/.ci/docker/ubuntu/Dockerfile @@ -42,7 +42,6 @@ COPY ./common/install_conda.sh install_conda.sh COPY ./common/common_utils.sh common_utils.sh COPY ./common/install_magma_conda.sh install_magma_conda.sh RUN bash ./install_conda.sh && rm install_conda.sh install_magma_conda.sh common_utils.sh /opt/conda/requirements-ci.txt /opt/conda/requirements-docs.txt -RUN if [ -n "${UNINSTALL_DILL}" ]; then pip uninstall -y dill; fi # Install gcc ARG GCC_VERSION @@ -68,36 +67,7 @@ ENV NCCL_INCLUDE_DIR="/usr/local/cuda/include/" ENV NCCL_LIB_DIR="/usr/local/cuda/lib64/" -# (optional) Install UCC -ARG UCX_COMMIT -ARG UCC_COMMIT ARG CUDA_VERSION -ENV UCX_COMMIT $UCX_COMMIT -ENV UCC_COMMIT $UCC_COMMIT -ENV UCX_HOME /usr -ENV UCC_HOME /usr -ADD ./common/install_ucc.sh install_ucc.sh -RUN if [ -n "${UCX_COMMIT}" ] && [ -n "${UCC_COMMIT}" ]; then bash ./install_ucc.sh; fi -RUN rm install_ucc.sh - -# (optional) Install vision packages like OpenCV -ARG VISION -COPY ./common/install_vision.sh ./common/cache_vision_models.sh ./common/common_utils.sh ./ -RUN if [ -n "${VISION}" ]; then bash ./install_vision.sh; fi -RUN rm install_vision.sh cache_vision_models.sh common_utils.sh -ENV INSTALLED_VISION ${VISION} - -# (optional) Install non-default Ninja version -ARG NINJA_VERSION -COPY ./common/install_ninja.sh install_ninja.sh -RUN if [ -n "${NINJA_VERSION}" ]; then bash ./install_ninja.sh; fi -RUN rm install_ninja.sh - -COPY ./common/install_openssl.sh install_openssl.sh -RUN bash ./install_openssl.sh -ENV OPENSSL_ROOT_DIR /opt/openssl -ENV OPENSSL_DIR /opt/openssl -RUN rm install_openssl.sh ARG INDUCTOR_BENCHMARKS COPY ./common/install_inductor_benchmark_deps.sh install_inductor_benchmark_deps.sh @@ -207,11 +177,6 @@ RUN rm install_openmpi.sh ARG BUILD_ENVIRONMENT ENV BUILD_ENVIRONMENT ${BUILD_ENVIRONMENT} -# Install LLVM dev version (Defined in the pytorch/builder github repository) -ARG SKIP_LLVM_SRC_BUILD_INSTALL -COPY --from=pytorch/llvm:9.0.1 /opt/llvm /opt/llvm -RUN if [ -n "${SKIP_LLVM_SRC_BUILD_INSTALL}" ]; then set -eu; rm -rf /opt/llvm; fi - # AWS specific CUDA build guidance ENV TORCH_NVCC_FLAGS "-Xfatbin -compress-all" ENV CUDA_PATH /usr/local/cuda diff --git a/.ci/libtorch/extract_libtorch_from_wheel.py b/.ci/libtorch/extract_libtorch_from_wheel.py new file mode 100644 index 0000000000000..423a1b82cb0b5 --- /dev/null +++ b/.ci/libtorch/extract_libtorch_from_wheel.py @@ -0,0 +1,334 @@ +#!/usr/bin/env python3 +"""Extract libtorch package from a PyTorch wheel. + +Creates a libtorch zip from a pre-built wheel by copying the C++ libraries, +headers, and CMake files. On Linux, optionally splits debug symbols from +libtorch_cpu.so into a separate debug zip. + +Usage: + python extract_libtorch_from_wheel.py \ + --wheel-dir DIR --output-dir DIR --platform linux|macos|windows +""" + +import argparse +import glob +import os +import re +import shutil +import subprocess +import sys +import zipfile +from pathlib import Path + + +def find_wheel(wheel_dir: str) -> Path: + wheels = glob.glob(os.path.join(wheel_dir, "*.whl")) + if not wheels: + raise FileNotFoundError(f"No .whl files found in {wheel_dir}") + if len(wheels) > 1: + raise RuntimeError(f"Multiple .whl files found in {wheel_dir}: {wheels}") + return Path(wheels[0]) + + +def parse_version_from_wheel(wheel_path: Path) -> str: + # Wheel filename format: {name}-{version}(-{build})?-{python}-{abi}-{platform}.whl + name = wheel_path.stem + parts = name.split("-") + if len(parts) < 3: + raise ValueError(f"Cannot parse version from wheel filename: {wheel_path.name}") + return parts[1] + + +def extract_wheel(wheel_path: Path, extract_dir: Path) -> Path: + with zipfile.ZipFile(wheel_path, "r") as zf: + zf.extractall(extract_dir) + # Find the torch directory + torch_dir = extract_dir / "torch" + if not torch_dir.is_dir(): + raise FileNotFoundError( + f"No 'torch' directory found in extracted wheel at {extract_dir}" + ) + return torch_dir + + +def should_exclude_lib(filename: str) -> bool: + """Return True for files that should not go into the libtorch package.""" + if filename.startswith("libtorch_python"): + return True + if re.match(r"_C\.cpython.*", filename): + return True + if filename.endswith((".py", ".pyc")): + return True + if filename == "__init__.py": + return True + return False + + +def _is_lib_file(name: str, platform: str) -> bool: + """Return True if the file looks like a library or header to include.""" + if platform == "linux": + return ".so" in name or name.endswith(".a") + elif platform == "macos": + return name.endswith((".dylib", ".a")) + elif platform == "windows": + return name.endswith((".dll", ".lib", ".pdb")) + return False + + +def copy_libraries(torch_dir: Path, libtorch_lib: Path, platform: str) -> None: + """Copy libraries from torch/lib/ to libtorch/lib/.""" + torch_lib = torch_dir / "lib" + if not torch_lib.is_dir(): + raise FileNotFoundError(f"torch/lib/ not found at {torch_lib}") + + for item in torch_lib.iterdir(): + if item.is_dir(): + # Copy subdirectories (e.g. libshm/) as-is + shutil.copytree(item, libtorch_lib / item.name, dirs_exist_ok=True) + continue + if should_exclude_lib(item.name): + continue + if _is_lib_file(item.name, platform): + shutil.copy2(item, libtorch_lib / item.name) + + # On macOS, also copy delocated dylibs from torch/.dylibs/ if present + if platform == "macos": + dylibs_dir = torch_dir / ".dylibs" + if dylibs_dir.is_dir(): + for item in dylibs_dir.iterdir(): + if item.suffix == ".dylib" and not should_exclude_lib(item.name): + shutil.copy2(item, libtorch_lib / item.name) + + +def copy_includes(torch_dir: Path, libtorch_include: Path) -> None: + torch_include = torch_dir / "include" + if not torch_include.is_dir(): + # Some older wheels might have include under torch/lib/include + torch_include = torch_dir / "lib" / "include" + if not torch_include.is_dir(): + raise FileNotFoundError("include/ not found in torch directory") + shutil.copytree(torch_include, libtorch_include, dirs_exist_ok=True) + + +def copy_cmake(torch_dir: Path, libtorch_share: Path) -> None: + torch_cmake = torch_dir / "share" / "cmake" + if not torch_cmake.is_dir(): + print(f"Warning: share/cmake/ not found at {torch_cmake}", file=sys.stderr) + return + cmake_dest = libtorch_share / "cmake" + shutil.copytree(torch_cmake, cmake_dest, dirs_exist_ok=True) + + +def copy_bin(torch_dir: Path, libtorch_bin: Path, platform: str) -> None: + """Copy binary executables (mainly relevant for Windows).""" + if platform == "windows": + torch_lib = torch_dir / "lib" + if torch_lib.is_dir(): + for item in torch_lib.iterdir(): + if item.suffix == ".dll" and not should_exclude_lib(item.name): + shutil.copy2(item, libtorch_bin / item.name) + + +def write_metadata(libtorch_dir: Path, version: str, git_hash: str) -> None: + (libtorch_dir / "build-version").write_text(version + "\n") + (libtorch_dir / "build-hash").write_text(git_hash + "\n") + + +def get_git_hash(torch_dir: Path) -> str: + """Read git_version from the wheel's torch/version.py.""" + version_file = torch_dir / "version.py" + if not version_file.exists(): + return "unknown" + from ast import literal_eval + + for line in version_file.read_text().splitlines(): + if line.strip().startswith("git_version"): + try: + return literal_eval(line.partition("=")[2].strip()) + except Exception: + pass + return "unknown" + + +def split_debug_symbols( + libtorch_dir: Path, output_dir: Path, zip_prefix: str, version: str +) -> None: + """Split debug symbols from libtorch_cpu.so (Linux only).""" + libtorch_cpu = libtorch_dir / "lib" / "libtorch_cpu.so" + if not libtorch_cpu.exists(): + print( + "Warning: libtorch_cpu.so not found, skipping debug symbol split", + file=sys.stderr, + ) + return + + debug_dir = libtorch_dir.parent / "debug" + debug_dir.mkdir(exist_ok=True) + dbg_file = debug_dir / "libtorch_cpu.so.dbg" + + # Copy to create debug file + shutil.copy2(libtorch_cpu, dbg_file) + + # Keep only debug symbols + subprocess.run( + ["strip", "--only-keep-debug", str(dbg_file)], + check=True, + ) + + # Strip debug info from release lib + subprocess.run( + ["strip", "--strip-debug", str(libtorch_cpu)], + check=True, + ) + + # Add debug link + subprocess.run( + ["objcopy", str(libtorch_cpu), f"--add-gnu-debuglink={dbg_file}"], + check=True, + cwd=str(libtorch_dir / "lib"), + ) + + # Extract CRC32 from the debug link section + try: + result = subprocess.run( + [ + "bash", + "-c", + f"objcopy --dump-section .gnu_debuglink=>(tail -c4 | od -t x4 -An | xargs echo) {libtorch_cpu}", + ], + capture_output=True, + text=True, + check=True, + ) + crc32 = result.stdout.strip() + except subprocess.CalledProcessError: + crc32 = "unknown" + + # Create debug zip + debug_zip = output_dir / f"debug-{zip_prefix}-{version}-{crc32}.zip" + with zipfile.ZipFile(debug_zip, "w", zipfile.ZIP_DEFLATED) as zf: + zf.write(dbg_file, "debug/libtorch_cpu.so.dbg") + + print(f"Debug symbols zip: {debug_zip}") + + +def create_libtorch_zip( + libtorch_dir: Path, + output_dir: Path, + zip_prefix: str, + version: str, +) -> Path: + zip_path = output_dir / f"{zip_prefix}-{version}.zip" + with zipfile.ZipFile(zip_path, "w", zipfile.ZIP_DEFLATED) as zf: + for root, dirs, files in os.walk(libtorch_dir): + for f in files: + filepath = Path(root) / f + arcname = filepath.relative_to(libtorch_dir.parent) + zf.write(filepath, arcname) + # Create latest symlink + latest_zip = output_dir / f"{zip_prefix}-latest.zip" + latest_zip.symlink_to(zip_path.name) + + print(f"Libtorch zip: {zip_path}") + print(f"Libtorch latest zip: {latest_zip}") + return zip_path + + +def compute_zip_prefix(platform: str, desired_cuda: str, libtorch_variant: str) -> str: + """Compute the zip filename prefix matching existing naming conventions. + + Linux: libtorch-shared-with-deps + macOS: libtorch-macos-arm64 + Windows: libtorch-win-shared-with-deps (or libtorch-win-arm64-shared-with-deps) + """ + if platform == "macos": + return "libtorch-macos-arm64" + elif platform == "windows": + return f"libtorch-win-{libtorch_variant}" + else: + return f"libtorch-{libtorch_variant}" + + +def main() -> None: + parser = argparse.ArgumentParser(description="Extract libtorch from PyTorch wheel") + parser.add_argument( + "--wheel-dir", required=True, help="Directory containing the .whl file" + ) + parser.add_argument( + "--output-dir", required=True, help="Directory for output zip files" + ) + parser.add_argument( + "--platform", + required=True, + choices=["linux", "macos", "windows"], + help="Target platform", + ) + parser.add_argument( + "--desired-cuda", + default="cpu", + help="CUDA variant (cpu, cu126, cu128, rocm7.1, etc.)", + ) + parser.add_argument( + "--libtorch-variant", + default="shared-with-deps", + help="Libtorch variant (shared-with-deps, etc.)", + ) + parser.add_argument( + "--git-hash", + default="", + help="Git hash to use for build-hash (auto-detected if not set)", + ) + args = parser.parse_args() + + wheel_dir = Path(args.wheel_dir) + output_dir = Path(args.output_dir) + output_dir.mkdir(parents=True, exist_ok=True) + + # Find and extract wheel + wheel_path = find_wheel(str(wheel_dir)) + version = parse_version_from_wheel(wheel_path) + print(f"Found wheel: {wheel_path}") + print(f"Version: {version}") + + extract_dir = wheel_dir / "_extract_tmp" + if extract_dir.exists(): + shutil.rmtree(extract_dir) + extract_dir.mkdir() + + try: + torch_dir = extract_wheel(wheel_path, extract_dir) + + # Create libtorch directory structure + libtorch_dir = extract_dir / "libtorch" + libtorch_dir.mkdir() + for subdir in ["lib", "bin", "include", "share"]: + (libtorch_dir / subdir).mkdir() + + # Copy components + copy_libraries(torch_dir, libtorch_dir / "lib", args.platform) + copy_includes(torch_dir, libtorch_dir / "include") + copy_cmake(torch_dir, libtorch_dir / "share") + copy_bin(torch_dir, libtorch_dir / "bin", args.platform) + + # Write metadata + git_hash = args.git_hash or get_git_hash(torch_dir) + write_metadata(libtorch_dir, version, git_hash) + + # Compute zip prefix + zip_prefix = compute_zip_prefix( + args.platform, args.desired_cuda, args.libtorch_variant + ) + + # Split debug symbols on Linux + if args.platform == "linux": + split_debug_symbols(libtorch_dir, output_dir, zip_prefix, version) + + # Create the zip + create_libtorch_zip(libtorch_dir, output_dir, zip_prefix, version) + + finally: + shutil.rmtree(extract_dir) + + +if __name__ == "__main__": + main() diff --git a/.ci/lumen_cli/cli/build_cli/register_build.py b/.ci/lumen_cli/cli/build_cli/register_build.py index 9f35a9c8165dc..277abce662bcd 100644 --- a/.ci/lumen_cli/cli/build_cli/register_build.py +++ b/.ci/lumen_cli/cli/build_cli/register_build.py @@ -1,4 +1,6 @@ -import argparse +from __future__ import annotations + +import argparse # noqa: TC003 import logging from cli.lib.common.cli_helper import register_targets, RichHelp, TargetSpec diff --git a/.ci/lumen_cli/cli/lib/common/cli_helper.py b/.ci/lumen_cli/cli/lib/common/cli_helper.py index 4086eb7d46e81..b19ad297ca1f0 100644 --- a/.ci/lumen_cli/cli/lib/common/cli_helper.py +++ b/.ci/lumen_cli/cli/lib/common/cli_helper.py @@ -3,6 +3,8 @@ """ +from __future__ import annotations + import argparse from abc import ABC, abstractmethod @@ -11,7 +13,7 @@ from collections.abc import Callable # Python 3.11+ from typing import Any, Required, TypedDict except ImportError: - from collections.abc import Callable + from collections.abc import Callable # noqa: TC003 from typing import Any, TypedDict from typing_extensions import Required # Fallback for Python <3.11 diff --git a/.ci/lumen_cli/cli/lib/common/docker_helper.py b/.ci/lumen_cli/cli/lib/common/docker_helper.py index b5f0a90e2d47a..1c96ab37e9219 100644 --- a/.ci/lumen_cli/cli/lib/common/docker_helper.py +++ b/.ci/lumen_cli/cli/lib/common/docker_helper.py @@ -2,8 +2,9 @@ Docker Utility helpers for CLI tasks. """ +from __future__ import annotations + import logging -from typing import Optional import docker from docker.errors import APIError, NotFound @@ -12,7 +13,7 @@ logger = logging.getLogger(__name__) # lazy singleton so we don't reconnect every call -_docker_client: Optional[docker.DockerClient] = None +_docker_client: docker.DockerClient | None = None def _get_client() -> docker.DockerClient: @@ -23,7 +24,7 @@ def _get_client() -> docker.DockerClient: def local_image_exists( - image_name: str, client: Optional[docker.DockerClient] = None + image_name: str, client: docker.DockerClient | None = None ) -> bool: """Return True if a local Docker image exists.""" if not image_name: diff --git a/.ci/lumen_cli/cli/lib/common/envs_helper.py b/.ci/lumen_cli/cli/lib/common/envs_helper.py index a654e7f18ed9f..06c33163711a1 100644 --- a/.ci/lumen_cli/cli/lib/common/envs_helper.py +++ b/.ci/lumen_cli/cli/lib/common/envs_helper.py @@ -2,11 +2,12 @@ Environment Variables and Dataclasses Utility helpers for CLI tasks. """ +from __future__ import annotations + import os from dataclasses import field, fields, is_dataclass, MISSING from pathlib import Path from textwrap import indent -from typing import Optional, Union from cli.lib.common.utils import str2bool @@ -18,9 +19,9 @@ def get_env(name: str, default: str = "") -> str: def env_path_optional( name: str, - default: Optional[Union[str, Path]] = None, + default: str | Path | None = None, resolve: bool = True, -) -> Optional[Path]: +) -> Path | None: """Get environment variable as optional Path.""" val = get_env(name) or default if not val: @@ -32,7 +33,7 @@ def env_path_optional( def env_path( name: str, - default: Optional[Union[str, Path]] = None, + default: str | Path | None = None, resolve: bool = True, ) -> Path: """Get environment variable as Path, raise if missing.""" @@ -61,7 +62,7 @@ def env_bool_field( def env_path_field( name: str, - default: Union[str, Path] = "", + default: str | Path = "", *, resolve: bool = True, ) -> Path: diff --git a/.ci/lumen_cli/cli/lib/common/path_helper.py b/.ci/lumen_cli/cli/lib/common/path_helper.py index 4f74aa6e509de..3102d66ff7131 100644 --- a/.ci/lumen_cli/cli/lib/common/path_helper.py +++ b/.ci/lumen_cli/cli/lib/common/path_helper.py @@ -1,15 +1,16 @@ """Path utility helpers for CLI tasks.""" +from __future__ import annotations + import logging import shutil from pathlib import Path -from typing import Union logger = logging.getLogger(__name__) -def get_path(path: Union[str, Path], resolve: bool = False) -> Path: +def get_path(path: str | Path, resolve: bool = False) -> Path: """Convert to Path object, optionally resolving to absolute path.""" if not path: raise ValueError("Path cannot be None or empty") @@ -17,14 +18,14 @@ def get_path(path: Union[str, Path], resolve: bool = False) -> Path: return result.resolve() if resolve else result -def ensure_dir_exists(path: Union[str, Path]) -> Path: +def ensure_dir_exists(path: str | Path) -> Path: """Create directory if it doesn't exist.""" path_obj = get_path(path) path_obj.mkdir(parents=True, exist_ok=True) return path_obj -def remove_dir(path: Union[str, Path, None]) -> None: +def remove_dir(path: str | Path | None) -> None: """Remove directory if it exists.""" if not path: return @@ -33,13 +34,13 @@ def remove_dir(path: Union[str, Path, None]) -> None: shutil.rmtree(path_obj) -def force_create_dir(path: Union[str, Path]) -> Path: +def force_create_dir(path: str | Path) -> Path: """Remove directory if exists, then create fresh empty directory.""" remove_dir(path) return ensure_dir_exists(path) -def copy(src: Union[str, Path], dst: Union[str, Path]) -> None: +def copy(src: str | Path, dst: str | Path) -> None: """Copy file or directory from src to dst.""" src_path = get_path(src, resolve=True) dst_path = get_path(dst, resolve=True) @@ -57,6 +58,6 @@ def copy(src: Union[str, Path], dst: Union[str, Path]) -> None: raise ValueError(f"Unsupported path type: {src_path}") -def is_path_exist(path: Union[str, Path, None]) -> bool: +def is_path_exist(path: str | Path | None) -> bool: """Check if path exists.""" return bool(path and get_path(path).exists()) diff --git a/.ci/lumen_cli/cli/lib/common/pip_helper.py b/.ci/lumen_cli/cli/lib/common/pip_helper.py index a53747e24d256..a0cb1e17840e1 100644 --- a/.ci/lumen_cli/cli/lib/common/pip_helper.py +++ b/.ci/lumen_cli/cli/lib/common/pip_helper.py @@ -1,11 +1,12 @@ +from __future__ import annotations + import glob import logging import shlex import shutil import sys -from collections.abc import Iterable +from collections.abc import Iterable # noqa: TC003 from importlib.metadata import PackageNotFoundError, version # noqa: UP035 -from typing import Optional, Union from cli.lib.common.utils import run_command @@ -17,8 +18,8 @@ def pip_install_packages( packages: Iterable[str] = (), env=None, *, - requirements: Optional[str] = None, - constraints: Optional[str] = None, + requirements: str | None = None, + constraints: str | None = None, prefer_uv: bool = False, ) -> None: use_uv = prefer_uv and shutil.which("uv") is not None @@ -37,14 +38,14 @@ def pip_install_packages( run_command(" ".join(map(shlex.quote, cmd)), env=env) -def pip_install_first_match(pattern: str, extras: Optional[str] = None, pref_uv=False): +def pip_install_first_match(pattern: str, extras: str | None = None, pref_uv=False): wheel = first_matching_pkg(pattern) target = f"{wheel}[{extras}]" if extras else wheel logger.info("Installing %s...", target) pip_install_packages([target], prefer_uv=pref_uv) -def run_python(args: Union[str, list[str]], env=None): +def run_python(args: str | list[str], env=None): """ Run the python in the current environment. """ diff --git a/.ci/lumen_cli/cli/lib/common/utils.py b/.ci/lumen_cli/cli/lib/common/utils.py index b03309810d986..0798d7d1d369b 100644 --- a/.ci/lumen_cli/cli/lib/common/utils.py +++ b/.ci/lumen_cli/cli/lib/common/utils.py @@ -2,6 +2,8 @@ General Utility helpers for CLI tasks. """ +from __future__ import annotations + import logging import os import shlex @@ -9,7 +11,6 @@ import sys from contextlib import contextmanager from pathlib import Path -from typing import Optional logger = logging.getLogger(__name__) @@ -19,8 +20,8 @@ def run_command( cmd: str, use_shell: bool = False, log_cmd: bool = True, - cwd: Optional[str] = None, - env: Optional[dict] = None, + cwd: str | None = None, + env: dict | None = None, check: bool = True, ) -> int: """Run a command with optional shell execution.""" @@ -61,7 +62,7 @@ def run_command( return proc.returncode -def str2bool(value: Optional[str]) -> bool: +def str2bool(value: str | None) -> bool: """Convert environment variables to boolean values.""" if not value: return False @@ -120,7 +121,7 @@ def working_directory(path: str): def get_wheels( output_dir: Path, - max_depth: Optional[int] = None, + max_depth: int | None = None, ) -> list[str]: """Return a list of wheels found in the given output directory.""" root = Path(output_dir) diff --git a/.ci/lumen_cli/cli/lib/core/torchtitan/lib.py b/.ci/lumen_cli/cli/lib/core/torchtitan/lib.py new file mode 100644 index 0000000000000..3b4a818ff04ac --- /dev/null +++ b/.ci/lumen_cli/cli/lib/core/torchtitan/lib.py @@ -0,0 +1,65 @@ +import logging +from pathlib import Path +from typing import Any + +import yaml +from cli.lib.common.git_helper import clone_external_repo +from cli.lib.common.utils import run_command, temp_environ, working_directory + + +logger = logging.getLogger(__name__) + +_TORCHTITAN_TEST_LIBRARY_PATH = Path(__file__).parent / "torchtitan_test_library.yaml" + + +def _load_torchtitan_test_library_yaml() -> dict[str, Any]: + if not _TORCHTITAN_TEST_LIBRARY_PATH.exists(): + raise FileNotFoundError( + f"torchtitan test library YAML not found: {_TORCHTITAN_TEST_LIBRARY_PATH}" + ) + with open(_TORCHTITAN_TEST_LIBRARY_PATH, encoding="utf-8") as f: + return yaml.safe_load(f) + + +def load_torchtitan_test_library() -> dict[str, Any]: + return _load_torchtitan_test_library_yaml() + + +def clone_torchtitan(dst: str = "torchtitan"): + _, commit = clone_external_repo( + target="torchtitan", + repo="https://github.com/pytorch/torchtitan.git", + dst=dst, + ) + return commit + + +def run_test_plan( + test_plan: str, + tests_map: dict[str, Any], +): + logger.info("Running torchtitan test plan: %s", test_plan) + if test_plan not in tests_map: + raise RuntimeError( + f"test plan '{test_plan}' not found in torchtitan test library" + ) + + tests = tests_map[test_plan] + title = tests.get("title", "unknown test") + logger.info("Running tests: %s", title) + + with ( + working_directory(tests.get("working_directory", "")), + temp_environ(tests.get("env_vars", {})), + ): + failures = [] + for step in tests["steps"]: + logger.info("Running step: %s", step) + code = run_command(cmd=step, check=False, use_shell=True) + if code != 0: + failures.append(step) + logger.info("Finished step: %s", step) + if failures: + logger.error("Failed steps: %s", failures) + raise RuntimeError(f"{len(failures)} test steps failed: {failures}") + logger.info("All tests passed for plan: %s", test_plan) diff --git a/.ci/lumen_cli/cli/lib/core/torchtitan/torchtitan_test.py b/.ci/lumen_cli/cli/lib/core/torchtitan/torchtitan_test.py new file mode 100644 index 0000000000000..d42fc93d6f649 --- /dev/null +++ b/.ci/lumen_cli/cli/lib/core/torchtitan/torchtitan_test.py @@ -0,0 +1,41 @@ +import logging +from typing import Any + +from cli.lib.common.cli_helper import BaseRunner +from cli.lib.common.pip_helper import pip_install_packages +from cli.lib.common.utils import working_directory +from cli.lib.core.torchtitan.lib import ( + clone_torchtitan, + load_torchtitan_test_library, + run_test_plan, +) + + +logger = logging.getLogger(__name__) + + +class TorchtitanTestRunner(BaseRunner): + def __init__(self, args: Any): + self.work_directory = "torchtitan" + self.test_plan = args.test_plan + + def prepare(self): + clone_torchtitan(dst=self.work_directory) + # torchao and torchcomms nightlies are required by torchtitan + pip_install_packages( + packages=[ + "--pre", + "torchao", + "torchcomms", + "--index-url", + "https://download.pytorch.org/whl/nightly/cu129", + ], + ) + with working_directory(self.work_directory): + pip_install_packages(packages=["-e", "."]) + pip_install_packages(packages=["pytest", "pytest-cov"]) + + def run(self): + self.prepare() + with working_directory(self.work_directory): + run_test_plan(self.test_plan, load_torchtitan_test_library()) diff --git a/.ci/lumen_cli/cli/lib/core/torchtitan/torchtitan_test_library.yaml b/.ci/lumen_cli/cli/lib/core/torchtitan/torchtitan_test_library.yaml new file mode 100644 index 0000000000000..eba45310855c3 --- /dev/null +++ b/.ci/lumen_cli/cli/lib/core/torchtitan/torchtitan_test_library.yaml @@ -0,0 +1,21 @@ +# torchtitan Test Library Configuration +# Each test plan maps to torchtitan's own test runners. +# When tests are added/removed in torchtitan, the daily pin bump picks them up. +# Test filtering and exclusions are managed in the torchtitan repo via +# scripts/ci/pytorch_ci_test_runner.sh. + +torchtitan_features_integration: + title: torchtitan Feature Integration Tests (8 GPU) + id: torchtitan_features_integration + env_vars: + NGPU: "8" + steps: + - scripts/ci/pytorch_ci_test_runner.sh feature_tests + +torchtitan_models_integration: + title: torchtitan Model Integration Tests (8 GPU) + id: torchtitan_models_integration + env_vars: + NGPU: "8" + steps: + - scripts/ci/pytorch_ci_test_runner.sh model_tests diff --git a/.ci/lumen_cli/cli/lib/core/vllm/disabled_vllm_tests.yaml b/.ci/lumen_cli/cli/lib/core/vllm/disabled_vllm_tests.yaml new file mode 100644 index 0000000000000..a350ced6be862 --- /dev/null +++ b/.ci/lumen_cli/cli/lib/core/vllm/disabled_vllm_tests.yaml @@ -0,0 +1,21 @@ +# Disabled vLLM tests for PyTorch CI. +# Node IDs are relative to vLLM's tests/ directory. +# Each entry requires 'test' (node ID) and 'issue' (tracking URL). +# +# To disable a test, copy one of the entries below into the disabled_tests list. +# +# Disable an entire test file (produces --ignore): +# - test: basic_correctness/test_basic_correctness.py +# issue: https://github.com/pytorch/pytorch/issues/12345 +# +# Disable a single test (produces --deselect): +# - test: basic_correctness/test_basic_correctness.py::test_something +# issue: https://github.com/pytorch/pytorch/issues/12345 +# +# Disable only in specific test plans via 'configs': +# - test: compile/test_fusion.py +# issue: https://github.com/pytorch/pytorch/issues/12345 +# configs: +# - vllm_pytorch_compilation_unit_tests + +disabled_tests: [] diff --git a/.ci/lumen_cli/cli/lib/core/vllm/lib.py b/.ci/lumen_cli/cli/lib/core/vllm/lib.py index 2532a97ebf345..f197080857f72 100644 --- a/.ci/lumen_cli/cli/lib/core/vllm/lib.py +++ b/.ci/lumen_cli/cli/lib/core/vllm/lib.py @@ -1,6 +1,11 @@ +from __future__ import annotations + +import json import logging import os +import re import textwrap +import urllib.request from pathlib import Path from typing import Any @@ -19,6 +24,8 @@ logger = logging.getLogger(__name__) _VLLM_TEST_LIBRARY_PATH = Path(__file__).parent / "vllm_test_library.yaml" +_DISABLED_VLLM_TESTS_PATH = Path(__file__).parent / "disabled_vllm_tests.yaml" +_DISABLED_VLLM_TESTS_ISSUE = 175899 def _load_vllm_test_library_yaml() -> dict[str, Any]: @@ -87,6 +94,92 @@ def check_parallelism(tests: Any, title: str, shard_id: int = 0, num_shards: int return True +def _load_disabled_vllm_tests_from_yaml() -> list[dict[str, Any]]: + if not _DISABLED_VLLM_TESTS_PATH.exists(): + return [] + with open(_DISABLED_VLLM_TESTS_PATH, encoding="utf-8") as f: + data = yaml.safe_load(f) + if not data or "disabled_tests" not in data: + return [] + entries = data["disabled_tests"] + if not entries: + return [] + for entry in entries: + if "test" not in entry or "issue" not in entry: + raise ValueError( + f"disabled_vllm_tests.yaml: each entry must have 'test' and 'issue' keys, got {entry}" + ) + return entries + + +def _parse_disabled_tests_from_issue_body(body: str) -> list[dict[str, Any]]: + match = re.search(r"```yaml\s*\n(.*?)```", body, re.DOTALL) + if not match: + return [] + block = match.group(1) + data = yaml.safe_load(block) + if not data or "disabled_tests" not in data: + return [] + return data["disabled_tests"] or [] + + +def _load_disabled_vllm_tests_from_github() -> list[dict[str, Any]]: + if not _DISABLED_VLLM_TESTS_ISSUE: + return [] + url = f"https://api.github.com/repos/pytorch/pytorch/issues/{_DISABLED_VLLM_TESTS_ISSUE}" + headers = {"Accept": "application/vnd.github.v3+json"} + token = os.environ.get("GITHUB_TOKEN") + if token: + headers["Authorization"] = f"token {token}" + try: + req = urllib.request.Request(url, headers=headers) + with urllib.request.urlopen(req, timeout=30) as resp: + issue = json.loads(resp.read()) + body = issue.get("body", "") or "" + entries = _parse_disabled_tests_from_issue_body(body) + # Filter out malformed entries — the issue body is user-editable + entries = [e for e in entries if "test" in e] + issue_url = issue.get("html_url", url) + for entry in entries: + entry.setdefault("issue", issue_url) + return entries + except Exception: + logger.warning( + "Failed to fetch disabled vLLM tests from GitHub issue #%d", + _DISABLED_VLLM_TESTS_ISSUE, + exc_info=True, + ) + return [] + + +def _load_disabled_vllm_tests() -> list[dict[str, Any]]: + yaml_entries = _load_disabled_vllm_tests_from_yaml() + github_entries = _load_disabled_vllm_tests_from_github() + seen = {e["test"] for e in yaml_entries} + merged = list(yaml_entries) + for entry in github_entries: + if entry["test"] not in seen: + seen.add(entry["test"]) + merged.append(entry) + return merged + + +def _build_disabled_test_flags( + disabled_tests: list[dict[str, Any]], test_plan: str +) -> str: + flags = [] + for entry in disabled_tests: + configs = entry.get("configs") + if configs and test_plan not in configs: + continue + node_id = entry["test"] + if "::" in node_id: + flags.append(f"--deselect={node_id}") + else: + flags.append(f"--ignore={node_id}") + return " ".join(flags) + + def run_test_plan( test_plan: str, test_target: str, @@ -110,6 +203,11 @@ def run_test_plan( if is_parallel: title = title.replace("%N", f"{shard_id}/{num_shards}") + disabled_tests = _load_disabled_vllm_tests() + disabled_flags = _build_disabled_test_flags(disabled_tests, test_plan) + if disabled_flags: + logger.info("Disabled test flags for %s: %s", test_plan, disabled_flags) + logger.info("Running tests: %s", title) if pkgs: logger.info("Installing packages: %s", pkgs) @@ -124,11 +222,14 @@ def run_test_plan( if is_parallel: step = replace_buildkite_placeholders(step, shard_id, num_shards) logger.info("Running parallel step: %s", step) - # Support retry with delay for all pytest commands, pytest-rerunfailures - # is already a dependency of vLLM. This is needed as a stop gap to reduce - # the number of requests to HF until #172300 can be landed to enable - # HF offline mode if "pytest" in step: + # Inject disabled test flags before rerun flags + if disabled_flags: + step = step.replace("pytest", f"pytest {disabled_flags}", 1) + # Support retry with delay for all pytest commands, pytest-rerunfailures + # is already a dependency of vLLM. This is needed as a stop gap to reduce + # the number of requests to HF until #172300 can be landed to enable + # HF offline mode. # Use a low retry count and a high delay value to lower the risk of # having a retry storm and make thing worse rerun_count = os.getenv( diff --git a/.ci/lumen_cli/cli/lib/core/vllm/vllm_build.py b/.ci/lumen_cli/cli/lib/core/vllm/vllm_build.py index 63e5f7a28de54..f0780f98541f3 100644 --- a/.ci/lumen_cli/cli/lib/core/vllm/vllm_build.py +++ b/.ci/lumen_cli/cli/lib/core/vllm/vllm_build.py @@ -1,9 +1,10 @@ +from __future__ import annotations + import logging import os import textwrap from dataclasses import dataclass from pathlib import Path -from typing import Optional from cli.lib.common.cli_helper import BaseRunner from cli.lib.common.docker_helper import local_image_exists @@ -235,7 +236,7 @@ def get_result_path(self, path): abs_path = get_path(path, resolve=True) return abs_path - def _get_torch_wheel_path_arg(self, torch_whl_dir: Optional[Path]) -> str: + def _get_torch_wheel_path_arg(self, torch_whl_dir: Path | None) -> str: if not torch_whl_dir: return "" return f"--build-arg TORCH_WHEELS_PATH={_VLLM_TEMP_FOLDER}" diff --git a/.ci/lumen_cli/cli/lib/core/vllm/vllm_test_library.yaml b/.ci/lumen_cli/cli/lib/core/vllm/vllm_test_library.yaml index f2f450b6f9004..402f2d8bf0e69 100644 --- a/.ci/lumen_cli/cli/lib/core/vllm/vllm_test_library.yaml +++ b/.ci/lumen_cli/cli/lib/core/vllm/vllm_test_library.yaml @@ -20,7 +20,7 @@ vllm_basic_models_test: - pytest -v -s models/test_registry.py - pytest -v -s models/test_utils.py - pytest -v -s models/test_vision.py - - pytest -v -s models/test_initialization.py + - HF_DATASETS_OFFLINE=0 TRANSFORMERS_OFFLINE=0 pytest -v -s models/test_initialization.py vllm_entrypoints_test: title: Entrypoints Test @@ -60,7 +60,7 @@ vllm_distributed_test_28_failure_test: VLLM_WORKER_MULTIPROC_METHOD: spawn num_gpus: 4 steps: - - pytest -v -s distributed/test_sequence_parallel.py + - pytest -v -s compile/correctness_e2e/test_sequence_parallel.py vllm_lora_28_failure_test: title: LoRA pytorch 2.8 failure test @@ -85,29 +85,28 @@ vllm_multi_model_test_28_failure_test: package_install: - git+https://github.com/TIGER-AI-Lab/Mantis.git steps: - - pytest -v -s models/multimodal/generation/test_voxtral.py -k 'not 5-128-half' - - HF_DATASETS_OFFLINE=0 TRANSFORMERS_OFFLINE=0 pytest -v -s models/multimodal/generation/test_voxtral.py -k 5-128-half + - HF_DATASETS_OFFLINE=0 TRANSFORMERS_OFFLINE=0 pytest -v -s models/multimodal/generation/test_voxtral.py - pytest -v -s models/multimodal/pooling vllm_pytorch_compilation_unit_tests: title: PyTorch Compilation Unit Tests id: vllm_pytorch_compilation_unit_tests steps: - - pytest -v -s compile/test_pass_manager.py - - pytest -v -s compile/test_fusion.py - - pytest -v -s compile/test_fusion_attn.py - - pytest -v -s compile/test_silu_mul_quant_fusion.py - - pytest -v -s compile/distributed/test_sequence_parallelism.py - - pytest -v -s compile/distributed/test_async_tp.py - - pytest -v -s compile/distributed/test_fusion_all_reduce.py + - pytest -v -s compile/passes/test_pass_manager.py + - pytest -v -s compile/passes/test_fusion.py + - pytest -v -s compile/passes/test_fusion_attn.py + - pytest -v -s compile/passes/test_silu_mul_quant_fusion.py + - pytest -v -s compile/passes/distributed/test_sequence_parallelism.py + - pytest -v -s compile/passes/distributed/test_async_tp.py + - pytest -v -s compile/passes/distributed/test_fusion_all_reduce.py - pytest -v -s compile/test_decorator.py vllm_language_model_test_extended_generation_28_failure_test: title: Language Models Test (Extended Generation) 2.8 release failure - id: vllm_languagde_model_test_extended_generation_28_failure_test + id: vllm_language_model_test_extended_generation_28_failure_test package_install: - --no-build-isolation - - git+https://github.com/Dao-AILab/causal-conv1d@v1.5.0.post8 + - git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0 steps: - pytest -v -s models/language/generation/test_mistral.py @@ -118,7 +117,7 @@ vllm_distributed_test_2_gpu_28_failure_test: VLLM_WORKER_MULTIPROC_METHOD: spawn num_gpus: 4 steps: - - pytest -v -s distributed/test_sequence_parallel.py + - pytest -v -s compile/correctness_e2e/test_sequence_parallel.py vllm_lora_test: title: LoRA Test %N diff --git a/.ci/lumen_cli/cli/test_cli/register_test.py b/.ci/lumen_cli/cli/test_cli/register_test.py index 2973341b83ed2..12088da4f2ef9 100644 --- a/.ci/lumen_cli/cli/test_cli/register_test.py +++ b/.ci/lumen_cli/cli/test_cli/register_test.py @@ -1,7 +1,10 @@ -import argparse +from __future__ import annotations + +import argparse # noqa: TC003 import logging from cli.lib.common.cli_helper import register_targets, RichHelp, TargetSpec +from cli.lib.core.torchtitan.torchtitan_test import TorchtitanTestRunner from cli.lib.core.vllm.vllm_test import VllmTestRunner @@ -13,8 +16,11 @@ "vllm": { "runner": VllmTestRunner, "help": "test vLLM with pytorch main", - } - # add yours ... + }, + "torchtitan": { + "runner": TorchtitanTestRunner, + "help": "test torchtitan with pytorch main", + }, } diff --git a/.ci/lumen_cli/pyproject.toml b/.ci/lumen_cli/pyproject.toml index b2ac379e34ab0..ce8cf59d99fda 100644 --- a/.ci/lumen_cli/pyproject.toml +++ b/.ci/lumen_cli/pyproject.toml @@ -6,7 +6,7 @@ dependencies = [ "GitPython==3.1.45", "docker==7.1.0", "pytest==7.3.2", - "uv==0.9.6" + "uv==0.11.6" ] [tool.setuptools] diff --git a/.ci/lumen_cli/tests/test_cli_helper.py b/.ci/lumen_cli/tests/test_cli_helper.py index 848f22d6be200..b839c8ddf62a6 100644 --- a/.ci/lumen_cli/tests/test_cli_helper.py +++ b/.ci/lumen_cli/tests/test_cli_helper.py @@ -1,3 +1,5 @@ +from __future__ import annotations + import argparse import io import unittest diff --git a/.ci/lumen_cli/tests/test_run_plan.py b/.ci/lumen_cli/tests/test_run_plan.py index 1ad3433cd6661..13d2695d21c0c 100644 --- a/.ci/lumen_cli/tests/test_run_plan.py +++ b/.ci/lumen_cli/tests/test_run_plan.py @@ -1,8 +1,11 @@ # tests/test_run_test_plan.py +from __future__ import annotations + import importlib +import json from contextlib import nullcontext from types import SimpleNamespace -from unittest.mock import MagicMock +from unittest.mock import MagicMock, patch as mock_patch import pytest @@ -29,13 +32,15 @@ def _get_check(c): def patch_module(monkeypatch): """ Patch helpers ('pip_install_packages', 'temp_environ', 'working_directory', - 'run_command', 'logger') inside the target module and expose them. + 'run_command', 'logger', '_load_disabled_vllm_tests') inside the target + module and expose them. """ module = importlib.import_module(MOD) # Create fakes/mocks pip_install_packages = MagicMock(name="pip_install_packages") run_command = MagicMock(name="run_command", return_value=0) + disabled_mock = MagicMock(name="_load_disabled_vllm_tests", return_value=[]) # temp_environ / working_directory: record calls but act as context managers temp_calls: list[dict] = [] @@ -51,6 +56,7 @@ def fake_temp_env(map: dict[str, str]): logger = SimpleNamespace( info=MagicMock(name="logger.info"), + warning=MagicMock(name="logger.warning"), error=MagicMock(name="logger.error"), ) @@ -64,6 +70,9 @@ def fake_temp_env(map: dict[str, str]): ) monkeypatch.setattr(module, "temp_environ", fake_temp_env, raising=True) monkeypatch.setattr(module, "logger", logger, raising=True) + monkeypatch.setattr( + module, "_load_disabled_vllm_tests", disabled_mock, raising=True + ) return SimpleNamespace( module=module, @@ -73,6 +82,7 @@ def fake_temp_env(map: dict[str, str]): temp_calls=temp_calls, workdir_calls=workdir_calls, logger=logger, + disabled_mock=disabled_mock, ) @@ -101,12 +111,10 @@ def test_success_runs_all_steps_and_uses_env_and_workdir(monkeypatch, patch_modu cmds = [_get_cmd(c) for c in calls] checks = [_get_check(c) for c in calls] - expected_cmds = [ - "export A=x && pytest -q", - "export B=y && pytest -q tests/unit", - ] - if cmds != expected_cmds: - raise AssertionError(f"Expected cmds={expected_cmds}, got cmds={cmds}") + if len(cmds) != 2: + raise AssertionError(f"Expected 2 commands, got {len(cmds)}: {cmds}") + if "pytest" not in cmds[0] or "pytest" not in cmds[1]: + raise AssertionError(f"Expected pytest in both commands, got {cmds}") if not all(chk is False for chk in checks): raise AssertionError(f"Expected all checks to be False, got checks={checks}") @@ -202,3 +210,313 @@ def test_custom_working_directory_used(patch_module): raise AssertionError( f"Expected workdir_calls=['examples/ci'], got {patch_module.workdir_calls}" ) + + +# -- Disabled vLLM test injection (integration tests) ------------------------- + +_SIMPLE_TESTS_MAP = { + "plan_a": { + "title": "Plan A", + "steps": ["pytest -v -s test_foo.py", "pytest -v -s test_bar.py"], + } +} + + +def _cmds(patch_module): + return [_get_cmd(c) for c in patch_module.run_command.call_args_list] + + +def test_deselect_injected_for_test_level(patch_module): + """Test-level node IDs (with ::) produce --deselect.""" + patch_module.disabled_mock.return_value = [ + { + "test": "test_foo.py::test_x", + "issue": "https://github.com/pytorch/pytorch/issues/175899", + } + ] + patch_module.run_test_plan("plan_a", "cpu", _SIMPLE_TESTS_MAP) + cmds = _cmds(patch_module) + if not any("--deselect=test_foo.py::test_x" in c for c in cmds): + raise AssertionError( + f"Expected --deselect=test_foo.py::test_x in commands, got {cmds}" + ) + + +def test_ignore_injected_for_file_level(patch_module): + """File-level node IDs (no ::) produce --ignore.""" + patch_module.disabled_mock.return_value = [ + { + "test": "test_foo.py", + "issue": "https://github.com/pytorch/pytorch/issues/175899", + } + ] + patch_module.run_test_plan("plan_a", "cpu", _SIMPLE_TESTS_MAP) + cmds = _cmds(patch_module) + if not any("--ignore=test_foo.py" in c for c in cmds): + raise AssertionError(f"Expected --ignore=test_foo.py in commands, got {cmds}") + + +def test_empty_disabled_list_no_modification(patch_module): + """No flags injected when disabled list is empty.""" + patch_module.run_test_plan("plan_a", "cpu", _SIMPLE_TESTS_MAP) + for cmd in _cmds(patch_module): + if "--ignore" in cmd or "--deselect" in cmd: + raise AssertionError(f"Expected no --ignore/--deselect flags, got {cmd}") + + +def test_non_pytest_steps_not_modified(patch_module): + """Only pytest steps get disabled flags; other commands are left alone.""" + tests_map = { + "mixed_steps": { + "title": "Mixed", + "steps": ["echo hello", "pytest -v test.py"], + } + } + patch_module.disabled_mock.return_value = [ + { + "test": "test.py::test_z", + "issue": "https://github.com/pytorch/pytorch/issues/175899", + } + ] + patch_module.run_test_plan("mixed_steps", "cpu", tests_map) + cmds = _cmds(patch_module) + if "--deselect" in cmds[0]: + raise AssertionError( + f"Non-pytest step should not have --deselect, got {cmds[0]}" + ) + if "--deselect=test.py::test_z" not in cmds[1]: + raise AssertionError( + f"Expected --deselect=test.py::test_z in pytest step, got {cmds[1]}" + ) + + +def test_disabled_and_rerun_flags_both_present(patch_module): + """Disabled flags and rerun flags compose into a single pytest invocation.""" + patch_module.disabled_mock.return_value = [ + { + "test": "skip.py::test_x", + "issue": "https://github.com/pytorch/pytorch/issues/175899", + } + ] + patch_module.run_test_plan("plan_a", "cpu", _SIMPLE_TESTS_MAP) + for cmd in _cmds(patch_module): + if "--deselect=skip.py::test_x" not in cmd: + raise AssertionError(f"Expected --deselect=skip.py::test_x, got {cmd}") + if "--reruns" not in cmd: + raise AssertionError(f"Expected --reruns in command, got {cmd}") + if cmd.count("pytest") != 1: + raise AssertionError(f"Expected exactly one 'pytest' token, got {cmd}") + + +# -- Disabled vLLM test helpers (unit tests) ----------------------------------- + + +class FakeResponse: + """Minimal fake for urllib.request.urlopen return value.""" + + def __init__(self, body): + self._data = json.dumps(body).encode() + + def read(self): + return self._data + + def __enter__(self): + return self + + def __exit__(self, *a): + pass + + +@pytest.fixture +def vllm_module(): + return importlib.import_module(MOD) + + +@pytest.mark.parametrize( + ("node_id", "expected_flag"), + ( + ("a.py::test_1", "--deselect=a.py::test_1"), + ("a.py", "--ignore=a.py"), + ), +) +def test_build_flags_single_entry(vllm_module, node_id, expected_flag): + """:: in node_id produces --deselect, otherwise --ignore.""" + flags = vllm_module._build_disabled_test_flags( + [{"test": node_id, "issue": "url"}], "plan" + ) + if flags != expected_flag: + raise AssertionError(f"Expected '{expected_flag}', got '{flags}'") + + +def test_build_flags_config_filter(vllm_module): + """Entries with matching configs are included; entries without configs apply to all.""" + entries = [ + {"test": "a.py", "issue": "url", "configs": ["plan_x"]}, + {"test": "b.py::test_1", "issue": "url"}, + ] + flags = vllm_module._build_disabled_test_flags(entries, "plan_x") + if "--ignore=a.py" not in flags: + raise AssertionError(f"Expected '--ignore=a.py' in '{flags}'") + if "--deselect=b.py::test_1" not in flags: + raise AssertionError(f"Expected '--deselect=b.py::test_1' in '{flags}'") + + +def test_build_flags_config_excludes(vllm_module): + """Entries with non-matching configs are excluded.""" + entries = [{"test": "a.py", "issue": "url", "configs": ["other"]}] + flags = vllm_module._build_disabled_test_flags(entries, "plan_x") + if flags != "": + raise AssertionError(f"Expected empty flags, got '{flags}'") + + +def test_parse_issue_body(vllm_module): + """Extracts disabled_tests from a ```yaml code block in issue body.""" + body = ( + "## Disabled vLLM Tests\n" + "```yaml\n" + "disabled_tests:\n" + " - test: foo.py::test_bar\n" + " - test: baz.py\n" + " configs:\n" + " - plan_a\n" + "```\n" + ) + entries = vllm_module._parse_disabled_tests_from_issue_body(body) + if len(entries) != 2: + raise AssertionError(f"Expected 2 entries, got {len(entries)}") + if entries[0]["test"] != "foo.py::test_bar": + raise AssertionError(f"Expected 'foo.py::test_bar', got '{entries[0]['test']}'") + if entries[1]["test"] != "baz.py": + raise AssertionError(f"Expected 'baz.py', got '{entries[1]['test']}'") + if entries[1]["configs"] != ["plan_a"]: + raise AssertionError(f"Expected ['plan_a'], got {entries[1]['configs']}") + + +def test_parse_issue_body_no_yaml(vllm_module): + """Returns [] when issue body has no yaml code block.""" + entries = vllm_module._parse_disabled_tests_from_issue_body("no yaml here") + if entries != []: + raise AssertionError(f"Expected [], got {entries}") + + +def test_load_yaml_valid_entries(tmp_path, vllm_module): + """Loads and validates entries from a well-formed YAML file.""" + yaml_file = tmp_path / "disabled.yaml" + yaml_file.write_text( + "disabled_tests:\n" + " - test: a.py\n" + " issue: https://github.com/pytorch/pytorch/issues/1\n" + " - test: b.py::test_x\n" + " issue: https://github.com/pytorch/pytorch/issues/2\n" + " configs:\n" + " - plan_a\n" + ) + with mock_patch.object(vllm_module, "_DISABLED_VLLM_TESTS_PATH", yaml_file): + entries = vllm_module._load_disabled_vllm_tests_from_yaml() + + if len(entries) != 2: + raise AssertionError(f"Expected 2 entries, got {len(entries)}") + if entries[0] != { + "test": "a.py", + "issue": "https://github.com/pytorch/pytorch/issues/175899", + }: + raise AssertionError(f"Unexpected first entry: {entries[0]}") + if entries[1]["test"] != "b.py::test_x": + raise AssertionError(f"Expected 'b.py::test_x', got '{entries[1]['test']}'") + if entries[1]["configs"] != ["plan_a"]: + raise AssertionError(f"Expected ['plan_a'], got {entries[1]['configs']}") + + +def test_load_yaml_missing_keys(tmp_path, vllm_module): + """Raises ValueError when YAML entry is missing required 'issue' key.""" + yaml_file = tmp_path / "disabled.yaml" + yaml_file.write_text("disabled_tests:\n - test: foo.py\n") + with mock_patch.object(vllm_module, "_DISABLED_VLLM_TESTS_PATH", yaml_file): + with pytest.raises(ValueError, match="must have 'test' and 'issue' keys"): + vllm_module._load_disabled_vllm_tests_from_yaml() + + +def test_load_yaml_missing_file(tmp_path, vllm_module): + """Returns [] when YAML file doesn't exist.""" + with mock_patch.object( + vllm_module, "_DISABLED_VLLM_TESTS_PATH", tmp_path / "nope.yaml" + ): + result = vllm_module._load_disabled_vllm_tests_from_yaml() + if result != []: + raise AssertionError(f"Expected [], got {result}") + + +def test_load_github_skipped_when_issue_unset(vllm_module): + """Returns [] immediately when _DISABLED_VLLM_TESTS_ISSUE is 0 (not configured).""" + with mock_patch.object(vllm_module, "_DISABLED_VLLM_TESTS_ISSUE", 0): + result = vllm_module._load_disabled_vllm_tests_from_github() + if result != []: + raise AssertionError(f"Expected [], got {result}") + + +def test_load_github_failure_returns_empty(vllm_module): + """Network errors are swallowed and return [].""" + + def _raise(*args, **kwargs): + raise OSError("network error") + + with ( + mock_patch.object(vllm_module, "_DISABLED_VLLM_TESTS_ISSUE", 175899), + mock_patch.object(vllm_module.urllib.request, "urlopen", _raise), + ): + result = vllm_module._load_disabled_vllm_tests_from_github() + if result != []: + raise AssertionError(f"Expected [], got {result}") + + +def test_load_github_filters_malformed_entries(vllm_module): + """Entries without 'test' key are filtered out; issue URL is auto-filled.""" + fake_resp = FakeResponse( + { + "body": "```yaml\ndisabled_tests:\n - bad_key: oops\n - test: good.py\n```", + "html_url": "https://github.com/pytorch/pytorch/issues/175899", + } + ) + with ( + mock_patch.object(vllm_module, "_DISABLED_VLLM_TESTS_ISSUE", 175899), + mock_patch.object( + vllm_module.urllib.request, "urlopen", lambda *a, **kw: fake_resp + ), + ): + entries = vllm_module._load_disabled_vllm_tests_from_github() + + if len(entries) != 1: + raise AssertionError(f"Expected 1 entry, got {len(entries)}") + if entries[0]["test"] != "good.py": + raise AssertionError(f"Expected 'good.py', got '{entries[0]['test']}'") + if entries[0]["issue"] != "https://github.com/pytorch/pytorch/issues/175899": + raise AssertionError(f"Expected issue URL, got '{entries[0]['issue']}'") + + +def test_deduplication_yaml_wins(vllm_module): + """YAML entries take precedence over GitHub entries with the same test key.""" + yaml_entries = [{"test": "a.py", "issue": "yaml-url"}] + github_entries = [ + {"test": "a.py", "issue": "github-url"}, + {"test": "b.py::test_1", "issue": "github-url"}, + ] + with ( + mock_patch.object( + vllm_module, + "_load_disabled_vllm_tests_from_yaml", + return_value=yaml_entries, + ), + mock_patch.object( + vllm_module, + "_load_disabled_vllm_tests_from_github", + return_value=github_entries, + ), + ): + merged = vllm_module._load_disabled_vllm_tests() + + if len(merged) != 2: + raise AssertionError(f"Expected 2 merged entries, got {len(merged)}") + if merged[0] != {"test": "a.py", "issue": "yaml-url"}: + raise AssertionError(f"Expected yaml entry to win, got {merged[0]}") + if merged[1] != {"test": "b.py::test_1", "issue": "github-url"}: + raise AssertionError(f"Unexpected second entry: {merged[1]}") diff --git a/.ci/magma/Makefile b/.ci/magma/Makefile index 4169aedd03fa5..848b567cb8992 100644 --- a/.ci/magma/Makefile +++ b/.ci/magma/Makefile @@ -16,6 +16,7 @@ DOCKER_RUN = set -eou pipefail; ${DOCKER_CMD} run --rm -i \ magma/build_magma.sh .PHONY: all +all: magma-cuda132 all: magma-cuda130 all: magma-cuda129 all: magma-cuda128 @@ -26,6 +27,12 @@ clean: $(RM) -r magma-* $(RM) -r output +.PHONY: magma-cuda132 +magma-cuda132: DESIRED_CUDA := 13.2 +magma-cuda132: CUDA_ARCH_LIST := -gencode arch=compute_80,code=sm_80 -gencode arch=compute_86,code=sm_86 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_100,code=sm_100 -gencode arch=compute_120,code=sm_120 +magma-cuda132: + $(DOCKER_RUN) + .PHONY: magma-cuda130 magma-cuda130: DESIRED_CUDA := 13.0 magma-cuda130: CUDA_ARCH_LIST := -gencode arch=compute_80,code=sm_80 -gencode arch=compute_86,code=sm_86 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_100,code=sm_100 -gencode arch=compute_120,code=sm_120 diff --git a/.ci/manywheel/build_common.sh b/.ci/manywheel/build_common.sh index d50bd623dace0..d76b75dea6850 100644 --- a/.ci/manywheel/build_common.sh +++ b/.ci/manywheel/build_common.sh @@ -91,11 +91,6 @@ export PYTORCH_BUILD_NUMBER=$build_number export CMAKE_LIBRARY_PATH="/opt/intel/lib:/lib:$CMAKE_LIBRARY_PATH" export CMAKE_INCLUDE_PATH="/opt/intel/include:$CMAKE_INCLUDE_PATH" -if [[ -e /opt/openssl ]]; then - export OPENSSL_ROOT_DIR=/opt/openssl - export CMAKE_INCLUDE_PATH="/opt/openssl/include":$CMAKE_INCLUDE_PATH -fi - mkdir -p /tmp/$WHEELHOUSE_DIR export PATCHELF_BIN=/usr/local/bin/patchelf @@ -118,6 +113,9 @@ retry pip install -qUr requirements-build.txt python setup.py clean retry pip install -qr requirements.txt case ${DESIRED_PYTHON} in + cp314*) + retry pip install -q --pre numpy==2.3.4 + ;; cp31*) retry pip install -q --pre numpy==2.1.0 ;; diff --git a/.ci/manywheel/build_cuda.sh b/.ci/manywheel/build_cuda.sh index 94bf6a6b4b26c..613301059f2a6 100644 --- a/.ci/manywheel/build_cuda.sh +++ b/.ci/manywheel/build_cuda.sh @@ -13,7 +13,7 @@ export ATEN_STATIC_CUDA=1 export USE_CUDA_STATIC_LINK=1 export INSTALL_TEST=0 # dont install test binaries into site-packages export USE_CUPTI_SO=0 -export USE_CUSPARSELT=${USE_CUSPARSELT:-1} # Enable if not disabled by libtorch build +export USE_CUSPARSELT=${USE_CUSPARSELT:-1} # Enable if not disabled by libtorch build. export USE_CUFILE=${USE_CUFILE:-1} export USE_SYSTEM_NCCL=1 export NCCL_INCLUDE_DIR="/usr/local/cuda/include/" @@ -109,13 +109,13 @@ TORCH_CUDA_ARCH_LIST="7.5;8.0;8.6;9.0;10.0" case ${CUDA_VERSION} in 12.6) TORCH_CUDA_ARCH_LIST="5.0;6.0;7.0;${TORCH_CUDA_ARCH_LIST//10.0/}" ;; # Only 12.6 includes legacy Maxwell/Pascal/Volta, -Hopper support 12.8) TORCH_CUDA_ARCH_LIST="${TORCH_CUDA_ARCH_LIST};12.0" ;; # +Blackwell support - 12.9) TORCH_CUDA_ARCH_LIST="${TORCH_CUDA_ARCH_LIST};12.0+PTX" # +Blackwell support + PTX for forward compatibility + 12.9) TORCH_CUDA_ARCH_LIST="${TORCH_CUDA_ARCH_LIST};12.0" # +Blackwell support + PTX for forward compatibility if [[ "$PACKAGE_TYPE" == "libtorch" ]]; then TORCH_CUDA_ARCH_LIST="${TORCH_CUDA_ARCH_LIST//8.6;/}" # Remove 8.6 for libtorch fi ;; - 13.0) - TORCH_CUDA_ARCH_LIST="${TORCH_CUDA_ARCH_LIST};$([[ "$ARCH" == "aarch64" ]] && echo "11.0;" || echo "")12.0+PTX" + 13.0|13.2) + TORCH_CUDA_ARCH_LIST="${TORCH_CUDA_ARCH_LIST};$([[ "$ARCH" == "aarch64" ]] && echo "11.0;" || echo "")12.0" export TORCH_NVCC_FLAGS="-compress-mode=size" export BUILD_BUNDLE_PTXAS=1 ;; diff --git a/.ci/manywheel/build_libtorch.sh b/.ci/manywheel/build_libtorch.sh index d78fbd5c3ed36..852ecf7500604 100644 --- a/.ci/manywheel/build_libtorch.sh +++ b/.ci/manywheel/build_libtorch.sh @@ -59,12 +59,6 @@ export PYTORCH_BUILD_NUMBER=$build_number export CMAKE_LIBRARY_PATH="/opt/intel/lib:/lib:$CMAKE_LIBRARY_PATH" export CMAKE_INCLUDE_PATH="/opt/intel/include:$CMAKE_INCLUDE_PATH" -# set OPENSSL_ROOT_DIR=/opt/openssl if it exists -if [[ -e /opt/openssl ]]; then - export OPENSSL_ROOT_DIR=/opt/openssl - export CMAKE_INCLUDE_PATH="/opt/openssl/include":$CMAKE_INCLUDE_PATH -fi - # If given a python version like 3.6m or 2.7mu, convert this to the format we # expect. The binary CI jobs pass in python versions like this; they also only # ever pass one python version, so we assume that DESIRED_PYTHON is not a list diff --git a/.ci/manywheel/build_rocm.sh b/.ci/manywheel/build_rocm.sh index bac56746f4501..fa5724dca25a7 100755 --- a/.ci/manywheel/build_rocm.sh +++ b/.ci/manywheel/build_rocm.sh @@ -97,20 +97,13 @@ ROCM_SO_FILES=( "libhipblaslt.so" "libhipsparselt.so" "libhiprtc.so" + "librocprofiler-sdk.so" + "librocprofiler-register.so" + "libhsa-amd-aqlprofile64.so" + "librocm-core.so" + "librocroller.so" ) -if [[ $ROCM_INT -ge 60100 ]]; then - ROCM_SO_FILES+=("librocprofiler-register.so") -fi - -if [[ $ROCM_INT -ge 60200 ]]; then - ROCM_SO_FILES+=("librocm-core.so") -fi - -if [[ $ROCM_INT -ge 70000 ]]; then - ROCM_SO_FILES+=("librocroller.so") -fi - OS_NAME=`awk -F= '/^NAME/{print $2}' /etc/os-release` if [[ "$OS_NAME" == *"CentOS Linux"* || "$OS_NAME" == *"AlmaLinux"* ]]; then LIBGOMP_PATH="/usr/lib64/libgomp.so.1" @@ -121,61 +114,23 @@ if [[ "$OS_NAME" == *"CentOS Linux"* || "$OS_NAME" == *"AlmaLinux"* ]]; then else LIBTINFO_PATH="/usr/lib64/libtinfo.so.6" fi + LIBDW_PATH="/usr/lib64/libdw.so.1" LIBDRM_PATH="/opt/amdgpu/lib64/libdrm.so.2" LIBDRM_AMDGPU_PATH="/opt/amdgpu/lib64/libdrm_amdgpu.so.1" - if [[ $ROCM_INT -ge 60100 && $ROCM_INT -lt 60300 ]]; then - # Below libs are direct dependencies of libhipsolver - LIBSUITESPARSE_CONFIG_PATH="/lib64/libsuitesparseconfig.so.4" - if [[ "$OS_NAME" == *"CentOS Linux"* ]]; then - LIBCHOLMOD_PATH="/lib64/libcholmod.so.2" - # Below libs are direct dependencies of libsatlas - LIBGFORTRAN_PATH="/lib64/libgfortran.so.3" - else - LIBCHOLMOD_PATH="/lib64/libcholmod.so.3" - # Below libs are direct dependencies of libsatlas - LIBGFORTRAN_PATH="/lib64/libgfortran.so.5" - fi - # Below libs are direct dependencies of libcholmod - LIBAMD_PATH="/lib64/libamd.so.2" - LIBCAMD_PATH="/lib64/libcamd.so.2" - LIBCCOLAMD_PATH="/lib64/libccolamd.so.2" - LIBCOLAMD_PATH="/lib64/libcolamd.so.2" - LIBSATLAS_PATH="/lib64/atlas/libsatlas.so.3" - # Below libs are direct dependencies of libsatlas - LIBQUADMATH_PATH="/lib64/libquadmath.so.0" - fi MAYBE_LIB64=lib64 elif [[ "$OS_NAME" == *"Ubuntu"* ]]; then LIBGOMP_PATH="/usr/lib/x86_64-linux-gnu/libgomp.so.1" LIBNUMA_PATH="/usr/lib/x86_64-linux-gnu/libnuma.so.1" LIBELF_PATH="/usr/lib/x86_64-linux-gnu/libelf.so.1" - if [[ $ROCM_INT -ge 50300 ]]; then - LIBTINFO_PATH="/lib/x86_64-linux-gnu/libtinfo.so.6" - else - LIBTINFO_PATH="/lib/x86_64-linux-gnu/libtinfo.so.5" - fi + LIBTINFO_PATH="/lib/x86_64-linux-gnu/libtinfo.so.6" + LIBDW_PATH="/usr/lib/x86_64-linux-gnu/libdw.so.1" LIBDRM_PATH="/usr/lib/x86_64-linux-gnu/libdrm.so.2" LIBDRM_AMDGPU_PATH="/usr/lib/x86_64-linux-gnu/libdrm_amdgpu.so.1" - if [[ $ROCM_INT -ge 60100 && $ROCM_INT -lt 60300 ]]; then - # Below libs are direct dependencies of libhipsolver - LIBCHOLMOD_PATH="/lib/x86_64-linux-gnu/libcholmod.so.3" - # Below libs are direct dependencies of libcholmod - LIBSUITESPARSE_CONFIG_PATH="/lib/x86_64-linux-gnu/libsuitesparseconfig.so.5" - LIBAMD_PATH="/lib/x86_64-linux-gnu/libamd.so.2" - LIBCAMD_PATH="/lib/x86_64-linux-gnu/libcamd.so.2" - LIBCCOLAMD_PATH="/lib/x86_64-linux-gnu/libccolamd.so.2" - LIBCOLAMD_PATH="/lib/x86_64-linux-gnu/libcolamd.so.2" - LIBMETIS_PATH="/lib/x86_64-linux-gnu/libmetis.so.5" - LIBLAPACK_PATH="/lib/x86_64-linux-gnu/liblapack.so.3" - LIBBLAS_PATH="/lib/x86_64-linux-gnu/libblas.so.3" - # Below libs are direct dependencies of libblas - LIBGFORTRAN_PATH="/lib/x86_64-linux-gnu/libgfortran.so.5" - LIBQUADMATH_PATH="/lib/x86_64-linux-gnu/libquadmath.so.0" - fi MAYBE_LIB64=lib fi OS_SO_PATHS=($LIBGOMP_PATH $LIBNUMA_PATH\ $LIBELF_PATH $LIBTINFO_PATH\ + $LIBDW_PATH\ $LIBDRM_PATH $LIBDRM_AMDGPU_PATH\ $LIBSUITESPARSE_CONFIG_PATH\ $LIBCHOLMOD_PATH $LIBAMD_PATH\ diff --git a/.ci/onnx/README.md b/.ci/onnx/README.md deleted file mode 100644 index 47739136aabdf..0000000000000 --- a/.ci/onnx/README.md +++ /dev/null @@ -1,12 +0,0 @@ -# Jenkins - -The scripts in this directory are the entrypoint for testing ONNX exporter. - -The environment variable `BUILD_ENVIRONMENT` is expected to be set to -the build environment you intend to test. It is a hint for the build -and test scripts to configure Caffe2 a certain way and include/exclude -tests. Docker images, they equal the name of the image itself. For -example: `py2-cuda9.0-cudnn7-ubuntu16.04`. The Docker images that are -built on Jenkins and are used in triggered builds already have this -environment variable set in their manifest. Also see -`./docker/jenkins/*/Dockerfile` and search for `BUILD_ENVIRONMENT`. diff --git a/.ci/onnx/common.sh b/.ci/onnx/common.sh deleted file mode 100644 index b8f912fbbb4e6..0000000000000 --- a/.ci/onnx/common.sh +++ /dev/null @@ -1,107 +0,0 @@ -#!/bin/bash - -set -ex - -source "$(dirname "${BASH_SOURCE[0]}")/../pytorch/common_utils.sh" - -LOCAL_DIR=$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd) -ROOT_DIR=$(cd "$LOCAL_DIR"/../.. && pwd) -TEST_DIR="$ROOT_DIR/test" -pytest_reports_dir="${TEST_DIR}/test-reports/python" - -# Figure out which Python to use -PYTHON="$(which python)" -if [[ "${BUILD_ENVIRONMENT}" =~ py((2|3)\.?[0-9]?\.?[0-9]?) ]]; then - PYTHON=$(which "python${BASH_REMATCH[1]}") -fi - -if [[ "${BUILD_ENVIRONMENT}" == *rocm* ]]; then - # HIP_PLATFORM is auto-detected by hipcc; unset to avoid build errors - unset HIP_PLATFORM -fi - -mkdir -p "$pytest_reports_dir" || true - -########################################## -# copied from .ci/pytorch/common_utils.sh -########################################## - -function get_pinned_commit() { - cat .github/ci_commit_pins/"${1}".txt -} - -function pip_install_whl() { - # This is used to install PyTorch and other build artifacts wheel locally - # without using any network connection - - # Convert the input arguments into an array - local args=("$@") - - # Check if the first argument contains multiple paths separated by spaces - if [[ "${args[0]}" == *" "* ]]; then - # Split the string by spaces into an array - IFS=' ' read -r -a paths <<< "${args[0]}" - # Loop through each path and install individually - for path in "${paths[@]}"; do - echo "Installing $path" - python3 -mpip install --no-index --no-deps "$path" - done - else - # Loop through each argument and install individually - for path in "${args[@]}"; do - echo "Installing $path" - python3 -mpip install --no-index --no-deps "$path" - done - fi -} - -function pip_build_and_install() { - local build_target=$1 - local wheel_dir=$2 - - local found_whl=0 - for file in "${wheel_dir}"/*.whl - do - if [[ -f "${file}" ]]; then - found_whl=1 - break - fi - done - - # Build the wheel if it doesn't exist - if [ "${found_whl}" == "0" ]; then - python3 -m pip wheel \ - --no-build-isolation \ - --no-deps \ - -w "${wheel_dir}" \ - "${build_target}" - fi - - for file in "${wheel_dir}"/*.whl - do - pip_install_whl "${file}" - done -} - -function install_torchvision() { - local orig_preload - local commit - commit=$(get_pinned_commit vision) - orig_preload=${LD_PRELOAD} - if [ -n "${LD_PRELOAD}" ]; then - # Silence dlerror to work-around glibc ASAN bug, see https://sourceware.org/bugzilla/show_bug.cgi?id=27653#c9 - echo 'char* dlerror(void) { return "";}'|gcc -fpic -shared -o "${HOME}/dlerror.so" -x c - - LD_PRELOAD=${orig_preload}:${HOME}/dlerror.so - fi - - if [[ "${BUILD_ENVIRONMENT}" == *cuda* ]]; then - # Not sure if both are needed, but why not - export FORCE_CUDA=1 - export WITH_CUDA=1 - fi - pip_build_and_install "git+https://github.com/pytorch/vision.git@${commit}" dist/vision - - if [ -n "${LD_PRELOAD}" ]; then - LD_PRELOAD=${orig_preload} - fi -} diff --git a/.ci/onnx/test.sh b/.ci/onnx/test.sh deleted file mode 100755 index 1f2a23b49dc45..0000000000000 --- a/.ci/onnx/test.sh +++ /dev/null @@ -1,29 +0,0 @@ -#!/bin/bash - -# shellcheck source=./common.sh -source "$(dirname "${BASH_SOURCE[0]}")/common.sh" - -# Workaround for dind-rootless userid mapping (https://github.com/pytorch/ci-infra/issues/96) -WORKSPACE_ORIGINAL_OWNER_ID=$(stat -c '%u' "/var/lib/jenkins/workspace") -cleanup_workspace() { - echo "sudo may print the following warning message that can be ignored. The chown command will still run." - echo " sudo: setrlimit(RLIMIT_STACK): Operation not permitted" - echo "For more details refer to https://github.com/sudo-project/sudo/issues/42" - sudo chown -R "$WORKSPACE_ORIGINAL_OWNER_ID" /var/lib/jenkins/workspace -} -# Disable shellcheck SC2064 as we want to parse the original owner immediately. -# shellcheck disable=SC2064 -trap_add cleanup_workspace EXIT -sudo chown -R jenkins /var/lib/jenkins/workspace -git config --global --add safe.directory /var/lib/jenkins/workspace - -if [[ "$BUILD_ENVIRONMENT" == *onnx* ]]; then - # TODO: This can be removed later once vision is also part of the Docker image - install_torchvision - # JIT C++ extensions require ninja, so put it into PATH. - export PATH="/var/lib/jenkins/.local/bin:$PATH" - # NB: ONNX test is fast (~15m) so it's ok to retry it few more times to avoid any flaky issue, we - # need to bring this to the standard PyTorch run_test eventually. The issue will be tracked in - # https://github.com/pytorch/pytorch/issues/98626 - "$ROOT_DIR/scripts/onnx/test.sh" -fi diff --git a/.circleci/scripts/binary_linux_test.sh b/.ci/pytorch/binary_linux_test.sh similarity index 98% rename from .circleci/scripts/binary_linux_test.sh rename to .ci/pytorch/binary_linux_test.sh index 3771ecc108f87..1180012458dcc 100755 --- a/.circleci/scripts/binary_linux_test.sh +++ b/.ci/pytorch/binary_linux_test.sh @@ -1,4 +1,5 @@ #!/bin/bash +# shellcheck disable=SC1091,SC2012,SC2154 OUTPUT_SCRIPT=${OUTPUT_SCRIPT:-/home/circleci/project/ci_test_script.sh} diff --git a/.circleci/scripts/binary_populate_env.sh b/.ci/pytorch/binary_populate_env.sh similarity index 98% rename from .circleci/scripts/binary_populate_env.sh rename to .ci/pytorch/binary_populate_env.sh index 74ad225db933b..53914914c8c93 100755 --- a/.circleci/scripts/binary_populate_env.sh +++ b/.ci/pytorch/binary_populate_env.sh @@ -1,4 +1,5 @@ #!/bin/bash +# shellcheck disable=SC2002,SC2004,SC2086,SC2129,SC2155 set -eux -o pipefail export TZ=UTC @@ -111,9 +112,8 @@ if [[ "$PACKAGE_TYPE" =~ .*wheel.* && -n "$PYTORCH_BUILD_VERSION" && "$PYTORCH_B fi fi -USE_GLOO_WITH_OPENSSL="ON" +USE_GLOO_WITH_OPENSSL="OFF" if [[ "$GPU_ARCH_TYPE" =~ .*aarch64.* ]]; then - USE_GLOO_WITH_OPENSSL="OFF" USE_GOLD_LINKER="OFF" fi diff --git a/.circleci/scripts/binary_upload.sh b/.ci/pytorch/binary_upload.sh similarity index 62% rename from .circleci/scripts/binary_upload.sh rename to .ci/pytorch/binary_upload.sh index d48077e112455..a0ca32be72852 100755 --- a/.circleci/scripts/binary_upload.sh +++ b/.ci/pytorch/binary_upload.sh @@ -1,5 +1,5 @@ #!/usr/bin/env bash - +# shellcheck disable=SC2231 set -euo pipefail PACKAGE_TYPE=${PACKAGE_TYPE:-wheel} @@ -62,17 +62,61 @@ s3_upload() { ) } +R2_UPLOAD=${R2_UPLOAD:-} +R2_BUCKET="s3://pytorch-downloads" +R2_ACCOUNT_ID=${R2_ACCOUNT_ID:-} +R2_ACCESS_KEY_ID=${R2_ACCESS_KEY_ID:-} +R2_SECRET_ACCESS_KEY=${R2_SECRET_ACCESS_KEY:-} + +r2_upload() { + if [[ -z "${R2_ACCOUNT_ID}" || -z "${R2_ACCESS_KEY_ID}" || -z "${R2_SECRET_ACCESS_KEY}" ]]; then + echo "WARNING: R2 credentials not configured, skipping R2 upload" + return + fi + local extension + local pkg_type + extension="$1" + pkg_type="$2" + r2_root_dir="${R2_BUCKET}/${pkg_type}/${UPLOAD_CHANNEL}" + if [[ -z ${UPLOAD_SUBFOLDER:-} ]]; then + r2_upload_dir="${r2_root_dir}/" + else + r2_upload_dir="${r2_root_dir}/${UPLOAD_SUBFOLDER}/" + fi + ( + for pkg in ${PKG_DIR}/*.${extension}; do + ( + set -x + shm_id=$(sha256sum "${pkg}" | awk '{print $1}') + AWS_ACCESS_KEY_ID="${R2_ACCESS_KEY_ID}" \ + AWS_SECRET_ACCESS_KEY="${R2_SECRET_ACCESS_KEY}" \ + AWS_SESSION_TOKEN="" \ + AWS_DEFAULT_REGION="auto" \ + ${AWS_S3_CP} --no-progress "${pkg}" "${r2_upload_dir}" \ + --metadata "checksum-sha256=${shm_id}" \ + --endpoint-url "https://${R2_ACCOUNT_ID}.r2.cloudflarestorage.com" + ) + done + ) +} + # Install dependencies (should be a no-op if previously installed) pip install -q awscli uv case "${PACKAGE_TYPE}" in libtorch) s3_upload "zip" "libtorch" + if [[ "${R2_UPLOAD}" == "true" ]]; then + r2_upload "zip" "libtorch" + fi BACKUP_DIR="libtorch/${UPLOAD_CHANNEL}/${UPLOAD_SUBFOLDER}" ;; # wheel can either refer to wheel/manywheel *wheel) s3_upload "whl" "whl" + if [[ "${R2_UPLOAD}" == "true" ]]; then + r2_upload "whl" "whl" + fi BACKUP_DIR="whl/${UPLOAD_CHANNEL}/${UPLOAD_SUBFOLDER}" ;; *) diff --git a/.circleci/scripts/binary_windows_build.sh b/.ci/pytorch/binary_windows_build.sh similarity index 97% rename from .circleci/scripts/binary_windows_build.sh rename to .ci/pytorch/binary_windows_build.sh index 59dbbb3d9b6a8..2f560233b1201 100644 --- a/.circleci/scripts/binary_windows_build.sh +++ b/.ci/pytorch/binary_windows_build.sh @@ -1,4 +1,5 @@ #!/bin/bash +# shellcheck disable=SC1090 set -eux -o pipefail source "${BINARY_ENV_FILE:-/c/w/env}" diff --git a/.circleci/scripts/binary_windows_test.sh b/.ci/pytorch/binary_windows_test.sh similarity index 93% rename from .circleci/scripts/binary_windows_test.sh rename to .ci/pytorch/binary_windows_test.sh index b8b82979caf48..f64a9447193bb 100644 --- a/.circleci/scripts/binary_windows_test.sh +++ b/.ci/pytorch/binary_windows_test.sh @@ -1,4 +1,5 @@ #!/bin/bash +# shellcheck disable=SC1090 set -eux -o pipefail source "${BINARY_ENV_FILE:-/c/w/env}" diff --git a/.ci/pytorch/build.sh b/.ci/pytorch/build.sh index eb3529a0c43f3..265a6f579f427 100755 --- a/.ci/pytorch/build.sh +++ b/.ci/pytorch/build.sh @@ -23,12 +23,6 @@ cmake --version echo "Environment variables:" env -# The sccache wrapped version of nvcc gets put in /opt/cache/lib in docker since -# there are some issues if it is always wrapped, so we need to add it to PATH -# during CI builds. -# https://github.com/pytorch/pytorch/blob/0b6c0898e6c352c8ea93daec854e704b41485375/.ci/docker/common/install_cache.sh#L97 -export PATH="/opt/cache/lib:$PATH" - if [[ "$BUILD_ENVIRONMENT" == *cuda* ]]; then # Use jemalloc during compilation to mitigate https://github.com/pytorch/pytorch/issues/116289 export LD_PRELOAD=/usr/lib/x86_64-linux-gnu/libjemalloc.so.2 @@ -41,16 +35,6 @@ if [[ "$BUILD_ENVIRONMENT" == *cuda13* ]]; then export USE_FBGEMM=0 fi -if [[ "$BUILD_ENVIRONMENT" == *cuda11* ]]; then - if [[ "$BUILD_ENVIRONMENT" != *clang* ]]; then - # TODO: there is a linking issue when building with UCC using clang, - # disable it for now and to be fix later. - # TODO: disable UCC temporarily to enable CUDA 12.1 in CI - export USE_UCC=1 - export USE_SYSTEM_UCC=1 - fi -fi - if [[ ${BUILD_ENVIRONMENT} == *"parallelnative"* ]]; then export ATEN_THREADING=NATIVE fi @@ -118,30 +102,6 @@ if [[ "$BUILD_ENVIRONMENT" == *riscv64* ]]; then fi -if [[ "$BUILD_ENVIRONMENT" == *libtorch* ]]; then - POSSIBLE_JAVA_HOMES=() - POSSIBLE_JAVA_HOMES+=(/usr/local) - POSSIBLE_JAVA_HOMES+=(/usr/lib/jvm/java-8-openjdk-amd64) - POSSIBLE_JAVA_HOMES+=(/Library/Java/JavaVirtualMachines/*.jdk/Contents/Home) - # Add the Windows-specific JNI - POSSIBLE_JAVA_HOMES+=("$PWD/.circleci/windows-jni/") - for JH in "${POSSIBLE_JAVA_HOMES[@]}" ; do - if [[ -e "$JH/include/jni.h" ]] ; then - # Skip if we're not on Windows but haven't found a JAVA_HOME - if [[ "$JH" == "$PWD/.circleci/windows-jni/" && "$OSTYPE" != "msys" ]] ; then - break - fi - echo "Found jni.h under $JH" - export JAVA_HOME="$JH" - export BUILD_JNI=ON - break - fi - done - if [ -z "$JAVA_HOME" ]; then - echo "Did not find jni.h" - fi -fi - # Use special scripts for Android builds if [[ "$BUILD_ENVIRONMENT" == *vulkan* ]]; then @@ -184,6 +144,7 @@ if [[ "$BUILD_ENVIRONMENT" == *xpu* ]]; then export USE_XCCL=1 export USE_MPI=0 export TORCH_XPU_ARCH_LIST=pvc + export USE_STATIC_MKL=1 fi # sccache will fail for CUDA builds if all cores are used for compiling @@ -214,9 +175,9 @@ if [[ "$BUILD_ENVIRONMENT" == *cuda* ]] && echo "${TORCH_CUDA_ARCH_LIST}" | tr ' export BUILD_CUSTOM_STEP="ninja -C build flash_attention -j ${J}" fi -if [[ "${BUILD_ENVIRONMENT}" == *clang* ]]; then - export CC=clang - export CXX=clang++ +# TODO: Removeme once all the wrappers are gone +if [[ "$BUILD_ENVIRONMENT" == *clang* ]] && [[ "$BUILD_ENVIRONMENT" == *cuda* ]]; then + sudo rm -f /opt/cache/bin/clang++ fi if [[ "$BUILD_ENVIRONMENT" == *-clang*-asan* ]]; then @@ -232,10 +193,6 @@ if [[ "${BUILD_ENVIRONMENT}" == *no-ops* ]]; then export USE_PER_OPERATOR_HEADERS=0 fi -if [[ "${BUILD_ENVIRONMENT}" == *-pch* ]]; then - export USE_PRECOMPILED_HEADERS=1 -fi - if [[ "${BUILD_ENVIRONMENT}" != *cuda* ]]; then export BUILD_STATIC_RUNTIME_BENCHMARK=ON fi @@ -264,187 +221,168 @@ if [[ "$BUILD_ENVIRONMENT" != *rocm* && "$BUILD_ENVIRONMENT" != *s390x* && "$BUI git config --global --add safe.directory /var/lib/jenkins/workspace fi -if [[ "$BUILD_ENVIRONMENT" == *-bazel-* ]]; then - set -e -o pipefail - - get_bazel - python3 tools/optional_submodules.py checkout_eigen +# check that setup.py would fail with bad arguments +echo "The next three invocations are expected to fail with invalid command error messages." +( ! get_exit_code python setup.py bad_argument ) +( ! get_exit_code python setup.py clean] ) +( ! get_exit_code python setup.py clean bad_argument ) + +if [[ "$BUILD_ENVIRONMENT" != *libtorch* ]]; then + # rocm builds fail when WERROR=1 + # XLA test build fails when WERROR=1 + # set only when building other architectures + # or building non-XLA tests. + if [[ "$BUILD_ENVIRONMENT" != *rocm* && "$BUILD_ENVIRONMENT" != *xla* && "$BUILD_ENVIRONMENT" != *riscv64* ]]; then + # TODO: Remove me and may be just focus on numpy-2.x testing + if [[ "$ANACONDA_PYTHON_VERSION" =~ ^3\.1[0-2]$ ]]; then + # Install numpy-2.0.2 for builds which are backward compatible with 1.X + # In relality it's only needed for numpy_2_x and vllm shards (where vllm depends on numpy-2) + python -mpip install numpy==2.0.2 + fi - # Leave 1 CPU free and use only up to 80% of memory to reduce the change of crashing - # the runner - BAZEL_MEM_LIMIT="--local_ram_resources=HOST_RAM*.8" - BAZEL_CPU_LIMIT="--local_cpu_resources=HOST_CPUS-1" + WERROR=1 python setup.py clean - if [[ "$CUDA_VERSION" == "cpu" ]]; then - # Build torch, the Python module, and tests for CPU-only - tools/bazel build --config=no-tty "${BAZEL_MEM_LIMIT}" "${BAZEL_CPU_LIMIT}" --config=cpu-only :torch :torch/_C.so :all_tests + WERROR=1 python -m build --wheel --no-isolation else - tools/bazel build --config=no-tty "${BAZEL_MEM_LIMIT}" "${BAZEL_CPU_LIMIT}" //... - fi -else - # check that setup.py would fail with bad arguments - echo "The next three invocations are expected to fail with invalid command error messages." - ( ! get_exit_code python setup.py bad_argument ) - ( ! get_exit_code python setup.py clean] ) - ( ! get_exit_code python setup.py clean bad_argument ) - - if [[ "$BUILD_ENVIRONMENT" != *libtorch* ]]; then - # rocm builds fail when WERROR=1 - # XLA test build fails when WERROR=1 - # set only when building other architectures - # or building non-XLA tests. - if [[ "$BUILD_ENVIRONMENT" != *rocm* && "$BUILD_ENVIRONMENT" != *xla* && "$BUILD_ENVIRONMENT" != *riscv64* ]]; then - # TODO: Remove me and may be just focus on numpy-2.x testing - if [[ "$ANACONDA_PYTHON_VERSION" =~ ^3\.1[0-2]$ ]]; then - # Install numpy-2.0.2 for builds which are backward compatible with 1.X - # In relality it's only needed for numpy_2_x and vllm shards (where vllm depends on numpy-2) - python -mpip install numpy==2.0.2 - fi - - WERROR=1 python setup.py clean - - WERROR=1 python -m build --wheel --no-isolation - else - python setup.py clean - if [[ "$BUILD_ENVIRONMENT" == *xla* ]]; then - source .ci/pytorch/install_cache_xla.sh - fi - python -m build --wheel --no-isolation - fi - pip_install_whl "$(echo dist/*.whl)" - if [[ "$BUILD_ENVIRONMENT" == *full-debug* ]]; then - # Regression test for https://github.com/pytorch/pytorch/issues/164297 - # Torch should be importable and that's about it - pushd /; python -c "import torch;print(torch.__config__.show(), torch.randn(5) + 1.7)"; popd + python setup.py clean + if [[ "$BUILD_ENVIRONMENT" == *xla* ]]; then + source .ci/pytorch/install_cache_xla.sh fi + python -m build --wheel --no-isolation + fi + pip_install_whl "$(echo dist/*.whl)" + if [[ "$BUILD_ENVIRONMENT" == *full-debug* ]]; then + # Regression test for https://github.com/pytorch/pytorch/issues/164297 + # Torch should be importable and that's about it + pushd /; python -c "import torch;print(torch.__config__.show(), torch.randn(5) + 1.7)"; popd + fi - if [[ "${BUILD_ADDITIONAL_PACKAGES:-}" == *vision* ]]; then - install_torchvision - fi + if [[ "${BUILD_ADDITIONAL_PACKAGES:-}" == *vision* ]]; then + install_torchvision + fi - if [[ "${BUILD_ADDITIONAL_PACKAGES:-}" == *audio* ]]; then - install_torchaudio - fi + if [[ "${BUILD_ADDITIONAL_PACKAGES:-}" == *audio* ]]; then + install_torchaudio + fi - if [[ "${BUILD_ADDITIONAL_PACKAGES:-}" == *torchrec* || "${BUILD_ADDITIONAL_PACKAGES:-}" == *fbgemm* ]]; then - install_torchrec_and_fbgemm - fi + if [[ "${BUILD_ADDITIONAL_PACKAGES:-}" == *torchrec* || "${BUILD_ADDITIONAL_PACKAGES:-}" == *fbgemm* ]]; then + install_torchrec_and_fbgemm + fi - if [[ "${BUILD_ADDITIONAL_PACKAGES:-}" == *torchao* ]]; then - install_torchao - fi + if [[ "${BUILD_ADDITIONAL_PACKAGES:-}" == *torchao* ]]; then + install_torchao + fi - if [[ "$BUILD_ENVIRONMENT" == *xpu* ]]; then - echo "Checking that xpu is compiled" - pushd dist/ - if python -c 'import torch; exit(0 if torch.xpu._is_compiled() else 1)'; then - echo "XPU support is compiled in." - else - echo "XPU support is NOT compiled in." - exit 1 - fi - popd + if [[ "$BUILD_ENVIRONMENT" == *xpu* ]]; then + echo "Checking that xpu is compiled" + pushd dist/ + if python -c 'import torch; exit(0 if torch.xpu._is_compiled() else 1)'; then + echo "XPU support is compiled in." + else + echo "XPU support is NOT compiled in." + exit 1 fi + popd + fi - # TODO: I'm not sure why, but somehow we lose verbose commands - set -x + # TODO: I'm not sure why, but somehow we lose verbose commands + set -x - assert_git_not_dirty - # Copy ninja build logs to dist folder - mkdir -p dist - if [ -f build/.ninja_log ]; then - cp build/.ninja_log dist - fi + assert_git_not_dirty + # Copy ninja build logs to dist folder + mkdir -p dist + if [ -f build/.ninja_log ]; then + cp build/.ninja_log dist + fi - if [[ "$BUILD_ENVIRONMENT" == *rocm* ]]; then - # remove sccache wrappers post-build; runtime compilation of MIOpen kernels does not yet fully support them - sudo rm -f /opt/cache/bin/cc - sudo rm -f /opt/cache/bin/c++ - sudo rm -f /opt/cache/bin/gcc - sudo rm -f /opt/cache/bin/g++ - # Restore original clang compilers that were backed up during sccache wrapping. - # Skip for theRock nightly: sccache wrapping is disabled, so no backup exists. - # theRock also uses ${ROCM_PATH}/lib/llvm/bin instead of /opt/rocm/llvm/bin. - if [[ -d /opt/rocm/llvm/bin ]]; then - pushd /opt/rocm/llvm/bin - if [[ -d original ]]; then - sudo mv original/clang . - sudo mv original/clang++ . - fi - sudo rm -rf original - popd + if [[ "$BUILD_ENVIRONMENT" == *rocm* ]]; then + # remove sccache wrappers post-build; runtime compilation of MIOpen kernels does not yet fully support them + sudo rm -f /opt/cache/bin/cc + sudo rm -f /opt/cache/bin/c++ + sudo rm -f /opt/cache/bin/gcc + sudo rm -f /opt/cache/bin/g++ + # Restore original clang compilers that were backed up during sccache wrapping. + # Skip for theRock nightly: sccache wrapping is disabled, so no backup exists. + # theRock also uses ${ROCM_PATH}/lib/llvm/bin instead of /opt/rocm/llvm/bin. + if [[ -d /opt/rocm/llvm/bin ]]; then + pushd /opt/rocm/llvm/bin + if [[ -d original ]]; then + sudo mv original/clang . + sudo mv original/clang++ . fi + sudo rm -rf original + popd fi - - CUSTOM_TEST_ARTIFACT_BUILD_DIR=${CUSTOM_TEST_ARTIFACT_BUILD_DIR:-"build/custom_test_artifacts"} - CUSTOM_TEST_USE_ROCM=$([[ "$BUILD_ENVIRONMENT" == *rocm* ]] && echo "ON" || echo "OFF") - CUSTOM_TEST_MODULE_PATH="${PWD}/cmake/public" - mkdir -pv "${CUSTOM_TEST_ARTIFACT_BUILD_DIR}" - - # Build custom operator tests. - CUSTOM_OP_BUILD="${CUSTOM_TEST_ARTIFACT_BUILD_DIR}/custom-op-build" - CUSTOM_OP_TEST="$PWD/test/custom_operator" - python --version - SITE_PACKAGES="$(python -c 'import site; print(";".join([x for x in site.getsitepackages()] + [x + "/torch" for x in site.getsitepackages()]))')" - - mkdir -p "$CUSTOM_OP_BUILD" - pushd "$CUSTOM_OP_BUILD" - cmake "$CUSTOM_OP_TEST" -DCMAKE_PREFIX_PATH="$SITE_PACKAGES" -DPython_EXECUTABLE="$(which python)" \ - -DCMAKE_MODULE_PATH="$CUSTOM_TEST_MODULE_PATH" -DUSE_ROCM="$CUSTOM_TEST_USE_ROCM" - make VERBOSE=1 - popd - assert_git_not_dirty - - # Build jit hook tests - JIT_HOOK_BUILD="${CUSTOM_TEST_ARTIFACT_BUILD_DIR}/jit-hook-build" - JIT_HOOK_TEST="$PWD/test/jit_hooks" - python --version - SITE_PACKAGES="$(python -c 'import site; print(";".join([x for x in site.getsitepackages()] + [x + "/torch" for x in site.getsitepackages()]))')" - mkdir -p "$JIT_HOOK_BUILD" - pushd "$JIT_HOOK_BUILD" - cmake "$JIT_HOOK_TEST" -DCMAKE_PREFIX_PATH="$SITE_PACKAGES" -DPython_EXECUTABLE="$(which python)" \ - -DCMAKE_MODULE_PATH="$CUSTOM_TEST_MODULE_PATH" -DUSE_ROCM="$CUSTOM_TEST_USE_ROCM" - make VERBOSE=1 - popd - assert_git_not_dirty - - # Build custom backend tests. - CUSTOM_BACKEND_BUILD="${CUSTOM_TEST_ARTIFACT_BUILD_DIR}/custom-backend-build" - CUSTOM_BACKEND_TEST="$PWD/test/custom_backend" - python --version - mkdir -p "$CUSTOM_BACKEND_BUILD" - pushd "$CUSTOM_BACKEND_BUILD" - cmake "$CUSTOM_BACKEND_TEST" -DCMAKE_PREFIX_PATH="$SITE_PACKAGES" -DPython_EXECUTABLE="$(which python)" \ - -DCMAKE_MODULE_PATH="$CUSTOM_TEST_MODULE_PATH" -DUSE_ROCM="$CUSTOM_TEST_USE_ROCM" - make VERBOSE=1 - popd - assert_git_not_dirty - else - # Test no-Python build - echo "Building libtorch" - - # This is an attempt to mitigate flaky libtorch build OOM error. By default, the build parallelization - # is set to be the number of CPU minus 2. So, let's try a more conservative value here. A 4xlarge has - # 16 CPUs - MAX_JOBS=$(nproc --ignore=4) - export MAX_JOBS - - # NB: Install outside of source directory (at the same level as the root - # pytorch folder) so that it doesn't get cleaned away prior to docker push. - BUILD_LIBTORCH_PY=$PWD/tools/build_libtorch.py - mkdir -p ../cpp-build/caffe2 - pushd ../cpp-build/caffe2 - WERROR=1 VERBOSE=1 DEBUG=1 python "$BUILD_LIBTORCH_PY" - popd fi + + CUSTOM_TEST_ARTIFACT_BUILD_DIR=${CUSTOM_TEST_ARTIFACT_BUILD_DIR:-"build/custom_test_artifacts"} + CUSTOM_TEST_USE_ROCM=$([[ "$BUILD_ENVIRONMENT" == *rocm* ]] && echo "ON" || echo "OFF") + CUSTOM_TEST_MODULE_PATH="${PWD}/cmake/public" + mkdir -pv "${CUSTOM_TEST_ARTIFACT_BUILD_DIR}" + + # Build custom operator tests. + CUSTOM_OP_BUILD="${CUSTOM_TEST_ARTIFACT_BUILD_DIR}/custom-op-build" + CUSTOM_OP_TEST="$PWD/test/custom_operator" + python --version + SITE_PACKAGES="$(python -c 'import site; print(";".join([x for x in site.getsitepackages()] + [x + "/torch" for x in site.getsitepackages()]))')" + + mkdir -p "$CUSTOM_OP_BUILD" + pushd "$CUSTOM_OP_BUILD" + cmake "$CUSTOM_OP_TEST" -DCMAKE_PREFIX_PATH="$SITE_PACKAGES" -DPython_EXECUTABLE="$(which python)" \ + -DCMAKE_MODULE_PATH="$CUSTOM_TEST_MODULE_PATH" -DUSE_ROCM="$CUSTOM_TEST_USE_ROCM" + make VERBOSE=1 + popd + assert_git_not_dirty + + # Build jit hook tests + JIT_HOOK_BUILD="${CUSTOM_TEST_ARTIFACT_BUILD_DIR}/jit-hook-build" + JIT_HOOK_TEST="$PWD/test/jit_hooks" + python --version + SITE_PACKAGES="$(python -c 'import site; print(";".join([x for x in site.getsitepackages()] + [x + "/torch" for x in site.getsitepackages()]))')" + mkdir -p "$JIT_HOOK_BUILD" + pushd "$JIT_HOOK_BUILD" + cmake "$JIT_HOOK_TEST" -DCMAKE_PREFIX_PATH="$SITE_PACKAGES" -DPython_EXECUTABLE="$(which python)" \ + -DCMAKE_MODULE_PATH="$CUSTOM_TEST_MODULE_PATH" -DUSE_ROCM="$CUSTOM_TEST_USE_ROCM" + make VERBOSE=1 + popd + assert_git_not_dirty + + # Build custom backend tests. + CUSTOM_BACKEND_BUILD="${CUSTOM_TEST_ARTIFACT_BUILD_DIR}/custom-backend-build" + CUSTOM_BACKEND_TEST="$PWD/test/custom_backend" + python --version + mkdir -p "$CUSTOM_BACKEND_BUILD" + pushd "$CUSTOM_BACKEND_BUILD" + cmake "$CUSTOM_BACKEND_TEST" -DCMAKE_PREFIX_PATH="$SITE_PACKAGES" -DPython_EXECUTABLE="$(which python)" \ + -DCMAKE_MODULE_PATH="$CUSTOM_TEST_MODULE_PATH" -DUSE_ROCM="$CUSTOM_TEST_USE_ROCM" + make VERBOSE=1 + popd + assert_git_not_dirty +else + # Test no-Python build + echo "Building libtorch" + + # This is an attempt to mitigate flaky libtorch build OOM error. By default, the build parallelization + # is set to be the number of CPU minus 2. So, let's try a more conservative value here. A 4xlarge has + # 16 CPUs + MAX_JOBS=$(nproc --ignore=4) + export MAX_JOBS + + BUILD_LIBTORCH_PY=$PWD/tools/build_libtorch.py + # Build outside the source tree so the artifacts don't interfere with + # the workspace. /tmp is writable on both EC2 and OSDC runners. + mkdir -p /tmp/cpp-build/caffe2 + pushd /tmp/cpp-build/caffe2 + WERROR=1 VERBOSE=1 DEBUG=1 python "$BUILD_LIBTORCH_PY" + popd fi -if [[ "$BUILD_ENVIRONMENT" != *libtorch* && "$BUILD_ENVIRONMENT" != *bazel* ]]; then +if [[ "$BUILD_ENVIRONMENT" != *libtorch* ]]; then # export test times so that potential sharded tests that'll branch off this build will use consistent data # don't do this for libtorch as libtorch is C++ only and thus won't have python tests run on its build PYTHONPATH=. python tools/stats/export_test_times.py fi -# don't do this for bazel or s390x or riscv64 as they don't use sccache -if [[ "$BUILD_ENVIRONMENT" != *s390x* && "$BUILD_ENVIRONMENT" != *riscv64* && "$BUILD_ENVIRONMENT" != *-bazel-* ]]; then +# don't do this for s390x or riscv64 as they don't use sccache +if [[ "$BUILD_ENVIRONMENT" != *s390x* && "$BUILD_ENVIRONMENT" != *riscv64* ]]; then print_sccache_stats fi diff --git a/.ci/pytorch/check_binary.sh b/.ci/pytorch/check_binary.sh index c8c89fe871fe3..9356970394e81 100755 --- a/.ci/pytorch/check_binary.sh +++ b/.ci/pytorch/check_binary.sh @@ -154,7 +154,7 @@ setup_link_flags () { TEST_CODE_DIR="$(dirname $(realpath ${BASH_SOURCE[0]}))/test_example_code" build_and_run_example_cpp () { setup_link_flags - g++ ${TEST_CODE_DIR}/$1.cpp -I${install_root}/include -I${install_root}/include/torch/csrc/api/include -std=gnu++17 -L${install_root}/lib ${REF_LIB} ${ADDITIONAL_LINKER_FLAGS} -ltorch $TORCH_CPU_LINK_FLAGS $TORCH_CUDA_LINK_FLAGS $C10_LINK_FLAGS -o $1 + g++ ${TEST_CODE_DIR}/$1.cpp -I${install_root}/include -I${install_root}/include/torch/csrc/api/include -std=gnu++20 -L${install_root}/lib ${REF_LIB} ${ADDITIONAL_LINKER_FLAGS} -ltorch $TORCH_CPU_LINK_FLAGS $TORCH_CUDA_LINK_FLAGS $C10_LINK_FLAGS -o $1 ./$1 } @@ -292,25 +292,6 @@ if [[ "$PACKAGE_TYPE" != 'libtorch' ]]; then popd fi -############################################################################### -# Check PyTorch supports TCP_TLS gloo transport -############################################################################### - -if [[ "$(uname)" == 'Linux' && "$PACKAGE_TYPE" != 'libtorch' ]]; then - GLOO_CHECK="import torch.distributed as dist -try: - dist.init_process_group('gloo', rank=0, world_size=1) -except RuntimeError as e: - print(e) -" - RESULT=`GLOO_DEVICE_TRANSPORT=TCP_TLS MASTER_ADDR=localhost MASTER_PORT=63945 python -c "$GLOO_CHECK"` - GLOO_TRANSPORT_IS_NOT_SUPPORTED='gloo transport is not supported' - if [[ "$RESULT" =~ "$GLOO_TRANSPORT_IS_NOT_SUPPORTED" ]]; then - echo "PyTorch doesn't support TLS_TCP transport, please build with USE_GLOO_WITH_OPENSSL=1" - exit 1 - fi -fi - ############################################################################### # Restore LD_LIBRARY_PATH to its original value ############################################################################### diff --git a/.ci/pytorch/common-build.sh b/.ci/pytorch/common-build.sh index 8ca9fdb34c77a..5ccf317711553 100644 --- a/.ci/pytorch/common-build.sh +++ b/.ci/pytorch/common-build.sh @@ -6,12 +6,6 @@ if [[ "$BUILD_ENVIRONMENT" != *win-* ]]; then # Save the absolute path in case later we chdir (as occurs in the gpu perf test) script_dir="$( cd "$(dirname "${BASH_SOURCE[0]}")" || exit ; pwd -P )" - if [[ "${BUILD_ENVIRONMENT}" == *-pch* ]]; then - # This is really weird, but newer sccache somehow produces broken binary - # see https://github.com/pytorch/pytorch/issues/139188 - sudo mv /opt/cache/bin/sccache-0.2.14a /opt/cache/bin/sccache - fi - if which sccache > /dev/null; then # Clear SCCACHE_BUCKET and SCCACHE_REGION if they are empty, otherwise # sccache will complain about invalid bucket configuration diff --git a/.ci/pytorch/common.sh b/.ci/pytorch/common.sh index 072b8da9b10c6..94d9629eac519 100644 --- a/.ci/pytorch/common.sh +++ b/.ci/pytorch/common.sh @@ -5,8 +5,8 @@ source "$(dirname "${BASH_SOURCE[0]}")/common_utils.sh" set -ex -o pipefail -# for ROCm environment variables -if [[ "${BUILD_ENVIRONMENT}" == *rocm* ]]; then +# Source ROCm environment variables (paths may vary between tarball/wheel installs) +if [[ "${BUILD_ENVIRONMENT}" == *rocm* ]] && [[ -f /etc/rocm_env.sh ]]; then # shellcheck disable=SC1091 source /etc/rocm_env.sh fi @@ -14,6 +14,19 @@ fi # Required environment variables: # $BUILD_ENVIRONMENT (should be set by your Docker image) +# Select compiler based on build environment name. Images that have both +# GCC and Clang installed default cc/c++ to Clang (via install_clang.sh), +# so we need to override when a gcc build is requested. +if [[ "${BUILD_ENVIRONMENT}" == *clang* ]]; then + export CC=clang + export CXX=clang++ +elif [[ "${BUILD_ENVIRONMENT}" == *gcc* ]]; then + export CC=gcc + export CXX=g++ + sudo update-alternatives --install /usr/bin/cc cc /usr/bin/gcc 100 + sudo update-alternatives --install /usr/bin/c++ c++ /usr/bin/g++ 100 +fi + # Figure out which Python to use for ROCm if [[ "${BUILD_ENVIRONMENT}" == *rocm* ]]; then # HIP_PLATFORM is auto-detected by hipcc; unset to avoid build errors diff --git a/.ci/pytorch/common_utils.sh b/.ci/pytorch/common_utils.sh index c4a92c997561e..354841db899f8 100644 --- a/.ci/pytorch/common_utils.sh +++ b/.ci/pytorch/common_utils.sh @@ -127,17 +127,6 @@ function get_exit_code() { return $retcode } -function get_bazel() { - # Download and use the cross-platform, dependency-free Python - # version of Bazelisk to fetch the platform specific version of - # Bazel to use from .bazelversion. - retry curl --location --output tools/bazel \ - https://raw.githubusercontent.com/bazelbuild/bazelisk/v1.23.0/bazelisk.py - shasum --algorithm=1 --check \ - <(echo '01df9cf7f08dd80d83979ed0d0666a99349ae93c tools/bazel') - chmod u+x tools/bazel -} - function install_monkeytype { # Install MonkeyType pip_install MonkeyType @@ -290,7 +279,7 @@ function install_torchrec_and_fbgemm() { function clone_pytorch_xla() { if [[ ! -d ./xla ]]; then - git clone --recursive --quiet https://github.com/pytorch/xla.git + git clone --recursive -b r2.12 https://github.com/pytorch/xla.git pushd xla # pin the xla hash so that we don't get broken by changes to xla git checkout "$(cat ../.github/ci_commit_pins/xla.txt)" @@ -307,25 +296,9 @@ function install_torchao() { } function install_flash_attn_cute() { - echo "Installing FlashAttention CuTe from GitHub..." - # Grab latest main til we have a pinned commit - local flash_attn_commit - flash_attn_commit=$(git ls-remote https://github.com/Dao-AILab/flash-attention.git HEAD | cut -f1) - - # Clone the repo to a temporary directory - rm -rf flash-attention-build - git clone --depth 1 --recursive https://github.com/Dao-AILab/flash-attention.git flash-attention-build - - pushd flash-attention-build - git checkout "${flash_attn_commit}" - - # Install only the 'cute' sub-directory - pip_install -e flash_attn/cute/ - popd - - # remove the local repo - rm -rf flash-attention-build - echo "FlashAttention CuTe installation complete." + echo "Installing FlashAttention 4 from PyPI..." + pip_install flash-attn-4==4.0.0b5 + echo "FlashAttention 4 installation complete." } function install_cutlass_dsl() { @@ -367,7 +340,7 @@ function install_cutlass_api() { git checkout "${cutlass_commit}" # Install cutlass_api with torch extras - pip_install -e "python/cutlass_api[torch]" + pip_install "python/cutlass_api[torch]" popd rm -rf cutlass-build diff --git a/.ci/pytorch/cpp_doc_push_script.sh b/.ci/pytorch/cpp_doc_push_script.sh index f085fa78bebe9..d0b4fd38826fa 100755 --- a/.ci/pytorch/cpp_doc_push_script.sh +++ b/.ci/pytorch/cpp_doc_push_script.sh @@ -1,7 +1,7 @@ #!/bin/bash # This is where the local pytorch install in the docker image is located -pt_checkout="/var/lib/jenkins/workspace" +pt_checkout="${GITHUB_WORKSPACE:-/var/lib/jenkins/workspace}" # Since we're cat-ing this file, we need to escape all $'s echo "cpp_doc_push_script.sh: Invoked with $*" @@ -60,6 +60,34 @@ time python tools/setup_helpers/generate_code.py \ pushd docs/cpp time make VERBOSE=1 html +# Run C++ API coverage check (allowlist-based + HTML formatting) +echo "Running C++ docs coverage check..." +python check_coverage.py --coverxygen || coverage_exit=$? + +# Generate coverxygen HTML report if coverxygen produced output +if [ -f coverxygen.info ] && command -v genhtml &> /dev/null; then + genhtml --no-function-coverage coverxygen.info -o build/html/_coverage \ + --title "PyTorch C++ API Doc Coverage" \ + --legend --highlight 2>/dev/null || true +fi + +# Copy coverage reports into the build output so they get uploaded +mkdir -p build/html/_coverage +cp -f cpp_coverage.txt cpp_html_issues.txt build/html/_coverage/ 2>/dev/null || true +cp -f coverxygen.info build/html/_coverage/ 2>/dev/null || true + +if [ "${coverage_exit:-0}" -ne 0 ]; then + echo "" + echo "========================================" + echo "C++ DOCS COVERAGE: HIGH-PRIORITY GAPS" + echo "========================================" + echo "" + cat cpp_coverage.txt + echo "" + echo "See the full coverage report at: _coverage/cpp_coverage.txt" + echo "See the HTML issues report at: _coverage/cpp_html_issues.txt" +fi + popd popd @@ -76,6 +104,7 @@ cp -r "${pt_checkout}"/docs/cpp/build/html/* . # Copy back _config.yml rm -rf _config.yml mv /tmp/cppdocs-sync/* . +touch .nojekyll # Make a new commit git add . || true diff --git a/.ci/pytorch/macos-build.sh b/.ci/pytorch/macos-build.sh index f2bcb486cf95b..3259c62149b61 100755 --- a/.ci/pytorch/macos-build.sh +++ b/.ci/pytorch/macos-build.sh @@ -48,7 +48,7 @@ if [[ ${BUILD_ENVIRONMENT} == *"distributed"* ]]; then else # Explicitly set USE_DISTRIBUTED=0 to align with the default build config on mac. This also serves as the sole CI config that tests # that building with USE_DISTRIBUTED=0 works at all. See https://github.com/pytorch/pytorch/issues/86448 - USE_DISTRIBUTED=0 USE_OPENMP=1 MACOSX_DEPLOYMENT_TARGET=11.0 WERROR=1 BUILD_TEST=OFF USE_PYTORCH_METAL=1 python -m build --wheel --no-isolation -C--build-option=--plat-name=macosx_11_0_arm64 + USE_DISTRIBUTED=0 USE_OPENMP=1 WERROR=1 BUILD_TEST=OFF USE_PYTORCH_METAL=1 python -m build --wheel --no-isolation fi if which sccache > /dev/null; then print_sccache_stats diff --git a/.ci/pytorch/macos-common.sh b/.ci/pytorch/macos-common.sh index 6826a52577a29..d21ee458f6c99 100755 --- a/.ci/pytorch/macos-common.sh +++ b/.ci/pytorch/macos-common.sh @@ -9,6 +9,6 @@ sysctl -a | grep machdep.cpu # These are required for both the build job and the test job. # In the latter to test cpp extensions. -export MACOSX_DEPLOYMENT_TARGET=11.1 +export MACOSX_DEPLOYMENT_TARGET=14.0 export CXX=clang++ export CC=clang diff --git a/.ci/pytorch/macos-test.sh b/.ci/pytorch/macos-test.sh index f6173f64224e2..332a52b11b10a 100755 --- a/.ci/pytorch/macos-test.sh +++ b/.ci/pytorch/macos-test.sh @@ -49,6 +49,8 @@ test_python_mps() { test_python_openreg() { setup_test_python + git submodule update --init --depth 1 third_party/googletest + time python test/run_test.py --openreg --verbose assert_git_not_dirty diff --git a/.ci/pytorch/python_doc_push_script.sh b/.ci/pytorch/python_doc_push_script.sh index 6bcd46c4815a6..dde476a9392ee 100755 --- a/.ci/pytorch/python_doc_push_script.sh +++ b/.ci/pytorch/python_doc_push_script.sh @@ -1,7 +1,7 @@ #!/bin/bash # This is where the local pytorch install in the docker image is located -pt_checkout="/var/lib/jenkins/workspace" +pt_checkout="${GITHUB_WORKSPACE:-/var/lib/jenkins/workspace}" source "$pt_checkout/.ci/pytorch/common_utils.sh" @@ -50,8 +50,14 @@ echo "install_path: $install_path version: $version" build_docs () { set +e - set -o pipefail - make "$1" 2>&1 | tee /tmp/docs_build.txt + # Don't pipe through tee: sphinx -j auto forks workers that inherit + # the pipe fd and hold it open after sphinx exits, causing tee to + # block forever. Write to a file and tail with --pid so it exits + # (after draining) when make finishes. + make "$1" > /tmp/docs_build.txt 2>&1 & + local make_pid=$! + tail -f --pid=$make_pid /tmp/docs_build.txt + wait $make_pid code=$? if [ $code -ne 0 ]; then set +x @@ -87,44 +93,36 @@ pushd docs if [ "$is_main_doc" = true ]; then build_docs html || exit $? - make coverage - # Now we have the coverage report, we need to make sure it is empty. - # Sphinx 7.2.6+ format: python.txt contains a statistics table with a TOTAL row - # showing the undocumented count in the third column. - # Example: | TOTAL | 99.83% | 2 | - # - # Also: see docs/source/conf.py for "coverage_ignore*" items, which should - # be documented then removed from there. - - # Extract undocumented count from TOTAL row in Sphinx 7.2.6 statistics table - # The table format is: | Module | Coverage | Undocumented | - # Extract the third column (undocumented count) from the TOTAL row - undocumented=$(grep "| TOTAL" build/coverage/python.txt | awk -F'|' '{print $4}' | tr -d ' ') - - if [ -z "$undocumented" ] || ! [[ "$undocumented" =~ ^[0-9]+$ ]]; then - echo coverage output not found - exit 1 - elif [ "$undocumented" -gt 0 ]; then - set +x # Disable command echoing for cleaner output - echo "" - echo "=====================" - echo "UNDOCUMENTED OBJECTS:" - echo "=====================" - echo "" - # Find the line number of the TOTAL row and print only what comes after it - total_line=$(grep -n "| TOTAL" build/coverage/python.txt | cut -d: -f1) - if [ -n "$total_line" ]; then - # Print only the detailed list (skip the statistics table) - tail -n +$((total_line + 2)) build/coverage/python.txt - else - # Fallback to showing entire file if TOTAL line not found - cat build/coverage/python.txt + # Coverage check is only needed on PR builds (push=false) to catch + # undocumented APIs early. Nightly/release builds (push=true) skip it + # since PRs already enforce coverage. + if [[ "${WITH_PUSH:-}" != true ]]; then + SPHINXOPTS="-WT --keep-going" make coverage + + undocumented=$(grep "| TOTAL" build/coverage/python.txt | awk -F'|' '{print $4}' | tr -d ' ') + + if [ -z "$undocumented" ] || ! [[ "$undocumented" =~ ^[0-9]+$ ]]; then + echo coverage output not found + exit 1 + elif [ "$undocumented" -gt 0 ]; then + set +x + echo "" + echo "=====================" + echo "UNDOCUMENTED OBJECTS:" + echo "=====================" + echo "" + total_line=$(grep -n "| TOTAL" build/coverage/python.txt | cut -d: -f1) + if [ -n "$total_line" ]; then + tail -n +$((total_line + 2)) build/coverage/python.txt + else + cat build/coverage/python.txt + fi + echo "" + echo "Make sure you've updated relevant .rsts in docs/source!" + echo "You can reproduce locally by running 'cd docs && make coverage && tail -n +\$((grep -n \"| TOTAL\" build/coverage/python.txt | cut -d: -f1) + 2)) build/coverage/python.txt'" + set -x + exit 1 fi - echo "" - echo "Make sure you've updated relevant .rsts in docs/source!" - echo "You can reproduce locally by running 'cd docs && make coverage && tail -n +\$((grep -n \"| TOTAL\" build/coverage/python.txt | cut -d: -f1) + 2)) build/coverage/python.txt'" - set -x # Re-enable command echoing - exit 1 fi else # skip coverage, format for stable or tags diff --git a/.ci/pytorch/smoke_test/check_binary_symbols.py b/.ci/pytorch/smoke_test/check_binary_symbols.py index 95f1e51a88a13..4a49be6db5730 100755 --- a/.ci/pytorch/smoke_test/check_binary_symbols.py +++ b/.ci/pytorch/smoke_test/check_binary_symbols.py @@ -1,4 +1,6 @@ #!/usr/bin/env python3 +from __future__ import annotations + import concurrent.futures import distutils.sysconfig import functools diff --git a/.ci/pytorch/smoke_test/check_wheel_tags.py b/.ci/pytorch/smoke_test/check_wheel_tags.py new file mode 100644 index 0000000000000..901304657b531 --- /dev/null +++ b/.ci/pytorch/smoke_test/check_wheel_tags.py @@ -0,0 +1,269 @@ +"""Validate wheel platform tags and macOS dylib minos. +Supports two modes: +1. Pre-install: reads .whl files from PYTORCH_FINAL_PACKAGE_DIR +2. Post-install: reads metadata from installed torch package (soft warnings) +- (macOS only) dylib minos matches the wheel platform tag +""" + +import os +import platform +import re +import subprocess +import sys +import tempfile +import zipfile +from pathlib import Path + + +EXPECTED_PLATFORM_TAGS: dict[str, str] = { + "linux": r"_x86_64$", + "linux-aarch64": r"_aarch64$", + "windows": r"^win_amd64$", + "win32": r"^win_amd64$", + "macos-arm64": r"^macosx_\d+_\d+_arm64$", + "darwin": r"^macosx_\d+_\d+_(arm64|x86_64)$", +} + + +def _extract_wheel_tags(whl_path: Path) -> list[str]: + """Extract Tag values from the WHEEL metadata file inside a .whl archive.""" + tags = [] + with zipfile.ZipFile(whl_path, "r") as zf: + wheel_files = [n for n in zf.namelist() if n.endswith("/WHEEL")] + if not wheel_files: + return tags + content = zf.read(wheel_files[0]).decode("utf-8") + for line in content.splitlines(): + if line.startswith("Tag:"): + tags.append(line.split(":", 1)[1].strip()) + return tags + + +def _extract_installed_wheel_tags(package: str = "torch") -> list[str]: + """Extract Tag values from an installed package's WHEEL metadata.""" + from importlib.metadata import distribution + + dist = distribution(package) + wheel_text = dist.read_text("WHEEL") + if not wheel_text: + return [] + tags = [] + for line in wheel_text.splitlines(): + if line.startswith("Tag:"): + tags.append(line.split(":", 1)[1].strip()) + return tags + + +def check_wheel_platform_tag() -> None: + """Validate that wheel Tags in WHEEL metadata match the expected platform. + + Mode 1: PYTORCH_FINAL_PACKAGE_DIR set → read .whl file (strict, raises on mismatch) + Mode 2: No wheel dir → read from installed torch package (soft, prints warnings) + """ + wheel_dir = os.getenv("PYTORCH_FINAL_PACKAGE_DIR", "") + + target_os = os.getenv("TARGET_OS", sys.platform) + if target_os == "linux" and platform.machine() == "aarch64": + target_os = "linux-aarch64" + expected_python = f"cp{sys.version_info.major}{sys.version_info.minor}" + import sysconfig + + abiflags = getattr(sys, "abiflags", "") + if not abiflags and ( + os.getenv("MATRIX_PYTHON_VERSION", "").endswith("t") + or bool(sysconfig.get_config_var("Py_GIL_DISABLED")) + or not getattr(sys, "_is_gil_enabled", lambda: True)() + ): + abiflags = "t" + expected_abi = f"cp{sys.version_info.major}{sys.version_info.minor}{abiflags}" + print(f"Expected ABI tag: {expected_abi}") + + platform_pattern = EXPECTED_PLATFORM_TAGS.get(target_os) + if not platform_pattern: + print( + f"No expected platform pattern for TARGET_OS={target_os}, " + "skipping wheel tag check" + ) + return + + # Mode 1: Read from .whl file + if wheel_dir and os.path.isdir(wheel_dir): + whls = list(Path(wheel_dir).glob("torch-*.whl")) + if not whls: + print(f"No torch wheel found in {wheel_dir}, skipping wheel tag check") + return + if len(whls) > 1: + raise RuntimeError( + f"Expected exactly one torch wheel in {wheel_dir}, " + f"found {len(whls)}: {[w.name for w in whls]}" + ) + whl = whls[0] + print(f"Checking wheel platform tag for: {whl.name}") + tags = _extract_wheel_tags(whl) + source = whl.name + else: + # Mode 2: Read from installed package (soft) + print("PYTORCH_FINAL_PACKAGE_DIR not set, reading from installed torch package") + try: + tags = _extract_installed_wheel_tags("torch") + source = "installed torch" + except Exception as e: + print(f"Could not read installed torch metadata: {e}, skipping") + return + + if not tags: + raise RuntimeError(f"No Tag found in WHEEL metadata of {source}") + + for tag_str in tags: + parts = tag_str.split("-") + if len(parts) != 3: + msg = ( + f"Malformed wheel tag '{tag_str}' in {source}, " + f"expected format: --" + ) + raise RuntimeError(msg) + + python_tag, abi_tag, platform_tag = parts + + print(f"Checking tag: {tag_str} (from {source})") + if python_tag != expected_python: + msg: str = ( + f"Python tag mismatch in {source}: " + f"got '{python_tag}', expected '{expected_python}'" + ) + raise RuntimeError(msg) + + if abi_tag != expected_abi: + msg = ( + f"ABI tag mismatch in {source}: " + f"got '{abi_tag}', expected '{expected_abi}'" + ) + raise RuntimeError(msg) + + if not re.search(platform_pattern, platform_tag): + msg = ( + f"Platform tag mismatch in {source}: " + f"got '{platform_tag}', expected pattern matching " + f"'{platform_pattern}' for TARGET_OS={target_os}" + ) + raise RuntimeError(msg) + + print(f"OK: Wheel tag(s) valid for {source}: {', '.join(tags)}") + + +def _check_dylibs_minos(dylibs: list, expected_minos: str, source: str) -> None: + mismatches = [] + for dylib in dylibs: + try: + result = subprocess.run( + ["otool", "-l", str(dylib)], + capture_output=True, + text=True, + timeout=30, + ) + except Exception: + continue + + minos = None + lines = result.stdout.splitlines() + for i, line in enumerate(lines): + s = line.strip() + if "LC_BUILD_VERSION" in s: + for j in range(i + 1, min(i + 6, len(lines))): + if lines[j].strip().startswith("minos"): + minos = lines[j].strip().split()[1] + break + break + if "LC_VERSION_MIN_MACOSX" in s: + for j in range(i + 1, min(i + 4, len(lines))): + if lines[j].strip().startswith("version"): + minos = lines[j].strip().split()[1] + break + break + + # A dylib with a lower minos than the wheel tag is safe (forward compatible). + # Only flag dylibs that require a *higher* macOS than the wheel claims to support. + if minos and tuple(int(x) for x in minos.split(".")) > tuple( + int(x) for x in expected_minos.split(".") + ): + mismatches.append( + f"{dylib.name}: minos={minos}, expected<={expected_minos}" + ) + + if mismatches: + raise RuntimeError( + f"minos/platform tag mismatch in {len(mismatches)} dylib(s):\n" + + "\n".join(f" {m}" for m in mismatches) + ) + print( + f"OK: All {len(dylibs)} dylib(s) have minos matching " + f"platform tag ({expected_minos}) for {source}" + ) + + +def check_mac_wheel_minos() -> None: + if sys.platform != "darwin": + return + + wheel_dir = os.getenv("PYTORCH_FINAL_PACKAGE_DIR", "") + + if wheel_dir and os.path.isdir(wheel_dir): + # Mode 1: extract dylibs from .whl file + whls = list(Path(wheel_dir).glob("*.whl")) + if not whls: + print(f"No .whl files in {wheel_dir}, skipping wheel minos check") + return + + macos_whl_re = re.compile(r"macosx_(\d+)_(\d+)_(\w+)\.whl$") + for whl in whls: + print(f"Checking wheel tag minos for: {whl.name}") + m = macos_whl_re.search(whl.name) + if not m: + print(f"No macOS platform tag in {whl.name}, skipping") + continue + expected_minos = f"{m.group(1)}.{m.group(2)}" + + with tempfile.TemporaryDirectory() as tmpdir: + with zipfile.ZipFile(whl, "r") as zf: + dylib_names = [n for n in zf.namelist() if n.endswith(".dylib")] + if not dylib_names: + print("No .dylib files in wheel, skipping minos check") + continue + for name in dylib_names: + zf.extract(name, tmpdir) + dylibs = list(Path(tmpdir).rglob("*.dylib")) + _check_dylibs_minos(dylibs, expected_minos, whl.name) + else: + # Mode 2: read from installed torch package + print("PYTORCH_FINAL_PACKAGE_DIR not set, checking installed torch dylibs") + try: + tags = _extract_installed_wheel_tags("torch") + except Exception as e: + print(f"Could not read installed torch metadata: {e}, skipping") + return + + expected_minos = None + for tag_str in tags: + m = re.search(r"macosx_(\d+)_(\d+)_\w+", tag_str) + if m: + expected_minos = f"{m.group(1)}.{m.group(2)}" + break + + if not expected_minos: + print("No macOS platform tag found in installed torch metadata, skipping") + return + + print(f"Expected minos from installed wheel tag: {expected_minos}") + + import torch + + torch_dir = Path(torch.__file__).parent + dylibs = list(torch_dir.rglob("*.dylib")) + if not dylibs: + raise RuntimeError("No .dylib files found in installed torch") + _check_dylibs_minos(dylibs, expected_minos, "installed torch") + + +if __name__ == "__main__": + check_wheel_platform_tag() + check_mac_wheel_minos() diff --git a/.ci/pytorch/smoke_test/smoke_test.py b/.ci/pytorch/smoke_test/smoke_test.py index d916f6b49a2f3..877218e71e307 100644 --- a/.ci/pytorch/smoke_test/smoke_test.py +++ b/.ci/pytorch/smoke_test/smoke_test.py @@ -1,3 +1,5 @@ +from __future__ import annotations + import argparse import importlib import json @@ -7,7 +9,8 @@ import sys from pathlib import Path from tempfile import NamedTemporaryFile -from typing import Optional + +from check_wheel_tags import check_mac_wheel_minos, check_wheel_platform_tag import torch import torch._dynamo @@ -24,6 +27,7 @@ package_type = os.getenv("MATRIX_PACKAGE_TYPE") target_os = os.getenv("TARGET_OS", sys.platform) BASE_DIR = Path(__file__).parent.parent.parent +PYTORCH_ROOT = BASE_DIR.parent is_cuda_system = gpu_arch_type == "cuda" NIGHTLY_ALLOWED_DELTA = 3 @@ -205,7 +209,7 @@ def test_cuda_gds_errors_captured() -> None: ) -def find_pypi_package_version(package: str) -> Optional[str]: +def find_pypi_package_version(package: str) -> str | None: from importlib import metadata dists = metadata.distributions() @@ -215,6 +219,89 @@ def find_pypi_package_version(package: str) -> Optional[str]: return None +def get_expected_cudnn_version_linux(cuda_version: str) -> str | None: + """Parse expected cuDNN version from generate_binary_build_matrix.py for Linux. + + Reads PYTORCH_EXTRA_INSTALL_REQUIREMENTS and extracts the cudnn version + for the given CUDA version (e.g. "12.6"). + """ + matrix_script = ( + PYTORCH_ROOT / ".github" / "scripts" / "generate_binary_build_matrix.py" + ) + if not matrix_script.exists(): + print(f"Warning: {matrix_script} not found, skipping cuDNN version check") + return None + + content = matrix_script.read_text() + # Match the full cudnn package version like nvidia-cudnn-cu12==9.10.2.21 + # and extract major.minor.patch (dropping the build number) + pattern = ( + rf'"{re.escape(cuda_version)}":\s*\(\s*' + r"[\s\S]*?nvidia-cudnn-cu\d+==(\d+\.\d+\.\d+)\.\d+" + ) + match = re.search(pattern, content) + if match: + return match.group(1) + return None + + +def get_expected_cudnn_version_windows(cuda_version: str) -> str | None: + """Parse expected cuDNN version from cuda_install.bat for Windows. + + Reads the batch file and extracts EXPECTED_CUDNN_VERSION for the given + CUDA version (e.g. "12.6" maps to CUDA_VER 126). + """ + bat_file = ( + PYTORCH_ROOT / ".ci" / "pytorch" / "windows" / "internal" / "cuda_install.bat" + ) + if not bat_file.exists(): + print(f"Warning: {bat_file} not found, skipping cuDNN version check") + return None + + content = bat_file.read_text() + # Convert "12.6" to "126" to match batch file's CUDA_VER format + cuda_ver_nodot = cuda_version.replace(".", "") + # Match: if %CUDA_VER% EQU 126 ( ... set EXPECTED_CUDNN_VERSION=9.10.2 ) + pattern = ( + rf"if %CUDA_VER% EQU {re.escape(cuda_ver_nodot)}\s*\(" + r"[\s\S]*?set EXPECTED_CUDNN_VERSION=(\d+\.\d+\.\d+)" + ) + match = re.search(pattern, content) + if match: + return match.group(1) + return None + + +def check_cudnn_version(cuda_version: str, actual_cudnn_version: str) -> None: + """Validate cuDNN version matches expected version from build config files.""" + if sys.platform in ["linux", "linux2"]: + expected = get_expected_cudnn_version_linux(cuda_version) + source = "generate_binary_build_matrix.py" + elif sys.platform == "win32": + expected = get_expected_cudnn_version_windows(cuda_version) + source = "cuda_install.bat" + else: + print(f"cuDNN version check not supported on platform {sys.platform}") + return + + if expected is None: + print( + f"Warning: Could not determine expected cuDNN version for CUDA {cuda_version} " + f"from {source}, skipping validation" + ) + return + + if not actual_cudnn_version.startswith(expected): + raise RuntimeError( + f"cuDNN version mismatch for CUDA {cuda_version}. " + f"Loaded: {actual_cudnn_version} Expected: {expected} (from {source})" + ) + print( + f"cuDNN version check passed: {actual_cudnn_version} matches " + f"expected {expected} from {source}" + ) + + def cudnn_to_version_str(cudnn_version: int) -> str: patch = int(cudnn_version % 10) minor = int((cudnn_version / 100) % 100) @@ -281,6 +368,20 @@ def smoke_test_cuda( torch_cudnn_version = cudnn_to_version_str(torch.backends.cudnn.version()) print(f"Torch cuDNN version: {torch_cudnn_version}") + torch_cudnn_compile_version = torch._C._cudnn.getCompileVersion() + print(f"Torch cuDNN compile-time version: {torch_cudnn_compile_version}") + torch_cudnn_runtime_version = tuple( + [int(x) for x in torch_cudnn_version.split(".")] + ) + if torch_cudnn_runtime_version != torch_cudnn_compile_version: + raise RuntimeError( + "cuDNN runtime version doesn't match comple version. " + f"Loaded: {torch_cudnn_runtime_version} " + f"Expected: {torch_cudnn_compile_version}" + ) + + check_cudnn_version(gpu_arch_ver, torch_cudnn_version) + if sys.platform in ["linux", "linux2"]: torch_nccl_version = ".".join(str(v) for v in torch.cuda.nccl.version()) print(f"Torch nccl; version: {torch_nccl_version}") @@ -538,6 +639,9 @@ def main() -> None: smoke_test_nvshmem() + check_wheel_platform_tag() + check_mac_wheel_minos() + if __name__ == "__main__": main() diff --git a/.ci/pytorch/test.sh b/.ci/pytorch/test.sh index 7bc94541a7558..0b10d1dda82b3 100755 --- a/.ci/pytorch/test.sh +++ b/.ci/pytorch/test.sh @@ -14,9 +14,10 @@ source "$(dirname "${BASH_SOURCE[0]}")/common.sh" # shellcheck source=./common-build.sh source "$(dirname "${BASH_SOURCE[0]}")/common-build.sh" -# Do not change workspace permissions for ROCm and s390x CI jobs -# as it can leave workspace with bad permissions for cancelled jobs -if [[ "$BUILD_ENVIRONMENT" != *rocm* && "$BUILD_ENVIRONMENT" != *s390x* && -d /var/lib/jenkins/workspace ]]; then +# Only change workspace permissions if passwordless sudo is available +# (e.g. ROCm and s390x CI jobs lack it, and changing permissions +# can leave the workspace in a bad state for cancelled jobs) +if sudo -n true 2>/dev/null && [[ -d /var/lib/jenkins/workspace ]]; then # Workaround for dind-rootless userid mapping (https://github.com/pytorch/ci-infra/issues/96) WORKSPACE_ORIGINAL_OWNER_ID=$(stat -c '%u' "/var/lib/jenkins/workspace") cleanup_workspace() { @@ -44,6 +45,16 @@ if [[ "$BUILD_ENVIRONMENT" == *cuda* ]]; then fi fi +# Remove onnxruntime if present to avoid interference with non-ONNX tests +if [[ "$TEST_CONFIG" != "onnx" ]]; then + pip uninstall -y onnxruntime 2>/dev/null || true +fi + +# Remove dill to test that serialization works without it +if [[ "$BUILD_ENVIRONMENT" == *py3.10-gcc11 ]]; then + pip uninstall -y dill 2>/dev/null || true +fi + echo "Environment variables:" env @@ -129,9 +140,7 @@ if [[ "${PYTORCH_TEST_RERUN_DISABLED_TESTS}" == "1" ]] || [[ "${CONTINUE_THROUGH fi # Get fully qualified path using realpath -if [[ "$BUILD_ENVIRONMENT" != *bazel* ]]; then - CUSTOM_TEST_ARTIFACT_BUILD_DIR=$(realpath "${CUSTOM_TEST_ARTIFACT_BUILD_DIR:-"build/custom_test_artifacts"}") -fi +CUSTOM_TEST_ARTIFACT_BUILD_DIR=$(realpath "${CUSTOM_TEST_ARTIFACT_BUILD_DIR:-"build/custom_test_artifacts"}") # Reduce set of tests to include when running run_test.py if [[ -n $TESTS_TO_INCLUDE ]]; then @@ -144,11 +153,34 @@ env echo "Testing pytorch" +# Set OMP_NUM_THREADS to nproc/4 on k8s ARC runners if not already set. +# +# We use nproc (cgroup-aware) rather than os.cpu_count() because on k8s (ARC) +# pods, os.cpu_count() returns the host's CPU count (e.g., 192) rather than +# the pod's cpuset allocation (e.g., 16). +# +# We use nproc/4 rather than nproc because OpenMP spin-waits at thread barriers. +# When thread count equals cpuset size (e.g., 16 threads on 16 CPUs), spinning +# barrier threads monopolize all CPUs and the OS must context-switch to let +# actual work complete. This causes ~5000x slowdowns on small tensor ops +# (e.g., aten::copy_ on 147KB: ~34ms instead of ~7us). Using nproc/4 leaves +# headroom for the main thread and for NUM_PROCS=3 parallel test processes. +if [[ -z "${OMP_NUM_THREADS:-}" ]] && [[ -n "${USE_ARC:-}" ]]; then + OMP_NUM_THREADS=$(( $(nproc) / 4 )) + # Floor of 4: low OMP_NUM_THREADS (1-2) changes floating-point reduction + # order, causing numerical mismatches in tests with tight tolerances + # (e.g., test_batchnorm_nhwc_cpu). + if [[ "$OMP_NUM_THREADS" -lt 4 ]]; then + OMP_NUM_THREADS=4 + fi + export OMP_NUM_THREADS +fi + export LANG=C.UTF-8 PR_NUMBER=${PR_NUMBER:-${CIRCLE_PR_NUMBER:-}} -if [[ -d "${HF_CACHE}" ]]; then +if [[ -d "${HF_CACHE}" && "$TEST_CONFIG" != "onnx" ]]; then export HF_HOME="${HF_CACHE}" fi @@ -222,12 +254,10 @@ if [[ "$BUILD_ENVIRONMENT" == *xpu* ]]; then timeout 30 xpu-smi discovery || true fi -if [[ "$BUILD_ENVIRONMENT" != *-bazel-* ]] ; then - # JIT C++ extensions require ninja (installed from requirements-ci.txt). - # ninja is installed in $HOME/.local/bin, e.g., /var/lib/jenkins/.local/bin for CI user jenkins - # but this script should be runnable by any user, including root - export PATH="$HOME/.local/bin:$PATH" -fi +# JIT C++ extensions require ninja (installed from requirements-ci.txt). +# ninja is installed in $HOME/.local/bin, e.g., /var/lib/jenkins/.local/bin for CI user jenkins +# but this script should be runnable by any user, including root +export PATH="$HOME/.local/bin:$PATH" if [[ "$BUILD_ENVIRONMENT" == *aarch64* ]]; then # TODO: revisit this once the CI is stabilized on aarch64 linux @@ -299,8 +329,7 @@ fi if [[ "$BUILD_ENVIRONMENT" == *-debug* ]]; then echo "We are in debug mode: $BUILD_ENVIRONMENT. Expect the python assertion to fail" (cd test && ! get_exit_code python -c "import torch; torch._C._crash_if_debug_asserts_fail(424242)") -elif [[ "$BUILD_ENVIRONMENT" != *-bazel-* ]]; then - # Noop when debug is disabled. Skip bazel jobs because torch isn't available there yet. +else echo "We are not in debug mode: $BUILD_ENVIRONMENT. Expect the assertion to pass" (cd test && python -c "import torch; torch._C._crash_if_debug_asserts_fail(424242)") fi @@ -311,12 +340,6 @@ elif [[ $TEST_CONFIG == 'nogpu_AVX512' ]]; then export ATEN_CPU_CAPABILITY=avx2 fi -if [[ "${TEST_CONFIG}" == "legacy_nvidia_driver" ]]; then - # Make sure that CUDA can be initialized - (cd test && python -c "import torch; torch.rand(2, 2, device='cuda')") - export USE_LEGACY_DRIVER=1 -fi - test_python_legacy_jit() { time python test/run_test.py --include test_jit_legacy test_jit_fuser_legacy --verbose assert_git_not_dirty @@ -365,8 +388,28 @@ test_python_smoke_b200() { inductor/test_flex_flash \ inductor/test_torchinductor \ inductor/test_nv_universal_gemm \ - $PYTHON_TEST_EXTRA_OPTION \ - --upload-artifacts-while-running + inductor/test_fused_attention \ + test_varlen_attention \ + $PYTHON_TEST_EXTRA_OPTION \ + --upload-artifacts-while-running + assert_git_not_dirty +} + + +test_python_smoke_xpu() { + # Smoke tests for XPU client + time python test/run_test.py --include test_transformers $PYTHON_TEST_EXTRA_OPTION --upload-artifacts-while-running + time test_xpu_sycl_tla_backend + assert_git_not_dirty +} + +test_dtensor() { + # Dynamically discover all test files under test/distributed/tensor/ + # so new tests are automatically picked up. + # shellcheck disable=SC2046 + time python test/run_test.py \ + --include $(find test/distributed/tensor -name 'test_*.py' -printf '%P\n' | sed 's|\.py$||; s|^|distributed/tensor/|' | sort | tr '\n' ' ') \ + --verbose $PYTHON_TEST_EXTRA_OPTION --upload-artifacts-while-running assert_git_not_dirty } @@ -393,6 +436,7 @@ test_h100_symm_mem() { export NVSHMEM_SYMMETRIC_SIZE=4G # Disable NVLink Switch features (not available on AWS H100 instances) export NVSHMEM_DISABLE_NVLS=1 + export NCCL_NVLS_ENABLE=0 _run_symm_mem_tests } @@ -402,10 +446,20 @@ test_b200_symm_mem() { test_h100_cutlass_backend() { # cutlass backend tests for H100 + git submodule update --init --depth 1 third_party/cutlass TORCHINDUCTOR_CUTLASS_DIR=$(realpath "./third_party/cutlass") python test/run_test.py --include inductor/test_cutlass_backend -k "not addmm" $PYTHON_TEST_EXTRA_OPTION --upload-artifacts-while-running TORCHINDUCTOR_CUTLASS_DIR=$(realpath "./third_party/cutlass") python test/run_test.py --include inductor/test_cutlass_evt $PYTHON_TEST_EXTRA_OPTION --upload-artifacts-while-running } +test_xpu_sycl_tla_backend() { + # Inductor sycl-tla backend tests for XPU + # shellcheck disable=SC1091 + source /opt/intel/oneapi/mkl/latest/env/vars.sh + sycl_tla_dir=$(realpath "./third_party/sycl-tla") + rm -rf "${sycl_tla_dir}" && git clone --depth 1 --single-branch -b v0.8 --quiet https://github.com/intel/sycl-tla.git "${sycl_tla_dir}" + TORCHINDUCTOR_CUTLASS_DIR=$(realpath "./third_party/sycl-tla") python test/run_test.py --include inductor/test_cutlass_backend -k "not addmm" $PYTHON_TEST_EXTRA_OPTION --upload-artifacts-while-running +} + test_lazy_tensor_meta_reference_disabled() { export TORCH_DISABLE_FUNCTIONALIZATION_META_REFERENCE=1 echo "Testing lazy tensor operations without meta reference" @@ -422,12 +476,16 @@ test_dynamo_core() { } test_dynamo_cpython() { + # Disable TD for cpython since it's pretty cheap to run the cpython tests (< 10 min) + # and if TD is enabled, only 25% of the tests will be executed + export NO_TD=1 time python test/run_test.py \ --include-cpython-tests \ --dynamo \ --verbose \ --upload-artifacts-while-running assert_git_not_dirty + unset NO_TD } test_dynamo_wrapped_shard() { @@ -452,6 +510,8 @@ test_dynamo_wrapped_shard() { } test_einops() { + pip install einops==0.5.0 + time python test/run_test.py --einops --verbose --upload-artifacts-while-running pip install einops==0.6.1 time python test/run_test.py --einops --verbose --upload-artifacts-while-running pip install einops==0.7.0 @@ -532,7 +592,7 @@ test_inductor_shard() { # Do not add --inductor for the following inductor unit tests, otherwise we will fail because of nested dynamo state python test/run_test.py \ - --include inductor/test_torchinductor inductor/test_torchinductor_opinfo inductor/test_aot_inductor \ + --include inductor/test_torchinductor inductor/test_torchinductor_opinfo inductor/test_aot_inductor inductor/test_cpu_select_algorithm \ --shard "$1" "$NUM_TEST_SHARDS" \ --verbose } @@ -590,7 +650,7 @@ test_inductor_cpp_wrapper_shard() { --shard "$1" "$NUM_TEST_SHARDS" \ --verbose python test/run_test.py \ - --include inductor/test_torchinductor inductor/test_max_autotune inductor/test_cpu_repro \ + --include inductor/test_torchinductor inductor/test_max_autotune inductor/test_cpu_repro inductor/test_triton_kernels \ --shard "$1" "$NUM_TEST_SHARDS" \ --verbose python test/run_test.py --inductor \ @@ -598,7 +658,12 @@ test_inductor_cpp_wrapper_shard() { -k 'take' \ --shard "$1" "$NUM_TEST_SHARDS" \ --verbose - + # Keep testing TORCHINDUCTOR_AUTOTUNE_AT_COMPILE_TIME=1 for the near future. + # Will drop this after AOTInductor also switches to lazy Triton compilation. + TORCHINDUCTOR_AUTOTUNE_AT_COMPILE_TIME=1 python test/run_test.py \ + --include inductor/test_torchinductor inductor/test_triton_kernels inductor/test_max_autotune \ + --shard "$1" "$NUM_TEST_SHARDS" \ + --verbose if [[ "${BUILD_ENVIRONMENT}" == *xpu* ]]; then python test/run_test.py \ --include inductor/test_mkldnn_pattern_matcher \ @@ -712,6 +777,10 @@ test_perf_for_dashboard() { TEST_REPORTS_DIR=$(pwd)/test/test-reports mkdir -p "$TEST_REPORTS_DIR" + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" ]]; then + mkdir -p "$TEST_REPORTS_DIR/profiler_traces" + fi + local suite="$1" shift @@ -771,52 +840,97 @@ test_perf_for_dashboard() { fi if [[ "$DASHBOARD_TAG" == *default-true* ]]; then + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" && "$target" == "performance" ]]; then + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${backend}_no_cudagraphs_${suite}_${dtype}_${mode}_${device}") + fi $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --backend "$backend" --disable-cudagraphs "$@" \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${backend}_no_cudagraphs_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi if [[ "$DASHBOARD_TAG" == *cudagraphs-true* ]]; then + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" && "$target" == "performance" ]]; then + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${backend}_with_cudagraphs_${suite}_${dtype}_${mode}_${device}") + fi $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --backend "$backend" "$@" \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${backend}_with_cudagraphs_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi if [[ "$DASHBOARD_TAG" == *dynamic-true* ]]; then + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" && "$target" == "performance" ]]; then + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${backend}_dynamic_${suite}_${dtype}_${mode}_${device}") + fi $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --backend "$backend" --dynamic-shapes \ --dynamic-batch-only "$@" \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${backend}_dynamic_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi if [[ "$DASHBOARD_TAG" == *cppwrapper-true* ]]; then + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" && "$target" == "performance" ]]; then + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${backend}_cpp_wrapper_${suite}_${dtype}_${mode}_${device}") + fi TORCHINDUCTOR_CPP_WRAPPER=1 $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --backend "$backend" --disable-cudagraphs "$@" \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${backend}_cpp_wrapper_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi if [[ "$DASHBOARD_TAG" == *freezing_cudagraphs-true* ]] && [[ "$mode" == "inference" ]]; then + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" && "$target" == "performance" ]]; then + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${backend}_with_cudagraphs_freezing_${suite}_${dtype}_${mode}_${device}") + fi $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --backend "$backend" "$@" --freezing \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${backend}_with_cudagraphs_freezing_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi if [[ "$DASHBOARD_TAG" == *freeze_autotune_cudagraphs-true* ]] && [[ "$mode" == "inference" ]]; then + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" && "$target" == "performance" ]]; then + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${backend}_with_cudagraphs_freezing_autotune_${suite}_${dtype}_${mode}_${device}") + fi TORCHINDUCTOR_MAX_AUTOTUNE=1 $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --backend "$backend" "$@" --freezing \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${backend}_with_cudagraphs_freezing_autotune_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi if [[ "$DASHBOARD_TAG" == *aotinductor-true* ]] && [[ "$mode" == "inference" ]]; then - if [[ "$target" == "accuracy" ]]; then # Also collect Export pass rate and display as a separate row + if [[ "$target" == "accuracy" ]]; then $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --export --disable-cudagraphs "$@" \ --output "$TEST_REPORTS_DIR/${backend}_export_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" && "$target" == "performance" ]]; then + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${backend}_aot_inductor_${suite}_${dtype}_${mode}_${device}") + fi $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --export-aot-inductor --disable-cudagraphs "$@" \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${backend}_aot_inductor_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi if [[ "$DASHBOARD_TAG" == *maxautotune-true* ]]; then + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" && "$target" == "performance" ]]; then + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${backend}_max_autotune_${suite}_${dtype}_${mode}_${device}") + fi TORCHINDUCTOR_MAX_AUTOTUNE=1 $TASKSET python "benchmarks/dynamo/$suite.py" \ "${target_flag[@]}" --"$mode" --"$dtype" --backend "$backend" "$@" \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${backend}_max_autotune_${suite}_${dtype}_${mode}_${device}_${target}.csv" fi + if [[ "$DASHBOARD_TAG" == *deterministic_perf-true* ]]; then + $TASKSET python "benchmarks/dynamo/$suite.py" \ + "${target_flag[@]}" --"$mode" --"$dtype" --backend "$backend" --disable-cudagraphs --deterministic "$@" \ + --output "$TEST_REPORTS_DIR/${backend}_deterministic_perf_${suite}_${dtype}_${mode}_${device}_${target}.csv" + fi done done } @@ -844,9 +958,15 @@ test_single_dynamo_benchmark() { fi if [[ "${TEST_CONFIG}" == *perf_compare* ]]; then + local profiler_trace_flags=() + if [[ "${EXPORT_PROFILER_TRACE:-0}" == "1" ]]; then + mkdir -p "$TEST_REPORTS_DIR/profiler_traces" + profiler_trace_flags=(--export-profiler-trace --profiler-trace-name "$TEST_REPORTS_DIR/profiler_traces/${name}_${suite}") + fi python "benchmarks/dynamo/$suite.py" \ --ci --performance --disable-cudagraphs --inductor \ "${DYNAMO_BENCHMARK_FLAGS[@]}" "$@" "${partition_flags[@]}" \ + "${profiler_trace_flags[@]}" \ --output "$TEST_REPORTS_DIR/${name}_${suite}.csv" elif [[ "${TEST_CONFIG}" == *perf* ]]; then test_perf_for_dashboard "$suite" \ @@ -900,6 +1020,49 @@ test_inductor_triton_cpu() { assert_git_not_dirty } +setup_torch_trace() { + if [[ "${ENABLE_TORCH_TRACE:-0}" != "1" ]]; then + return + fi + local trace_dir="${RUNNER_TEMP:-/tmp}/torch_traces" + mkdir -p "$trace_dir" + export TORCH_TRACE="$trace_dir" + echo "TORCH_TRACE enabled: writing structured trace logs to $trace_dir" +} + +collect_tlparse_output() { + if [[ "${ENABLE_TORCH_TRACE:-0}" != "1" ]]; then + return + fi + local trace_dir="${RUNNER_TEMP:-/tmp}/torch_traces" + local test_reports_dir + test_reports_dir=$(pwd)/test/test-reports + + if [[ ! -d "$trace_dir" ]] || [[ -z "$(ls -A "$trace_dir" 2>/dev/null)" ]]; then + echo "No torch trace files found in $trace_dir, skipping tlparse" + return + fi + + echo "Collecting tlparse output from $trace_dir" + + # Install tlparse if not already available + if ! command -v tlparse &>/dev/null; then + pip install tlparse 2>/dev/null || { + echo "Warning: failed to install tlparse, skipping HTML generation" + return + } + fi + + # Run tlparse to generate HTML report + mkdir -p "$test_reports_dir/tlparse_output" + tlparse -o "$test_reports_dir/tlparse_output/" --no-browser --overwrite "$trace_dir" 2>&1 || { + echo "Warning: tlparse failed to generate HTML output" + return + } + + echo "TLParse output generated in $test_reports_dir/tlparse_output/" +} + test_dynamo_benchmark() { # Usage: test_dynamo_benchmark huggingface 0 TEST_REPORTS_DIR=$(pwd)/test/test-reports @@ -982,6 +1145,147 @@ test_inductor_torchbench_smoketest_perf() { done } +test_unbacked_parity_smoketest() { + # Check that unbacked batch-only has performance parity with backed batch-only + # Fails if any model regresses >THRESHOLD% consistently across 3 retries + TEST_REPORTS_DIR=$(pwd)/test/test-reports + mkdir -p "$TEST_REPORTS_DIR" + + local THRESHOLD=1.0 + local MAX_RETRIES=3 + local MODELS="MobileBertForMaskedLM|DistilBertForMaskedLM|DistillGPT2|T5Small" + + # Issue 6: Write per-run output files for post-failure debugging + run_comparison() { + local run_num=$1 + local output_file="$TEST_REPORTS_DIR/unbacked_parity_results_run${run_num}.txt" + python benchmarks/dynamo/huggingface.py \ + --compare-backed-unbacked \ + --performance --inference --inductor --device cuda \ + --filter "$MODELS" 2>&1 | tee "$output_file" + } + + check_regressions() { + local run_num=$1 + local output_file="$TEST_REPORTS_DIR/unbacked_parity_results_run${run_num}.txt" + # Parse the comparison table and check for regressions > threshold + # Returns 0 if regressions found, 1 if no regressions + local regressions=() + while IFS= read -r line; do + # Issue 3: Broadened regex to match model names with hyphens, slashes, dots + # Match lines like: " ModelName 10.000 10.500 +5.0%" + if [[ "$line" =~ ^[[:space:]]+([A-Za-z0-9_./-]+)[[:space:]]+([0-9.]+)[[:space:]]+([0-9.]+)[[:space:]]+\+([0-9.]+)% ]]; then + local model="${BASH_REMATCH[1]}" + local diff="${BASH_REMATCH[4]}" + # Nit: Use awk instead of bc -l to avoid dependency on bc + if awk "BEGIN{exit !($diff > $THRESHOLD)}"; then + regressions+=("$model:+${diff}%") + fi + fi + done < "$output_file" + + if [[ ${#regressions[@]} -gt 0 ]]; then + echo "Regressions found: ${regressions[*]}" + return 0 + fi + return 1 + } + + check_failures() { + local run_num=$1 + local output_file="$TEST_REPORTS_DIR/unbacked_parity_results_run${run_num}.txt" + # Issue 2: Check for any model failure — not just paired failures. + # Specifically flags when unbacked fails but backed succeeds (regression signal). + # Returns 0 if failures found, 1 if no failures + local current_model="" + local backed_failed=false + local unbacked_failed=false + local both_failures=() + local unbacked_only_failures=() + + # Append a sentinel header so the loop naturally evaluates the last real model + while IFS= read -r line; do + if [[ "$line" =~ ^---[[:space:]]+([A-Za-z0-9_./-]+)[[:space:]]+--- ]]; then + if [[ -n "$current_model" ]]; then + if $backed_failed && $unbacked_failed; then + both_failures+=("$current_model") + elif $unbacked_failed && ! $backed_failed; then + unbacked_only_failures+=("$current_model") + fi + fi + current_model="${BASH_REMATCH[1]}" + backed_failed=false + unbacked_failed=false + elif [[ "$line" =~ backed.*FAILED|backed.*TIMEOUT|backed.*ERROR ]]; then + backed_failed=true + elif [[ "$line" =~ unbacked.*FAILED|unbacked.*TIMEOUT|unbacked.*ERROR ]]; then + unbacked_failed=true + fi + done < <(cat "$output_file"; echo "--- END ---") + + local has_failures=false + if [[ ${#both_failures[@]} -gt 0 ]]; then + echo "❌ FAILURES DETECTED: Both backed and unbacked failed for: ${both_failures[*]}" + has_failures=true + fi + if [[ ${#unbacked_only_failures[@]} -gt 0 ]]; then + echo "❌ FAILURES DETECTED: Unbacked failed (but backed succeeded) for: ${unbacked_only_failures[*]}" + has_failures=true + fi + + if $has_failures; then + return 0 + fi + return 1 + } + + # Run initial comparison + echo "=== Run 1/$MAX_RETRIES ===" + run_comparison 1 + + # Check for failures first + if check_failures 1; then + echo "❌ Test failed: Models failed to run (see above for details)" + exit 1 + fi + + # Check for regressions + if ! check_regressions 1; then + echo "✅ PASSED: No regressions above ${THRESHOLD}% threshold" + exit 0 + fi + + # Regression detected - retry to confirm + local regression_count=1 + for ((retry=2; retry<=MAX_RETRIES; retry++)); do + echo "" + echo "=== Retry $retry/$MAX_RETRIES (potential regression detected) ===" + run_comparison "$retry" + + # Issue 4: Also check for failures on retries (e.g., intermittent OOM) + if check_failures "$retry"; then + echo "❌ Test failed: Models failed on retry $retry (see above for details)" + exit 1 + fi + + if check_regressions "$retry"; then + ((regression_count++)) + fi + done + + # Check if regression was consistent (majority of runs) + local required=$((MAX_RETRIES / 2 + 1)) + if [[ $regression_count -ge $required ]]; then + echo "" + echo "❌ REGRESSION CONFIRMED: Detected in $regression_count/$MAX_RETRIES runs (threshold: ${THRESHOLD}%)" + exit 1 + else + echo "" + echo "✅ PASSED: Regressions were not consistent ($regression_count/$MAX_RETRIES runs, needed $required)" + exit 0 + fi +} + test_inductor_set_cpu_affinity(){ JEMALLOC_LIB="$(find /usr/lib -name libjemalloc.so.2)" export LD_PRELOAD="$JEMALLOC_LIB":"$LD_PRELOAD" @@ -1156,6 +1460,18 @@ test_libtorch_jit() { popd } +test_libtorch_profiler() { + echo "Testing profiler C++ tests" + export CPP_TESTS_DIR="${TORCH_BIN_DIR}" + export LD_LIBRARY_PATH="${TORCH_LIB_DIR}:${LD_LIBRARY_PATH}" + + # Run E2E test first (needs clean Kineto state) + python test/run_test.py --cpp --verbose -i cpp/test_privateuse1_profiler -k "EndToEndProfiling" + + # Run all other tests + python test/run_test.py --cpp --verbose -i cpp/test_privateuse1_profiler -k "not EndToEndProfiling" +} + test_libtorch_api() { # Start background download MNIST_DIR="${PWD}/test/cpp/api/mnist" @@ -1625,74 +1941,6 @@ EOF assert_git_not_dirty } -test_bazel() { - set -e -o pipefail - - # bazel test needs sccache setup. - # shellcheck source=./common-build.sh - source "$(dirname "${BASH_SOURCE[0]}")/common-build.sh" - - get_bazel - - if [[ "$CUDA_VERSION" == "cpu" ]]; then - # Test //c10/... without Google flags and logging libraries. The - # :all_tests target in the subsequent Bazel invocation tests - # //c10/... with the Google libraries. - tools/bazel test --config=cpu-only --test_timeout=480 --test_output=all --test_tag_filters=-gpu-required --test_filter=-*CUDA \ - --no//c10:use_gflags --no//c10:use_glog //c10/... - - tools/bazel test --config=cpu-only --test_timeout=480 --test_output=all --test_tag_filters=-gpu-required --test_filter=-*CUDA :all_tests - else - # Increase the test timeout to 480 like CPU tests because modules_test frequently timeout - tools/bazel test --test_timeout=480 --test_output=errors \ - //:any_test \ - //:autograd_test \ - //:dataloader_test \ - //:dispatch_test \ - //:enum_test \ - //:expanding_array_test \ - //:fft_test \ - //:functional_test \ - //:grad_mode_test \ - //:inference_mode_test \ - //:init_test \ - //:jit_test \ - //:memory_test \ - //:meta_tensor_test \ - //:misc_test \ - //:moduledict_test \ - //:modulelist_test \ - //:modules_test \ - //:namespace_test \ - //:nested_test \ - //:nn_utils_test \ - //:operations_test \ - //:ordered_dict_test \ - //:parallel_benchmark_test \ - //:parameterdict_test \ - //:parameterlist_test \ - //:sequential_test \ - //:serialize_test \ - //:special_test \ - //:static_test \ - //:support_test \ - //:tensor_flatten_test \ - //:tensor_indexing_test \ - //:tensor_options_cuda_test \ - //:tensor_options_test \ - //:tensor_test \ - //:torch_dist_autograd_test \ - //:torch_include_test \ - //:transformer_test \ - //:test_bazel \ - //c10/cuda/test:test \ - //c10/test:core_tests \ - //c10/test:typeid_test \ - //c10/test:util/ssize_test \ - //c10/test:util_base_tests - fi -} - test_benchmarks() { if [[ "$BUILD_ENVIRONMENT" == *cuda* && $TEST_CONFIG != *nogpu* ]]; then pip_install "pytest-benchmark==3.2.3" @@ -1763,34 +2011,6 @@ test_executorch() { assert_git_not_dirty } -test_linux_aarch64() { - python test/run_test.py --include test_modules test_utils test_mkldnn test_mkldnn_fusion test_openmp test_torch test_dynamic_shapes \ - test_transformers test_multiprocessing test_numpy_interop test_autograd test_binary_ufuncs test_complex test_spectral_ops \ - test_foreach test_reductions test_unary_ufuncs test_tensor_creation_ops test_ops profiler/test_memory_profiler \ - distributed/elastic/timer/api_test distributed/elastic/timer/local_timer_example distributed/elastic/timer/local_timer_test \ - test_linalg \ - --shard "$SHARD_NUMBER" "$NUM_TEST_SHARDS" --verbose - - # Dynamo tests - python test/run_test.py --include dynamo/test_compile dynamo/test_backends dynamo/test_comptime dynamo/test_config \ - dynamo/test_functions dynamo/test_fx_passes_pre_grad dynamo/test_interop dynamo/test_model_output dynamo/test_modules \ - dynamo/test_optimizers dynamo/test_recompile_ux dynamo/test_recompiles \ - --shard "$SHARD_NUMBER" "$NUM_TEST_SHARDS" --verbose - - # Inductor tests - python test/run_test.py --include inductor/test_torchinductor inductor/test_benchmark_fusion inductor/test_codecache \ - inductor/test_config inductor/test_control_flow inductor/test_coordinate_descent_tuner inductor/test_fx_fusion \ - inductor/test_group_batch_fusion inductor/test_inductor_freezing inductor/test_inductor_utils \ - inductor/test_inplacing_pass inductor/test_kernel_benchmark inductor/test_layout_optim \ - inductor/test_max_autotune inductor/test_memory_planning inductor/test_metrics inductor/test_multi_kernel inductor/test_pad_mm \ - inductor/test_pattern_matcher inductor/test_perf inductor/test_profiler inductor/test_select_algorithm inductor/test_smoke \ - inductor/test_split_cat_fx_passes inductor/test_compile inductor/test_torchinductor \ - inductor/test_torchinductor_codegen_dynamic_shapes inductor/test_torchinductor_dynamic_shapes inductor/test_memory \ - inductor/test_triton_cpu_backend inductor/test_triton_extension_backend inductor/test_mkldnn_pattern_matcher inductor/test_cpu_cpp_wrapper \ - inductor/test_cpu_select_algorithm \ - --shard "$SHARD_NUMBER" "$NUM_TEST_SHARDS" --verbose -} - test_operator_benchmark() { TEST_REPORTS_DIR=$(pwd)/test/test-reports mkdir -p "$TEST_REPORTS_DIR" @@ -1853,15 +2073,19 @@ test_attention_microbenchmark() { } test_openreg() { + git submodule update --init --depth 1 third_party/googletest python test/run_test.py --openreg --verbose assert_git_not_dirty } -if ! [[ "${BUILD_ENVIRONMENT}" == *libtorch* || "${BUILD_ENVIRONMENT}" == *-bazel-* ]]; then +if ! [[ "${BUILD_ENVIRONMENT}" == *libtorch* ]]; then (cd test && python -c "import torch; print(torch.__config__.show())") (cd test && python -c "import torch; print(torch.__config__.parallel_info())") fi -if [[ "${TEST_CONFIG}" == *numpy_2* ]]; then +if [[ "${TEST_CONFIG}" == "onnx" ]]; then + install_torchvision + "$(dirname "${BASH_SOURCE[0]}")/../../scripts/onnx/test.sh" +elif [[ "${TEST_CONFIG}" == *numpy_2* ]]; then # Install numpy-2.0.2 and compatible scipy & numba versions # Force re-install of pandas to avoid error where pandas checks numpy version from initial install and fails upon import TMP_PANDAS_VERSION=$(python -c "import pandas; print(pandas.__version__)" 2>/dev/null) @@ -1871,8 +2095,6 @@ if [[ "${TEST_CONFIG}" == *numpy_2* ]]; then python -m pip install --pre numpy==2.0.2 scipy==1.13.1 numba==0.60.0 fi python test/run_test.py --include dynamo/test_functions.py dynamo/test_unspec.py test_binary_ufuncs.py test_fake_tensor.py test_linalg.py test_numpy_interop.py test_tensor_creation_ops.py test_torch.py torch_np/test_basic.py -elif [[ "${BUILD_ENVIRONMENT}" == *aarch64* && "${TEST_CONFIG}" == 'default' ]]; then - test_linux_aarch64 elif [[ "${TEST_CONFIG}" == *backward* ]]; then test_forward_backward_compatibility # Do NOT add tests after bc check tests, see its comment. @@ -1887,6 +2109,9 @@ elif [[ "$TEST_CONFIG" == *vllm* ]]; then (cd .ci/lumen_cli && python -m pip install -e .) python -m cli.run test external vllm --test-plan "$TEST_CONFIG" --shard-id "$SHARD_NUMBER" --num-shards "$NUM_TEST_SHARDS" +elif [[ "$TEST_CONFIG" == *torchtitan* ]]; then + (cd .ci/lumen_cli && python -m pip install -e .) + python -m cli.run test external torchtitan --test-plan "$TEST_CONFIG" --shard-id "$SHARD_NUMBER" --num-shards "$NUM_TEST_SHARDS" elif [[ "${TEST_CONFIG}" == *executorch* ]]; then test_executorch elif [[ "$TEST_CONFIG" == 'jit_legacy' ]]; then @@ -1920,7 +2145,9 @@ elif [[ "${TEST_CONFIG}" == *operator_microbenchmark* ]]; then elif [[ "${TEST_CONFIG}" == *attention_microbenchmark* ]]; then test_attention_microbenchmark elif [[ "${TEST_CONFIG}" == *inductor_distributed* ]]; then + setup_torch_trace test_inductor_distributed + collect_tlparse_output elif [[ "${TEST_CONFIG}" == *inductor-halide* ]]; then test_inductor_halide elif [[ "${TEST_CONFIG}" == *inductor-pallas* ]]; then @@ -1934,17 +2161,19 @@ elif [[ "${TEST_CONFIG}" == *aoti_cross_compile_for_windows* ]]; then elif [[ "${TEST_CONFIG}" == *huggingface* ]]; then install_torchvision id=$((SHARD_NUMBER-1)) - test_dynamo_benchmark huggingface "$id" + setup_torch_trace + if [[ "${TEST_CONFIG}" == *unbacked_parity* ]]; then + test_unbacked_parity_smoketest + else + test_dynamo_benchmark huggingface "$id" + fi + collect_tlparse_output elif [[ "${TEST_CONFIG}" == *timm* ]]; then install_torchvision - TIMM_PIN="$(< .ci/docker/ci_commit_pins/timm.txt)" - export HF_HOME="${HF_HOME}/timm_${TIMM_PIN}" - if [[ "${TRANSFORMERS_OFFLINE:-1}" == "0" ]]; then - python benchmarks/dynamo/timm_models.py --download-only \ - && touch "${HF_HOME}/.timm_cache_complete" - fi id=$((SHARD_NUMBER-1)) + setup_torch_trace test_dynamo_benchmark timm_models "$id" + collect_tlparse_output elif [[ "${TEST_CONFIG}" == cachebench ]]; then install_torchaudio install_torchvision @@ -1975,19 +2204,27 @@ elif [[ "${TEST_CONFIG}" == *torchbench* ]]; then LIBTBB_PATH="$(find "$(dirname "$(which python)")/../lib/" -name libtbb.so.12)" export LD_PRELOAD="$LIBTBB_PATH":"$LD_PRELOAD" fi + setup_torch_trace PYTHONPATH=/torchbench test_dynamo_benchmark torchbench "$id" + collect_tlparse_output fi elif [[ "${TEST_CONFIG}" == *inductor_cpp_wrapper* ]]; then install_torchvision + setup_torch_trace PYTHONPATH=/torchbench test_inductor_cpp_wrapper_shard "$SHARD_NUMBER" if [[ "$SHARD_NUMBER" -eq "1" ]]; then test_inductor_aoti_cpp fi + collect_tlparse_output elif [[ "${TEST_CONFIG}" == *inductor_core* ]]; then + setup_torch_trace test_inductor_core + collect_tlparse_output elif [[ "${TEST_CONFIG}" == *inductor* ]]; then install_torchvision + setup_torch_trace test_inductor_shard "${SHARD_NUMBER}" + collect_tlparse_output elif [[ "${TEST_CONFIG}" == *einops* ]]; then test_einops elif [[ "${TEST_CONFIG}" == *dynamo_core* ]]; then @@ -2022,27 +2259,25 @@ elif [[ "${SHARD_NUMBER}" == 2 && $NUM_TEST_SHARDS -gt 1 ]]; then test_custom_script_ops test_custom_backend test_torch_function_benchmark + test_libtorch_profiler elif [[ "${SHARD_NUMBER}" -gt 2 ]]; then # Handle arbitrary number of shards install_torchvision test_python_shard "$SHARD_NUMBER" elif [[ "${BUILD_ENVIRONMENT}" == *vulkan* ]]; then test_vulkan -elif [[ "${BUILD_ENVIRONMENT}" == *-bazel-* ]]; then - test_bazel elif [[ "${BUILD_ENVIRONMENT}" == *-mobile-lightweight-dispatch* ]]; then test_libtorch elif [[ "${TEST_CONFIG}" = docs_test ]]; then test_docs_test -elif [[ "${BUILD_ENVIRONMENT}" == *xpu* ]]; then - install_torchvision - test_python - test_aten - test_xpu_bin elif [[ "${TEST_CONFIG}" == smoke ]]; then test_python_smoke elif [[ "${TEST_CONFIG}" == smoke_b200 ]]; then test_python_smoke_b200 +elif [[ "${TEST_CONFIG}" == smoke_xpu ]]; then + test_python_smoke_xpu +elif [[ "${TEST_CONFIG}" == dtensor ]]; then + test_dtensor elif [[ "${TEST_CONFIG}" == h100_distributed ]]; then test_h100_distributed elif [[ "${TEST_CONFIG}" == "h100-symm-mem" ]]; then diff --git a/.ci/pytorch/test_example_code/CMakeLists.txt b/.ci/pytorch/test_example_code/CMakeLists.txt index e87f37ae61fb4..c89a728f27649 100644 --- a/.ci/pytorch/test_example_code/CMakeLists.txt +++ b/.ci/pytorch/test_example_code/CMakeLists.txt @@ -1,4 +1,4 @@ -cmake_minimum_required(VERSION 3.0 FATAL_ERROR) +cmake_minimum_required(VERSION 3.10 FATAL_ERROR) project(simple-torch-test) find_package(Torch REQUIRED) @@ -8,7 +8,7 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${TORCH_CXX_FLAGS}") add_executable(simple-torch-test simple-torch-test.cpp) target_include_directories(simple-torch-test PRIVATE ${TORCH_INCLUDE_DIRS}) target_link_libraries(simple-torch-test "${TORCH_LIBRARIES}") -set_property(TARGET simple-torch-test PROPERTY CXX_STANDARD 17) +set_property(TARGET simple-torch-test PROPERTY CXX_STANDARD 20) find_package(CUDAToolkit 11.8) diff --git a/.ci/pytorch/win-test-helpers/test_openreg.bat b/.ci/pytorch/win-test-helpers/test_openreg.bat index 0470057daf641..db3c0758f44e6 100644 --- a/.ci/pytorch/win-test-helpers/test_openreg.bat +++ b/.ci/pytorch/win-test-helpers/test_openreg.bat @@ -6,6 +6,8 @@ if not errorlevel 0 ( exit /b ) +git submodule update --init --depth 1 third_party/googletest + pushd test echo Run openreg tests diff --git a/.ci/pytorch/windows/arm64/build_libtorch.bat b/.ci/pytorch/windows/arm64/build_libtorch.bat index 1ac14ff697730..33080af122c73 100644 --- a/.ci/pytorch/windows/arm64/build_libtorch.bat +++ b/.ci/pytorch/windows/arm64/build_libtorch.bat @@ -27,8 +27,8 @@ where cl.exe :: change to source directory cd %PYTORCH_ROOT% -:: copy libuv.dll -copy %libuv_ROOT%\lib\Release\uv.dll torch\lib\uv.dll +:: copy libuv.dll (cmake installs the dll to bin/, not lib/Release/) +copy %libuv_ROOT%\bin\uv.dll torch\lib\uv.dll :: create virtual environment python -m venv .venv diff --git a/.ci/pytorch/windows/arm64/build_pytorch.bat b/.ci/pytorch/windows/arm64/build_pytorch.bat index b5c2ef65b84ad..7d10b26339d25 100644 --- a/.ci/pytorch/windows/arm64/build_pytorch.bat +++ b/.ci/pytorch/windows/arm64/build_pytorch.bat @@ -5,6 +5,7 @@ set CMAKE_BUILD_TYPE=%BUILD_TYPE% set CMAKE_C_COMPILER_LAUNCHER=sccache set CMAKE_CXX_COMPILER_LAUNCHER=sccache set libuv_ROOT=%DEPENDENCIES_DIR%\libuv\install +set INSTALL_TEST=0 set MSSdk=1 if defined PYTORCH_BUILD_VERSION ( set PYTORCH_BUILD_VERSION=%PYTORCH_BUILD_VERSION% @@ -27,8 +28,8 @@ where cl.exe :: change to source directory cd %PYTORCH_ROOT% -:: copy libuv.dll -copy %libuv_ROOT%\lib\Release\uv.dll torch\lib\uv.dll +:: copy libuv.dll (cmake installs the dll to bin/, not lib/Release/) +copy %libuv_ROOT%\bin\uv.dll torch\lib\uv.dll :: create virtual environment python -m venv .venv @@ -57,4 +58,4 @@ sccache --show-stats if %errorlevel% neq 0 ( echo "Failed on build_pytorch. (exitcode = %errorlevel%)" exit /b 1 -) \ No newline at end of file +) diff --git a/.ci/pytorch/windows/arm64/smoke_test.bat b/.ci/pytorch/windows/arm64/smoke_test.bat index 2c90f6158062f..202f6eb8cc9c5 100644 --- a/.ci/pytorch/windows/arm64/smoke_test.bat +++ b/.ci/pytorch/windows/arm64/smoke_test.bat @@ -40,10 +40,10 @@ set INCLUDE=%INCLUDE%;%install_root%\include;%install_root%\include\torch\csrc\a set LIB=%LIB%;%install_root%\lib set PATH=%PATH%;%install_root%\lib -cl %PYTORCH_ROOT%\.ci\pytorch\test_example_code\simple-torch-test.cpp c10.lib torch_cpu.lib /EHsc /std:c++17 +cl %PYTORCH_ROOT%\.ci\pytorch\test_example_code\simple-torch-test.cpp c10.lib torch_cpu.lib /EHsc /std:c++20 if ERRORLEVEL 1 exit /b 1 .\simple-torch-test.exe if ERRORLEVEL 1 exit /b 1 -:end \ No newline at end of file +:end diff --git a/.ci/pytorch/windows/internal/cuda_config.bat b/.ci/pytorch/windows/internal/cuda_config.bat index 352eb8a3391bf..69059e2e46ee1 100644 --- a/.ci/pytorch/windows/internal/cuda_config.bat +++ b/.ci/pytorch/windows/internal/cuda_config.bat @@ -22,6 +22,10 @@ if "%CUDA_VER%"=="126" ( set "CUDA_DOTTED_VERSION=13.0" set "CUDA_ARCH_LIST=7.5;8.0;8.6;9.0;10.0;12.0" set "VISION_GENCODE=-gencode=arch=compute_75,code=sm_75 -gencode=arch=compute_80,code=compute_80 -gencode=arch=compute_86,code=compute_86 -gencode=arch=compute_90,code=compute_90 -gencode=arch=compute_100,code=compute_100 -gencode=arch=compute_120,code=compute_120" +) else if "%CUDA_VER%"=="132" ( + set "CUDA_DOTTED_VERSION=13.2" + set "CUDA_ARCH_LIST=7.5;8.0;8.6;9.0;10.0;12.0" + set "VISION_GENCODE=-gencode=arch=compute_75,code=sm_75 -gencode=arch=compute_80,code=compute_80 -gencode=arch=compute_86,code=compute_86 -gencode=arch=compute_90,code=compute_90 -gencode=arch=compute_100,code=compute_100 -gencode=arch=compute_120,code=compute_120" ) else ( echo Unknown CUDA version: %CUDA_VER% exit /b 1 diff --git a/.ci/pytorch/windows/internal/cuda_install.bat b/.ci/pytorch/windows/internal/cuda_install.bat index 1349d3e661f55..c1050edecc0b9 100644 --- a/.ci/pytorch/windows/internal/cuda_install.bat +++ b/.ci/pytorch/windows/internal/cuda_install.bat @@ -17,16 +17,17 @@ set /a CUDA_VER=%CUDA_VERSION% set CUDA_VER_MAJOR=%CUDA_VERSION:~0,-1% set CUDA_VER_MINOR=%CUDA_VERSION:~-1,1% set CUDA_VERSION_STR=%CUDA_VER_MAJOR%.%CUDA_VER_MINOR% -set CUDNN_FOLDER="cuda" -set CUDNN_LIB_FOLDER="lib\x64" +set CUDNN_FOLDER=cuda +set CUDNN_LIB_FOLDER=lib\x64 -:: Skip all of this if we already have cuda installed -if exist "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\bin\nvcc.exe" goto set_cuda_env_vars +:: If CUDA is already installed, skip CUDA installation but still verify cuDNN +if exist "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\bin\nvcc.exe" goto check_cudnn if %CUDA_VER% EQU 126 goto cuda126 if %CUDA_VER% EQU 128 goto cuda128 if %CUDA_VER% EQU 129 goto cuda129 if %CUDA_VER% EQU 130 goto cuda130 +if %CUDA_VER% EQU 132 goto cuda132 echo CUDA %CUDA_VERSION_STR% is not supported exit /b 1 @@ -34,110 +35,53 @@ exit /b 1 goto cuda_common :cuda126 - set CUDA_INSTALL_EXE=cuda_12.6.2_560.94_windows.exe -if not exist "%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" ( - curl -k -L "https://ossci-windows.s3.amazonaws.com/%CUDA_INSTALL_EXE%" --output "%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" & REM @lint-ignore - if errorlevel 1 exit /b 1 - set "CUDA_SETUP_FILE=%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" - set "ARGS=cuda_profiler_api_12.6 thrust_12.6 nvcc_12.6 cuobjdump_12.6 nvprune_12.6 nvprof_12.6 cupti_12.6 cublas_12.6 cublas_dev_12.6 cudart_12.6 cufft_12.6 cufft_dev_12.6 curand_12.6 curand_dev_12.6 cusolver_12.6 cusolver_dev_12.6 cusparse_12.6 cusparse_dev_12.6 npp_12.6 npp_dev_12.6 nvrtc_12.6 nvrtc_dev_12.6 nvml_dev_12.6 nvjitlink_12.6 nvtx_12.6" -) - -set CUDNN_FOLDER=cudnn-windows-x86_64-9.5.0.50_cuda12-archive -set CUDNN_LIB_FOLDER="lib" -set "CUDNN_INSTALL_ZIP=%CUDNN_FOLDER%.zip" -if not exist "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" ( - curl -k -L "http://s3.amazonaws.com/ossci-windows/%CUDNN_INSTALL_ZIP%" --output "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" & REM @lint-ignore - if errorlevel 1 exit /b 1 - set "CUDNN_SETUP_FILE=%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" -) - -@REM cuDNN 8.3+ required zlib to be installed on the path -echo Installing ZLIB dlls -curl -k -L "http://s3.amazonaws.com/ossci-windows/zlib123dllx64.zip" --output "%SRC_DIR%\temp_build\zlib123dllx64.zip" -7z x "%SRC_DIR%\temp_build\zlib123dllx64.zip" -o"%SRC_DIR%\temp_build\zlib" -xcopy /Y "%SRC_DIR%\temp_build\zlib\dll_x64\*.dll" "C:\Windows\System32" - -goto cuda_common +set "ARGS=cuda_profiler_api_12.6 thrust_12.6 nvcc_12.6 cuobjdump_12.6 nvprune_12.6 nvprof_12.6 cupti_12.6 cublas_12.6 cublas_dev_12.6 cudart_12.6 cufft_12.6 cufft_dev_12.6 curand_12.6 curand_dev_12.6 cusolver_12.6 cusolver_dev_12.6 cusparse_12.6 cusparse_dev_12.6 npp_12.6 npp_dev_12.6 nvrtc_12.6 nvrtc_dev_12.6 nvml_dev_12.6 nvjitlink_12.6 nvtx_12.6" +set CUDNN_FOLDER=cudnn-windows-x86_64-9.10.2.21_cuda12-archive +goto cuda_download :cuda128 - set CUDA_INSTALL_EXE=cuda_12.8.0_571.96_windows.exe -if not exist "%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" ( - curl -k -L "https://ossci-windows.s3.amazonaws.com/%CUDA_INSTALL_EXE%" --output "%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" & REM @lint-ignore - if errorlevel 1 exit /b 1 - set "CUDA_SETUP_FILE=%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" - set "ARGS=cuda_profiler_api_12.8 thrust_12.8 nvcc_12.8 cuobjdump_12.8 nvprune_12.8 nvprof_12.8 cupti_12.8 cublas_12.8 cublas_dev_12.8 cudart_12.8 cufft_12.8 cufft_dev_12.8 curand_12.8 curand_dev_12.8 cusolver_12.8 cusolver_dev_12.8 cusparse_12.8 cusparse_dev_12.8 npp_12.8 npp_dev_12.8 nvrtc_12.8 nvrtc_dev_12.8 nvml_dev_12.8 nvjitlink_12.8 nvtx_12.8" -) - -set CUDNN_FOLDER=cudnn-windows-x86_64-9.7.0.66_cuda12-archive -set CUDNN_LIB_FOLDER="lib" -set "CUDNN_INSTALL_ZIP=%CUDNN_FOLDER%.zip" -if not exist "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" ( - curl -k -L "http://s3.amazonaws.com/ossci-windows/%CUDNN_INSTALL_ZIP%" --output "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" & REM @lint-ignore - if errorlevel 1 exit /b 1 - set "CUDNN_SETUP_FILE=%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" -) - -@REM cuDNN 8.3+ required zlib to be installed on the path -echo Installing ZLIB dlls -curl -k -L "http://s3.amazonaws.com/ossci-windows/zlib123dllx64.zip" --output "%SRC_DIR%\temp_build\zlib123dllx64.zip" -7z x "%SRC_DIR%\temp_build\zlib123dllx64.zip" -o"%SRC_DIR%\temp_build\zlib" -xcopy /Y "%SRC_DIR%\temp_build\zlib\dll_x64\*.dll" "C:\Windows\System32" - -goto cuda_common +set "ARGS=cuda_profiler_api_12.8 thrust_12.8 nvcc_12.8 cuobjdump_12.8 nvprune_12.8 nvprof_12.8 cupti_12.8 cublas_12.8 cublas_dev_12.8 cudart_12.8 cufft_12.8 cufft_dev_12.8 curand_12.8 curand_dev_12.8 cusolver_12.8 cusolver_dev_12.8 cusparse_12.8 cusparse_dev_12.8 npp_12.8 npp_dev_12.8 nvrtc_12.8 nvrtc_dev_12.8 nvml_dev_12.8 nvjitlink_12.8 nvtx_12.8" +set CUDNN_FOLDER=cudnn-windows-x86_64-9.20.0.48_cuda12-archive +goto cuda_download :cuda129 - set CUDA_INSTALL_EXE=cuda_12.9.1_576.57_windows.exe -if not exist "%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" ( - curl -k -L "https://ossci-windows.s3.amazonaws.com/%CUDA_INSTALL_EXE%" --output "%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" & REM @lint-ignore - if errorlevel 1 exit /b 1 - set "CUDA_SETUP_FILE=%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" - set "ARGS=cuda_profiler_api_12.9 thrust_12.9 nvcc_12.9 cuobjdump_12.9 nvprune_12.9 nvprof_12.9 cupti_12.9 cublas_12.9 cublas_dev_12.9 cudart_12.9 cufft_12.9 cufft_dev_12.9 curand_12.9 curand_dev_12.9 cusolver_12.9 cusolver_dev_12.9 cusparse_12.9 cusparse_dev_12.9 npp_12.9 npp_dev_12.9 nvrtc_12.9 nvrtc_dev_12.9 nvml_dev_12.9 nvjitlink_12.9 nvtx_12.9" -) - -set CUDNN_FOLDER=cudnn-windows-x86_64-9.10.2.21_cuda12-archive -set CUDNN_LIB_FOLDER="lib" -set "CUDNN_INSTALL_ZIP=%CUDNN_FOLDER%.zip" -if not exist "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" ( - curl -k -L "http://s3.amazonaws.com/ossci-windows/%CUDNN_INSTALL_ZIP%" --output "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" & REM @lint-ignore - if errorlevel 1 exit /b 1 - set "CUDNN_SETUP_FILE=%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" -) - -@REM cuDNN 8.3+ required zlib to be installed on the path -echo Installing ZLIB dlls -curl -k -L "http://s3.amazonaws.com/ossci-windows/zlib123dllx64.zip" --output "%SRC_DIR%\temp_build\zlib123dllx64.zip" -7z x "%SRC_DIR%\temp_build\zlib123dllx64.zip" -o"%SRC_DIR%\temp_build\zlib" -xcopy /Y "%SRC_DIR%\temp_build\zlib\dll_x64\*.dll" "C:\Windows\System32" - -goto cuda_common +set "ARGS=cuda_profiler_api_12.9 thrust_12.9 nvcc_12.9 cuobjdump_12.9 nvprune_12.9 nvprof_12.9 cupti_12.9 cublas_12.9 cublas_dev_12.9 cudart_12.9 cufft_12.9 cufft_dev_12.9 curand_12.9 curand_dev_12.9 cusolver_12.9 cusolver_dev_12.9 cusparse_12.9 cusparse_dev_12.9 npp_12.9 npp_dev_12.9 nvrtc_12.9 nvrtc_dev_12.9 nvml_dev_12.9 nvjitlink_12.9 nvtx_12.9" +set CUDNN_FOLDER=cudnn-windows-x86_64-9.20.0.48_cuda12-archive +goto cuda_download :cuda130 - set CUDA_INSTALL_EXE=cuda_13.0.0_windows.exe +set "ARGS=" +set CUDNN_FOLDER=cudnn-windows-x86_64-9.20.0.48_cuda13-archive +goto cuda_download + +:cuda132 +set CUDA_INSTALL_EXE=cuda_13.2.1_windows.exe +set "ARGS=" +set CUDNN_FOLDER=cudnn-windows-x86_64-9.20.0.48_cuda13-archive +goto cuda_download + +:: Common download logic for CUDA toolkit, cuDNN, and ZLIB +:cuda_download +set CUDNN_LIB_FOLDER=lib +set "CUDNN_INSTALL_ZIP=%CUDNN_FOLDER%.zip" + if not exist "%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" ( curl -k -L "https://ossci-windows.s3.amazonaws.com/%CUDA_INSTALL_EXE%" --output "%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" & REM @lint-ignore if errorlevel 1 exit /b 1 set "CUDA_SETUP_FILE=%SRC_DIR%\temp_build\%CUDA_INSTALL_EXE%" - set "ARGS=" ) -set CUDNN_FOLDER=cudnn-windows-x86_64-9.12.0.46_cuda13-archive -set CUDNN_LIB_FOLDER="lib" -set "CUDNN_INSTALL_ZIP=%CUDNN_FOLDER%.zip" if not exist "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" ( curl -k -L "http://s3.amazonaws.com/ossci-windows/%CUDNN_INSTALL_ZIP%" --output "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" & REM @lint-ignore if errorlevel 1 exit /b 1 set "CUDNN_SETUP_FILE=%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" ) -@REM cuDNN 8.3+ required zlib to be installed on the path -echo Installing ZLIB dlls -curl -k -L "http://s3.amazonaws.com/ossci-windows/zlib123dllx64.zip" --output "%SRC_DIR%\temp_build\zlib123dllx64.zip" -7z x "%SRC_DIR%\temp_build\zlib123dllx64.zip" -o"%SRC_DIR%\temp_build\zlib" -xcopy /Y "%SRC_DIR%\temp_build\zlib\dll_x64\*.dll" "C:\Windows\System32" +call :install_zlib goto cuda_common @@ -189,9 +133,12 @@ if not exist "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_ echo Installing cuDNN... 7z x %CUDNN_SETUP_FILE% -o"%SRC_DIR%\temp_build\cudnn" - xcopy /Y "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\bin\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\bin" - xcopy /Y "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\%CUDNN_LIB_FOLDER%\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\lib\x64" - xcopy /Y "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\include\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\include" + xcopy /Y /S "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\bin\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\bin\" + if exist "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\bin\x64\*.*" ( + xcopy /Y "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\bin\x64\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\bin\" + ) + xcopy /Y /S "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\%CUDNN_LIB_FOLDER%\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\lib\x64\" + xcopy /Y /S "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\include\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\include\" echo Installing GPU driver DLLs 7z x %SRC_DIR%\temp_build\gpu_driver_dlls.zip -o"C:\Windows\System32" @@ -211,6 +158,85 @@ if not exist "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_ goto set_cuda_env_vars +:check_cudnn +:: When CUDA is pre-installed on the AMI, cuDNN may still be missing. +:: Set the correct cuDNN variables for the CUDA version, then install if needed. + +set CUDNN_LIB_FOLDER=lib +if %CUDA_VER% EQU 126 ( + set CUDNN_FOLDER=cudnn-windows-x86_64-9.10.2.21_cuda12-archive + set EXPECTED_CUDNN_VERSION=9.10.2 +) +if %CUDA_VER% EQU 128 ( + set CUDNN_FOLDER=cudnn-windows-x86_64-9.20.0.48_cuda12-archive + set EXPECTED_CUDNN_VERSION=9.20.0 +) +if %CUDA_VER% EQU 129 ( + set CUDNN_FOLDER=cudnn-windows-x86_64-9.20.0.48_cuda12-archive + set EXPECTED_CUDNN_VERSION=9.20.0 +) +if %CUDA_VER% EQU 130 ( + set CUDNN_FOLDER=cudnn-windows-x86_64-9.20.0.48_cuda13-archive + set EXPECTED_CUDNN_VERSION=9.20.0 +) +if %CUDA_VER% EQU 132 ( + set CUDNN_FOLDER=cudnn-windows-x86_64-9.20.0.48_cuda13-archive + set EXPECTED_CUDNN_VERSION=9.20.0 +) +set "CUDNN_INSTALL_ZIP=%CUDNN_FOLDER%.zip" + +set "CUDNN_VERSION_FILE=%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\include\cudnn_version.h" + +if not exist "%CUDNN_VERSION_FILE%" ( + echo cuDNN not found, installing %CUDNN_FOLDER%... + goto install_cudnn +) + +for /f "tokens=3" %%a in ('findstr /C:"#define CUDNN_MAJOR " "%CUDNN_VERSION_FILE%"') do set INSTALLED_MAJOR=%%a +for /f "tokens=3" %%a in ('findstr /C:"#define CUDNN_MINOR " "%CUDNN_VERSION_FILE%"') do set INSTALLED_MINOR=%%a +for /f "tokens=3" %%a in ('findstr /C:"#define CUDNN_PATCHLEVEL " "%CUDNN_VERSION_FILE%"') do set INSTALLED_PATCHLEVEL=%%a +set "INSTALLED_CUDNN_VERSION=%INSTALLED_MAJOR%.%INSTALLED_MINOR%.%INSTALLED_PATCHLEVEL%" + +if "%INSTALLED_CUDNN_VERSION%" == "%EXPECTED_CUDNN_VERSION%" ( + echo cuDNN %INSTALLED_CUDNN_VERSION% already installed at %ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR% + goto set_cuda_env_vars +) + +echo cuDNN version mismatch: installed %INSTALLED_CUDNN_VERSION%, expected %EXPECTED_CUDNN_VERSION%. Reinstalling... + +:: Remove old cuDNN DLLs so they don't shadow the new version at runtime. +:: AMI-installed cuDNN places DLLs directly in bin\, while newer archives +:: use bin\x64\. Without cleanup the old DLLs in bin\ are found first. +del /Q "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\bin\cudnn*.dll" 2>nul + +:install_cudnn + +if not exist "%SRC_DIR%\temp_build" mkdir "%SRC_DIR%\temp_build" + +curl -k -L "http://s3.amazonaws.com/ossci-windows/%CUDNN_INSTALL_ZIP%" --output "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" & REM @lint-ignore +if errorlevel 1 exit /b 1 + +7z x "%SRC_DIR%\temp_build\%CUDNN_INSTALL_ZIP%" -o"%SRC_DIR%\temp_build\cudnn" +if errorlevel 1 ( + echo Failed to extract cuDNN archive %CUDNN_INSTALL_ZIP% + exit /b 1 +) +echo Listing extracted cuDNN archive contents: +dir /S /B "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%" +xcopy /Y /S "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\bin\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\bin\" +:: Newer cuDNN archives place DLLs under bin\x64\. Flatten them into bin\ +:: so they are found via PATH (which only includes bin\, not bin\x64\). +if exist "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\bin\x64\*.*" ( + xcopy /Y "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\bin\x64\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\bin\" +) +xcopy /Y /S "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\%CUDNN_LIB_FOLDER%\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\lib\x64\" +xcopy /Y /S "%SRC_DIR%\temp_build\cudnn\%CUDNN_FOLDER%\include\*.*" "%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\include\" + +call :install_zlib + +echo Cleaning temp files +rd /s /q "%SRC_DIR%\temp_build" || ver > nul + :set_cuda_env_vars echo Setting up environment... @@ -218,3 +244,13 @@ set "PATH=%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%\b set "CUDA_PATH=%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%" set "CUDA_PATH_V%CUDA_VER_MAJOR%_%CUDA_VER_MINOR%=%ProgramFiles%\NVIDIA GPU Computing Toolkit\CUDA\v%CUDA_VERSION_STR%" set "NVTOOLSEXT_PATH=%ProgramFiles%\NVIDIA Corporation\NvToolsExt" + +goto :eof + +@REM cuDNN 8.3+ requires zlib to be installed on the path +:install_zlib +echo Installing ZLIB dlls +curl -k -L "http://s3.amazonaws.com/ossci-windows/zlib123dllx64.zip" --output "%SRC_DIR%\temp_build\zlib123dllx64.zip" +7z x "%SRC_DIR%\temp_build\zlib123dllx64.zip" -o"%SRC_DIR%\temp_build\zlib" +xcopy /Y "%SRC_DIR%\temp_build\zlib\dll_x64\*.dll" "C:\Windows\System32" +goto :eof diff --git a/.ci/pytorch/windows/internal/smoke_test.bat b/.ci/pytorch/windows/internal/smoke_test.bat index f671a9d0e0abb..8cce63a81693d 100644 --- a/.ci/pytorch/windows/internal/smoke_test.bat +++ b/.ci/pytorch/windows/internal/smoke_test.bat @@ -1,10 +1,15 @@ set SRC_DIR=%~dp0 +set TARGET_OS=windows pushd %SRC_DIR%\.. if not "%CUDA_VERSION%" == "cpu" if not "%CUDA_VERSION%" == "xpu" call internal\driver_update.bat if errorlevel 1 exit /b 1 +echo "Check if CUDA and CUDNN versions need to be updated" +call internal\cuda_install.bat +if errorlevel 1 exit /b 1 + if "%CUDA_VERSION%" == "xpu" ( call internal\xpu_install.bat if errorlevel 1 exit /b 1 @@ -94,6 +99,10 @@ echo Checking that basic CNN works %PYTHON_EXEC% %PYTORCH_ROOT%\.ci\pytorch\test_example_code\cnn_smoke.py if ERRORLEVEL 1 exit /b 1 +echo Running smoke_test.py +%PYTHON_EXEC% %PYTORCH_ROOT%\.ci\pytorch\smoke_test\smoke_test.py --package=torchonly --torch-compile-check disabled --runtime-error-check disabled +if ERRORLEVEL 1 exit /b 1 + goto end :libtorch @@ -131,13 +140,13 @@ set INCLUDE=%INCLUDE%;%install_root%\include;%install_root%\include\torch\csrc\a set LIB=%LIB%;%install_root%\lib set PATH=%PATH%;%install_root%\lib -cl %PYTORCH_ROOT%\.ci\pytorch\test_example_code\simple-torch-test.cpp c10.lib torch_cpu.lib /EHsc /std:c++17 +cl %PYTORCH_ROOT%\.ci\pytorch\test_example_code\simple-torch-test.cpp c10.lib torch_cpu.lib /EHsc /std:c++20 if ERRORLEVEL 1 exit /b 1 .\simple-torch-test.exe if ERRORLEVEL 1 exit /b 1 -cl %PYTORCH_ROOT%\.ci\pytorch\test_example_code\check-torch-mkl.cpp c10.lib torch_cpu.lib /EHsc /std:c++17 +cl %PYTORCH_ROOT%\.ci\pytorch\test_example_code\check-torch-mkl.cpp c10.lib torch_cpu.lib /EHsc /std:c++20 if ERRORLEVEL 1 exit /b 1 .\check-torch-mkl.exe @@ -148,7 +157,7 @@ if "%NVIDIA_GPU_EXISTS%" == "0" ( goto end ) -cl %PYTORCH_ROOT%\.ci\pytorch\test_example_code\check-torch-cuda.cpp torch_cpu.lib c10.lib torch_cuda.lib /EHsc /std:c++17 /link /INCLUDE:?warp_size@cuda@at@@YAHXZ +cl %PYTORCH_ROOT%\.ci\pytorch\test_example_code\check-torch-cuda.cpp torch_cpu.lib c10.lib torch_cuda.lib /EHsc /std:c++20 /link /INCLUDE:?warp_size@cuda@at@@YAHXZ .\check-torch-cuda.exe if ERRORLEVEL 1 exit /b 1 diff --git a/.ci/wheel/build_wheel.sh b/.ci/wheel/build_wheel.sh index afd1faf2a5f7c..2563d5ba31765 100755 --- a/.ci/wheel/build_wheel.sh +++ b/.ci/wheel/build_wheel.sh @@ -97,7 +97,7 @@ fi whl_tmp_dir="${MAC_PACKAGE_WORK_DIR}/dist" mkdir -p "$whl_tmp_dir" -mac_version='macosx-11.0-arm64' +mac_version='macosx-14.0-arm64' libtorch_arch='arm64' # Create a consistent wheel package name to rename the wheel to @@ -125,27 +125,19 @@ popd export TH_BINARY_BUILD=1 export INSTALL_TEST=0 # dont install test binaries into site-packages -export MACOSX_DEPLOYMENT_TARGET=11.0 +export MACOSX_DEPLOYMENT_TARGET=14.0 EXTRA_CONDA_INSTALL_FLAGS="" CONDA_ENV_CREATE_FLAGS="" RENAME_WHEEL=false VERIFY_WHEELNAME=true case $desired_python in - 3.14t) - echo "Using 3.14 deps" - NUMPY_PINNED_VERSION="==2.1.0" - ;; - 3.14) - echo "Using 3.14t deps" - NUMPY_PINNED_VERSION="==2.1.0" - ;; - 3.13t) - echo "Using 3.13t deps" - NUMPY_PINNED_VERSION="==2.1.0" + 3.14*) + echo "Using ${desired_python} deps" + NUMPY_PINNED_VERSION="==2.3.4" ;; - 3.13) - echo "Using 3.13 deps" + 3.13*) + echo "Using ${desired_python} deps" NUMPY_PINNED_VERSION="==2.1.0" ;; 3.12) @@ -179,6 +171,7 @@ retry pip install "${PINNED_PACKAGES[@]}" -r "${pytorch_rootdir}/requirements.tx if [[ -d "/opt/llvm-openmp" ]]; then export OMP_PREFIX=/opt/llvm-openmp else + echo "libomp not found, installing via brew" retry brew install libomp fi @@ -217,7 +210,7 @@ if [[ -z "$BUILD_PYTHONLESS" && $RENAME_WHEEL == true ]]; then # Copy the whl to a final destination before tests are run echo "Renaming Wheel file: $wheel_filename_gen to $wheel_filename_new" cp "$whl_tmp_dir/$wheel_filename_gen" "$PYTORCH_FINAL_PACKAGE_DIR/$wheel_filename_new" -elif [[ $RENAME_WHEEL == false ]]; then +elif [[ -z "$BUILD_PYTHONLESS" && $RENAME_WHEEL == false ]]; then echo "Copying Wheel file: $wheel_filename_gen to $PYTORCH_FINAL_PACKAGE_DIR" cp "$whl_tmp_dir/$wheel_filename_gen" "$PYTORCH_FINAL_PACKAGE_DIR/$wheel_filename_gen" if [[ "$VERIFY_WHEELNAME" == "true" && "$wheel_filename_gen" != "$wheel_filename_new" ]]; then diff --git a/.circleci/.gitignore b/.circleci/.gitignore deleted file mode 100644 index c2153925851f5..0000000000000 --- a/.circleci/.gitignore +++ /dev/null @@ -1,2 +0,0 @@ -*.svg -*.png diff --git a/.circleci/README.md b/.circleci/README.md deleted file mode 100644 index 24dde8b47666f..0000000000000 --- a/.circleci/README.md +++ /dev/null @@ -1,4 +0,0 @@ -Warning -======= - -PyTorch migration from CircleCI to github actions has been completed. All continuous integration & deployment workflows are defined in `.github/workflows` folder diff --git a/.circleci/codegen_validation/compare_normalized_yaml.sh b/.circleci/codegen_validation/compare_normalized_yaml.sh deleted file mode 100755 index af2a8268d3c55..0000000000000 --- a/.circleci/codegen_validation/compare_normalized_yaml.sh +++ /dev/null @@ -1,17 +0,0 @@ -#!/bin/bash -xe - - -YAML_FILENAME=verbatim-sources/workflows-pytorch-ge-config-tests.yml -DIFF_TOOL=meld - - -# Allows this script to be invoked from any directory: -cd $(dirname "$0") - -pushd .. - - -$DIFF_TOOL $YAML_FILENAME <(./codegen_validation/normalize_yaml_fragment.py < $YAML_FILENAME) - - -popd diff --git a/.circleci/codegen_validation/normalize_yaml_fragment.py b/.circleci/codegen_validation/normalize_yaml_fragment.py deleted file mode 100755 index 232eaa833b932..0000000000000 --- a/.circleci/codegen_validation/normalize_yaml_fragment.py +++ /dev/null @@ -1,26 +0,0 @@ -#!/usr/bin/env python3 - -import os -import sys - -import yaml - - -# Need to import modules that lie on an upward-relative path -sys.path.append(os.path.dirname(sys.path[0])) - -import cimodel.lib.miniyaml as miniyaml - - -def regurgitate(depth, use_pyyaml_formatter=False): - data = yaml.safe_load(sys.stdin) - - if use_pyyaml_formatter: - output = yaml.dump(data, sort_keys=True) - sys.stdout.write(output) - else: - miniyaml.render(sys.stdout, data, depth) - - -if __name__ == "__main__": - regurgitate(3) diff --git a/.circleci/codegen_validation/overwrite_with_normalized.sh b/.circleci/codegen_validation/overwrite_with_normalized.sh deleted file mode 100755 index 7665984cedb69..0000000000000 --- a/.circleci/codegen_validation/overwrite_with_normalized.sh +++ /dev/null @@ -1,15 +0,0 @@ -#!/bin/bash -xe - -YAML_FILENAME=$1 - -# Allows this script to be invoked from any directory: -cd $(dirname "$0") - -pushd .. - -TEMP_FILENAME=$(mktemp) - -cat $YAML_FILENAME | ./codegen_validation/normalize_yaml_fragment.py > $TEMP_FILENAME -mv $TEMP_FILENAME $YAML_FILENAME - -popd diff --git a/.circleci/scripts/README.md b/.circleci/scripts/README.md deleted file mode 100644 index c06504fd43042..0000000000000 --- a/.circleci/scripts/README.md +++ /dev/null @@ -1,4 +0,0 @@ -All the scripts in this directory are callable from `~/workspace/.circleci/scripts/foo.sh`. -Don't try to call them as `.circleci/scripts/foo.sh`, that won't -(necessarily) work. See Note [Workspace for CircleCI scripts] in -job-specs-setup.yml for more details. diff --git a/.circleci/scripts/driver_update.bat b/.circleci/scripts/driver_update.bat deleted file mode 100644 index fb87743666213..0000000000000 --- a/.circleci/scripts/driver_update.bat +++ /dev/null @@ -1,8 +0,0 @@ -set "DRIVER_DOWNLOAD_LINK=https://s3.amazonaws.com/ossci-windows/452.39-data-center-tesla-desktop-win10-64bit-international.exe" -curl --retry 3 --retry-all-errors -kL %DRIVER_DOWNLOAD_LINK% --output 452.39-data-center-tesla-desktop-win10-64bit-international.exe -if errorlevel 1 exit /b 1 - -start /wait 452.39-data-center-tesla-desktop-win10-64bit-international.exe -s -noreboot -if errorlevel 1 exit /b 1 - -del 452.39-data-center-tesla-desktop-win10-64bit-international.exe || ver > NUL diff --git a/.circleci/scripts/publish_android_snapshot.sh b/.circleci/scripts/publish_android_snapshot.sh deleted file mode 100755 index 352b9d54f871a..0000000000000 --- a/.circleci/scripts/publish_android_snapshot.sh +++ /dev/null @@ -1,46 +0,0 @@ -#!/usr/bin/env bash -# DO NOT ADD 'set -x' not to reveal CircleCI secret context environment variables -set -eu -o pipefail - -export ANDROID_NDK_HOME=/opt/ndk -export ANDROID_HOME=/opt/android/sdk - -export GRADLE_VERSION=6.8.3 -export GRADLE_HOME=/opt/gradle/gradle-$GRADLE_VERSION -export GRADLE_PATH=$GRADLE_HOME/bin/gradle - -echo "BUILD_ENVIRONMENT:$BUILD_ENVIRONMENT" -ls -la ~/workspace - -GRADLE_PROPERTIES=~/workspace/android/gradle.properties - -IS_SNAPSHOT="$(grep 'VERSION_NAME=[0-9\.]\+-SNAPSHOT' "$GRADLE_PROPERTIES")" -echo "IS_SNAPSHOT:$IS_SNAPSHOT" - -if [ -z "$IS_SNAPSHOT" ]; then - echo "Error: version is not snapshot." -elif [ -z "$SONATYPE_NEXUS_USERNAME" ]; then - echo "Error: missing env variable SONATYPE_NEXUS_USERNAME." -elif [ -z "$SONATYPE_NEXUS_PASSWORD" ]; then - echo "Error: missing env variable SONATYPE_NEXUS_PASSWORD." -elif [ -z "$ANDROID_SIGN_KEY" ]; then - echo "Error: missing env variable ANDROID_SIGN_KEY." -elif [ -z "$ANDROID_SIGN_PASS" ]; then - echo "Error: missing env variable ANDROID_SIGN_PASS." -else - GRADLE_LOCAL_PROPERTIES=~/workspace/android/local.properties - rm -f $GRADLE_LOCAL_PROPERTIES - - echo "sdk.dir=/opt/android/sdk" >> $GRADLE_LOCAL_PROPERTIES - echo "ndk.dir=/opt/ndk" >> $GRADLE_LOCAL_PROPERTIES - - echo "SONATYPE_NEXUS_USERNAME=${SONATYPE_NEXUS_USERNAME}" >> $GRADLE_PROPERTIES - echo "mavenCentralRepositoryUsername=${SONATYPE_NEXUS_USERNAME}" >> $GRADLE_PROPERTIES - echo "SONATYPE_NEXUS_PASSWORD=${SONATYPE_NEXUS_PASSWORD}" >> $GRADLE_PROPERTIES - echo "mavenCentralRepositoryPassword=${SONATYPE_NEXUS_PASSWORD}" >> $GRADLE_PROPERTIES - - echo "signing.keyId=${ANDROID_SIGN_KEY}" >> $GRADLE_PROPERTIES - echo "signing.password=${ANDROID_SIGN_PASS}" >> $GRADLE_PROPERTIES - - $GRADLE_PATH -p ~/workspace/android/ uploadArchives -fi diff --git a/.circleci/windows-jni/include/jni.h b/.circleci/windows-jni/include/jni.h deleted file mode 100644 index 1fdc6f3ad69d3..0000000000000 --- a/.circleci/windows-jni/include/jni.h +++ /dev/null @@ -1,1131 +0,0 @@ -/* - * Copyright (C) 2006 The Android Open Source Project - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -/* - * JNI specification, as defined by Sun: - * http://java.sun.com/javase/6/docs/technotes/guides/jni/spec/jniTOC.html - * - * Everything here is expected to be VM-neutral. - */ - -#ifndef JNI_H_ -#define JNI_H_ - -#include -#include - -/* Primitive types that match up with Java equivalents. */ -typedef uint8_t jboolean; /* unsigned 8 bits */ -typedef int8_t jbyte; /* signed 8 bits */ -typedef uint16_t jchar; /* unsigned 16 bits */ -typedef int16_t jshort; /* signed 16 bits */ -typedef int32_t jint; /* signed 32 bits */ -typedef int64_t jlong; /* signed 64 bits */ -typedef float jfloat; /* 32-bit IEEE 754 */ -typedef double jdouble; /* 64-bit IEEE 754 */ - -/* "cardinal indices and sizes" */ -typedef jint jsize; - -#ifdef __cplusplus -/* - * Reference types, in C++ - */ -class _jobject {}; -class _jclass : public _jobject {}; -class _jstring : public _jobject {}; -class _jarray : public _jobject {}; -class _jobjectArray : public _jarray {}; -class _jbooleanArray : public _jarray {}; -class _jbyteArray : public _jarray {}; -class _jcharArray : public _jarray {}; -class _jshortArray : public _jarray {}; -class _jintArray : public _jarray {}; -class _jlongArray : public _jarray {}; -class _jfloatArray : public _jarray {}; -class _jdoubleArray : public _jarray {}; -class _jthrowable : public _jobject {}; - -typedef _jobject* jobject; -typedef _jclass* jclass; -typedef _jstring* jstring; -typedef _jarray* jarray; -typedef _jobjectArray* jobjectArray; -typedef _jbooleanArray* jbooleanArray; -typedef _jbyteArray* jbyteArray; -typedef _jcharArray* jcharArray; -typedef _jshortArray* jshortArray; -typedef _jintArray* jintArray; -typedef _jlongArray* jlongArray; -typedef _jfloatArray* jfloatArray; -typedef _jdoubleArray* jdoubleArray; -typedef _jthrowable* jthrowable; -typedef _jobject* jweak; - - -#else /* not __cplusplus */ - -/* - * Reference types, in C. - */ -typedef void* jobject; -typedef jobject jclass; -typedef jobject jstring; -typedef jobject jarray; -typedef jarray jobjectArray; -typedef jarray jbooleanArray; -typedef jarray jbyteArray; -typedef jarray jcharArray; -typedef jarray jshortArray; -typedef jarray jintArray; -typedef jarray jlongArray; -typedef jarray jfloatArray; -typedef jarray jdoubleArray; -typedef jobject jthrowable; -typedef jobject jweak; - -#endif /* not __cplusplus */ - -struct _jfieldID; /* opaque structure */ -typedef struct _jfieldID* jfieldID; /* field IDs */ - -struct _jmethodID; /* opaque structure */ -typedef struct _jmethodID* jmethodID; /* method IDs */ - -struct JNIInvokeInterface; - -typedef union jvalue { - jboolean z; - jbyte b; - jchar c; - jshort s; - jint i; - jlong j; - jfloat f; - jdouble d; - jobject l; -} jvalue; - -typedef enum jobjectRefType { - JNIInvalidRefType = 0, - JNILocalRefType = 1, - JNIGlobalRefType = 2, - JNIWeakGlobalRefType = 3 -} jobjectRefType; - -typedef struct { - const char* name; - const char* signature; - void* fnPtr; -} JNINativeMethod; - -struct _JNIEnv; -struct _JavaVM; -typedef const struct JNINativeInterface* C_JNIEnv; - -#if defined(__cplusplus) -typedef _JNIEnv JNIEnv; -typedef _JavaVM JavaVM; -#else -typedef const struct JNINativeInterface* JNIEnv; -typedef const struct JNIInvokeInterface* JavaVM; -#endif - -/* - * Table of interface function pointers. - */ -struct JNINativeInterface { - void* reserved0; - void* reserved1; - void* reserved2; - void* reserved3; - - jint (*GetVersion)(JNIEnv *); - - jclass (*DefineClass)(JNIEnv*, const char*, jobject, const jbyte*, - jsize); - jclass (*FindClass)(JNIEnv*, const char*); - - jmethodID (*FromReflectedMethod)(JNIEnv*, jobject); - jfieldID (*FromReflectedField)(JNIEnv*, jobject); - /* spec doesn't show jboolean parameter */ - jobject (*ToReflectedMethod)(JNIEnv*, jclass, jmethodID, jboolean); - - jclass (*GetSuperclass)(JNIEnv*, jclass); - jboolean (*IsAssignableFrom)(JNIEnv*, jclass, jclass); - - /* spec doesn't show jboolean parameter */ - jobject (*ToReflectedField)(JNIEnv*, jclass, jfieldID, jboolean); - - jint (*Throw)(JNIEnv*, jthrowable); - jint (*ThrowNew)(JNIEnv *, jclass, const char *); - jthrowable (*ExceptionOccurred)(JNIEnv*); - void (*ExceptionDescribe)(JNIEnv*); - void (*ExceptionClear)(JNIEnv*); - void (*FatalError)(JNIEnv*, const char*); - - jint (*PushLocalFrame)(JNIEnv*, jint); - jobject (*PopLocalFrame)(JNIEnv*, jobject); - - jobject (*NewGlobalRef)(JNIEnv*, jobject); - void (*DeleteGlobalRef)(JNIEnv*, jobject); - void (*DeleteLocalRef)(JNIEnv*, jobject); - jboolean (*IsSameObject)(JNIEnv*, jobject, jobject); - - jobject (*NewLocalRef)(JNIEnv*, jobject); - jint (*EnsureLocalCapacity)(JNIEnv*, jint); - - jobject (*AllocObject)(JNIEnv*, jclass); - jobject (*NewObject)(JNIEnv*, jclass, jmethodID, ...); - jobject (*NewObjectV)(JNIEnv*, jclass, jmethodID, va_list); - jobject (*NewObjectA)(JNIEnv*, jclass, jmethodID, jvalue*); - - jclass (*GetObjectClass)(JNIEnv*, jobject); - jboolean (*IsInstanceOf)(JNIEnv*, jobject, jclass); - jmethodID (*GetMethodID)(JNIEnv*, jclass, const char*, const char*); - - jobject (*CallObjectMethod)(JNIEnv*, jobject, jmethodID, ...); - jobject (*CallObjectMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jobject (*CallObjectMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - jboolean (*CallBooleanMethod)(JNIEnv*, jobject, jmethodID, ...); - jboolean (*CallBooleanMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jboolean (*CallBooleanMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - jbyte (*CallByteMethod)(JNIEnv*, jobject, jmethodID, ...); - jbyte (*CallByteMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jbyte (*CallByteMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - jchar (*CallCharMethod)(JNIEnv*, jobject, jmethodID, ...); - jchar (*CallCharMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jchar (*CallCharMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - jshort (*CallShortMethod)(JNIEnv*, jobject, jmethodID, ...); - jshort (*CallShortMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jshort (*CallShortMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - jint (*CallIntMethod)(JNIEnv*, jobject, jmethodID, ...); - jint (*CallIntMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jint (*CallIntMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - jlong (*CallLongMethod)(JNIEnv*, jobject, jmethodID, ...); - jlong (*CallLongMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jlong (*CallLongMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - jfloat (*CallFloatMethod)(JNIEnv*, jobject, jmethodID, ...); - jfloat (*CallFloatMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jfloat (*CallFloatMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - jdouble (*CallDoubleMethod)(JNIEnv*, jobject, jmethodID, ...); - jdouble (*CallDoubleMethodV)(JNIEnv*, jobject, jmethodID, va_list); - jdouble (*CallDoubleMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - void (*CallVoidMethod)(JNIEnv*, jobject, jmethodID, ...); - void (*CallVoidMethodV)(JNIEnv*, jobject, jmethodID, va_list); - void (*CallVoidMethodA)(JNIEnv*, jobject, jmethodID, jvalue*); - - jobject (*CallNonvirtualObjectMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jobject (*CallNonvirtualObjectMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jobject (*CallNonvirtualObjectMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - jboolean (*CallNonvirtualBooleanMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jboolean (*CallNonvirtualBooleanMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jboolean (*CallNonvirtualBooleanMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - jbyte (*CallNonvirtualByteMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jbyte (*CallNonvirtualByteMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jbyte (*CallNonvirtualByteMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - jchar (*CallNonvirtualCharMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jchar (*CallNonvirtualCharMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jchar (*CallNonvirtualCharMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - jshort (*CallNonvirtualShortMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jshort (*CallNonvirtualShortMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jshort (*CallNonvirtualShortMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - jint (*CallNonvirtualIntMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jint (*CallNonvirtualIntMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jint (*CallNonvirtualIntMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - jlong (*CallNonvirtualLongMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jlong (*CallNonvirtualLongMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jlong (*CallNonvirtualLongMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - jfloat (*CallNonvirtualFloatMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jfloat (*CallNonvirtualFloatMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jfloat (*CallNonvirtualFloatMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - jdouble (*CallNonvirtualDoubleMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - jdouble (*CallNonvirtualDoubleMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - jdouble (*CallNonvirtualDoubleMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - void (*CallNonvirtualVoidMethod)(JNIEnv*, jobject, jclass, - jmethodID, ...); - void (*CallNonvirtualVoidMethodV)(JNIEnv*, jobject, jclass, - jmethodID, va_list); - void (*CallNonvirtualVoidMethodA)(JNIEnv*, jobject, jclass, - jmethodID, jvalue*); - - jfieldID (*GetFieldID)(JNIEnv*, jclass, const char*, const char*); - - jobject (*GetObjectField)(JNIEnv*, jobject, jfieldID); - jboolean (*GetBooleanField)(JNIEnv*, jobject, jfieldID); - jbyte (*GetByteField)(JNIEnv*, jobject, jfieldID); - jchar (*GetCharField)(JNIEnv*, jobject, jfieldID); - jshort (*GetShortField)(JNIEnv*, jobject, jfieldID); - jint (*GetIntField)(JNIEnv*, jobject, jfieldID); - jlong (*GetLongField)(JNIEnv*, jobject, jfieldID); - jfloat (*GetFloatField)(JNIEnv*, jobject, jfieldID); - jdouble (*GetDoubleField)(JNIEnv*, jobject, jfieldID); - - void (*SetObjectField)(JNIEnv*, jobject, jfieldID, jobject); - void (*SetBooleanField)(JNIEnv*, jobject, jfieldID, jboolean); - void (*SetByteField)(JNIEnv*, jobject, jfieldID, jbyte); - void (*SetCharField)(JNIEnv*, jobject, jfieldID, jchar); - void (*SetShortField)(JNIEnv*, jobject, jfieldID, jshort); - void (*SetIntField)(JNIEnv*, jobject, jfieldID, jint); - void (*SetLongField)(JNIEnv*, jobject, jfieldID, jlong); - void (*SetFloatField)(JNIEnv*, jobject, jfieldID, jfloat); - void (*SetDoubleField)(JNIEnv*, jobject, jfieldID, jdouble); - - jmethodID (*GetStaticMethodID)(JNIEnv*, jclass, const char*, const char*); - - jobject (*CallStaticObjectMethod)(JNIEnv*, jclass, jmethodID, ...); - jobject (*CallStaticObjectMethodV)(JNIEnv*, jclass, jmethodID, va_list); - jobject (*CallStaticObjectMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - jboolean (*CallStaticBooleanMethod)(JNIEnv*, jclass, jmethodID, ...); - jboolean (*CallStaticBooleanMethodV)(JNIEnv*, jclass, jmethodID, - va_list); - jboolean (*CallStaticBooleanMethodA)(JNIEnv*, jclass, jmethodID, - jvalue*); - jbyte (*CallStaticByteMethod)(JNIEnv*, jclass, jmethodID, ...); - jbyte (*CallStaticByteMethodV)(JNIEnv*, jclass, jmethodID, va_list); - jbyte (*CallStaticByteMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - jchar (*CallStaticCharMethod)(JNIEnv*, jclass, jmethodID, ...); - jchar (*CallStaticCharMethodV)(JNIEnv*, jclass, jmethodID, va_list); - jchar (*CallStaticCharMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - jshort (*CallStaticShortMethod)(JNIEnv*, jclass, jmethodID, ...); - jshort (*CallStaticShortMethodV)(JNIEnv*, jclass, jmethodID, va_list); - jshort (*CallStaticShortMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - jint (*CallStaticIntMethod)(JNIEnv*, jclass, jmethodID, ...); - jint (*CallStaticIntMethodV)(JNIEnv*, jclass, jmethodID, va_list); - jint (*CallStaticIntMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - jlong (*CallStaticLongMethod)(JNIEnv*, jclass, jmethodID, ...); - jlong (*CallStaticLongMethodV)(JNIEnv*, jclass, jmethodID, va_list); - jlong (*CallStaticLongMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - jfloat (*CallStaticFloatMethod)(JNIEnv*, jclass, jmethodID, ...); - jfloat (*CallStaticFloatMethodV)(JNIEnv*, jclass, jmethodID, va_list); - jfloat (*CallStaticFloatMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - jdouble (*CallStaticDoubleMethod)(JNIEnv*, jclass, jmethodID, ...); - jdouble (*CallStaticDoubleMethodV)(JNIEnv*, jclass, jmethodID, va_list); - jdouble (*CallStaticDoubleMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - void (*CallStaticVoidMethod)(JNIEnv*, jclass, jmethodID, ...); - void (*CallStaticVoidMethodV)(JNIEnv*, jclass, jmethodID, va_list); - void (*CallStaticVoidMethodA)(JNIEnv*, jclass, jmethodID, jvalue*); - - jfieldID (*GetStaticFieldID)(JNIEnv*, jclass, const char*, - const char*); - - jobject (*GetStaticObjectField)(JNIEnv*, jclass, jfieldID); - jboolean (*GetStaticBooleanField)(JNIEnv*, jclass, jfieldID); - jbyte (*GetStaticByteField)(JNIEnv*, jclass, jfieldID); - jchar (*GetStaticCharField)(JNIEnv*, jclass, jfieldID); - jshort (*GetStaticShortField)(JNIEnv*, jclass, jfieldID); - jint (*GetStaticIntField)(JNIEnv*, jclass, jfieldID); - jlong (*GetStaticLongField)(JNIEnv*, jclass, jfieldID); - jfloat (*GetStaticFloatField)(JNIEnv*, jclass, jfieldID); - jdouble (*GetStaticDoubleField)(JNIEnv*, jclass, jfieldID); - - void (*SetStaticObjectField)(JNIEnv*, jclass, jfieldID, jobject); - void (*SetStaticBooleanField)(JNIEnv*, jclass, jfieldID, jboolean); - void (*SetStaticByteField)(JNIEnv*, jclass, jfieldID, jbyte); - void (*SetStaticCharField)(JNIEnv*, jclass, jfieldID, jchar); - void (*SetStaticShortField)(JNIEnv*, jclass, jfieldID, jshort); - void (*SetStaticIntField)(JNIEnv*, jclass, jfieldID, jint); - void (*SetStaticLongField)(JNIEnv*, jclass, jfieldID, jlong); - void (*SetStaticFloatField)(JNIEnv*, jclass, jfieldID, jfloat); - void (*SetStaticDoubleField)(JNIEnv*, jclass, jfieldID, jdouble); - - jstring (*NewString)(JNIEnv*, const jchar*, jsize); - jsize (*GetStringLength)(JNIEnv*, jstring); - const jchar* (*GetStringChars)(JNIEnv*, jstring, jboolean*); - void (*ReleaseStringChars)(JNIEnv*, jstring, const jchar*); - jstring (*NewStringUTF)(JNIEnv*, const char*); - jsize (*GetStringUTFLength)(JNIEnv*, jstring); - /* JNI spec says this returns const jbyte*, but that's inconsistent */ - const char* (*GetStringUTFChars)(JNIEnv*, jstring, jboolean*); - void (*ReleaseStringUTFChars)(JNIEnv*, jstring, const char*); - jsize (*GetArrayLength)(JNIEnv*, jarray); - jobjectArray (*NewObjectArray)(JNIEnv*, jsize, jclass, jobject); - jobject (*GetObjectArrayElement)(JNIEnv*, jobjectArray, jsize); - void (*SetObjectArrayElement)(JNIEnv*, jobjectArray, jsize, jobject); - - jbooleanArray (*NewBooleanArray)(JNIEnv*, jsize); - jbyteArray (*NewByteArray)(JNIEnv*, jsize); - jcharArray (*NewCharArray)(JNIEnv*, jsize); - jshortArray (*NewShortArray)(JNIEnv*, jsize); - jintArray (*NewIntArray)(JNIEnv*, jsize); - jlongArray (*NewLongArray)(JNIEnv*, jsize); - jfloatArray (*NewFloatArray)(JNIEnv*, jsize); - jdoubleArray (*NewDoubleArray)(JNIEnv*, jsize); - - jboolean* (*GetBooleanArrayElements)(JNIEnv*, jbooleanArray, jboolean*); - jbyte* (*GetByteArrayElements)(JNIEnv*, jbyteArray, jboolean*); - jchar* (*GetCharArrayElements)(JNIEnv*, jcharArray, jboolean*); - jshort* (*GetShortArrayElements)(JNIEnv*, jshortArray, jboolean*); - jint* (*GetIntArrayElements)(JNIEnv*, jintArray, jboolean*); - jlong* (*GetLongArrayElements)(JNIEnv*, jlongArray, jboolean*); - jfloat* (*GetFloatArrayElements)(JNIEnv*, jfloatArray, jboolean*); - jdouble* (*GetDoubleArrayElements)(JNIEnv*, jdoubleArray, jboolean*); - - void (*ReleaseBooleanArrayElements)(JNIEnv*, jbooleanArray, - jboolean*, jint); - void (*ReleaseByteArrayElements)(JNIEnv*, jbyteArray, - jbyte*, jint); - void (*ReleaseCharArrayElements)(JNIEnv*, jcharArray, - jchar*, jint); - void (*ReleaseShortArrayElements)(JNIEnv*, jshortArray, - jshort*, jint); - void (*ReleaseIntArrayElements)(JNIEnv*, jintArray, - jint*, jint); - void (*ReleaseLongArrayElements)(JNIEnv*, jlongArray, - jlong*, jint); - void (*ReleaseFloatArrayElements)(JNIEnv*, jfloatArray, - jfloat*, jint); - void (*ReleaseDoubleArrayElements)(JNIEnv*, jdoubleArray, - jdouble*, jint); - - void (*GetBooleanArrayRegion)(JNIEnv*, jbooleanArray, - jsize, jsize, jboolean*); - void (*GetByteArrayRegion)(JNIEnv*, jbyteArray, - jsize, jsize, jbyte*); - void (*GetCharArrayRegion)(JNIEnv*, jcharArray, - jsize, jsize, jchar*); - void (*GetShortArrayRegion)(JNIEnv*, jshortArray, - jsize, jsize, jshort*); - void (*GetIntArrayRegion)(JNIEnv*, jintArray, - jsize, jsize, jint*); - void (*GetLongArrayRegion)(JNIEnv*, jlongArray, - jsize, jsize, jlong*); - void (*GetFloatArrayRegion)(JNIEnv*, jfloatArray, - jsize, jsize, jfloat*); - void (*GetDoubleArrayRegion)(JNIEnv*, jdoubleArray, - jsize, jsize, jdouble*); - - /* spec shows these without const; some jni.h do, some don't */ - void (*SetBooleanArrayRegion)(JNIEnv*, jbooleanArray, - jsize, jsize, const jboolean*); - void (*SetByteArrayRegion)(JNIEnv*, jbyteArray, - jsize, jsize, const jbyte*); - void (*SetCharArrayRegion)(JNIEnv*, jcharArray, - jsize, jsize, const jchar*); - void (*SetShortArrayRegion)(JNIEnv*, jshortArray, - jsize, jsize, const jshort*); - void (*SetIntArrayRegion)(JNIEnv*, jintArray, - jsize, jsize, const jint*); - void (*SetLongArrayRegion)(JNIEnv*, jlongArray, - jsize, jsize, const jlong*); - void (*SetFloatArrayRegion)(JNIEnv*, jfloatArray, - jsize, jsize, const jfloat*); - void (*SetDoubleArrayRegion)(JNIEnv*, jdoubleArray, - jsize, jsize, const jdouble*); - - jint (*RegisterNatives)(JNIEnv*, jclass, const JNINativeMethod*, - jint); - jint (*UnregisterNatives)(JNIEnv*, jclass); - jint (*MonitorEnter)(JNIEnv*, jobject); - jint (*MonitorExit)(JNIEnv*, jobject); - jint (*GetJavaVM)(JNIEnv*, JavaVM**); - - void (*GetStringRegion)(JNIEnv*, jstring, jsize, jsize, jchar*); - void (*GetStringUTFRegion)(JNIEnv*, jstring, jsize, jsize, char*); - - void* (*GetPrimitiveArrayCritical)(JNIEnv*, jarray, jboolean*); - void (*ReleasePrimitiveArrayCritical)(JNIEnv*, jarray, void*, jint); - - const jchar* (*GetStringCritical)(JNIEnv*, jstring, jboolean*); - void (*ReleaseStringCritical)(JNIEnv*, jstring, const jchar*); - - jweak (*NewWeakGlobalRef)(JNIEnv*, jobject); - void (*DeleteWeakGlobalRef)(JNIEnv*, jweak); - - jboolean (*ExceptionCheck)(JNIEnv*); - jobject (*NewDirectByteBuffer)(JNIEnv*, void*, jlong); - - void* (*GetDirectBufferAddress)(JNIEnv*, jobject); - jlong (*GetDirectBufferCapacity)(JNIEnv*, jobject); - - /* added in JNI 1.6 */ - jobjectRefType (*GetObjectRefType)(JNIEnv*, jobject); -}; - -/* - * C++ object wrapper. - * - * This is usually overlaid on a C struct whose first element is a - * JNINativeInterface*. We rely somewhat on compiler behavior. - */ -struct _JNIEnv { - /* do not rename this; it does not seem to be entirely opaque */ - const struct JNINativeInterface* functions; - -#if defined(__cplusplus) - jint GetVersion() - { return functions->GetVersion(this); } - - jclass DefineClass(const char *name, jobject loader, const jbyte* buf, - jsize bufLen) - { return functions->DefineClass(this, name, loader, buf, bufLen); } - - jclass FindClass(const char* name) - { return functions->FindClass(this, name); } - - jmethodID FromReflectedMethod(jobject method) - { return functions->FromReflectedMethod(this, method); } - - jfieldID FromReflectedField(jobject field) - { return functions->FromReflectedField(this, field); } - - jobject ToReflectedMethod(jclass cls, jmethodID methodID, jboolean isStatic) - { return functions->ToReflectedMethod(this, cls, methodID, isStatic); } - - jclass GetSuperclass(jclass clazz) - { return functions->GetSuperclass(this, clazz); } - - jboolean IsAssignableFrom(jclass clazz1, jclass clazz2) - { return functions->IsAssignableFrom(this, clazz1, clazz2); } - - jobject ToReflectedField(jclass cls, jfieldID fieldID, jboolean isStatic) - { return functions->ToReflectedField(this, cls, fieldID, isStatic); } - - jint Throw(jthrowable obj) - { return functions->Throw(this, obj); } - - jint ThrowNew(jclass clazz, const char* message) - { return functions->ThrowNew(this, clazz, message); } - - jthrowable ExceptionOccurred() - { return functions->ExceptionOccurred(this); } - - void ExceptionDescribe() - { functions->ExceptionDescribe(this); } - - void ExceptionClear() - { functions->ExceptionClear(this); } - - void FatalError(const char* msg) - { functions->FatalError(this, msg); } - - jint PushLocalFrame(jint capacity) - { return functions->PushLocalFrame(this, capacity); } - - jobject PopLocalFrame(jobject result) - { return functions->PopLocalFrame(this, result); } - - jobject NewGlobalRef(jobject obj) - { return functions->NewGlobalRef(this, obj); } - - void DeleteGlobalRef(jobject globalRef) - { functions->DeleteGlobalRef(this, globalRef); } - - void DeleteLocalRef(jobject localRef) - { functions->DeleteLocalRef(this, localRef); } - - jboolean IsSameObject(jobject ref1, jobject ref2) - { return functions->IsSameObject(this, ref1, ref2); } - - jobject NewLocalRef(jobject ref) - { return functions->NewLocalRef(this, ref); } - - jint EnsureLocalCapacity(jint capacity) - { return functions->EnsureLocalCapacity(this, capacity); } - - jobject AllocObject(jclass clazz) - { return functions->AllocObject(this, clazz); } - - jobject NewObject(jclass clazz, jmethodID methodID, ...) - { - va_list args; - va_start(args, methodID); - jobject result = functions->NewObjectV(this, clazz, methodID, args); - va_end(args); - return result; - } - - jobject NewObjectV(jclass clazz, jmethodID methodID, va_list args) - { return functions->NewObjectV(this, clazz, methodID, args); } - - jobject NewObjectA(jclass clazz, jmethodID methodID, jvalue* args) - { return functions->NewObjectA(this, clazz, methodID, args); } - - jclass GetObjectClass(jobject obj) - { return functions->GetObjectClass(this, obj); } - - jboolean IsInstanceOf(jobject obj, jclass clazz) - { return functions->IsInstanceOf(this, obj, clazz); } - - jmethodID GetMethodID(jclass clazz, const char* name, const char* sig) - { return functions->GetMethodID(this, clazz, name, sig); } - -#define CALL_TYPE_METHOD(_jtype, _jname) \ - _jtype Call##_jname##Method(jobject obj, jmethodID methodID, ...) \ - { \ - _jtype result; \ - va_list args; \ - va_start(args, methodID); \ - result = functions->Call##_jname##MethodV(this, obj, methodID, \ - args); \ - va_end(args); \ - return result; \ - } - -#define CALL_TYPE_METHODV(_jtype, _jname) \ - _jtype Call##_jname##MethodV(jobject obj, jmethodID methodID, \ - va_list args) \ - { return functions->Call##_jname##MethodV(this, obj, methodID, args); } - -#define CALL_TYPE_METHODA(_jtype, _jname) \ - _jtype Call##_jname##MethodA(jobject obj, jmethodID methodID, \ - jvalue* args) \ - { return functions->Call##_jname##MethodA(this, obj, methodID, args); } - -#define CALL_TYPE(_jtype, _jname) \ - CALL_TYPE_METHOD(_jtype, _jname) \ - CALL_TYPE_METHODV(_jtype, _jname) \ - CALL_TYPE_METHODA(_jtype, _jname) - CALL_TYPE(jobject, Object) - CALL_TYPE(jboolean, Boolean) - CALL_TYPE(jbyte, Byte) - CALL_TYPE(jchar, Char) - CALL_TYPE(jshort, Short) - CALL_TYPE(jint, Int) - CALL_TYPE(jlong, Long) - CALL_TYPE(jfloat, Float) - CALL_TYPE(jdouble, Double) - - void CallVoidMethod(jobject obj, jmethodID methodID, ...) - { - va_list args; - va_start(args, methodID); - functions->CallVoidMethodV(this, obj, methodID, args); - va_end(args); - } - - void CallVoidMethodV(jobject obj, jmethodID methodID, va_list args) - { functions->CallVoidMethodV(this, obj, methodID, args); } - - void CallVoidMethodA(jobject obj, jmethodID methodID, jvalue* args) - { functions->CallVoidMethodA(this, obj, methodID, args); } -#define CALL_NONVIRT_TYPE_METHOD(_jtype, _jname) \ - _jtype CallNonvirtual##_jname##Method(jobject obj, jclass clazz, \ - jmethodID methodID, ...) \ - { \ - _jtype result; \ - va_list args; \ - va_start(args, methodID); \ - result = functions->CallNonvirtual##_jname##MethodV(this, obj, \ - clazz, methodID, args); \ - va_end(args); \ - return result; \ - } -#define CALL_NONVIRT_TYPE_METHODV(_jtype, _jname) \ - _jtype CallNonvirtual##_jname##MethodV(jobject obj, jclass clazz, \ - jmethodID methodID, va_list args) \ - { return functions->CallNonvirtual##_jname##MethodV(this, obj, clazz, \ - methodID, args); } -#define CALL_NONVIRT_TYPE_METHODA(_jtype, _jname) \ - _jtype CallNonvirtual##_jname##MethodA(jobject obj, jclass clazz, \ - jmethodID methodID, jvalue* args) \ - { return functions->CallNonvirtual##_jname##MethodA(this, obj, clazz, \ - methodID, args); } -#define CALL_NONVIRT_TYPE(_jtype, _jname) \ - CALL_NONVIRT_TYPE_METHOD(_jtype, _jname) \ - CALL_NONVIRT_TYPE_METHODV(_jtype, _jname) \ - CALL_NONVIRT_TYPE_METHODA(_jtype, _jname) - CALL_NONVIRT_TYPE(jobject, Object) - CALL_NONVIRT_TYPE(jboolean, Boolean) - CALL_NONVIRT_TYPE(jbyte, Byte) - CALL_NONVIRT_TYPE(jchar, Char) - CALL_NONVIRT_TYPE(jshort, Short) - CALL_NONVIRT_TYPE(jint, Int) - CALL_NONVIRT_TYPE(jlong, Long) - CALL_NONVIRT_TYPE(jfloat, Float) - CALL_NONVIRT_TYPE(jdouble, Double) - void CallNonvirtualVoidMethod(jobject obj, jclass clazz, - jmethodID methodID, ...) - { - va_list args; - va_start(args, methodID); - functions->CallNonvirtualVoidMethodV(this, obj, clazz, methodID, args); - va_end(args); - } - void CallNonvirtualVoidMethodV(jobject obj, jclass clazz, - jmethodID methodID, va_list args) - { functions->CallNonvirtualVoidMethodV(this, obj, clazz, methodID, args); } - void CallNonvirtualVoidMethodA(jobject obj, jclass clazz, - jmethodID methodID, jvalue* args) - { functions->CallNonvirtualVoidMethodA(this, obj, clazz, methodID, args); } - jfieldID GetFieldID(jclass clazz, const char* name, const char* sig) - { return functions->GetFieldID(this, clazz, name, sig); } - jobject GetObjectField(jobject obj, jfieldID fieldID) - { return functions->GetObjectField(this, obj, fieldID); } - jboolean GetBooleanField(jobject obj, jfieldID fieldID) - { return functions->GetBooleanField(this, obj, fieldID); } - jbyte GetByteField(jobject obj, jfieldID fieldID) - { return functions->GetByteField(this, obj, fieldID); } - jchar GetCharField(jobject obj, jfieldID fieldID) - { return functions->GetCharField(this, obj, fieldID); } - jshort GetShortField(jobject obj, jfieldID fieldID) - { return functions->GetShortField(this, obj, fieldID); } - jint GetIntField(jobject obj, jfieldID fieldID) - { return functions->GetIntField(this, obj, fieldID); } - jlong GetLongField(jobject obj, jfieldID fieldID) - { return functions->GetLongField(this, obj, fieldID); } - jfloat GetFloatField(jobject obj, jfieldID fieldID) - { return functions->GetFloatField(this, obj, fieldID); } - jdouble GetDoubleField(jobject obj, jfieldID fieldID) - { return functions->GetDoubleField(this, obj, fieldID); } - void SetObjectField(jobject obj, jfieldID fieldID, jobject value) - { functions->SetObjectField(this, obj, fieldID, value); } - void SetBooleanField(jobject obj, jfieldID fieldID, jboolean value) - { functions->SetBooleanField(this, obj, fieldID, value); } - void SetByteField(jobject obj, jfieldID fieldID, jbyte value) - { functions->SetByteField(this, obj, fieldID, value); } - void SetCharField(jobject obj, jfieldID fieldID, jchar value) - { functions->SetCharField(this, obj, fieldID, value); } - void SetShortField(jobject obj, jfieldID fieldID, jshort value) - { functions->SetShortField(this, obj, fieldID, value); } - void SetIntField(jobject obj, jfieldID fieldID, jint value) - { functions->SetIntField(this, obj, fieldID, value); } - void SetLongField(jobject obj, jfieldID fieldID, jlong value) - { functions->SetLongField(this, obj, fieldID, value); } - void SetFloatField(jobject obj, jfieldID fieldID, jfloat value) - { functions->SetFloatField(this, obj, fieldID, value); } - void SetDoubleField(jobject obj, jfieldID fieldID, jdouble value) - { functions->SetDoubleField(this, obj, fieldID, value); } - jmethodID GetStaticMethodID(jclass clazz, const char* name, const char* sig) - { return functions->GetStaticMethodID(this, clazz, name, sig); } - -#define CALL_STATIC_TYPE_METHOD(_jtype, _jname) \ - _jtype CallStatic##_jname##Method(jclass clazz, jmethodID methodID, \ - ...) \ - { \ - _jtype result; \ - va_list args; \ - va_start(args, methodID); \ - result = functions->CallStatic##_jname##MethodV(this, clazz, \ - methodID, args); \ - va_end(args); \ - return result; \ - } -#define CALL_STATIC_TYPE_METHODV(_jtype, _jname) \ - _jtype CallStatic##_jname##MethodV(jclass clazz, jmethodID methodID, \ - va_list args) \ - { return functions->CallStatic##_jname##MethodV(this, clazz, methodID, \ - args); } -#define CALL_STATIC_TYPE_METHODA(_jtype, _jname) \ - _jtype CallStatic##_jname##MethodA(jclass clazz, jmethodID methodID, \ - jvalue* args) \ - { return functions->CallStatic##_jname##MethodA(this, clazz, methodID, \ - args); } - -#define CALL_STATIC_TYPE(_jtype, _jname) \ - CALL_STATIC_TYPE_METHOD(_jtype, _jname) \ - CALL_STATIC_TYPE_METHODV(_jtype, _jname) \ - CALL_STATIC_TYPE_METHODA(_jtype, _jname) - CALL_STATIC_TYPE(jobject, Object) - CALL_STATIC_TYPE(jboolean, Boolean) - CALL_STATIC_TYPE(jbyte, Byte) - CALL_STATIC_TYPE(jchar, Char) - CALL_STATIC_TYPE(jshort, Short) - CALL_STATIC_TYPE(jint, Int) - CALL_STATIC_TYPE(jlong, Long) - CALL_STATIC_TYPE(jfloat, Float) - CALL_STATIC_TYPE(jdouble, Double) - void CallStaticVoidMethod(jclass clazz, jmethodID methodID, ...) - { - va_list args; - va_start(args, methodID); - functions->CallStaticVoidMethodV(this, clazz, methodID, args); - va_end(args); - } - void CallStaticVoidMethodV(jclass clazz, jmethodID methodID, va_list args) - { functions->CallStaticVoidMethodV(this, clazz, methodID, args); } - void CallStaticVoidMethodA(jclass clazz, jmethodID methodID, jvalue* args) - { functions->CallStaticVoidMethodA(this, clazz, methodID, args); } - - jfieldID GetStaticFieldID(jclass clazz, const char* name, const char* sig) - { return functions->GetStaticFieldID(this, clazz, name, sig); } - - jobject GetStaticObjectField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticObjectField(this, clazz, fieldID); } - jboolean GetStaticBooleanField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticBooleanField(this, clazz, fieldID); } - jbyte GetStaticByteField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticByteField(this, clazz, fieldID); } - jchar GetStaticCharField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticCharField(this, clazz, fieldID); } - jshort GetStaticShortField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticShortField(this, clazz, fieldID); } - jint GetStaticIntField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticIntField(this, clazz, fieldID); } - jlong GetStaticLongField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticLongField(this, clazz, fieldID); } - jfloat GetStaticFloatField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticFloatField(this, clazz, fieldID); } - jdouble GetStaticDoubleField(jclass clazz, jfieldID fieldID) - { return functions->GetStaticDoubleField(this, clazz, fieldID); } - - void SetStaticObjectField(jclass clazz, jfieldID fieldID, jobject value) - { functions->SetStaticObjectField(this, clazz, fieldID, value); } - void SetStaticBooleanField(jclass clazz, jfieldID fieldID, jboolean value) - { functions->SetStaticBooleanField(this, clazz, fieldID, value); } - void SetStaticByteField(jclass clazz, jfieldID fieldID, jbyte value) - { functions->SetStaticByteField(this, clazz, fieldID, value); } - void SetStaticCharField(jclass clazz, jfieldID fieldID, jchar value) - { functions->SetStaticCharField(this, clazz, fieldID, value); } - void SetStaticShortField(jclass clazz, jfieldID fieldID, jshort value) - { functions->SetStaticShortField(this, clazz, fieldID, value); } - void SetStaticIntField(jclass clazz, jfieldID fieldID, jint value) - { functions->SetStaticIntField(this, clazz, fieldID, value); } - void SetStaticLongField(jclass clazz, jfieldID fieldID, jlong value) - { functions->SetStaticLongField(this, clazz, fieldID, value); } - void SetStaticFloatField(jclass clazz, jfieldID fieldID, jfloat value) - { functions->SetStaticFloatField(this, clazz, fieldID, value); } - void SetStaticDoubleField(jclass clazz, jfieldID fieldID, jdouble value) - { functions->SetStaticDoubleField(this, clazz, fieldID, value); } - - jstring NewString(const jchar* unicodeChars, jsize len) - { return functions->NewString(this, unicodeChars, len); } - - jsize GetStringLength(jstring string) - { return functions->GetStringLength(this, string); } - - const jchar* GetStringChars(jstring string, jboolean* isCopy) - { return functions->GetStringChars(this, string, isCopy); } - - void ReleaseStringChars(jstring string, const jchar* chars) - { functions->ReleaseStringChars(this, string, chars); } - - jstring NewStringUTF(const char* bytes) - { return functions->NewStringUTF(this, bytes); } - - jsize GetStringUTFLength(jstring string) - { return functions->GetStringUTFLength(this, string); } - - const char* GetStringUTFChars(jstring string, jboolean* isCopy) - { return functions->GetStringUTFChars(this, string, isCopy); } - - void ReleaseStringUTFChars(jstring string, const char* utf) - { functions->ReleaseStringUTFChars(this, string, utf); } - - jsize GetArrayLength(jarray array) - { return functions->GetArrayLength(this, array); } - - jobjectArray NewObjectArray(jsize length, jclass elementClass, - jobject initialElement) - { return functions->NewObjectArray(this, length, elementClass, - initialElement); } - - jobject GetObjectArrayElement(jobjectArray array, jsize index) - { return functions->GetObjectArrayElement(this, array, index); } - - void SetObjectArrayElement(jobjectArray array, jsize index, jobject value) - { functions->SetObjectArrayElement(this, array, index, value); } - - jbooleanArray NewBooleanArray(jsize length) - { return functions->NewBooleanArray(this, length); } - jbyteArray NewByteArray(jsize length) - { return functions->NewByteArray(this, length); } - jcharArray NewCharArray(jsize length) - { return functions->NewCharArray(this, length); } - jshortArray NewShortArray(jsize length) - { return functions->NewShortArray(this, length); } - jintArray NewIntArray(jsize length) - { return functions->NewIntArray(this, length); } - jlongArray NewLongArray(jsize length) - { return functions->NewLongArray(this, length); } - jfloatArray NewFloatArray(jsize length) - { return functions->NewFloatArray(this, length); } - jdoubleArray NewDoubleArray(jsize length) - { return functions->NewDoubleArray(this, length); } - - jboolean* GetBooleanArrayElements(jbooleanArray array, jboolean* isCopy) - { return functions->GetBooleanArrayElements(this, array, isCopy); } - jbyte* GetByteArrayElements(jbyteArray array, jboolean* isCopy) - { return functions->GetByteArrayElements(this, array, isCopy); } - jchar* GetCharArrayElements(jcharArray array, jboolean* isCopy) - { return functions->GetCharArrayElements(this, array, isCopy); } - jshort* GetShortArrayElements(jshortArray array, jboolean* isCopy) - { return functions->GetShortArrayElements(this, array, isCopy); } - jint* GetIntArrayElements(jintArray array, jboolean* isCopy) - { return functions->GetIntArrayElements(this, array, isCopy); } - jlong* GetLongArrayElements(jlongArray array, jboolean* isCopy) - { return functions->GetLongArrayElements(this, array, isCopy); } - jfloat* GetFloatArrayElements(jfloatArray array, jboolean* isCopy) - { return functions->GetFloatArrayElements(this, array, isCopy); } - jdouble* GetDoubleArrayElements(jdoubleArray array, jboolean* isCopy) - { return functions->GetDoubleArrayElements(this, array, isCopy); } - - void ReleaseBooleanArrayElements(jbooleanArray array, jboolean* elems, - jint mode) - { functions->ReleaseBooleanArrayElements(this, array, elems, mode); } - void ReleaseByteArrayElements(jbyteArray array, jbyte* elems, - jint mode) - { functions->ReleaseByteArrayElements(this, array, elems, mode); } - void ReleaseCharArrayElements(jcharArray array, jchar* elems, - jint mode) - { functions->ReleaseCharArrayElements(this, array, elems, mode); } - void ReleaseShortArrayElements(jshortArray array, jshort* elems, - jint mode) - { functions->ReleaseShortArrayElements(this, array, elems, mode); } - void ReleaseIntArrayElements(jintArray array, jint* elems, - jint mode) - { functions->ReleaseIntArrayElements(this, array, elems, mode); } - void ReleaseLongArrayElements(jlongArray array, jlong* elems, - jint mode) - { functions->ReleaseLongArrayElements(this, array, elems, mode); } - void ReleaseFloatArrayElements(jfloatArray array, jfloat* elems, - jint mode) - { functions->ReleaseFloatArrayElements(this, array, elems, mode); } - void ReleaseDoubleArrayElements(jdoubleArray array, jdouble* elems, - jint mode) - { functions->ReleaseDoubleArrayElements(this, array, elems, mode); } - - void GetBooleanArrayRegion(jbooleanArray array, jsize start, jsize len, - jboolean* buf) - { functions->GetBooleanArrayRegion(this, array, start, len, buf); } - void GetByteArrayRegion(jbyteArray array, jsize start, jsize len, - jbyte* buf) - { functions->GetByteArrayRegion(this, array, start, len, buf); } - void GetCharArrayRegion(jcharArray array, jsize start, jsize len, - jchar* buf) - { functions->GetCharArrayRegion(this, array, start, len, buf); } - void GetShortArrayRegion(jshortArray array, jsize start, jsize len, - jshort* buf) - { functions->GetShortArrayRegion(this, array, start, len, buf); } - void GetIntArrayRegion(jintArray array, jsize start, jsize len, - jint* buf) - { functions->GetIntArrayRegion(this, array, start, len, buf); } - void GetLongArrayRegion(jlongArray array, jsize start, jsize len, - jlong* buf) - { functions->GetLongArrayRegion(this, array, start, len, buf); } - void GetFloatArrayRegion(jfloatArray array, jsize start, jsize len, - jfloat* buf) - { functions->GetFloatArrayRegion(this, array, start, len, buf); } - void GetDoubleArrayRegion(jdoubleArray array, jsize start, jsize len, - jdouble* buf) - { functions->GetDoubleArrayRegion(this, array, start, len, buf); } - - void SetBooleanArrayRegion(jbooleanArray array, jsize start, jsize len, - const jboolean* buf) - { functions->SetBooleanArrayRegion(this, array, start, len, buf); } - void SetByteArrayRegion(jbyteArray array, jsize start, jsize len, - const jbyte* buf) - { functions->SetByteArrayRegion(this, array, start, len, buf); } - void SetCharArrayRegion(jcharArray array, jsize start, jsize len, - const jchar* buf) - { functions->SetCharArrayRegion(this, array, start, len, buf); } - void SetShortArrayRegion(jshortArray array, jsize start, jsize len, - const jshort* buf) - { functions->SetShortArrayRegion(this, array, start, len, buf); } - void SetIntArrayRegion(jintArray array, jsize start, jsize len, - const jint* buf) - { functions->SetIntArrayRegion(this, array, start, len, buf); } - void SetLongArrayRegion(jlongArray array, jsize start, jsize len, - const jlong* buf) - { functions->SetLongArrayRegion(this, array, start, len, buf); } - void SetFloatArrayRegion(jfloatArray array, jsize start, jsize len, - const jfloat* buf) - { functions->SetFloatArrayRegion(this, array, start, len, buf); } - void SetDoubleArrayRegion(jdoubleArray array, jsize start, jsize len, - const jdouble* buf) - { functions->SetDoubleArrayRegion(this, array, start, len, buf); } - - jint RegisterNatives(jclass clazz, const JNINativeMethod* methods, - jint nMethods) - { return functions->RegisterNatives(this, clazz, methods, nMethods); } - - jint UnregisterNatives(jclass clazz) - { return functions->UnregisterNatives(this, clazz); } - - jint MonitorEnter(jobject obj) - { return functions->MonitorEnter(this, obj); } - - jint MonitorExit(jobject obj) - { return functions->MonitorExit(this, obj); } - - jint GetJavaVM(JavaVM** vm) - { return functions->GetJavaVM(this, vm); } - - void GetStringRegion(jstring str, jsize start, jsize len, jchar* buf) - { functions->GetStringRegion(this, str, start, len, buf); } - - void GetStringUTFRegion(jstring str, jsize start, jsize len, char* buf) - { return functions->GetStringUTFRegion(this, str, start, len, buf); } - - void* GetPrimitiveArrayCritical(jarray array, jboolean* isCopy) - { return functions->GetPrimitiveArrayCritical(this, array, isCopy); } - - void ReleasePrimitiveArrayCritical(jarray array, void* carray, jint mode) - { functions->ReleasePrimitiveArrayCritical(this, array, carray, mode); } - - const jchar* GetStringCritical(jstring string, jboolean* isCopy) - { return functions->GetStringCritical(this, string, isCopy); } - - void ReleaseStringCritical(jstring string, const jchar* carray) - { functions->ReleaseStringCritical(this, string, carray); } - - jweak NewWeakGlobalRef(jobject obj) - { return functions->NewWeakGlobalRef(this, obj); } - - void DeleteWeakGlobalRef(jweak obj) - { functions->DeleteWeakGlobalRef(this, obj); } - - jboolean ExceptionCheck() - { return functions->ExceptionCheck(this); } - - jobject NewDirectByteBuffer(void* address, jlong capacity) - { return functions->NewDirectByteBuffer(this, address, capacity); } - - void* GetDirectBufferAddress(jobject buf) - { return functions->GetDirectBufferAddress(this, buf); } - - jlong GetDirectBufferCapacity(jobject buf) - { return functions->GetDirectBufferCapacity(this, buf); } - - /* added in JNI 1.6 */ - jobjectRefType GetObjectRefType(jobject obj) - { return functions->GetObjectRefType(this, obj); } -#endif /*__cplusplus*/ -}; - - -/* - * JNI invocation interface. - */ -struct JNIInvokeInterface { - void* reserved0; - void* reserved1; - void* reserved2; - jint (*DestroyJavaVM)(JavaVM*); - jint (*AttachCurrentThread)(JavaVM*, JNIEnv**, void*); - jint (*DetachCurrentThread)(JavaVM*); - jint (*GetEnv)(JavaVM*, void**, jint); - jint (*AttachCurrentThreadAsDaemon)(JavaVM*, JNIEnv**, void*); -}; - -/* - * C++ version. - */ -struct _JavaVM { - const struct JNIInvokeInterface* functions; - -#if defined(__cplusplus) - jint DestroyJavaVM() - { return functions->DestroyJavaVM(this); } - jint AttachCurrentThread(JNIEnv** p_env, void* thr_args) - { return functions->AttachCurrentThread(this, p_env, thr_args); } - jint DetachCurrentThread() - { return functions->DetachCurrentThread(this); } - jint GetEnv(void** env, jint version) - { return functions->GetEnv(this, env, version); } - jint AttachCurrentThreadAsDaemon(JNIEnv** p_env, void* thr_args) - { return functions->AttachCurrentThreadAsDaemon(this, p_env, thr_args); } -#endif /*__cplusplus*/ -}; - -struct JavaVMAttachArgs { - jint version; /* must be >= JNI_VERSION_1_2 */ - const char* name; /* NULL or name of thread as modified UTF-8 str */ - jobject group; /* global ref of a ThreadGroup object, or NULL */ -}; -typedef struct JavaVMAttachArgs JavaVMAttachArgs; - -/* - * JNI 1.2+ initialization. (As of 1.6, the pre-1.2 structures are no - * longer supported.) - */ -typedef struct JavaVMOption { - const char* optionString; - void* extraInfo; -} JavaVMOption; - -typedef struct JavaVMInitArgs { - jint version; /* use JNI_VERSION_1_2 or later */ - jint nOptions; - JavaVMOption* options; - jboolean ignoreUnrecognized; -} JavaVMInitArgs; - -#ifdef __cplusplus -extern "C" { -#endif -/* - * VM initialization functions. - * - * Note these are the only symbols exported for JNI by the VM. - */ -jint JNI_GetDefaultJavaVMInitArgs(void*); -jint JNI_CreateJavaVM(JavaVM**, JNIEnv**, void*); -jint JNI_GetCreatedJavaVMs(JavaVM**, jsize, jsize*); - -#define JNIIMPORT -// To match the JNIEXPORT on Windows -#define JNIEXPORT __declspec(dllexport) -#define JNICALL - -/* - * Prototypes for functions exported by loadable shared libs. These are - * called by JNI, not provided by JNI. - */ -JNIEXPORT jint JNI_OnLoad(JavaVM* vm, void* reserved); -JNIEXPORT void JNI_OnUnload(JavaVM* vm, void* reserved); - -#ifdef __cplusplus -} -#endif - -/* - * Manifest constants. - */ -#define JNI_FALSE 0 -#define JNI_TRUE 1 - -#define JNI_VERSION_1_1 0x00010001 -#define JNI_VERSION_1_2 0x00010002 -#define JNI_VERSION_1_4 0x00010004 -#define JNI_VERSION_1_6 0x00010006 - -#define JNI_OK (0) /* no error */ -#define JNI_ERR (-1) /* generic error */ -#define JNI_EDETACHED (-2) /* thread detached from the VM */ -#define JNI_EVERSION (-3) /* JNI version error */ - -#define JNI_COMMIT 1 /* copy content, do not free buffer */ -#define JNI_ABORT 2 /* free buffer w/o copying back */ - -#endif /* JNI_H_ */ diff --git a/.claude/skills/aoti-debug/SKILL.md b/.claude/skills/aoti-debug/SKILL.md index 56a17e08393e9..a790d1b15c26e 100644 --- a/.claude/skills/aoti-debug/SKILL.md +++ b/.claude/skills/aoti-debug/SKILL.md @@ -7,6 +7,22 @@ description: Debug AOTInductor (AOTI) errors and crashes. Use when encountering This skill helps diagnose and fix common AOTInductor issues. +## Error Pattern Routing + +**Check the error message and route to the appropriate sub-guide:** + +### Triton Index Out of Bounds +If the error matches this pattern: +``` +Assertion `index out of bounds: 0 <= tmpN < ksM` failed +``` +**→ Follow the guide in `triton-index-out-of-bounds.md`** + +### All Other Errors +Continue with the sections below. + +--- + ## First Step: Always Check Device and Shape Matching **For ANY AOTI error (segfault, exception, crash, wrong output), ALWAYS check these first:** diff --git a/.claude/skills/aoti-debug/triton-index-out-of-bounds.md b/.claude/skills/aoti-debug/triton-index-out-of-bounds.md new file mode 100644 index 0000000000000..fee827e32feb2 --- /dev/null +++ b/.claude/skills/aoti-debug/triton-index-out-of-bounds.md @@ -0,0 +1,265 @@ +# AOTI Triton Index Out of Bounds Debug Guide + +This guide helps debug AOTI Triton kernel assertion errors with the `index out of bounds` pattern. + +## Error Pattern + +This guide applies when you see errors like: + +``` +/var/tmp/torchinductor_*/.../*.py:NN: unknown: block: [X,Y,Z], thread: [X,Y,Z] +Assertion `index out of bounds: 0 <= tmpN < ksM` failed. +``` + +### Key Information from Error + +| Field | Value | Meaning | +|-------|-------|---------| +| File Path | `/var/tmp/torchinductor_*/*.py` | Generated Triton kernel file (runtime) | +| Line Number | `:NN` | Line in the generated kernel where assertion failed | +| Block/Thread | `[X,Y,Z]` | CUDA block and thread indices | +| Assertion | `0 <= tmpN < ksM` | Index `tmpN` must be within bounds `[0, ksM)` | + +### Understanding the Assertion + +- `tmpN`: A computed index value in the Triton kernel +- `ksM`: A dynamic kernel size parameter (runtime value) +- The assertion fails when `tmpN < 0` or `tmpN >= ksM` + +--- + +## Step 1: Collect AOTI Package + +You need access to the AOTI package that was compiled. This is typically a `.pt2` package or extracted archive containing a `wrapper.cpp` file. + +**Key File**: `*.wrapper.cpp` contains: +- All Triton kernel source code (embedded as comments) +- Kernel launch configurations +- Input/output tensor mappings +- Dynamic shape variable definitions + +--- + +## Step 2: Locate the Failing Kernel in C++ Wrapper + +### Search for the Assertion Pattern + +Extract the assertion pattern from the error (e.g., `tmp18 < ks0`) and search: + +```bash +# Search for the specific assertion +grep -n "tmpN < ksM" /path/to/*.wrapper.cpp + +# Get context around the assertion (80 lines before, 20 after) +grep -n -B80 -A20 "tmpN < ksM" /path/to/*.wrapper.cpp +``` + +### Find the Full Kernel Definition + +The kernel is embedded as a Python docstring comment in the C++ wrapper: + +```cpp + /* + async_compile.triton('triton_red_fused_...', ''' + import triton + import triton.language as tl + ... + def triton_red_fused_...(in_ptr0, out_ptr1, ks0, xnumel, r0_numel, ...): +``` + +--- + +## Step 3: Understand the Kernel Logic + +Analyze the code path leading to the assertion. Common patterns that cause index out of bounds: + +### Pattern: Empty Tensor with ks0 = 0 + +When a dynamic shape `ks0 = 0`: +1. `tmp13 = (-1) + 0 = -1` +2. Index wrapping logic produces `-1` +3. Assertion `0 <= -1 < 0` fails + +### Example Kernel Pattern + +```python +tmp13 = (-1) + ks0 # ks0 - 1 +tmp14 = tl.where(tmp12, tmp10, tmp13) # if condition: use tmp10, else: ks0-1 +tmp15 = ks0 +tmp16 = tmp14 + tmp15 # wrap-around for negative indices +tmp17 = tmp14 < 0 +tmp18 = tl.where(tmp17, tmp16, tmp14) # if negative: add ks0 + +# ASSERTION: 0 <= tmp18 < ks0 +tl.device_assert(((0 <= tmp18) & (tmp18 < ks0)), "index out of bounds") +``` + +--- + +## Step 4: Identify the Dynamic Shape Variable + +### Find Where the Kernel is Called + +```bash +grep -n "call_triton_KERNEL_NAME" /path/to/*.wrapper.cpp +``` + +### Example Output + +```cpp +call_triton_red_fused_...(arg1415_1, buf696, s607, 1L, s13, ...); +``` + +### Parameter Mapping + +| Parameter | Value | Meaning | +|-----------|-------|---------| +| `in_ptr0` | `arg1415_1` | Input tensor | +| `out_ptr1` | `buf696` | Output buffer | +| `ks0` | `s607` | **Dynamic shape - this is the failing bound** | + +### Find the Definition of the Shape Variable + +```bash +grep -n "int64_t s607 = " /path/to/*.wrapper.cpp +``` + +This shows which input tensor dimension defines the shape: + +```cpp +int64_t s607 = arg1416_1_size[0]; +``` + +--- + +## Step 5: Trace Back to Model Input + +### Find Input Index + +Inputs are numbered sequentially. Find which input the argument corresponds to: + +```bash +grep -n 'inputs_info_\[INDEX\].name = "argNNN_1"' /path/to/*.wrapper.cpp +``` + +### Check Input Constraints + +```bash +grep -n "argNNN_1_size\[0\]" /path/to/*.wrapper.cpp +``` + +Look for guards like: +```cpp +if (arg_size[0] > 230400) { // Upper bound check only - no lower bound! +``` + +**Common Issue**: Upper bound checks exist but no lower bound checks for `>= 1`. + +--- + +## Step 6: Map to Model Code + +### Use Source Node Comments + +The C++ wrapper includes comments showing which PyTorch operations generated each kernel: + +```bash +grep -n -B5 "call_triton_KERNEL_NAME" /path/to/*.wrapper.cpp | grep "Source Nodes" +``` + +### Example Output + +```cpp +// Topologically Sorted Source Nodes: [slice_1, sub_89, cumsum, ge_231, where_2, index_copy] +``` + +### Map Operations to Python Code + +| ATen Operation | Python Code Pattern | +|----------------|---------------------| +| `cumsum` | `torch.cumsum(tensor, dim=0)` | +| `sub` | `idx - 1` | +| `ge` | `idx >= 0` | +| `where` | `torch.where(condition, ...)` | +| `index_copy` | `tensor.index_copy(0, indices, source)` | + +--- + +## Root Cause Analysis + +### Common Root Causes + +1. **Empty tensor at runtime**: A jagged/variable-length tensor has size 0 at runtime but wasn't tested during compilation +2. **Missing lower bound guards**: AOTI only generates upper bound checks, not lower bound checks +3. **Edge case not in sample inputs**: Sample inputs during AOTI export never included the edge case + +--- + +## Fix Recommendations + +### Option 1: Add Guard in Forward Method + +```python +def forward(self, lengths: torch.Tensor, ...) -> torch.Tensor: + if lengths.numel() == 0: + device = lengths.device + return torch.empty(0, self.output_dim, device=device) + # ... rest of method +``` + +### Option 2: Fix the Specific Operation + +Add handling for empty tensors in the problematic operation: + +```python +def process_events(self, lengths: torch.Tensor, ...): + if lengths.numel() == 0: + return torch.empty(0, self.emb_dim, device=lengths.device) + # ... rest of method +``` + +### Option 3: Include Edge Cases in AOTI Export + +Ensure sample inputs during AOTI export include: +- Empty tensors (size 0) +- Minimum size tensors (size 1) +- Maximum expected sizes + +--- + +## Useful Commands Summary + +### Searching in AOTI Wrapper + +```bash +# Find kernel by assertion pattern +grep -n "tmpN < ksM" *.wrapper.cpp + +# Get full kernel context +grep -n -B80 -A20 "ASSERTION_PATTERN" *.wrapper.cpp + +# Find kernel call site +grep -n "call_KERNEL_NAME" *.wrapper.cpp + +# Find dynamic shape definition +grep -n "int64_t SHAPE_VAR = " *.wrapper.cpp + +# Find input mapping +grep -n 'inputs_info_\[INDEX\].name' *.wrapper.cpp + +# Find size constraints +grep -n "SHAPE_VAR_size\[0\]" *.wrapper.cpp +``` + +### Environment Variables for Debugging + +```bash +# Enable debug output during torch.compile +export TORCH_COMPILE_DEBUG=1 + +# Save generated kernels to persistent location +export TORCHINDUCTOR_CACHE_DIR=/path/to/save/kernels + +# Enable CUDA launch blocking for accurate stack traces +export CUDA_LAUNCH_BLOCKING=1 +``` diff --git a/.claude/skills/document-public-apis/SKILL.md b/.claude/skills/document-public-apis/SKILL.md new file mode 100644 index 0000000000000..211b444fc5461 --- /dev/null +++ b/.claude/skills/document-public-apis/SKILL.md @@ -0,0 +1,342 @@ +--- +name: document-public-apis +description: Document undocumented public APIs in PyTorch by removing functions from coverage_ignore_functions and coverage_ignore_classes in docs/source/conf.py, running Sphinx coverage, and adding the appropriate autodoc directives to the correct .md or .rst doc files. Use when a user asks to remove functions from conf.py ignore lists. +--- + +# Document Public APIs + +This skill documents undocumented public APIs in PyTorch by removing entries from the coverage ignore lists in `docs/source/conf.py` and adding Sphinx autodoc directives (e.g., `autosummary`, `currentmodule`, `autoclass`, `automodule`) to the corresponding `.md` or `.rst` doc source files in `docs/source/`. + +**"Documenting" means adding autodoc directives to doc source files — NEVER modifying Python source code.** Do not add or edit docstrings in `.py` files. Your only job is to add the correct directive to the correct doc file. + +**IMPORTANT: Before adding any function to the sphinx doctree, verify it has a real docstring.** Use a quick Python check (e.g., `python -c "from torch.module import func; print(bool(func.__doc__))"`) to confirm the function has actual documentation content — not just an empty docstring or a bare `.. warning:: This API is experimental` stub. Functions without meaningful docstrings should be left in the `coverage_ignore_functions`/`coverage_ignore_classes` lists. Adding undocumented functions to the doctree creates empty or near-empty pages that degrade documentation quality. + +## Overview + +`docs/source/conf.py` contains two lists that suppress Sphinx coverage warnings for undocumented APIs: + +- `coverage_ignore_functions`: undocumented functions +- `coverage_ignore_classes`: undocumented classes + +Entries are organized by **module comment groups**. Each group has a module label comment followed by the function/class names that belong to that module: + +```python +coverage_ignore_functions = [ + # torch.ao.quantization.fx.convert <-- module label comment + "convert", # <-- entries belonging to this module + "convert_custom_module", + "convert_standalone_module", + "convert_weighted_module", + # torch.ao.quantization.fx.fuse <-- next module group + "fuse", + # torch.nn.functional + "assert_int_or_pair", # looks unintentionally public <-- entry with inline comment + "constant", # deprecated <-- entry with inline comment +] +``` + +There are two kinds of comments: +- **Module label comments** (`# torch.ao.quantization.fx.convert`): these label which module the entries below belong to. They appear on their own line before a group of entries. +- **Inline comments** (`# deprecated`, `# documented as adaptive_max_pool1d`): these appear after a string entry on the same line and explain *why* the entry is in the ignore list. + +The module label comment directly tells you: +1. Which module the functions belong to +2. Where to add them in the docs (e.g., `# torch.ao.quantization.fx.convert` → the functions go under `torch.ao.quantization.fx.convert` in the doc file) + +## Instructions + +Each invocation of this skill processes **one batch** of module groups. Pick one or more complete module groups from the ignore lists, document their functions, and verify. + +### Step 1: Select module groups to document + +Read `docs/source/conf.py` and select one or more **complete module groups** to document. A module group is a module label comment and all entries beneath it up to the next module label comment. Process entire groups — never split a group across batches. + +For example, selecting the `torch.ao.quantization.fx.convert` group means taking all of: + +```python +# torch.ao.quantization.fx.convert +"convert", +"convert_custom_module", +"convert_standalone_module", +"convert_weighted_module", +``` + +Work through the lists top-to-bottom. Choose enough groups to make meaningful progress (aim for 5–15 functions total, but always include complete groups even if that means going slightly over). + +**Check inline comments before including an entry.** Some entries have inline comments that indicate they should not be documented: + +- `# deprecated` — The function is deprecated. Leave it in the ignore list. +- `# documented as ` — Already documented under a different name. Leave it. +- `# looks unintentionally public` — Probably not meant to be public API. Leave it. +- `# legacy helper for ...` — Same as deprecated. Leave it. +- `# utility function` - Leave it. + +If a module group has a **mix** of regular entries and entries with inline comments, still process the group — but only comment out the regular entries. Leave entries with inline comments untouched in the ignore list. + +### Step 1b: Verify functions have actual docstrings + +For each function selected in Step 1, check that it has a meaningful docstring by running: + +```bash +python -c "from torch.module.path import func_name; doc = func_name.__doc__; print('HAS DOC' if doc and len(doc.strip()) > 80 else 'NO DOC'); print(repr(doc[:120]) if doc else 'None')" +``` + +A function has a **meaningful docstring** if it has real descriptive content — not just: +- `None` or empty string +- Only a `.. warning:: This API is experimental` stub with no description +- Only a one-line auto-generated signature + +**Functions without meaningful docstrings must stay in the ignore list.** Remove them from your batch. If an entire module group has no functions with docstrings, skip the whole group. + +### Step 2: Present the batch to the user + +**Before making any edits**, present the selected module groups and their functions to the user. Indicate which functions passed the docstring check and which were excluded. Show them organized by module: + +``` +Module: torch.ao.quantization.fx.convert + - convert + - convert_custom_module + - convert_standalone_module + - convert_weighted_module + +Module: torch.ao.quantization.fx.fuse + - fuse +``` + +Then use the `AskUserQuestion` tool to let the user confirm, with options like: +- "Proceed with this batch" +- "Skip some entries" (user can specify which to remove) +- "Pick a different batch" + +### Step 3: Comment out entries in conf.py + +After the user confirms, edit `docs/source/conf.py` and **comment out** (do not delete) the selected entries. Use a `#` prefix on each string entry line: + +```python +# torch.ao.quantization.fx.convert +# "convert", +# "convert_custom_module", +# "convert_standalone_module", +# "convert_weighted_module", +``` + +This preserves the original entries so they can be restored if verification fails. + +### Step 4: Run Sphinx coverage + +```bash +cd docs && make coverage +``` + +**Ignore the terminal output of `make coverage`.** It often contains unrelated tracebacks and errors from Sphinx extensions (e.g., `onnx_ir`, `katex`, `sphinxcontrib`) that have nothing to do with coverage. The only thing that matters is whether `docs/build/coverage/python.txt` was generated. Read that file to see the specific undocumented APIs. + +The format of `python.txt` lists each undocumented API as: + +``` +torch.ao.quantization.fx.convert + * convert + * convert_custom_module + * convert_standalone_module + * convert_weighted_module +``` + +**Not all commented-out functions will appear in `python.txt`.** Some may already be documented elsewhere. This is fine — only add directives for functions that actually appear in `python.txt`. + +If `make coverage` fails due to missing dependencies, first run: + +```bash +cd docs && pip install -r requirements.txt +``` + +### Step 5: Add documentation directives + +For each function listed in `python.txt`, use the **module label comment** from `conf.py` to determine where it should be added. The module comment gives you the full module path, which maps to a doc source file and a section within that file. + +#### Finding the correct doc file + +The module comment maps to a doc source file in `docs/source/`. When unsure, search for other functions from the same module: + +```bash +grep -rn "torch.module_name" docs/source/*.md docs/source/*.rst +``` + +Or list candidate files: + +```bash +ls docs/source/*module_name* +``` + +If no doc file exists for a submodule, check whether a parent module's doc file has a section for it (e.g., `backends.md` has sections for `torch.backends.cuda`, `torch.backends.cudnn`, etc.). If not, add a new section to the parent file following existing patterns. + +#### Adding the directives + +**Read the target doc file first** and match the exact patterns already used there. Do not invent new patterns or use bare `autofunction` with fully qualified names — always use the proper hierarchical structure with `automodule`, `currentmodule`, and short names. Do not use `. py:module::` since that just suppresses errors and doesn't actually document the function. Look at other files that match the target file's format (e.g., `.md` vs. `.rst`) under `docs/source/` to see examples. + +There are two file formats. Match the one used in the target file. + +**Pattern A — MyST Markdown files (`.md`):** Used in files like `accelerator.md`, `backends.md`, `cuda.md`. + +The hierarchical structure uses `automodule` to register the module, `currentmodule` to set context, then short names: + +```markdown +## torch.ao.quantization.fx.convert + +```{eval-rst} +.. automodule:: torch.ao.quantization.fx.convert +``` + +```{eval-rst} +.. currentmodule:: torch.ao.quantization.fx.convert +``` + +```{eval-rst} +.. autofunction:: convert +``` + +```{eval-rst} +.. autofunction:: convert_custom_module +``` +``` + +For `autosummary` blocks (used in some files instead of individual directives): + +```markdown +```{eval-rst} +.. autosummary:: + :toctree: generated + :nosignatures: + + existing_function + your_new_function +`` ` +``` + +For classes: + +```markdown +```{eval-rst} +.. autoclass:: YourClass + :members: +`` ` +``` + +**Pattern B — reStructuredText files (`.rst`):** Used in files like `torch.rst`, `nn.rst`. + +Same hierarchical structure without the markdown fences: + +```rst +torch.ao.quantization.fx.convert +--------------------------------- + +.. automodule:: torch.ao.quantization.fx.convert + +.. currentmodule:: torch.ao.quantization.fx.convert + +.. autosummary:: + :toctree: generated + :nosignatures: + + convert + convert_custom_module + convert_standalone_module + convert_weighted_module +``` + +For individual directives: + +```rst +.. automodule:: torch.submodule + +.. currentmodule:: torch.submodule + +.. autofunction:: function_name + +.. autoclass:: ClassName + :members: +``` + +**Key rules:** +- The module label comment from `conf.py` (e.g., `# torch.ao.quantization.fx.convert`) tells you exactly which `automodule` and `currentmodule` to use. +- Always set `.. automodule::` and `.. currentmodule::` before documenting functions from a module. +- Use **short names** (e.g., `convert`, not `torch.ao.quantization.fx.convert.convert`) after `currentmodule` is set. +- If the module already has an `automodule`/`currentmodule` in the file, don't add another — just add your function under the existing one. +- Match whichever style the file already uses (`autosummary` blocks vs. individual `autofunction` directives). + +#### Placing in the right section + +Read the target doc file and find the appropriate section. If the module already has a section (e.g., `## torch.backends.cuda` in `backends.md`), add the functions there. If no section exists yet, create one following the existing section patterns in the file. Group all functions from the same module group together. + +### Step 6: Verify with coverage + +Run coverage again: + +```bash +cd docs && make coverage +``` + +Ignore the terminal output — only read `docs/build/coverage/python.txt`. Verification passes when `python.txt` contains **zero undocumented functions across ALL modules**. It should only have the statistics table with 100% coverage and 0 undocumented for every module. For example: + +``` +Undocumented Python objects +=========================== + +Statistics +---------- + ++---------------------------+----------+--------------+ +| Module | Coverage | Undocumented | ++===========================+==========+==============+ +| torch | 100.00% | 0 | ++---------------------------+----------+--------------+ +| torch.accelerator | 100.00% | 0 | ++---------------------------+----------+--------------+ +``` + +If any module shows undocumented functions (coverage below 100% or undocumented count > 0), verification has failed. + +**If verification succeeds (zero undocumented across all modules):** Go to Step 7. + +**If verification fails (any undocumented functions remain):** Read `docs/build/coverage/python.txt` to see which functions are still listed as undocumented. Common issues include: + +- Wrong doc file: the function was added to the wrong `.md`/`.rst` file. Move the directive to the correct file. +- Wrong directive type: e.g., used `autofunction` for a class, or `autoclass` for a function. Fix the directive. +- Wrong module path in the directive: e.g., `torch.foo.bar` should be `torch.foo.baz.bar`. Correct the qualified name. +- Function added to an `autosummary` block with the wrong `currentmodule`: make sure the `.. currentmodule::` directive above the block matches. +- Missing `automodule` for a submodule that hasn't been registered yet. Add a `.. automodule:: torch.submodule` directive before documenting functions from that submodule. + +Fix the doc directive based on the error, then re-run `make coverage`. Repeat until verification passes. + +If a function still fails after multiple attempts, **stop and show the error to the user.** Present the function name and the error, then use the `AskUserQuestion` tool with options like: +- "Uncomment it to restore to ignore list (skip for now)" +- "Try a different approach" +- "Investigate further" + +### Step 7: Report progress + +**Present a progress summary to the user** showing: + +- Which module groups were processed and how many functions were documented +- Which functions were skipped or restored to the ignore list (and why) +- How many entries remain in `coverage_ignore_functions` and `coverage_ignore_classes` + +### Step 8: Clean up commented-out entries in conf.py + +Now that verification has passed, delete the commented-out string entries from Step 3. These are lines that start with `# "` inside `coverage_ignore_functions` and `coverage_ignore_classes`. Commented-out string entries always contain **quotes** — that's how you distinguish them from module label comments: + +```python +# "disable_global_flags", <-- commented-out string entry (has quotes) → DELETE +# torch.backends <-- module label comment (no quotes) → KEEP if it has active entries +``` + +Also delete any module label comments that no longer have active entries beneath them (i.e., all their entries were either commented out and now deleted, or had inline comments and were left in place but the module label is otherwise empty). + +## Important notes + +- **Follow the steps exactly as written.** The `make coverage` step is the primary verification for correct Sphinx directives, and Step 1b's docstring check ensures you only document functions that have real content. +- **Never modify Python source files (`.py`).** This skill only edits `docs/source/conf.py` and doc source files (`.md`/`.rst`) in `docs/source/`. Do not add or edit docstrings. The only reason to inspect Python modules is in Step 1b to check whether a docstring exists — never to modify source code. +- Entries are commented out in Step 3, verified in Step 6, and cleaned up in Step 8 after verification passes. Never delete uncommented entries directly. +- **Read inline comments** on entries before deciding to document them. Entries marked `# deprecated`, `# documented as ...`, `# looks unintentionally public`, or `# legacy helper` should stay in the ignore list. +- The `coverage_ignore_functions` list uses bare function names (not fully qualified), so the same name can appear multiple times for different modules. Use the module label comment above each entry to identify which module it belongs to. Be careful during Step 8 cleanup to only delete the correct commented-out lines — commented-out string entries have **quotes** (`# "func_name",`), module label comments do not. +- Always match the existing style of the target doc file — don't mix `.md` style directives into `.rst` files or vice versa. +- **Use the module label comment** (e.g., `# torch.ao.quantization.fx.convert`) as the primary guide for both the `automodule`/`currentmodule` directives and for finding the right section in the doc file. +- Always process complete module groups — never split a group across invocations. diff --git a/.claude/skills/metal-kernel/SKILL.md b/.claude/skills/metal-kernel/SKILL.md index 75e6684b73a40..18ba4f64570ea 100644 --- a/.claude/skills/metal-kernel/SKILL.md +++ b/.claude/skills/metal-kernel/SKILL.md @@ -325,6 +325,83 @@ python test/test_mps.py -k test_output_match_my_op python test/test_mps.py ``` +## Debugging Metal Kernels with `torch.mps.compile_shader` + +Use `torch.mps.compile_shader` to JIT-compile and test individual Metal kernels in isolation. This is invaluable for debugging multi-kernel pipelines where you need to verify each stage independently. + +### Basic Usage + +```python +import torch + +source = ''' +#include +using namespace metal; + +kernel void my_kernel( + const device float* input [[buffer(0)]], + device float* output [[buffer(1)]], + uint tid [[thread_position_in_grid]]) { + output[tid] = input[tid] * 2.0; +} +''' + +lib = torch.mps.compile_shader(source) + +inp = torch.tensor([1.0, 2.0, 3.0], device='mps') +out = torch.zeros(3, device='mps') +lib.my_kernel(inp, out, threads=[3, 1, 1], group_size=[3, 1, 1]) +torch.mps.synchronize() +print(out) # tensor([2., 4., 6.], device='mps:0') +``` + +### Dispatch Semantics + +`compile_shader` uses **`dispatchThreads`** semantics (same as `mtl_dispatch1DJob` in PyTorch): +- `threads=[N, 1, 1]` — total number of threads (NOT threadgroups) +- `group_size=[G, 1, 1]` — threads per threadgroup + +This differs from the `dispatchThreadgroups` API used by some host-side code. To match `dispatchThreadgroups:MTLSizeMake(num_tgs, num_slices, 1) threadsPerThreadgroup:MTLSizeMake(TG_SIZE, 1, 1)`: + +```python +# Equivalent compile_shader call: +lib.kernel(args..., + threads=[num_tgs * TG_SIZE, num_slices, 1], + group_size=[TG_SIZE, 1, 1]) +``` + +### Constant Buffer Parameters + +Pass scalar constants as single-element tensors: + +```python +slice_size = torch.tensor([1024], dtype=torch.int32, device='mps') +lib.my_kernel(data, output, slice_size, threads=[1024, 1, 1], group_size=[256, 1, 1]) +``` + +### Debugging Strategy for Multi-Kernel Pipelines + +When a pipeline of kernels (e.g., histogram → prefix_sum → scatter) produces wrong results, test each kernel individually and verify its output against a Python/NumPy reference: + +```python +# 1. Run GPU kernel +lib.histogram(keys, hist, ..., threads=[N, 1, 1], group_size=[256, 1, 1]) +torch.mps.synchronize() + +# 2. Compute reference in Python +ref_hist = compute_histogram_cpu(keys.cpu().numpy(), ...) + +# 3. Compare +assert np.array_equal(hist.cpu().numpy(), ref_hist), "Histogram mismatch!" +``` + +This isolates which kernel in the pipeline is broken, rather than debugging the entire pipeline at once. + +### Common Pitfalls + +- **Wrong `threads` count** — `threads` is total threads, not threadgroups. For 5 threadgroups of 256, use `threads=[1280, 1, 1]`. +- **Threadgroup memory** — `compile_shader` doesn't support `[[threadgroup(N)]]` parameters directly. If your kernel needs threadgroup memory, restructure to use `threadgroup` arrays declared inside the kernel body instead. + ## Checklist - [ ] Added MPS dispatch to `native_functions.yaml` diff --git a/.claude/skills/pr-review/SKILL.md b/.claude/skills/pr-review/SKILL.md index 020ad83671bbd..01d68fb6a7610 100644 --- a/.claude/skills/pr-review/SKILL.md +++ b/.claude/skills/pr-review/SKILL.md @@ -5,7 +5,7 @@ description: Review PyTorch pull requests for code quality, test coverage, secur # PyTorch PR Review Skill -Review PyTorch pull requests focusing on what CI cannot check: code quality, test coverage adequacy, security vulnerabilities, and backward compatibility. Linting, formatting, type checking, and import ordering are handled by CI. +Review PyTorch pull requests focusing on what CI cannot check: code quality, test coverage adequacy, security vulnerabilities, and backward compatibility. ## Usage Modes @@ -80,60 +80,79 @@ For local branch reviews: ### GitHub Actions Mode -When invoked via workflow, PR data is passed as context. The PR number or diff will be available in the prompt. +When invoked via `@claude /pr-review` on a GitHub PR, the action pre-fetches PR +metadata and injects it into the prompt. Detect this mode by the presence of +``, ``, and `` tags in the prompt. -## Review Workflow +The prompt already contains: +- PR metadata (title, author, branch names, additions/deletions, file count) +- PR body/description +- All comments and review comments (with file/line references) +- List of changed files with paths and change types + +Use git commands to get the diff and commit history. The base branch name is in the +prompt context (look for `PR Branch: -> ` or the `baseBranch` field). + +```bash +# Get the full diff against the base branch +git diff origin/...HEAD + +# Get diff stats +git diff --stat origin/...HEAD + +# Get commit history for this PR +git log origin/..HEAD --oneline + +# If the base branch ref is not available, fetch it first +git fetch origin --depth=1 +``` + +Do NOT use `gh` CLI commands in this mode -- only git commands are available. +All PR metadata, comments, and reviews are already in the prompt context; +only the diff and commit log need to be fetched via git. + +## Review Philosophy + +A single line of code can have deep cross-cutting implications: a missing device guard causes silent data corruption on multi-GPU, a missing `Composite` dispatch key breaks every out-of-tree backend, a manual dtype check instead of `TensorIterator` silently skips type promotion. **Treat every line as potentially load-bearing.** -### Step 1: Fetch PR Information +1. **Investigate, don't guess** — When uncertain whether a checklist item applies, spawn a sub-agent to read the relevant code. A reviewer who guesses wrong provides negative value. +2. **Review the design, not just the implementation** — A PR can have perfectly correct implementation of a bad design. Question side-channel communication, on/off private flags, and demand concrete interface documentation for new contracts between components. +3. **Focus on what CI cannot check** — Don't comment on formatting, linting, type errors, or CI failures. Focus on design quality, interface correctness, thread safety, BC implications, test adequacy, and pattern adherence. +4. **Everything is a must-fix** — There are no "nits." If it's worth mentioning, it's worth fixing. Every inconsistency degrades the codebase over time. +5. **Be specific and actionable** — Reference file paths and line numbers. Name the function/class/file the author should use. +6. **Match the immediate context** — Read how similar features are already implemented in the same file. Pattern mismatches within a file are always wrong. +7. **Assume competence** — The author knows PyTorch; explain only non-obvious context. +8. **No repetition** — Each observation appears in exactly one section of the review output. -For local mode, use `gh` commands to get: -1. PR metadata (title, description, author) -2. List of changed files -3. Full diff of changes -4. Existing comments/reviews -5. Fetch associated issue information when applicable +### Using sub-agents -### Step 2: Analyze Changes +The review checklist is large. You cannot hold the full context of every infrastructure system in your head. **Spawn sub-agents** to investigate whether checklist items apply: read surrounding code, infrastructure the PR should be using, or tests that should exist. Spawn them in parallel for independent areas. A typical medium PR should spawn 3-8 sub-agents. -Read through the diff systematically: +## Review Workflow + +### Step 1: Understand Context + +Before reviewing, build understanding of what the PR touches and why: 1. Identify the purpose of the change from title/description/issue 2. Group changes by type (new code, tests, config, docs) 3. Note the scope of changes (files affected, lines changed) +4. Spawn sub-agents to read the unchanged code surrounding each significantly changed file to understand existing patterns and infrastructure -### Step 3: Deep Review +### Step 2: Deep Review -Perform thorough line-by-line analysis using the review checklist. See [review-checklist.md](review-checklist.md) for detailed criteria covering: -- Code quality and design -- Testing adequacy -- Security considerations -- Performance implications -- Any behavior change not expected by author +Go through **every changed line** in the diff and evaluate it against the review checklist in [review-checklist.md](review-checklist.md). -### Step 4: Check Backward Compatibility +### Step 3: Check Backward Compatibility -Evaluate BC implications. See [bc-guidelines.md](bc-guidelines.md) for: -- What constitutes a BC-breaking change -- Required deprecation patterns -- Common BC pitfalls +Evaluate BC implications per [bc-guidelines.md](bc-guidelines.md). For non-trivial BC questions, spawn a sub-agent to search for existing callers of the modified API. -### Step 5: Formulate Review +### Step 4: Formulate Review -Structure your review with actionable feedback organized by category. - -## Review Areas - -| Area | Focus | Reference | -|------|-------|-----------| -| Code Quality | Abstractions, patterns, complexity | [review-checklist.md](review-checklist.md) | -| API Design | New patterns, flag-based access, broader implications | [review-checklist.md](review-checklist.md) | -| Testing | Coverage, patterns, edge cases | [review-checklist.md](review-checklist.md) | -| Security | Injection, credentials, input handling | [review-checklist.md](review-checklist.md) | -| Performance | Regressions, device handling, memory | [review-checklist.md](review-checklist.md) | -| BC | Breaking changes, deprecation | [bc-guidelines.md](bc-guidelines.md) | +Structure your review with actionable feedback organized by category. Every finding should be traceable to a specific line in the diff. ## Output Format -Structure your review as follows: +Structure your review as follows. **Omit sections where you have no findings** — don't write "No concerns" for every empty section. Only include sections with actual observations. ```markdown ## PR Review: # @@ -144,25 +163,29 @@ Structure your review as follows: Brief overall assessment of the changes (1-2 sentences). ### Code Quality -[Issues and suggestions, or "No concerns" if none] +[Issues and suggestions] -### API Design -[Flag new patterns, internal-access flags, or broader implications if any. Otherwise omit this section.] +### Infrastructure +[Flag any checklist items from the PyTorch Infrastructure section that apply. +Reference the specific infrastructure the PR should be using.] ### Testing -- [ ] Tests exist for new functionality -- [ ] Edge cases covered -- [ ] Tests follow PyTorch patterns (TestCase, assertEqual) -[Additional testing feedback] +[Testing adequacy findings — missing OpInfo usage, non-device-generic tests, etc.] + +### API Design +[Flag new patterns, internal-access flags, or broader implications if any.] ### Security -[Issues if any, or "No security concerns identified"] +[Issues if any] + +### Thread Safety +[Threading concerns if any] ### Backward Compatibility -[BC concerns if any, or "No BC-breaking changes"] +[BC concerns if any] ### Performance -[Performance concerns if any, or "No performance concerns"] +[Performance concerns if any] ### Recommendation **Approve** / **Request Changes** / **Needs Discussion** @@ -185,19 +208,13 @@ When requested, add file-specific feedback with line references: - `torch/nn/modules/linear.py:78` - This allocation could be moved outside the loop ``` -## Key Principles - -1. **No repetition** - Each observation appears in exactly one section. Never repeat the same issue, concern, or suggestion across multiple sections. If an issue spans categories (e.g., a security issue that also affects performance), place it in the most relevant section only. -2. **Focus on what CI cannot check** - Don't comment on formatting, linting, or type errors -3. **Be specific** - Reference file paths and line numbers -4. **Be actionable** - Provide concrete suggestions, not vague concerns -5. **Be proportionate** - Minor issues shouldn't block, but note them -6. **Assume competence** - The author knows PyTorch; explain only non-obvious context - ## Files to Reference -When reviewing, consult these project files for context: +When reviewing, consult these project files for context — read them rather than relying on memory, as they change frequently: - `CLAUDE.md` - Coding style philosophy and testing patterns - `CONTRIBUTING.md` - PR requirements and review process - `torch/testing/_internal/common_utils.py` - Test patterns and utilities - `torch/testing/_internal/opinfo/core.py` - OpInfo test framework +- `aten/src/ATen/native/native_functions.yaml` - Operator declarations (for checking tags, dispatch keys, structured kernels) +- `tools/autograd/derivatives.yaml` - Backward formulas (for checking if an op should register here) +- `aten/src/ATen/native/tags.yaml` - Operator semantic tags diff --git a/.claude/skills/pr-review/review-checklist.md b/.claude/skills/pr-review/review-checklist.md index 3f4f8929d7a74..7c588027cbdca 100644 --- a/.claude/skills/pr-review/review-checklist.md +++ b/.claude/skills/pr-review/review-checklist.md @@ -7,7 +7,10 @@ This checklist covers areas that CI cannot check. Skip items related to linting, ### Abstractions and Design - [ ] **Clear abstractions** - State management is explicit; no dynamic attribute setting/getting -- [ ] **Match existing patterns** - Code follows architectural patterns already in the codebase +- [ ] **No side-channel communication** - If behavior changes based on a hidden flag or dynamically-set attribute, the interface itself should change instead (different function signature, different class, different code path). Side-channel patterns (set a private flag in one place, check it in another via `getattr`) create undocumented behavioral modes +- [ ] **Proper interface, not on/off flags** - A private boolean that switches between two fundamentally different behaviors should be two separate code paths or a proper interface change, not a flag +- [ ] **Interface documentation** - New internal calling conventions, protocols, or contracts between components must have concrete documentation: what the caller provides, what the callee receives, what invariants hold, and cleanup responsibilities. Motivational comments ("this allows X") are not interface documentation +- [ ] **Match existing patterns in the same file** - Before accepting new code in a file, read how similar features are already implemented in that same file. If the file uses class attributes for boolean flags, new boolean flags must use class attributes. If the file uses a specific setter pattern, new setters must use the same pattern - [ ] **No over-engineering** - Only requested changes are made; no speculative features - [ ] **No premature abstraction** - Helpers and utilities are only created when reused; three similar lines is better than a one-use helper - [ ] **No trivial helpers** - Avoid 1-2 LOC helper functions used only once (unless significantly improves readability) @@ -29,14 +32,128 @@ When a PR introduces new API patterns, carefully evaluate the broader implicatio - [ ] **Useful comments only** - Comments explain non-obvious context that cannot be inferred locally. For large comment use the `# Note [Good note title]` and `See Note [Good note title]` to write larger comments that can be referenced from multiple places in the codebase. - [ ] **No backward-compatibility hacks** - Unused code is deleted completely, not renamed with underscores or marked with "removed" comments - [ ] **Appropriate complexity** - Solutions are as simple as possible for the current requirements +- [ ] **Documentation shows correct patterns only** - Docs and markdown files should show the right way to do things directly, not anti-patterns followed by corrections. Code examples must have correct indentation, names, and syntax + +### Initialization and Module Design + +- [ ] **No fragile init ordering** - If multiple imports/calls must happen in a specific undocumented order, flag the design. Dependencies should be explicit or combined into a single entry point +- [ ] **Idempotent global state** - Registries and global lists that accumulate entries must handle multiple calls safely (no duplicate registration, clear cleanup story) + +## PyTorch Infrastructure + +When a PR touches code in the scope of any item below, **stop and investigate** whether the established infrastructure should be used. -### Common Issues to Flag +### C++ Kernel Infrastructure -- Dynamic `setattr`/`getattr` for state management (prefer explicit class members) -- Unused imports, variables, or dead code paths -- Copy-pasted code that could be a shared helper -- Magic numbers without explanation -- Overly defensive error handling for impossible cases +- [ ] **TensorIterator** — PR adds or modifies C++ kernel code that iterates over tensor data (raw pointers, `at::parallel_for`, manual contiguity checks, manual output reshape/resize) +- [ ] **DispatchStub** — PR adds C++ kernel code with manual `if (device_type == kCPU) ... else if (device_type == kCUDA)` dispatch instead of using `DECLARE_DISPATCH` / `DEFINE_DISPATCH` / `REGISTER_DISPATCH` from `aten/src/ATen/native/DispatchStub.h` +- [ ] **Structured Kernels** — PR adds a new ATen operator with separate hand-written functional, inplace, and out= variants instead of using `structured: True` + `structured_delegate` in `native_functions.yaml` to generate boilerplate +- [ ] **TORCH_CHECK variants** — PR uses generic `TORCH_CHECK` for conditions that have a more specific variant: `ValueError` → `TORCH_CHECK_VALUE`, `IndexError` → `TORCH_CHECK_INDEX`, `TypeError` → `TORCH_CHECK_TYPE`, `NotImplementedError` → `TORCH_CHECK_NOT_IMPLEMENTED` +- [ ] **AT_DISPATCH macros** — PR manually switches on `dtype` with `if (dtype == kFloat) ... else if (dtype == kDouble)` instead of using `AT_DISPATCH_FLOATING_TYPES`, `AT_DISPATCH_ALL_TYPES_AND`, or the `AT_DISPATCH_SWITCH` / `AT_DISPATCH_CASE` pattern from `aten/src/ATen/Dispatch.h` +- [ ] **Device guards (RAII)** — PR manually saves/restores device context (`cudaSetDevice` + try/catch) instead of using `DeviceGuard` or `OptionalDeviceGuard` from `c10/core/DeviceGuard.h`. **Note:** Operators registered in `native_functions.yaml` get automatic `DeviceGuard` insertion from codegen (controlled by `device_guard: True`, the default) — do NOT flag missing device guards for these ops unless they explicitly set `device_guard: False` +- [ ] **Memory format propagation** — PR allocates output tensors with `at::empty(shape, options)` (defaulting to contiguous) without calling `input.suggest_memory_format()` to preserve ChannelsLast or other input formats +- [ ] **Subclass-safe tensor allocation** — PR uses `at::empty(shape, input.options())` instead of `input.new_empty(shape)` or `at::empty_like(input)`, which don't propagate tensor subclass metadata +- [ ] **TORCH_LIBRARY operator registration** — PR registers operators using manual dispatcher calls instead of `TORCH_LIBRARY` / `TORCH_LIBRARY_IMPL` macros from `torch/library.h` +- [ ] **TORCH_WARN_DEPRECATION** — PR uses `TORCH_WARN` for deprecation notices instead of `TORCH_WARN_DEPRECATION` which issues a proper `DeprecationWarning` + +### CUDA & Device Management + +- [ ] **C10_CUDA_CHECK** — PR calls raw CUDA APIs (`cudaMalloc`, `cudaMemcpy`, etc.) without wrapping in `C10_CUDA_CHECK()` from `c10/cuda/CUDAException.h` +- [ ] **C10_CUDA_KERNEL_LAUNCH_CHECK** — PR launches CUDA kernels with `<<<>>>` syntax but doesn't follow with `C10_CUDA_KERNEL_LAUNCH_CHECK()` immediately after to detect launch errors early +- [ ] **CUDAStreamGuard** — PR manually manages CUDA streams (`cudaStreamCreate`/`cudaStreamDestroy`) instead of using `CUDAStreamGuard` or getting streams from `at::cuda::getCurrentCUDAStream()` / `getStreamFromPool()` +- [ ] **CUDAEvent synchronization** — PR uses `cudaDeviceSynchronize()` or `cudaStreamSynchronize()` for cross-stream ordering instead of `CUDAEvent::record()` + `CUDAEvent::block()` which avoids unnecessary full synchronization +- [ ] **recordStream for allocator** — PR uses tensors on non-default CUDA streams without calling `c10::cuda::CUDACachingAllocator::recordStream()` to prevent premature memory reuse +- [ ] **CUDA graph compatibility** — PR adds host-GPU synchronization, unpinned memory transfers, or other graph-unsafe operations without checking `currentStreamCaptureStatusMayInitCtx()` to detect CUDA graph capture mode +- [ ] **AcceleratorHooksInterface** — PR adds device-specific `#ifdef USE_CUDA` blocks in generic code instead of using `AcceleratorHooksInterface` from `aten/src/ATen/detail/AcceleratorHooksInterface.h` for device-agnostic behavior +- [ ] **DeviceGuardImplInterface** — PR implements custom device management without going through `DeviceGuardImplInterface` from `c10/core/impl/DeviceGuardImplInterface.h`, bypassing the standard device abstraction layer + +### Operator Registration & Codegen + +- [ ] **native_functions.yaml** — PR adds a new ATen operator by writing manual C++ bindings and Python wrappers instead of declaring it in `aten/src/ATen/native/native_functions.yaml` and letting codegen produce the boilerplate +- [ ] **Operator tags** — PR adds an operator to `native_functions.yaml` without appropriate tags from `tags.yaml` (e.g., `pointwise`, `reduction`, `view_copy`, `core`, `pt2_compliant_tag`) +- [ ] **Missing Composite fallback** — PR adds a new operator to `native_functions.yaml` with only backend-specific dispatch keys (e.g., `CPU`, `CUDA`) but no `CompositeImplicitAutograd` or `CompositeExplicitAutograd` fallback. Without a Composite entry, the op will fail on all backends that don't have an explicit registration (XLA, MPS, HPU, PrivateUse1, etc.). Every new op should either have a Composite implementation or a clear justification for why it can only work on specific backends +- [ ] **Meta function registration** — PR adds a new operator without a meta (shape-only) implementation, blocking `torch.compile` and `torch.export`. Meta implementations can be registered in Python via `@register_meta` from `torch/_meta_registrations.py`, via `torch.library.impl(..., "Meta")`, or in C++ as a structured kernel with a `meta` dispatch key or via any `Composite` dispatch key in `native_functions.yaml` (since Composite kernels automatically work on Meta tensors) +- [ ] **Fake tensor implementation** — PR adds a custom op (registered via `torch.library`) without a fake implementation. Custom ops need `@register_fake` / `my_op.register_fake()` for `torch.compile` to trace through the op. For C++ ops registered via `native_functions.yaml`, the meta kernel serves this purpose. For Python `torch.library` custom ops, use `@torch.library.register_fake("mylib::my_op")` or `@my_op.register_fake` to provide a shape/dtype-only implementation. The fake impl receives `FakeImplCtx` with `ctx.new_dynamic_size()` for data-dependent output shapes +- [ ] **Schema annotations** — PR defines operator schemas without proper alias annotations (`Tensor(a)`, `Tensor(a!)`) for view and in-place ops, which breaks functionalization and autograd's alias tracking + +### Autograd + +- [ ] **derivatives.yaml** — PR writes a custom `autograd.Function` subclass for an operation that should have its backward formula registered in `tools/autograd/derivatives.yaml` (the centralized backward formula registry for ATen ops) +- [ ] **setup_context pattern** — PR writes `autograd.Function` with `forward(ctx, ...)` (legacy pattern) instead of separated `forward(...)` + `setup_context(ctx, inputs, output)` which is required for functorch compatibility (vmap, grad) +- [ ] **ctx.save_for_backward** — PR saves tensors in `autograd.Function` via `ctx.my_tensor = tensor` instead of `ctx.save_for_backward(tensor)`, causing memory leaks by keeping tensors alive longer than needed +- [ ] **gradcheck testing** — PR adds custom backward logic but doesn't test it with `torch.autograd.gradcheck()` / `gradgradcheck()` which verify numerical correctness of gradients via finite differences +- [ ] **Forward-mode AD** — PR adds a new differentiable op with backward formula in `derivatives.yaml` but doesn't add a `result:` entry for forward-mode AD (JVP). Can often use `auto_element_wise` or `auto_linear` for automatic generation +- [ ] **register_autograd for custom ops** — PR writes a full `autograd.Function` subclass for a custom op registered via `torch.library` instead of using the simpler `@my_op.register_autograd(backward, setup_context=...)` API +- [ ] **Vmap rule for custom ops** — PR adds a custom op or `autograd.Function` without a vmap rule (`generate_vmap_rule = True` or manual `vmap()` static method), breaking `torch.vmap` support + +### Python Utilities + +- [ ] **__torch_function__ support** — PR adds a new Python-level function that takes tensors but doesn't check `has_torch_function()` / call `handle_torch_function()`, breaking tensor subclass dispatch +- [ ] **Pytree registration** — PR manually flattens/unflattens custom container types (dataclasses, named tuples) instead of registering them with `torch.utils._pytree.register_pytree_node()` or `register_dataclass()` +- [ ] **tree_map** — PR manually walks nested structures of tensors with recursive functions instead of using `torch.utils._pytree.tree_map()` +- [ ] **_DecoratorContextManager** — PR implements a context manager that should also work as a decorator but doesn't inherit from `torch.utils._contextlib._DecoratorContextManager` +- [ ] **Deprecation utilities** — PR deprecates a function using ad-hoc `warnings.warn()` calls instead of PyTorch's deprecation infrastructure (`lazy_deprecated_import` for module-level, `TORCH_WARN_DEPRECATION` for C++) +- [ ] **No print statements** — PR adds `print()` calls for debugging or diagnostics. Use `torch._logging` utilities instead (`getArtifactLogger`, `LazyString`, `warning_once`). For the `torch.compile` stack specifically, use `trace_structured()` for structured artifacts that integrate with `tlparse` for production debugging. No bare `print()` should ever land in production code +- [ ] **torch.backends context** — PR manually saves/restores backend flags (`cudnn.deterministic`, etc.) instead of using the `torch.backends.cudnn.flags()` context manager + +### nn Module Patterns + +- [ ] **ModuleList / ModuleDict** — PR stores submodules in plain Python `list` or `dict` instead of `nn.ModuleList` or `nn.ModuleDict`, causing them to be invisible to `parameters()`, `to()`, `state_dict()`, etc. +- [ ] **nn.init methods** — PR manually initializes weights with `self.weight.data.normal_(0, 0.01)` instead of using `torch.nn.init.kaiming_uniform_()`, `xavier_uniform_()`, etc., which handle fan-in/fan-out calculations correctly +- [ ] **Parametrization framework** — PR implements custom weight reparameterization via forward pre-hooks (the deprecated pattern) instead of using `torch.nn.utils.parametrize.register_parametrization()` +- [ ] **_load_from_state_dict versioning** — PR changes a module's parameter layout without implementing `_load_from_state_dict()` for backward-compatible loading of old checkpoints (see BatchNorm's `_version = 2` pattern) +- [ ] **clip_grad_norm_** — PR manually computes gradient norms and clips in training loops instead of using `torch.nn.utils.clip_grad_norm_()` or `clip_grad_value_()` +- [ ] **LazyModule pattern** — PR implements deferred parameter initialization with manual shape inference in `forward()` instead of using `LazyModuleMixin` with `UninitializedParameter` + +### Dynamo / Inductor / Compile + +- [ ] **@register_lowering** — PR adds Inductor support for an op by modifying core lowering code instead of using `@register_lowering(aten.my_op)` from `torch/_inductor/lowering.py` with automatic type promotion and broadcasting +- [ ] **Inductor decompositions** — PR writes a full Inductor lowering for a complex op that can be decomposed into simpler already-lowered ops via `@register_decomposition` in `torch/_inductor/decomposition.py` +- [ ] **CustomGraphPass** — PR writes ad-hoc FX graph iteration for Inductor optimization instead of implementing `CustomGraphPass` (with `__call__` and `uuid()`) from `torch/_inductor/custom_graph_pass.py` +- [ ] **config.patch** — PR manually saves/restores Dynamo or Inductor config values in tests instead of using `torch._dynamo.config.patch()` as a decorator or context manager +- [ ] **Graph break hints** — PR calls `unimplemented()` in Dynamo without providing `gb_type`, `explanation`, or `hints` (like `SUPPORTABLE`, `FUNDAMENTAL`), making it hard for users to understand and fix graph breaks +- [ ] **Dynamo trace rules** — PR adds manual skip/inline logic in Dynamo variable tracking instead of updating `manual_torch_name_rule_map`, `MOD_INLINELIST`, or `MOD_SKIPLIST` in `torch/_dynamo/trace_rules.py` +- [ ] **torch.compile compatibility** — PR adds a new op or modifies an existing one without verifying it works under `torch.compile` (should test with `pt2_compliant_tag` and run opcheck) + +### FX / Export + +- [ ] **FX PassBase** — PR writes a custom FX graph transformation with manual graph walking instead of inheriting from `PassBase` (with `requires()`, `call()`, `ensures()`) from `torch/fx/passes/infra/pass_base.py` +- [ ] **FX PassManager** — PR manually orders and applies multiple FX passes instead of using `PassManager` with `this_before_that_pass_constraint` from `torch/fx/passes/infra/pass_manager.py` +- [ ] **FX Interpreter** — PR manually iterates FX graph nodes and tracks values in a dict instead of subclassing `torch.fx.Interpreter` which provides structured `run_node()` / `call_function()` / `call_module()` overrides +- [ ] **Subgraph rewriter** — PR manually matches and replaces graph patterns instead of using `replace_pattern()` from `torch/fx/subgraph_rewriter.py` +- [ ] **ShapeProp** — PR manually executes FX graphs to annotate shapes on nodes instead of using `ShapeProp(gm).propagate(*args)` from `torch/fx/passes/shape_prop.py` +- [ ] **torch.export dynamic shapes** — PR hard-codes tensor shapes in export constraints instead of using `Dim(name, min, max)` and `dims()` from `torch/export/dynamic_shapes.py` +- [ ] **make_fx** — PR manually initializes `torch.fx.Tracer` for proxy-based tracing instead of using `make_fx(f, tracing_mode="symbolic")` from `torch/fx/experimental/proxy_tensor.py` + +### Type Promotion & Dtypes + +- [ ] **elementwise_dtypes / TensorIterator** — PR manually implements type promotion logic for elementwise ops. In Python, use `elementwise_dtypes()` from `torch/_prims_common/` with the appropriate `ELEMENTWISE_TYPE_PROMOTION_KIND`. In C++, use `TensorIteratorConfig` which handles type promotion automatically: call `.promote_inputs_to_common_dtype(true)` and `.cast_common_dtype_to_outputs(true)` on the config builder, then `TensorIterator` computes `common_dtype()` for the kernel and handles all input/output casting. Kernels should operate on `iter.common_dtype()` via `AT_DISPATCH` rather than manually checking and promoting dtypes +- [ ] **result_type** — PR manually resolves output dtype from mixed-dtype inputs instead of using `torch.result_type()` (Python) or `at::result_type()` / `update_result_type_state()` (C++) +- [ ] **Complex dtype handling** — PR manually maps between complex and real dtypes (e.g., `complex64` to `float32`) instead of using `corresponding_real_dtype()` / `corresponding_complex_dtype()` from `torch/_prims_common/` +- [ ] **promoteTypes** — PR writes manual dtype promotion tables instead of using `c10::promoteTypes(a, b)` from `c10/core/ScalarType.h` + +### Serialization + +- [ ] **weights_only=False** — PR adds `torch.load(..., weights_only=False)`, explicitly opting out of safe deserialization. `weights_only=True` is already the default; setting it to `False` enables arbitrary code execution via pickle and is almost never the right thing to do. Flag this and ask the author to register safe globals via `torch.serialization.add_safe_globals()` instead +- [ ] **safe_globals** — PR adds new types to serialization that should be loadable with `weights_only=True` but doesn't register them via `torch.serialization.add_safe_globals()` +- [ ] **skip_data context** — PR implements metadata-only checkpoint inspection by reading full tensors instead of using `torch.serialization.skip_data()` context manager + +### Distributed + +- [ ] **DeviceMesh** — PR manually creates multiple `ProcessGroup`s for multi-dimensional parallelism (TP + DP) instead of using `DeviceMesh` from `torch/distributed/device_mesh.py` which manages this automatically +- [ ] **Distributed testing** — PR spawns multiple real processes for distributed unit tests instead of using `MultiThreadedPG` from `torch/testing/_internal/distributed/multi_threaded_pg.py` for single-process testing + +### Tensor Subclasses + +- [ ] **_make_wrapper_subclass** — PR creates tensor subclasses by calling `torch.Tensor.__new__()` directly instead of using `torch.Tensor._make_wrapper_subclass()` which properly sets up the subclass wrapper +- [ ] **__tensor_flatten__ / __tensor_unflatten__** — PR adds a tensor subclass without implementing `__tensor_flatten__()` and `__tensor_unflatten__()`, breaking serialization and `torch.compile` support + +### Miscellaneous + +- [ ] **torch._check** — PR uses `assert` or `if not cond: raise` in Python op implementations instead of `torch._check()` / `torch._check_is_size()` which work correctly with meta tensors and symbolic shapes +- [ ] **C++ extension building** — PR uses raw `setuptools` or `distutils` for building C++ extensions instead of `torch.utils.cpp_extension.CppExtension` / `CUDAExtension` / `load_inline()` which handle compiler flags, ABI, and includes +- [ ] **register_package for custom devices** — PR adds custom device serialization handling by monkey-patching `torch.save`/`torch.load` instead of using `torch.serialization.register_package()` to register location tag and restore functions +- [ ] **@register_backend** — PR adds a `torch.compile` backend by manually modifying internal dispatch tables instead of using `@torch._dynamo.backends.registry.register_backend(name=...)` ## Testing @@ -48,41 +165,29 @@ When a PR introduces new API patterns, carefully evaluate the broader implicatio ### Test Patterns -- [ ] **Use OpInfo** - Any testing for an operator or a cross cutting feature must be done via OpInfo +- [ ] **Proper module ownership** - Test files must have a real `# Owner(s): ["module: ..."]` label, not `"module: unknown"`. The author should create a new module label if needed and add themselves as owner +- [ ] **Use OpInfo** - Any testing for an operator or a cross-cutting feature must be done via OpInfo. Flag manual tests (e.g., `assertEqual(a + b, expected)`) for operators that already have OpInfo entries — these are redundant and will rot. When a PR adds dtype/device support to an operator, the testing should come from existing OpInfo infrastructure automatically (e.g., by adding the dtype to the operator's OpInfo `dtypes`), not from new manual tests. Likewise, a test checking a specific behavior for a single operator should not be a standalone test — the OpInfo infrastructure for that test category should be updated to cover the behavior across all applicable operators +- [ ] **Use ModuleInfo** - Manual forward/backward tests for `nn.Module` subclasses should use `ModuleInfo` from `torch/testing/_internal/common_modules.py` and the `@modules` decorator instead of hand-written per-module tests - [ ] **Use TestCase** - Tests inherit from `torch.testing._internal.common_utils.TestCase` - [ ] **Use run_tests** - Test file ends with `if __name__ == "__main__": run_tests()` -- [ ] **Use assertEqual for tensors** - Tensor comparisons use `assertEqual`, not raw assertions +- [ ] **Use assertEqual for tensors** - Tensor comparisons use `assertEqual`, not raw assertions or `torch.allclose` +- [ ] **Device generic** - Any test checking compute result should happen in a device-generic test class (taking device as an argument) via `instantiate_device_type_tests`. Device-specific tests should be very rare and in device-specific test files +- [ ] **Use @dtypes** - PR writes separate test methods per dtype or manual `for dtype in [...]` loops instead of using the `@dtypes(...)` decorator from `common_device_type.py` +- [ ] **Use @parametrize** - PR duplicates test methods that differ only in a parameter instead of using `@parametrize` from `common_utils.py` +- [ ] **Use @ops for operator tests** - PR writes manual per-operator test iterations instead of using the `@ops(op_db)` decorator which automatically parametrizes tests over OpInfo entries +- [ ] **Use make_tensor** - PR creates test tensors with `torch.rand(shape)` (implicit CPU, implicit dtype) instead of `make_tensor(shape, device=device, dtype=dtype)` from `torch.testing` which enforces explicit device/dtype +- [ ] **Use common dtype groups** - PR manually lists dtypes like `[torch.float32, torch.float64]` instead of using helpers like `floating_types()`, `all_types_and_complex()`, etc. from `common_dtype.py` +- [ ] **Use toleranceOverride** - PR hard-codes tolerance values in individual assertions instead of using `@toleranceOverride` / `@precisionOverride` decorators which set per-dtype tolerances +- [ ] **Use DecorateInfo for OpInfo skips** - PR adds `@skipIf` conditionals inside OpInfo test methods instead of using `DecorateInfo` in the OpInfo's `skips` or `decorators` tuple +- [ ] **Use largeTensorTest** - PR manually checks free memory before large-tensor tests instead of using `@largeTensorTest("4 GB")` decorator from `common_device_type.py` - [ ] **Descriptive test names** - Test method names describe what is being tested -- [ ] **Device generic** - Any test checking compute result should happen in a Device-generic test class (taking device as an argument). Device-specific test should be very rare and in device-specific test files. ### Test Quality - [ ] **Edge cases covered** - Tests include boundary conditions, empty inputs, error cases -- [ ] **Error conditions tested** - Expected exceptions are tested with `assertRaises` or `assertRaisesRegex` -- [ ] **No duplicated test logic** - Similar tests share a private helper method (e.g., `_test_foo(config)`) called from individual tests with different configs - -**Example of good test structure:** -```python -def _test_feature_with_config(self, flag, expected_shape): - """Shared test logic called by device-specific tests.""" - x = torch.randn(10) - result = my_feature(x, flag) - self.assertEqual(result.shape, expected_shape) - -def test_feature_enabled(self): - self._test_feature_with_config(True, (10, 10)) - -def test_feature_disabled(self): - self._test_feature_with_config(False, (10, 5)) -``` - -### Common Testing Issues - -- Tests that only check the happy path without error cases -- Duplicated test code that should be a parameterized helper -- Tests that don't clean up resources (files, CUDA memory) -- Flaky tests (timing-dependent, order-dependent, golden value) -- Tests that skip without clear justification +- [ ] **Error conditions tested** - Expected exceptions are tested with `assertRaisesRegex`, not bare `assertRaises`. `assertRaisesRegex` verifies both the exception type and message, catching cases where the right exception is raised for the wrong reason. Bare `assertRaises` should be flagged — always require a message pattern match +- [ ] **No duplicated test logic** - Similar tests share a private helper method called from individual tests with different configs +- [ ] **Use weakref for lifetime testing** - PR uses `sys.getrefcount()` to test whether objects are kept alive. Use `weakref.ref()` instead — create a weak reference, delete the strong references, then check if the weakref is dead (`wr() is None`). `sys.getrefcount` is a CPython implementation detail that varies across versions and is fragile ## Security @@ -105,6 +210,43 @@ When reviewing changes to PyTorch APIs and user-facing code: - [ ] **Distributed primitives** - `torch.distributed`, RPC, and TCPStore have no auth/encryption and accept connections from anywhere; they are for internal networks only, not untrusted environments - [ ] **No new pickle usage** - Avoid adding `pickle.load` or `torch.load` without `weights_only=True` on paths that could receive untrusted data +## Thread Safety & Concurrency + +### Python Threading + +- [ ] **No unprotected shared mutable state** - Shared data structures accessed from multiple threads are protected by locks or are inherently thread-safe +- [ ] **Lock ordering** - When multiple locks are acquired, ordering is consistent to avoid deadlocks +- [ ] **No GIL-reliant correctness** - Code that mutates shared state should not rely on the GIL for thread safety, since the GIL may not be present in free-threaded builds + +### C++ Threading + +- [ ] **No data races** - Shared mutable state is protected by mutexes or uses atomics with appropriate memory ordering +- [ ] **RAII lock guards** - Prefer `std::lock_guard` or `std::unique_lock` over manual `lock()`/`unlock()` to ensure exception-safe unlocking +- [ ] **No lock-order inversions** - When acquiring multiple locks, a consistent global ordering is followed +- [ ] **Correct atomic memory ordering** - `std::memory_order_relaxed` is only used when ordering with other operations is genuinely unnecessary; default to `seq_cst` or use `acquire`/`release` pairs + +### CPython C API Thread Safety + +This is particularly important for PyTorch's autograd, which has multi-threaded C++ code calling into the CPython C API. + +- [ ] **GIL held for Python object access** - Any code that touches `PyObject*` (incref, decref, attribute access, container mutation) must hold the GIL. When releasing the GIL for long-running C++ work (`Py_BEGIN_ALLOW_THREADS`), verify no Python objects are accessed in that region +- [ ] **Borrowed references across GIL release** - Borrowed references (`PyTuple_GET_ITEM`, `PyList_GET_ITEM`) become unsafe if the GIL is released and reacquired, since another thread may have mutated the container +- [ ] **Decref-before-update hazard** - When replacing an item in a container (tuple, list, dict), update the container slot first, then `Py_DECREF` the old value. Decref can trigger `__del__` finalizers that re-enter and observe the container in an inconsistent state. Without the GIL (free-threaded builds), this is also a data race. This is **always** a must-fix — even if "safe in practice" because of refcount guarantees, the pattern is wrong and breaks under NoGIL. The correct pattern costs nothing extra + +### Free-Threaded Python (NoGIL, PEP 703) + +CPython 3.13t+ can run without the GIL. Code that was previously safe under the GIL may have races in free-threaded builds: + +- [ ] **No implicit GIL serialization assumptions** - Code paths that assume only one thread can execute Python at a time are broken under NoGIL. Look for shared mutable state accessed from C extensions without explicit locking +- [ ] **Raw `PyTuple_SET_ITEM` / `PyList_SET_ITEM`** - These are raw slot writes with no memory ordering guarantees. In free-threaded builds, concurrent reads from other threads may see stale or torn values. Consider whether the data structure could be accessed concurrently and whether atomic operations or the thread-safe API alternatives are needed +- [ ] **Module-level mutable state in C extensions** - Global/static `PyObject*` variables or C-level caches accessed from multiple threads need synchronization in NoGIL builds + +### PyTorch-Specific Concurrency + +- [ ] **Autograd engine multi-threading** - The autograd engine runs node `apply()` methods from worker threads. Code in custom autograd node implementations must be safe for concurrent execution across different nodes, and must hold the GIL when accessing Python objects +- [ ] **CUDA stream synchronization** - Operations across different CUDA streams require explicit synchronization (`cudaStreamSynchronize`, `cudaEventRecord`/`cudaStreamWaitEvent`). Missing synchronization can cause silent data corruption +- [ ] **DataLoader worker safety** - Objects shared between the main process and DataLoader worker processes (or threads) must be fork-safe or use appropriate IPC mechanisms + ## Performance ### Obvious Regressions @@ -125,9 +267,7 @@ When reviewing changes to PyTorch APIs and user-facing code: - [ ] **Efficient data structures** - Appropriate containers for access patterns - [ ] **Gradient memory** - Proper use of `no_grad()`, `detach()` to avoid unnecessary graph retention -### Common Performance Issues +### Profiling & Benchmarking -- Creating new tensors inside training loops instead of pre-allocating -- Synchronous CUDA operations where async would work -- Keeping computation graph alive longer than needed -- Redundant clones or copies +- [ ] **Use torch.profiler** - PR adds manual `time.time()` instrumentation instead of using `torch.profiler.profile()` context manager with `schedule()` and `tensorboard_trace_handler()` +- [ ] **Use torch.utils.benchmark.Timer** - PR benchmarks with `time.time()` loops instead of `torch.utils.benchmark.Timer` which handles warmup, statistics, and proper CUDA synchronization diff --git a/.claude/skills/pt2-bug-basher/SKILL.md b/.claude/skills/pt2-bug-basher/SKILL.md new file mode 100644 index 0000000000000..5c30f9002f79c --- /dev/null +++ b/.claude/skills/pt2-bug-basher/SKILL.md @@ -0,0 +1,273 @@ +--- +name: pt2-bug-basher +disable-model-invocation: true +description: Debug PyTorch 2 compiler stack failures including Dynamo graph breaks, Inductor codegen errors, AOTAutograd crashes, and accuracy mismatches. Use when encountering torch.compile errors, BackendCompilerFailed exceptions, recompilation issues, Triton kernel failures, FX graph problems, or when the user mentions debugging PT2, Dynamo, Inductor, or compiled model issues. +--- + +# PT2 Bug Basher + +Debug test failures and runtime errors in the PyTorch 2 compiler stack (Dynamo, Inductor, AOTAutograd, FX graphs). + +## Workflow Summary + +1. **Environment check** -- Ask the user which conda environment to use. Verify it is active by checking `$CONDA_DEFAULT_ENV`. Then run `python -c "import torch; print(torch.__version__)"` to confirm torch is importable and report the version. If the environment is not active or torch cannot be imported, stop and ask the user to activate the correct environment before proceeding. +2. **Reproduce** -- Get a consistent reproduction of the failure +3. **Minimize** -- Reduce the repro to the smallest possible standalone case. Strip away unrelated model logic, use minimal tensor shapes, and isolate the specific op or pattern that triggers the bug. +4. **Add a unit test** -- **Do this BEFORE diving into code search or root cause investigation.** Add a failing test to the codebase that captures the bug. Place it in a specific, topic-appropriate test file (e.g., `test/dynamo/test_repros.py`, `test/inductor/test_torchinductor.py`, `test/export/test_export.py`). **Avoid `test/dynamo/test_misc.py`** — it is already oversized; find a more specific test file that matches the area of the bug. Use `torch.testing._internal.common_utils.TestCase` and `run_tests`. The test must fail before the fix and pass after. Having the test first keeps you grounded — you know exactly what "fixed" looks like before you start exploring the codebase. +5. **Validate on main** -- Use `EnterWorktree` to create a worktree checked out at `main`. Copy the new test file into the worktree and run the test there to confirm it **fails** on main. If the test passes on main, stop — the test may not be capturing the right bug, or the bug may already be fixed. Exit the worktree with `ExitWorktree` (action: remove) and return to the working branch before continuing. +6. **Gather logs** -- Run with appropriate `TORCH_LOGS` settings +7. **Classify** -- Use the [Error Triage](#error-triage) table to identify the category +8. **Inspect artifacts** -- Check FX graphs, IR, and generated code via `TORCH_COMPILE_DEBUG=1` +9. **Identify root cause** -- Trace from the error back through the compilation pipeline +10. **Fix** -- Apply the fix +11. **Verify** -- Run the new unit test AND nearby related existing tests (e.g., if you changed how `is_exporting` works, also run the existing `test_is_exporting` export test). Use `pytest -k` to quickly run related tests by name. The task is not complete until all pass. +12. **Self-review** -- Use the `/pr-review` skill to review your own changes before presenting them. Fix any issues it flags. +13. **Celebrate** -- Summarize the changes: explain the root cause, what was changed and why, and which tests were added/verified. Then tell the user the bug is squashed. Include a fun, varied motivational message or easter egg to keep spirits high (e.g., a pun, a quote, an ASCII art bug getting squashed). Keep it short and different each time. + +## Investigation Strategy + +### Prefer direct tools over meta_codesearch + +Use `Grep`, `Glob`, and `Read` directly for code exploration. **Do not spawn `meta_codesearch` agents** — they are slow and expensive. The [Architectural Knowledge](#architectural-knowledge) and [Key Source Files](#key-source-files) sections below should give you enough context to know where to look. A targeted `Grep` for a function name is always faster. + +### Know which compilation mode you're in + +Before reading implementation code, determine the compilation mode. These share code but diverge in important ways: +- **`torch.compile`** -- Dynamo + Inductor. `tx.export=False`, no `_compiling_state_context()`. +- **`torch.export` (strict)** -- `tx.export=True`, `_compiling_state_context()` active. +- **`torch.export` (non-strict, **the default**)** -- Uses Dynamo via `fullgraph_capture` but `tx.export` may differ from strict. `_compiling_state_context()` active. Check `torch._export.config.use_new_tracer_experimental` — it changes which code path is used. + +### Distinguish trace-time vs runtime + +Many PT2 bugs come from confusing these two: +- **Trace-time**: Inside Dynamo's symbolic interpreter. Dynamo intercepts function calls and may constant-fold them (e.g., `is_exporting()` → `ConstantVariable(True)`). +- **Runtime**: Real tensors, real Python calls, module-level flags like `torch.compiler._is_exporting_flag`. + +When debugging, add temporary `print()` statements directly in the source file rather than monkey-patching from outside — dispatch chains make monkey-patching unreliable. + +## Gathering Information + +Pick the right diagnostic tool based on the error category: + +- **Quick overview**: `TORCH_LOGS="+dynamo,graph_breaks,recompiles" python your_script.py` +- **Full debug artifacts**: `TORCH_COMPILE_DEBUG=1 python your_script.py` — creates `torch_compile_debug/` with FX graphs, Inductor IR, and generated code +- **Generated code only**: `TORCH_LOGS="output_code" python your_script.py` +- **Structured tracing**: `TORCH_TRACE=/path/to/trace python your_script.py` then `tlparse /path/to/trace` +- **Single-threaded (for pdb)**: `TORCHINDUCTOR_COMPILE_THREADS=1 python your_script.py` + +## Error Triage + +Classify the failure using the error message and traceback: + +| Error Pattern | Category | Jump To | +|---|---|---| +| `Unsupported: ...` or `graph break` in logs | Graph break | [Graph Breaks](#graph-breaks) | +| `BackendCompilerFailed` | Inductor/backend crash | [Backend Failures](#backend-compiler-failures) | +| `RecompileError` or `cache_size_limit` | Recompilation | [Recompilation](#recompilation-issues) | +| Accuracy mismatch / wrong numerical output | Accuracy | [Accuracy](#accuracy-issues) | +| `InternalTorchDynamoError` | Dynamo bug | [Internal Errors](#internal-dynamo-errors) | +| Segfault or CUDA IMA | Runtime crash | [Runtime Crashes](#runtime-crashes) | +| Triton assertion / index out of bounds | Triton kernel bug | [Triton Failures](#triton-kernel-failures) | + +## Debugging by Category + +### Graph Breaks + +Graph breaks split the compiled graph into smaller subgraphs, often causing performance regressions or unexpected behavior. + +**Diagnosis:** +```bash +TORCH_LOGS="graph_breaks" python your_script.py +``` + +**Key files:** +- `torch/_dynamo/exc.py` -- `Unsupported` exception class +- `torch/_dynamo/variables/` -- where most graph break decisions happen + +**Common causes:** +- Unsupported Python constructs (data-dependent control flow, unsupported builtins) +- Tensor operations that can't be traced (in-place ops on inputs, unsupported dtypes) +- Calls to non-traceable functions + +**Fix approach:** +1. Read the graph break message to identify the unsupported operation +2. Check if there's a decomposition or supported alternative +3. If the operation genuinely can't be traced, consider `torch._dynamo.allow_in_graph` or restructuring user code + +### Backend Compiler Failures + +`BackendCompilerFailed` means Inductor (or another backend) crashed during compilation. + +**Diagnosis:** +```bash +TORCHDYNAMO_REPRO_AFTER=aot TORCHDYNAMO_REPRO_LEVEL=2 python your_script.py +``` + +This generates `minifier_launcher.py` that isolates the minimal failing graph. + +**Key files:** +- `torch/_dynamo/repro/after_aot.py` -- repro/minifier for post-AOT failures +- `torch/_inductor/` -- the backend itself + +**Fix approach:** +1. Run the minifier to get a minimal reproduction +2. Inspect the FX graph (`TORCH_COMPILE_DEBUG=1`) to understand what ops are involved +3. Check if it's a lowering issue (`torch/_inductor/lowering.py`), scheduling issue, or codegen issue +4. Look at the generated output code if the error is in codegen + +### Recompilation Issues + +Excessive recompilation happens when guards are too specific, causing cache misses. + +**Diagnosis:** +```bash +TORCH_LOGS="recompiles,recompiles_verbose,guards" python your_script.py +``` + +**Key config:** +- `torch._dynamo.config.recompile_limit` (default: 8) +- `torch._dynamo.config.fail_on_recompile_limit_hit` -- set to `True` to get a hard error + +**Common causes:** +- Changing tensor shapes without marking them dynamic +- Python scalar values that change between calls +- Global state mutations between calls + +**Fix approach:** +1. Read the recompilation reason from logs +2. Identify the failing guard +3. Either mark the relevant dimension as dynamic with `torch._dynamo.mark_dynamic()` or fix the source of guard instability + +### Accuracy Issues + +The compiled model produces different numerical results than eager mode. + +**Diagnosis:** +```bash +TORCHDYNAMO_REPRO_AFTER=aot TORCHDYNAMO_REPRO_LEVEL=4 python your_script.py +``` + +This compares compiled vs. eager with an fp64 reference and dumps a repro if accuracy fails. + +**Key utilities:** +- `torch/_dynamo/debug_utils.py` -- `same_two_models()`, `backend_accuracy_fails()`, `cast_to_fp64()` +- `torch._dynamo.config.repro_tolerance` (default: 1e-3) + +**Fix approach:** +1. Get the minimal failing graph from the minifier +2. Compare eager vs. compiled output at fp64 precision +3. Binary search through ops to find the diverging operation +4. Check for known numerical issues (reduction order, fused kernels, dtype promotions) + +### Internal Dynamo Errors + +`InternalTorchDynamoError` indicates a bug in Dynamo itself. + +**Diagnosis:** +```bash +TORCHDYNAMO_VERBOSE=1 python your_script.py +# or equivalently: +TORCH_LOGS="+dynamo" python your_script.py +``` + +**Key files:** +- `torch/_dynamo/symbolic_convert.py` -- bytecode interpreter +- `torch/_dynamo/variables/` -- variable tracking system +- `torch/_dynamo/guards.py` -- guard generation + +**Fix approach:** +1. Get the full stack trace with `TORCHDYNAMO_VERBOSE=1` +2. Identify which bytecode instruction or variable type caused the crash +3. Create a minimal repro (the error message often includes a minifier path) +4. Debug with `TORCHINDUCTOR_COMPILE_THREADS=1` and pdb if needed + +### Runtime Crashes + +Segfaults and CUDA illegal memory access errors during execution of compiled code. + +**Diagnosis (make crash deterministic):** +```bash +PYTORCH_NO_CUDA_MEMORY_CACHING=1 CUDA_LAUNCH_BLOCKING=1 python your_script.py +``` + +**For CUDA IMA, add NaN checks:** +```bash +TORCHINDUCTOR_NAN_ASSERTS=1 python your_script.py +``` + +**For Inductor-level sync debugging:** +```python +torch._inductor.config.triton.debug_sync_kernel = True # sync after every kernel +torch._inductor.config.triton.debug_sync_graph = True # sync before/after graph +``` + +**Fix approach:** +1. Make the crash deterministic with `PYTORCH_NO_CUDA_MEMORY_CACHING=1 CUDA_LAUNCH_BLOCKING=1` +2. Check if it's an input mismatch (shapes, devices, dtypes) +3. Inspect the generated kernel code with `TORCH_LOGS="output_code"` +4. Use `TORCHINDUCTOR_NAN_ASSERTS=1` to find the first kernel producing bad values +5. Check for dynamic shapes issues (historically a common source of IMA) + +### Triton Kernel Failures + +Triton assertion failures or index-out-of-bounds in generated kernels. + +**Diagnosis:** +```bash +TORCH_LOGS="output_code,schedule" python your_script.py +``` + +**Key files:** +- `torch/_inductor/codegen/triton.py` -- Triton codegen +- `torch/_inductor/scheduler.py` -- kernel fusion decisions + +**Fix approach:** +1. Get the generated Triton kernel from `output_code` logs +2. Check index computations for off-by-one or wrong stride calculations +3. Look at the IR (`TORCH_COMPILE_DEBUG=1`) to trace back to the FX op +4. Check if fusion decisions created invalid index combinations + +## Key Source Files + +| File | Purpose | +|---|---| +| `torch/_dynamo/exc.py` | Exception hierarchy and error formatting | +| `torch/_dynamo/debug_utils.py` | Minifier support, accuracy checking, input serialization | +| `torch/_dynamo/repro/after_dynamo.py` | Repro/minifier for Dynamo-stage failures | +| `torch/_dynamo/repro/after_aot.py` | Repro/minifier for post-AOTAutograd failures | +| `torch/_dynamo/repro/aoti.py` | Repro/minifier for AOTI failures | +| `torch/_dynamo/config.py` | Dynamo config (repro levels, recompile limits) | +| `torch/_dynamo/variables/torch.py` | Torch function handling, tracing state functions | +| `torch/_dynamo/variables/higher_order_ops.py` | HOP tracing (cond, map, etc.) | +| `torch/_dynamo/symbolic_convert.py` | Bytecode interpreter, InstructionTranslator | +| `torch/_dynamo/convert_frame.py` | Frame compilation, `fullgraph_capture` entry point | +| `torch/_dynamo/functional_export.py` | New export tracer (`_dynamo_graph_capture_for_export`) | +| `torch/_dynamo/eval_frame.py` | `torch._dynamo.export`, `optimize_assert` | +| `torch/_export/_trace.py` | Export pipeline (`_export`, `_strict_export`, `_non_strict_export`, `_export_to_aten_ir`) | +| `torch/_export/utils.py` | `_compiling_state_context()` | +| `torch/compiler/__init__.py` | `is_compiling()`, `is_exporting()`, runtime flags | +| `torch/_higher_order_ops/cond.py` | `torch.cond` implementation and proxy tracing | +| `torch/_higher_order_ops/utils.py` | `reenter_make_fx` for HOP branch tracing | +| `torch/_inductor/config.py` | Inductor config (debug flags, trace settings) | +| `torch/_inductor/debug.py` | DebugContext, graph visualization, IR logging | +| `torch/_logging/_registrations.py` | All registered log aliases and artifacts | + +## Using the Minifier + +The minifier reduces a failing graph to the smallest reproduction: + +```bash +# Step 1: Generate the minifier launcher +TORCHDYNAMO_REPRO_AFTER=aot TORCHDYNAMO_REPRO_LEVEL=2 python your_script.py + +# Step 2: Run the minifier +python minifier_launcher.py minify + +# Step 3: Run the minimized repro +python minifier_launcher.py run +``` + +For accuracy issues, use level 4: +```bash +TORCHDYNAMO_REPRO_AFTER=aot TORCHDYNAMO_REPRO_LEVEL=4 python your_script.py +``` diff --git a/.claude/skills/scrub-issue/SKILL.md b/.claude/skills/scrub-issue/SKILL.md new file mode 100644 index 0000000000000..aeb5e0d6839d4 --- /dev/null +++ b/.claude/skills/scrub-issue/SKILL.md @@ -0,0 +1,362 @@ +--- +name: scrub-issue +disable-model-invocation: true +description: Fetch, analyze, reproduce, and minimize GitHub issue reproductions. Use when asked to check if an issue reproduces, minimize a repro, analyze a bug report, or scrub/triage an issue for reproducibility. +--- + +# Minimize Issue Reproduction + +Fetch a GitHub issue, evaluate whether it has a reasonable repro, check if it +still reproduces, and systematically minimize the repro to the smallest possible +self-contained script. + +## Tools + +Assume the current environment is correct and run `python` directly. Only use +`conda run -n ` for version bisection (step 5a) where you need to +temporarily use a different environment. Use the Bash tool's `timeout` +parameter to enforce timeouts when running repro scripts. + +- `gh issue view --repo pytorch/pytorch` to fetch the issue body +- `gh issue view --comments --repo pytorch/pytorch` to fetch comments +- `python + * + * + * + * - Serve locally using: python3 -m http.server 8888 + * - Open http://localhost:8888 in your browser + * + * 2. WHAT TO TEST: + * - Ensure ALL tabs/views render correctly and switch properly + * - Verify BOTH interaction modes work: + * * Click mode: stack traces appear on click + * * Hover mode: stack traces appear on mouseover + * - Test zoom and brush controls for timeline navigation + * - Verify memory allocation blocks are rendered and interactive + * + * 3. TEST DATA REQUIREMENTS: + * - DO NOT just test with small dummy .pickle files + * - Use realistic, decent-sized .pickle files (10-100+ MB range) + * - Large files stress-test rendering performance and memory handling + * - Test with snapshots from real model training/inference runs + * + * 4. COMMON ISSUES TO WATCH FOR: + * - Performance degradation with large snapshots + * - Stack trace popups not appearing or positioning incorrectly + * - Tab switching not updating the visualization properly + * - Zoom/brush state not persisting across interactions + * + * ================================================================================ + */ + 'use strict'; import * as d3 from "https://cdn.jsdelivr.net/npm/d3@7/+esm"; @@ -5,6 +63,9 @@ import {axisLeft} from "https://cdn.jsdelivr.net/npm/d3-axis@3/+esm"; import {scaleLinear} from "https://cdn.jsdelivr.net/npm/d3-scale@4/+esm"; import {zoom, zoomIdentity} from "https://cdn.jsdelivr.net/npm/d3-zoom@3/+esm"; import {brushX} from "https://cdn.jsdelivr.net/npm/d3-brush@3/+esm"; +import {process_alloc_data, isPrivatePoolId, formatSize, formatAddr, + elideRepeats, frameFilter, format_user_metadata, + format_forward_frames, format_frames} from "./process_alloc_data.js"; // Global configuration for trace interaction mode // 'hover' = show trace on hover (default) @@ -48,8 +109,8 @@ function version_space() { }; } -function Segment(addr, size, stream, frames, version, user_metadata) { - return {addr, size, stream, version, frames, user_metadata}; +function Segment(addr, size, stream, frames, version, user_metadata, segment_pool_id) { + return {addr, size, stream, version, frames, user_metadata, segment_pool_id}; } function Block(addr, size, requested_size, frames, free_requested, version, user_metadata) { @@ -115,25 +176,6 @@ function EventSelector(outer, events, stack_info, memory_view) { return es; } -function formatSize(num, showBytes = true) { - const orig = num; - // https://stackoverflow.com/questions/1094841/get-human-readable-version-of-file-size - const units = ['', 'Ki', 'Mi', 'Gi', 'Ti', 'Pi', 'Ei', 'Zi']; - for (const unit of units) { - if (Math.abs(num) < 1024.0) { - if (showBytes) { - return `${num.toFixed(1)}${unit}B (${orig} bytes)`; - } - return `${num.toFixed(1)}${unit}B`; - } - num /= 1024.0; - } - return `${num.toFixed(1)}YiB`; -} -function formatAddr(event) { - const prefix = event.action.startsWith('segment') ? 's\'' : 'b\''; - return `${prefix}${event.addr.toString(16)}_${event.version}`; -} function formatEvent(event) { const stream = event.stream === null ? '' : `\n (stream ${event.stream})`; @@ -246,6 +288,7 @@ function MemoryView(outer, stack_info, snapshot, device) { seg.frames || [], seg.version, seg.user_metadata, + seg.segment_pool_id, ), ); for (const b of seg.blocks) { @@ -490,13 +533,17 @@ function MemoryView(outer, stack_info, snapshot, device) { const user_metadata_str = format_user_metadata(t.user_metadata); const frames_str = format_frames(t.frames); const forward_frames_str = format_forward_frames(t.forward_frames); + let pool_str = ''; + if (isPrivatePoolId(t.segment_pool_id)) { + pool_str = `, pool_id (${t.segment_pool_id[0]}, ${t.segment_pool_id[1]})`; + } return ( `s${t.addr.toString(16)}_${t.version}: segment ${formatSize( t.size, )} allocated, ` + `${formatSize(free)} free${internal} (stream ${ t.stream - })\n` + + }${pool_str})\n` + (user_metadata_str ? user_metadata_str + '\n' : '') + frames_str + forward_frames_str @@ -564,11 +611,15 @@ function MemoryView(outer, stack_info, snapshot, device) { const user_metadata_str = format_user_metadata(t.user_metadata); const frames_str = format_frames(t.frames); const forward_frames_str = format_forward_frames(t.forward_frames); + let pool_str = ''; + if (isPrivatePoolId(t.segment?.segment_pool_id)) { + pool_str = `, pool_id (${t.segment.segment_pool_id[0]}, ${t.segment.segment_pool_id[1]})`; + } return ( `b${t.addr.toString(16)}_${t.version} ` + `${formatSize(t.requested_size)} allocation${requested} (stream ${ t.segment.stream - })\n` + + }${pool_str})\n` + (user_metadata_str ? user_metadata_str + '\n' : '') + frames_str + forward_frames_str @@ -778,6 +829,7 @@ function annotate_snapshot(snapshot) { } } b.version = snapshot.block_version(b.addr, false); + b.segment_pool_id = seg.segment_pool_id; // Note [BigInt and Number Safe Arithmetic] // Device pointer addresses may be represented as either Number or BigInt. // Use explicit conversions to perform arithmetic safely and avoid mixing @@ -794,364 +846,6 @@ function annotate_snapshot(snapshot) { } } -function elideRepeats(frames) { - const result = []; - const length = frames.length; - for (let i = 0; i < length; ) { - let j = i + 1; - const f = frames[i]; - while (j < length && f === frames[j]) { - j++; - } - switch (j - i) { - case 1: - result.push(f); - break; - case 2: - result.push(f, f); - break; - default: - result.push(f, ``); - break; - } - i = j; - } - return result; -} -function frameFilter({name, filename}) { - const omitFunctions = [ - 'unwind::unwind', - 'CapturedTraceback::gather', - 'gather_with_cpp', - '_start', - '__libc_start_main', - 'PyEval_', - 'PyObject_', - 'PyFunction_', - ]; - - const omitFilenames = [ - 'core/boxing', - '/Register', - '/Redispatch', - 'pythonrun.c', - 'Modules/main.c', - 'Objects/call.c', - 'Objects/methodobject.c', - 'pycore_ceval.h', - 'ceval.c', - 'cpython/abstract.h', - ]; - - for (const of of omitFunctions) { - if (name.includes(of)) { - return false; - } - } - - for (const of of omitFilenames) { - if (filename.includes(of)) { - return false; - } - } - - return true; -} - -function format_user_metadata(user_metadata) { - if (!user_metadata) { - return ''; - } - // Handle string metadata - if (typeof user_metadata === 'string') { - return `User Metadata:\n ${user_metadata}`; - } - // Handle object metadata - if (typeof user_metadata === 'object' && Object.keys(user_metadata).length === 0) { - return ''; - } - const metadata_lines = Object.entries(user_metadata) - .map(([key, value]) => ` ${key}: ${value}`); - return 'User Metadata:\n' + metadata_lines.join('\n'); -} - -function format_forward_frames(forward_frames) { - if (!forward_frames || forward_frames.length === 0) { - return ''; - } - // forward_frames is a list of strings (each string is a frame line from the forward pass) - // Each frame string already includes newlines, so we just join them directly - let frames_str = forward_frames.join(''); - // Ensure we don't have a trailing newline that could cause display issues - frames_str = frames_str.trimEnd(); - return `\n\n=== Forward Pass Stack Trace (where this tensor was created) ===\n${frames_str}`; -} - -function format_frames(frames) { - if (frames.length === 0) { - return ( - `This block has no frames. Potential causes:\n` + - `1) This block was allocated before _record_memory_history was enabled.\n` + - `2) The context or stacks passed to _record_memory_history does not include this block. Consider changing context to 'state', 'alloc', or 'all', or changing stacks to 'all'.\n` + - `3) This event occurred during backward, which has no python frames, and memory history did not include C++ frames. Use stacks='all' to record both C++ and python frames.` - ); - } - const frame_strings = frames - .filter(frameFilter) - .map(f => { - let frame_str = `${f.filename}:${f.line}:${f.name}`; - - // Add FX debug information if available - if (f.fx_node_op || f.fx_node_name || f.fx_node_target) { - const fx_parts = []; - if (f.fx_node_name) fx_parts.push(`node=${f.fx_node_name}`); - if (f.fx_node_op) fx_parts.push(`op=${f.fx_node_op}`); - if (f.fx_node_target) fx_parts.push(`target=${f.fx_node_target}`); - frame_str += `\n >> FX: ${fx_parts.join(', ')}`; - } - - if (f.fx_original_trace) { - frame_str += `\n >> Original Model Code:`; - const original_lines = f.fx_original_trace.trim().split('\n'); - // Show all lines of the original trace - for (const line of original_lines) { - frame_str += `\n ${line}`; - } - } - - return frame_str; - }); - return elideRepeats(frame_strings).join('\n'); -} - -function process_alloc_data(snapshot, device, plot_segments, max_entries) { - const elements = []; - const initially_allocated = []; - const actions = []; - const addr_to_alloc = {}; - - const alloc = plot_segments ? 'segment_alloc' : 'alloc'; - const [free, free_completed] = plot_segments - ? ['segment_free', 'segment_free'] - : ['free', 'free_completed']; - for (const e of snapshot.device_traces[device]) { - switch (e.action) { - case alloc: - elements.push(e); - addr_to_alloc[e.addr] = elements.length - 1; - actions.push(elements.length - 1); - break; - case free: - case free_completed: - if (e.addr in addr_to_alloc) { - actions.push(addr_to_alloc[e.addr]); - delete addr_to_alloc[e.addr]; - } else { - elements.push(e); - initially_allocated.push(elements.length - 1); - actions.push(elements.length - 1); - } - break; - default: - break; - } - } - for (const seg of snapshot.segments) { - if (seg.device !== device) { - continue; - } - if (plot_segments) { - if (!(seg.address in addr_to_alloc)) { - const element = { - action: 'alloc', - addr: seg.address, - size: seg.total_size, - frames: [], - stream: seg.stream, - version: seg.version, - }; - elements.push(element); - initially_allocated.push(elements.length - 1); - } - } else { - for (const b of seg.blocks) { - if (b.state === 'active_allocated' && !(b.addr in addr_to_alloc)) { - const element = { - action: 'alloc', - addr: b.addr, - size: b.requested_size, - frames: b.frames, - stream: seg.stream, - version: b.version, - }; - elements.push(element); - initially_allocated.push(elements.length - 1); - } - } - } - } - initially_allocated.reverse(); - // if there are no actions, the graph will be blank, - // but if there are existing allocations we do not want to hide them - // by having just one allocate action it will show a flat graph with all segments - if (actions.length === 0 && initially_allocated.length > 0) { - actions.push(initially_allocated.pop()); - } - - const current = []; - const current_data = []; - const data = []; - let max_size = 0; - - let total_mem = 0; - let total_summarized_mem = 0; - let timestep = 0; - - const max_at_time = []; - - const summarized_mem = { - elem: 'summarized', - timesteps: [], - offsets: [total_mem], - size: [], - color: 0, - }; - const summarized_elems = {}; - - function advance(n) { - summarized_mem.timesteps.push(timestep); - summarized_mem.offsets.push(total_mem); - summarized_mem.size.push(total_summarized_mem); - timestep += n; - for (let i = 0; i < n; i++) { - max_at_time.push(total_mem + total_summarized_mem); - } - } - - const sizes = elements - .map((x, i) => [x.size, i]) - .sort(([x, _xi], [y, _yi]) => y - x); - - const draw_elem = {}; - for (const [_s, e] of sizes.slice(0, max_entries)) { - draw_elem[e] = true; - } - - function add_allocation(elem) { - const element_obj = elements[elem]; - const size = element_obj.size; - current.push(elem); - let color = elem; - if (snapshot.categories.length > 0) { - color = snapshot.categories.indexOf(element_obj.category || 'unknown'); - } - const e = { - elem, - timesteps: [timestep], - offsets: [total_mem], - size, - color, - }; - current_data.push(e); - data.push(e); - total_mem += size; - element_obj.max_allocated_mem = total_mem + total_summarized_mem; - } - - for (const elem of initially_allocated) { - if (elem in draw_elem) { - add_allocation(elem); - } else { - total_summarized_mem += elements[elem].size; - summarized_elems[elem] = true; - } - } - - for (const elem of actions) { - const size = elements[elem].size; - if (!(elem in draw_elem)) { - if (elem in summarized_elems) { - advance(1); - total_summarized_mem -= size; - summarized_elems[elem] = null; - } else { - total_summarized_mem += size; - summarized_elems[elem] = true; - advance(1); - } - continue; - } - const idx = current.findLastIndex(x => x === elem); - // first time we see an action we add it - // second time we remove it - if (idx === -1) { - add_allocation(elem); - advance(1); - } else { - advance(1); - const removed = current_data[idx]; - removed.timesteps.push(timestep); - removed.offsets.push(removed.offsets.at(-1)); - current.splice(idx, 1); - current_data.splice(idx, 1); - - if (idx < current.length) { - for (let j = idx; j < current.length; j++) { - const e = current_data[j]; - e.timesteps.push(timestep); - e.offsets.push(e.offsets.at(-1)); - e.timesteps.push(timestep + 3); - e.offsets.push(e.offsets.at(-1) - size); - } - advance(3); - } - total_mem -= size; - } - max_size = Math.max(total_mem + total_summarized_mem, max_size); - } - - for (const elem of current_data) { - elem.timesteps.push(timestep); - elem.offsets.push(elem.offsets.at(-1)); - } - data.push(summarized_mem); - - return { - max_size, - allocations_over_time: data, - max_at_time, - summarized_mem, - elements_length: elements.length, - context_for_id: id => { - const elem = elements[id]; - let text = `Addr: ${formatAddr(elem)}`; - text = `${text}, Size: ${formatSize(elem.size)} allocation`; - text = `${text}, Total memory used after allocation: ${formatSize( - elem.max_allocated_mem, - )}`; - const context = elem?.compile_context ?? 'None'; - text = `${text}, Compile context: ${context}`; - if (elem.stream !== null) { - text = `${text}, stream ${elem.stream}`; - } - if (elem.timestamp !== null) { - var d = new Date(elem.time_us / 1000); - text = `${text}, timestamp ${d}`; - } - if (!elem.action.includes('alloc')) { - text = `${text}\nalloc not recorded, stack trace for free:`; - } - const user_metadata_str = format_user_metadata(elem.user_metadata); - if (user_metadata_str) { - text = `${text}\n${user_metadata_str}`; - } - text = `${text}\n${format_frames(elem.frames)}`; - text = `${text}${format_forward_frames(elem.forward_frames)}`; - return text; - }, - }; -} - function MemoryPlot( svg, data, @@ -1213,7 +907,8 @@ function MemoryPlot( .enter() .append('polygon') .attr('points', format_points) - .attr('fill', d => colors[d.color % colors.length]); + .attr('fill', d => colors[d.color % colors.length]) + .attr('opacity', d => d.opacity ?? 1); const axis = plot_coordinate_space.append('g').call(yaxis); @@ -1272,11 +967,34 @@ function MemoryPlot( function ContextViewer(text, data) { let current_selected = null; + function restore_search_highlight(d) { + if (!d) return; + const addr = d.attr('data-search-match') === 'true'; + const frame = d.attr('data-frame-match') === 'true'; + if (addr && frame) { + d.attr('stroke', '#ff00ff') + .attr('stroke-width', 3) + .attr('stroke-dasharray', '6,3') + .attr('vector-effect', 'non-scaling-stroke'); + } else if (addr) { + d.attr('stroke', 'red') + .attr('stroke-width', 2) + .attr('stroke-dasharray', null) + .attr('vector-effect', 'non-scaling-stroke'); + } else if (frame) { + d.attr('stroke', '#2196F3') + .attr('stroke-width', 2) + .attr('stroke-dasharray', null) + .attr('vector-effect', 'non-scaling-stroke'); + } + } + return { default_selected: null, set_selected: d => { if (current_selected !== null) { - current_selected.attr('stroke', null).attr('stroke-width', null); + current_selected.attr('stroke', null).attr('stroke-width', null).attr('stroke-dasharray', null); + restore_search_highlight(current_selected); } if (d === null) { text.text(''); @@ -1287,6 +1005,10 @@ function ContextViewer(text, data) { 'Small tensors that were not plotted to cutdown on render time.\n' + 'Use detail slider to see smaller allocations.', ); + } else if (typeof dd.elem === 'string' && dd.elem.startsWith('pool:')) { + const pool_key = dd.elem.slice(5); + const capacity = Array.isArray(dd.size) ? dd.size.at(-1) : dd.size; + text.text(`Private Pool (${pool_key}): capacity ${formatSize(capacity)}`); } else { text.text(`${dd.elem} ${data.context_for_id(dd.elem)}`); } @@ -1387,13 +1109,20 @@ function create_trace_view( device, plot_segments = false, max_entries = 15000, + include_private_inactive = false, ) { const left_pad = 70; - const data = process_alloc_data(snapshot, device, plot_segments, max_entries); + const data = process_alloc_data(snapshot, device, plot_segments, max_entries, include_private_inactive); dst.selectAll('svg').remove(); dst.selectAll('div').remove(); max_entries = Math.min(max_entries, data.elements_length); + if (include_private_inactive) { + dst.append('div') + .attr('style', 'padding: 4px 8px; background: #fff3cd; border: 1px solid #ffc107; font-size: 13px; margin-bottom: 4px;') + .text('Note: Private pool memory (the gray bar) is shown as allocated until the pool\'s segment is freed. ' + + 'This view requires that MemPools are not deleted before torch.cuda.memory._snapshot() is called.'); + } const d = dst.append('div'); d.append('input') .attr('type', 'range') @@ -1401,12 +1130,28 @@ function create_trace_view( .attr('max', data.elements_length) .attr('value', max_entries) .on('change', function () { - create_trace_view(dst, snapshot, device, plot_segments, this.value); + create_trace_view(dst, snapshot, device, plot_segments, this.value, include_private_inactive); }); d.append('label').text( `Detail: ${max_entries} of ${data.elements_length} entries`, ); + d.append('span').text(' | '); + const search_input = d.append('input') + .attr('type', 'text') + .attr('placeholder', 'Search address (hex)...') + .attr('style', 'width: 180px; margin-left: 4px; font-family: monospace;'); + const search_label = d.append('label') + .attr('style', 'margin-left: 4px;'); + + d.append('span').text(' | '); + const frame_input = d.append('input') + .attr('type', 'text') + .attr('placeholder', 'Search stack frame...') + .attr('style', 'width: 200px; margin-left: 4px; font-family: monospace;'); + const frame_label = d.append('label') + .attr('style', 'margin-left: 4px;'); + const grid_container = dst .append('div') .attr( @@ -1443,6 +1188,61 @@ function create_trace_view( ); const delegate = ContextViewer(context_div.append('pre').text('none'), data); plot.set_delegate(delegate); + + function apply_search_highlights() { + const addr_query = search_input.node().value.toLowerCase().trim(); + const frame_query = frame_input.node().value.toLowerCase().trim(); + const polygons = plot_svg.selectAll('polygon'); + let addr_matches = 0; + let frame_matches = 0; + polygons.each(function () { + const dd = d3.select(this).datum(); + if (!dd || typeof dd.elem !== 'number') { + d3.select(this) + .attr('data-search-match', null) + .attr('data-frame-match', null); + return; + } + const ctx = data.context_for_id(dd.elem); + const ctx_lower = ctx.toLowerCase(); + const addr_hit = addr_query && ctx_lower.includes(addr_query); + const frame_hit = frame_query && ctx_lower.includes(frame_query); + d3.select(this) + .attr('data-search-match', addr_hit ? 'true' : null) + .attr('data-frame-match', frame_hit ? 'true' : null); + if (addr_hit && frame_hit) { + d3.select(this) + .attr('stroke', '#ff00ff') + .attr('stroke-width', 3) + .attr('stroke-dasharray', '6,3') + .attr('vector-effect', 'non-scaling-stroke'); + } else if (addr_hit) { + d3.select(this) + .attr('stroke', 'red') + .attr('stroke-width', 2) + .attr('stroke-dasharray', null) + .attr('vector-effect', 'non-scaling-stroke'); + } else if (frame_hit) { + d3.select(this) + .attr('stroke', '#2196F3') + .attr('stroke-width', 2) + .attr('stroke-dasharray', null) + .attr('vector-effect', 'non-scaling-stroke'); + } else { + d3.select(this) + .attr('stroke', null) + .attr('stroke-width', null) + .attr('stroke-dasharray', null); + } + if (addr_hit) addr_matches++; + if (frame_hit) frame_matches++; + }); + search_label.text(addr_query ? `${addr_matches} match${addr_matches !== 1 ? 'es' : ''}` : ''); + frame_label.text(frame_query ? `${frame_matches} match${frame_matches !== 1 ? 'es' : ''}` : ''); + } + + search_input.on('input', apply_search_highlights); + frame_input.on('input', apply_search_highlights); } function create_settings_view(dst, snapshot, device) { @@ -1764,6 +1564,8 @@ function decode_base64(input) { const kinds = { 'Active Memory Timeline': create_trace_view, + 'Allocated Memory (incl. Private Pools)': (dst, snapshot, device) => + create_trace_view(dst, snapshot, device, false, 15000, true), 'Allocator State History': create_segment_view, 'Active Cached Segment Timeline': (dst, snapshot, device) => create_trace_view(dst, snapshot, device, true), diff --git a/torch/utils/viz/_cycles.py b/torch/utils/viz/_cycles.py index df4bf34db2114..2f68cf75dee4d 100644 --- a/torch/utils/viz/_cycles.py +++ b/torch/utils/viz/_cycles.py @@ -467,7 +467,6 @@ def to_html(nodes): if n.context is None: continue s = _listener_template.format(id=str(i + 1), stack=escape(f'{n.label}:\n{n.context}')) - # pyrefly: ignore [bad-argument-type] listeners.append(s) dot = to_dot(nodes) return _template.replace('$DOT', repr(dot)).replace('$LISTENERS', '\n'.join(listeners)) diff --git a/torch/utils/viz/process_alloc_data.js b/torch/utils/viz/process_alloc_data.js new file mode 100644 index 0000000000000..6e5feda1046ea --- /dev/null +++ b/torch/utils/viz/process_alloc_data.js @@ -0,0 +1,902 @@ +// Pure data-processing functions for PyTorch memory visualization. +// Extracted from MemoryViz.js so they can be tested independently (no d3/DOM deps). +// +// This file is the single source of truth for these functions: +// - MemoryViz.js imports them via ESM: import {...} from "./process_alloc_data.js" +// - Node.js tests load this file by stripping the export line and eval-ing +// +// TRACE EVENT ACTIONS (from c10/core/CachingDeviceAllocator.h TraceEntry::Action): +// +// "alloc" - Sub-allocation returned to user from the caching allocator. +// Recorded in alloc_found_block() (CUDACachingAllocator.cpp:1834). +// +// "free_requested" - User code called free (tensor out of scope). The block may not +// be immediately returned to the free pool if it's in use on +// another stream via record_stream. +// Recorded in free() (CUDACachingAllocator.cpp:2123). +// +// "free_completed" - Block actually returned to the allocator's free pool. For simple +// cases this fires immediately after free_requested. For cross-stream +// blocks, deferred until CUDA events confirm all streams are done. +// Recorded in free_block() (CUDACachingAllocator.cpp:3148). +// +// "segment_alloc" - New segment allocated from OS via cudaMalloc (or cuMemCreate for +// expandable segments). +// Recorded in alloc_from_expandable_segment() (CUDACachingAllocator.cpp:3548). +// +// "segment_free" - Segment returned to OS via cudaFree. Happens during empty_cache() +// or defragmentation. Only for non-expandable segments. +// Recorded in release_block() (CUDACachingAllocator.cpp:3686). +// +// "segment_map" - Physical pages mapped into an expandable segment via cuMemMap. +// The segment grows. Only with expandable segments enabled. +// Recorded in alloc_from_expandable_segment() (CUDACachingAllocator.cpp:3092). +// +// "segment_unmap" - Physical pages unmapped from an expandable segment via cuMemUnmap. +// Virtual address range retained, physical memory returned to OS. +// Only with expandable segments. Causes "pool_id unknown" for any +// trace events whose addresses fall in the unmapped range, since +// the segment no longer exists at snapshot time. +// Recorded in unmap_block() (CUDACachingAllocator.cpp:3790). +// +// "snapshot" - A call to torch.cuda.memory._snapshot(). Timestamp marker to +// correlate trace events with snapshot state. addr=0. +// Recorded in snapshot() (CUDACachingAllocator.cpp:2689). +// +// "oom" - Allocator failed to satisfy an allocation after all retries. +// addr=device_free (bytes free on GPU), size=requested allocation. +// Recorded in malloc() (CUDACachingAllocator.cpp:1629). +// +// HOW SEGMENT EVENTS ARE USED IN VISUALIZATION: +// +// The snapshot pickle contains two separate data sources: +// 1. device_traces - Ring buffer of TraceEntry actions (alloc, free, segment_map, etc.) +// 2. segments - Point-in-time dump of all segments/blocks at _snapshot() time +// +// Block-level views ("Active Memory Timeline", "Allocated Memory (incl. Private Pools)"): +// - process_alloc_data matches "alloc" and "free_completed" from device_traces. +// - segment_alloc/segment_free/segment_map/segment_unmap are ignored (skipped in switch). +// - The segments snapshot is used only to resolve pool_id via find_pool_id(). +// +// Segment-level view ("Active Cached Segment Timeline"): +// - process_alloc_data is called with plot_segments=true. +// - Matches "segment_alloc" and "segment_free" instead of alloc/free. +// - segment_map/segment_unmap are NOT matched (they don't appear in the switch). +// - Segments from the snapshot that weren't seen in the trace are added as +// initially_allocated (Phase 2). +// +// Allocator State History ("Allocator State History"): +// - EventSelector lists ALL trace events including segment_map/segment_unmap. +// - MemoryView renders the segment/block layout from the segments snapshot. +// - Clicking an event in the list redraws the layout at that point in time. +// +// Ring buffer overflow: +// - All trace event types share the same ring buffer. When it overflows, older +// events are overwritten. The allocator_settings.trace_alloc_overflowed flag +// indicates this happened, and trace_alloc_max_entries gives the buffer size. +// - Segment snapshot data (segments array) is NOT affected by ring buffer overflow. +// - The segment snapshot is always complete regardless of overflow. + +/** + * Returns true if pool_id represents a private (user-created) memory pool, + * as opposed to the default pool [0, 0]. + * + * @param {number[]|null} pool_id - Two-element array [owner_id, pool_id] from + * the CUDA caching allocator. The default pool is [0, 0]; any other non-null + * value is a private pool (e.g. FSDP's MemPool). + * @returns {boolean} + */ +function isPrivatePoolId(pool_id) { + return pool_id && !(pool_id[0] === 0 && pool_id[1] === 0); +} + +/** + * Formats a byte count as a human-readable string (e.g. "1.5GiB (1610612736 bytes)"). + * + * @param {number} num - Size in bytes. + * @param {boolean} [showBytes=true] - Whether to include the raw byte count in parentheses. + * @returns {string} + */ +function formatSize(num, showBytes = true) { + const orig = num; + // https://stackoverflow.com/questions/1094841/get-human-readable-version-of-file-size + const units = ['', 'Ki', 'Mi', 'Gi', 'Ti', 'Pi', 'Ei', 'Zi']; + for (const unit of units) { + if (Math.abs(num) < 1024.0) { + if (showBytes) { + return `${num.toFixed(1)}${unit}B (${orig} bytes)`; + } + return `${num.toFixed(1)}${unit}B`; + } + num /= 1024.0; + } + return `${num.toFixed(1)}YiB`; +} + +/** + * Formats a trace event's address as a display string like "b'7f4c00000_3". + * Segment-level events get an "s'" prefix, block-level events get "b'". + * + * @param {{action: string, addr: number|BigInt, version: number}} event + * @returns {string} + */ +function formatAddr(event) { + const prefix = event.action.startsWith('segment') ? 's\'' : 'b\''; + return `${prefix}${event.addr.toString(16)}_${event.version}`; +} + +/** + * Collapses consecutive duplicate strings in an array. If a string appears + * N > 2 times in a row, it's replaced with [str, ""]. + * Used to compress repetitive stack frames in the display. + * + * @param {string[]} frames + * @returns {string[]} + */ +function elideRepeats(frames) { + const result = []; + const length = frames.length; + for (let i = 0; i < length; ) { + let j = i + 1; + const f = frames[i]; + while (j < length && f === frames[j]) { + j++; + } + switch (j - i) { + case 1: + result.push(f); + break; + case 2: + result.push(f, f); + break; + default: + result.push(f, ``); + break; + } + i = j; + } + return result; +} + +/** + * Returns false for stack frames that are internal runtime noise + * (e.g. Python interpreter internals, C++ dispatch machinery). + * Used as a filter predicate on frame arrays. + * + * @param {{name: string, filename: string}} frame + * @returns {boolean} + */ +function frameFilter({name, filename}) { + const omitFunctions = [ + 'unwind::unwind', + 'CapturedTraceback::gather', + 'gather_with_cpp', + '_start', + '__libc_start_main', + 'PyEval_', + 'PyObject_', + 'PyFunction_', + ]; + + const omitFilenames = [ + 'core/boxing', + '/Register', + '/Redispatch', + 'pythonrun.c', + 'Modules/main.c', + 'Objects/call.c', + 'Objects/methodobject.c', + 'pycore_ceval.h', + 'ceval.c', + 'cpython/abstract.h', + ]; + + for (const of of omitFunctions) { + if (name.includes(of)) { + return false; + } + } + + for (const of of omitFilenames) { + if (filename.includes(of)) { + return false; + } + } + + return true; +} + +/** + * Formats user-attached metadata (from torch.cuda.memory._record_memory_history) + * as a display string. Returns '' if no metadata is present. + * + * @param {string|Object|null|undefined} user_metadata + * @returns {string} + */ +function format_user_metadata(user_metadata) { + if (!user_metadata) { + return ''; + } + if (typeof user_metadata === 'string') { + return `User Metadata:\n ${user_metadata}`; + } + if (typeof user_metadata === 'object' && Object.keys(user_metadata).length === 0) { + return ''; + } + const metadata_lines = Object.entries(user_metadata) + .map(([key, value]) => ` ${key}: ${value}`); + return 'User Metadata:\n' + metadata_lines.join('\n'); +} + +/** + * Formats the forward-pass stack trace (captured via torch.autograd) as a + * display string showing where a tensor was originally created. + * + * @param {string[]|null|undefined} forward_frames + * @returns {string} + */ +function format_forward_frames(forward_frames) { + if (!forward_frames || forward_frames.length === 0) { + return ''; + } + let frames_str = forward_frames.join(''); + frames_str = frames_str.trimEnd(); + return `\n\n=== Forward Pass Stack Trace (where this tensor was created) ===\n${frames_str}`; +} + +/** + * Formats an array of stack frames into a human-readable string. + * Filters out runtime noise via frameFilter, annotates FX graph debug info + * when available, and collapses consecutive duplicate frames. + * + * @param {{filename: string, line: number, name: string, + * fx_node_op?: string, fx_node_name?: string, + * fx_node_target?: string, fx_original_trace?: string}[]} frames + * @returns {string} + */ +function format_frames(frames) { + if (frames.length === 0) { + return ( + `This block has no frames. Potential causes:\n` + + `1) This block was allocated before _record_memory_history was enabled.\n` + + `2) The context or stacks passed to _record_memory_history does not include this block. Consider changing context to 'state', 'alloc', or 'all', or changing stacks to 'all'.\n` + + `3) This event occurred during backward, which has no python frames, and memory history did not include C++ frames. Use stacks='all' to record both C++ and python frames.\n` + + `4) This block was reconstructed from the allocator's segment snapshot (not from a trace event). The snapshot records which blocks exist at the moment _snapshot() is called, but does not carry stack frames. This typically happens for blocks that were allocated before tracing started and never freed, or for inactive blocks in private memory pools.\n` + + `5) The original alloc event was evicted from the trace ring buffer (older entries are overwritten when the buffer is full). Increase the max_entries argument to _record_memory_history to retain more events.` + ); + } + const frame_strings = frames + .filter(frameFilter) + .map(f => { + let frame_str = `${f.filename}:${f.line}:${f.name}`; + + if (f.fx_node_op || f.fx_node_name || f.fx_node_target) { + const fx_parts = []; + if (f.fx_node_name) fx_parts.push(`node=${f.fx_node_name}`); + if (f.fx_node_op) fx_parts.push(`op=${f.fx_node_op}`); + if (f.fx_node_target) fx_parts.push(`target=${f.fx_node_target}`); + frame_str += `\n >> FX: ${fx_parts.join(', ')}`; + } + + if (f.fx_original_trace) { + frame_str += `\n >> Original Model Code:`; + const original_lines = f.fx_original_trace.trim().split('\n'); + for (const line of original_lines) { + frame_str += `\n ${line}`; + } + } + + return frame_str; + }); + return elideRepeats(frame_strings).join('\n'); +} + +/** + * Transforms a memory snapshot into a stacked-area timeline suitable for + * rendering by MemoryPlot. This is the core data-processing function behind + * the "Active Memory Timeline" and "Allocated Memory (incl. Private Pools)" + * visualization tabs. + * + * HIGH-LEVEL ALGORITHM: + * + * 1. TRACE EVENT MATCHING: Scans device_traces to pair alloc events with their + * corresponding free_completed events (by address). Events whose matching + * alloc was lost (e.g. ring buffer wrap) become "initially_allocated" — + * blocks assumed to exist at the start of the trace. + * + * 2. SEGMENT SNAPSHOT: Supplements trace data with the current segment state. + * Blocks marked active_allocated (or inactive in private pools, when + * include_private_inactive=true) that weren't seen in the trace are also + * added as initially_allocated. + * + * 3. DETAIL LIMITING: Only the largest max_entries elements get individual + * rectangles in the plot. Smaller elements are aggregated into a single + * "summarized" band to keep rendering fast. + * + * 4. STACKED AREA CONSTRUCTION: Replays alloc/free events in order, building + * a stacked-area dataset where each element has timesteps, y-offsets, and + * a size. Elements are stacked bottom-to-top; frees remove from the stack + * and shift elements above downward. + * + * 5. PRIVATE POOL ENVELOPES (include_private_inactive=true): Each private pool + * (e.g. FSDP's MemPool) gets a single gray "envelope" rectangle whose + * height is the pool's high-water mark. Active blocks within the pool are + * rendered as colored stripes inside the envelope. The envelope only grows + * (never shrinks), representing reserved capacity. + * + * Initially-allocated private pool blocks are PRE-LOADED into pool state + * so that when their free event appears in the trace, they are correctly + * recognized as frees (not misinterpreted as new allocations). + * + * NOTE ON FREE EVENT MATCHING: The C++ allocator emits 'free_requested' and + * 'free_completed' for each deallocation. This function matches against 'free' + * (which no longer appears in modern traces — effectively dead code) and + * 'free_completed'. Only free_completed does the actual matching. Matching + * both 'free_requested' AND 'free_completed' would cause double-processing + * since they share the same address. + * + * @param {Object} snapshot - Memory snapshot from torch.cuda.memory._snapshot(). + * @param {Object[]} snapshot.segments - Current allocator segment state. + * @param {Object[][]} snapshot.device_traces - Per-device arrays of trace events. + * Each event has {action, addr, size, frames, stream, segment_pool_id?, ...}. + * @param {string[]} snapshot.categories - Category names for color-coding. + * @param {number} device - Device index into snapshot.device_traces. + * @param {boolean} plot_segments - If true, plot segment-level (cudaMalloc) + * events instead of sub-allocation events. + * @param {number} max_entries - Maximum number of elements to render individually. + * Elements beyond this limit are aggregated into the "summarized" band. + * @param {boolean} [include_private_inactive=false] - If true, include inactive + * blocks from private pools and render pool envelopes. Used by the + * "Allocated Memory (incl. Private Pools)" tab. + * + * @returns {{ + * max_size: number, + * allocations_over_time: Object[], + * max_at_time: number[], + * summarized_mem: Object, + * elements_length: number, + * context_for_id: function(number): string + * }} + * - max_size: peak total memory observed during the action replay (used for + * y-axis scaling). Note: this is only updated inside the action loop, so + * the initial state from initially_allocated may not be reflected here + * (use max_at_time for the true peak). + * - allocations_over_time: array of stacked-area data objects, each with + * {elem, timesteps[], offsets[], size, color}. + * - max_at_time: total memory at each timestep (for minimap rendering). + * - summarized_mem: the aggregated band for small elements. + * - elements_length: total number of unique allocation elements. + * - context_for_id: function that returns a human-readable description + * string for a given element index (address, size, stack trace, etc.). + */ +function process_alloc_data(snapshot, device, plot_segments, max_entries, include_private_inactive = false) { + const elements = []; + // Contains two types of blocks + // 1. free without alloc in trace + // 2. actively allocated in segments, but no matching alloc in trace + const initially_allocated = []; + const actions = []; + const addr_to_alloc = {}; + + const device_segments = snapshot.segments + .filter(s => s.device === device) + .sort((a, b) => { + if (a.address === b.address) return 0; + return a.address < b.address ? -1 : 1; + }); + + // Binary search to find which segment contains a given address. + function find_pool_id(addr) { + let left = 0; + let right = device_segments.length - 1; + while (left <= right) { + const mid = Math.floor((left + right) / 2); + const seg = device_segments[mid]; + const seg_end = seg.address + (typeof seg.address === "bigint" ? BigInt(seg.total_size) : seg.total_size); + if (addr < seg.address) { + right = mid - 1; + } else if (addr >= seg_end) { + left = mid + 1; + } else { + return seg.segment_pool_id; + } + } + return null; + } + + const alloc = plot_segments ? 'segment_alloc' : 'alloc'; + const [free, free_completed] = plot_segments + ? ['segment_free', 'segment_free'] + : ['free', 'free_completed']; + for (const e of snapshot.device_traces[device]) { + switch (e.action) { + case alloc: + elements.push(e); + addr_to_alloc[e.addr] = elements.length - 1; + actions.push(elements.length - 1); + break; + case free: + case free_completed: + if (e.addr in addr_to_alloc) { + // Matched: reuse the element from the alloc event + actions.push(addr_to_alloc[e.addr]); + delete addr_to_alloc[e.addr]; + } else { + // Unmatched free: alloc happened before recording (or was evicted + // from the ring buffer). Create a new element from the free event; + // its stack trace will show the free site, not the alloc site. + elements.push(e); + initially_allocated.push(elements.length - 1); + actions.push(elements.length - 1); + } + break; + default: + break; + } + } + + // --- Phase 2: Add elements from the snapshot --- + for (const seg of snapshot.segments) { + if (seg.device !== device) { + continue; + } + if (plot_segments) { + if (!(seg.address in addr_to_alloc)) { + const element = { + action: 'alloc', + addr: seg.address, + size: seg.total_size, + frames: [], + stream: seg.stream, + version: seg.version, + }; + elements.push(element); + initially_allocated.push(elements.length - 1); + } + } else { + for (const b of seg.blocks) { + const addr = b.addr ?? b.address; + if (b.state === 'active_allocated' && !(addr in addr_to_alloc)) { + const element = { + action: 'alloc', + addr, + size: b.requested_size, + frames: b.frames, + stream: seg.stream, + version: b.version, + segment_pool_id: seg.segment_pool_id, + ghost: true, + }; + elements.push(element); + initially_allocated.push(elements.length - 1); + } + } + } + } + + // Resolve pool IDs for trace elements by looking up which segment they fall in + for (const elem of elements) { + if (!elem.segment_pool_id) { + elem.segment_pool_id = find_pool_id(elem.addr); + } + } + + initially_allocated.reverse(); + // If there are no trace actions but there are existing allocations, + // show a flat graph with the initial state + if (actions.length === 0 && initially_allocated.length > 0) { + actions.push(initially_allocated.pop()); + } + + // --- Phase 3: Build the stacked-area timeline --- + const current = []; // stack of element indices (bottom to top) + const current_data = []; // parallel array of visualization data objects + const data = []; // all data objects (including completed ones) + let max_size = 0; + + let total_mem = 0; + let total_summarized_mem = 0; + let timestep = 0; + + const max_at_time = []; + + const summarized_mem = { + elem: 'summarized', + timesteps: [], + offsets: [total_mem], + size: [], + color: 0, + }; + const summarized_elems = {}; + + // Record the current memory state and advance time by n steps + function advance(n) { + summarized_mem.timesteps.push(timestep); + summarized_mem.offsets.push(total_mem); + summarized_mem.size.push(total_summarized_mem); + timestep += n; + for (let i = 0; i < n; i++) { + max_at_time.push(total_mem + total_summarized_mem); + } + } + + // Only render the largest max_entries elements individually; + // everything else goes into the summarized band + const sizes = elements + .map((x, i) => [x.size, i]) + .sort(([x, _xi], [y, _yi]) => y - x); + + const draw_elem = {}; + for (const [_s, e] of sizes.slice(0, max_entries)) { + draw_elem[e] = true; + } + + // Push an element onto the memory stack + function add_allocation(elem) { + const element_obj = elements[elem]; + const size = element_obj.size; + current.push(elem); + let color = elem; + if (snapshot.categories.length > 0) { + color = snapshot.categories.indexOf(element_obj.category || 'unknown'); + } + const e = { + elem, + timesteps: [timestep], + offsets: [total_mem], + size, + color, + }; + if (element_obj.ghost) e.ghost = true; + current_data.push(e); + data.push(e); + total_mem += size; + element_obj.max_allocated_mem = total_mem + total_summarized_mem; + } + + // --- Pool envelope tracking (only when include_private_inactive=true) --- + // Each private pool gets a gray envelope whose height = high-water mark. + // Active blocks are rendered as colored stripes within the envelope. + const pools = {}; + const pool_active_elems = {}; + + function get_pool_key(elem_idx) { + const pid = elements[elem_idx].segment_pool_id; + if (!isPrivatePoolId(pid)) return null; + const stream = elements[elem_idx].stream; + return `${pid[0]},${pid[1]},s${stream}`; + } + + function get_or_create_pool(pool_key) { + if (!(pool_key in pools)) { + pools[pool_key] = { + max: 0, active: 0, envelope_data: null, + block_stack: [], // [{elem, size, inner_offset, stripe_data}] + }; + } + return pools[pool_key]; + } + + function elem_color(elem_idx) { + if (snapshot.categories.length > 0) { + return snapshot.categories.indexOf(elements[elem_idx].category || 'unknown'); + } + return elem_idx; + } + + function shift_pool_stripes(pool, delta) { + for (const block of pool.block_stack) { + const s = block.stripe_data; + s.timesteps.push(timestep); + s.offsets.push(s.offsets.at(-1)); + s.timesteps.push(timestep + 3); + s.offsets.push(s.offsets.at(-1) + delta); + } + } + + // Animate shifting all elements above idx by delta (used when an element + // is inserted or removed from the middle of the stack) + function shift_elements_above(idx, delta) { + for (let j = idx; j < current.length; j++) { + const e = current_data[j]; + e.timesteps.push(timestep); + e.offsets.push(e.offsets.at(-1)); + e.timesteps.push(timestep + 3); + e.offsets.push(e.offsets.at(-1) + delta); + if (Array.isArray(e.size)) { + e.size.push(e.size.at(-1)); + e.size.push(e.size.at(-1)); + } + const pk = typeof current[j] === 'string' && current[j].startsWith('pool:') + ? current[j].slice(5) : null; + if (pk && pk in pools) { + shift_pool_stripes(pools[pk], delta); + } + } + } + + // Grow a pool envelope to accommodate new_active bytes. + // The envelope only grows (never shrinks) — it represents reserved capacity. + function grow_pool_envelope(pool, pool_key, new_active) { + if (new_active <= pool.max) return; + const delta = new_active - pool.max; + pool.max = new_active; + const env = pool.envelope_data; + env.timesteps.push(timestep); + env.offsets.push(env.offsets.at(-1)); + env.size.push(env.size.at(-1)); + env.timesteps.push(timestep + 3); + env.offsets.push(env.offsets.at(-1)); + env.size.push(pool.max); + const pidx = current.indexOf(`pool:${pool_key}`); + if (pidx >= 0) { + shift_elements_above(pidx + 1, delta); + } + total_mem += delta; + advance(3); + } + + // --- Process initially_allocated elements --- + // Private pool blocks are pre-loaded into pool state at timestep 0 (no + // animation) so that their envelope starts at the correct initial size + // and free events are correctly recognized as frees. + for (const elem of initially_allocated) { + if (include_private_inactive && get_pool_key(elem)) { + const pk = get_pool_key(elem); + const size = elements[elem].size; + const pool = get_or_create_pool(pk); + pool_active_elems[elem] = pk; + + if (pool.envelope_data === null) { + const env = { + elem: `pool:${pk}`, + timesteps: [0], + offsets: [total_mem], + size: [0], + color: 9, + }; + pool.envelope_data = env; + current.push(`pool:${pk}`); + current_data.push(env); + data.push(env); + } + + const inner_offset = pool.active; + pool.active += size; + + // Grow envelope directly without animation — these blocks pre-exist + if (pool.active > pool.max) { + const delta = pool.active - pool.max; + pool.max = pool.active; + const env = pool.envelope_data; + env.size[env.size.length - 1] = pool.max; + total_mem += delta; + const pidx = current.indexOf(`pool:${pk}`); + if (pidx >= 0) { + for (let j = pidx + 1; j < current.length; j++) { + const e = current_data[j]; + e.offsets[e.offsets.length - 1] += delta; + } + } + } + + const stripe = { + elem, + timesteps: [0], + offsets: [pool.envelope_data.offsets.at(-1) + inner_offset], + size, + color: elem_color(elem), + opacity: 0.5, + ghost: elements[elem].ghost || false, + }; + pool.block_stack.push({elem, size, inner_offset, stripe_data: stripe}); + data.push(stripe); + continue; + } + if (elem in draw_elem) { + add_allocation(elem); + } else { + total_summarized_mem += elements[elem].size; + summarized_elems[elem] = true; + } + } + + // Fix up pool stripe offsets — stripes are not in current_data so they + // don't get shifted when other pools grow during initially_allocated + // processing. Recompute from the envelope's final offset. + for (const pk in pools) { + const p = pools[pk]; + if (!p.envelope_data) continue; + const env_offset = p.envelope_data.offsets.at(-1); + for (const block of p.block_stack) { + const s = block.stripe_data; + for (let i = 0; i < s.offsets.length; i++) { + s.offsets[i] = env_offset + block.inner_offset; + } + } + } + + // --- Replay alloc/free actions to build the timeline --- + for (const elem of actions) { + const size = elements[elem].size; + const pool_key = include_private_inactive ? get_pool_key(elem) : null; + + if (pool_key) { + // --- Private pool element --- + if (!(elem in pool_active_elems)) { + // Pool alloc: add to pool, grow envelope if needed + pool_active_elems[elem] = pool_key; + const pool = get_or_create_pool(pool_key); + + if (pool.envelope_data === null) { + const env = { + elem: `pool:${pool_key}`, + timesteps: [timestep], + offsets: [total_mem], + size: [0], + color: 9, + }; + pool.envelope_data = env; + current.push(`pool:${pool_key}`); + current_data.push(env); + data.push(env); + } + + const inner_offset = pool.active; + pool.active += size; + + if (pool.active > pool.max) { + grow_pool_envelope(pool, pool_key, pool.active); + } + + const stripe = { + elem, + timesteps: [timestep], + offsets: [pool.envelope_data.offsets.at(-1) + inner_offset], + size, + color: elem_color(elem), + opacity: 0.5, + }; + pool.block_stack.push({elem, size, inner_offset, stripe_data: stripe}); + data.push(stripe); + advance(1); + elements[elem].max_allocated_mem = total_mem + total_summarized_mem; + } else { + // Pool free: end stripe, shift stripes above down within the pool. + // The envelope stays at its high-water mark (never shrinks). + const pool = pools[pool_key]; + const block_idx = pool.block_stack.findIndex(b => b.elem === elem); + if (block_idx >= 0) { + advance(1); + const block = pool.block_stack[block_idx]; + block.stripe_data.timesteps.push(timestep); + block.stripe_data.offsets.push(block.stripe_data.offsets.at(-1)); + + pool.block_stack.splice(block_idx, 1); + pool.active -= size; + + if (block_idx < pool.block_stack.length) { + for (let j = block_idx; j < pool.block_stack.length; j++) { + const b = pool.block_stack[j]; + b.inner_offset -= size; + const s = b.stripe_data; + s.timesteps.push(timestep); + s.offsets.push(s.offsets.at(-1)); + s.timesteps.push(timestep + 3); + s.offsets.push(pool.envelope_data.offsets.at(-1) + b.inner_offset); + } + advance(3); + } + } else { + pool.active -= size; + advance(1); + } + delete pool_active_elems[elem]; + } + max_size = Math.max(total_mem + total_summarized_mem, max_size); + continue; + } + + // --- Non-pool element --- + if (!(elem in draw_elem)) { + // Too small to render individually — goes into the summarized band + if (elem in summarized_elems) { + advance(1); + total_summarized_mem -= size; + summarized_elems[elem] = null; + } else { + total_summarized_mem += size; + summarized_elems[elem] = true; + advance(1); + } + continue; + } + const idx = current.findLastIndex(x => x === elem); + if (idx === -1) { + // First appearance → alloc + add_allocation(elem); + advance(1); + } else { + // Second appearance → free: remove from stack, shift elements above down + advance(1); + const removed = current_data[idx]; + removed.timesteps.push(timestep); + removed.offsets.push(removed.offsets.at(-1)); + current.splice(idx, 1); + current_data.splice(idx, 1); + + if (idx < current.length) { + shift_elements_above(idx, -size); + advance(3); + } + total_mem -= size; + } + max_size = Math.max(total_mem + total_summarized_mem, max_size); + } + + // --- Finalize: close all still-active elements --- + for (const elem of current_data) { + elem.timesteps.push(timestep); + elem.offsets.push(elem.offsets.at(-1)); + if (Array.isArray(elem.size)) { + elem.size.push(elem.size.at(-1)); + } + } + for (const pk in pools) { + for (const block of pools[pk].block_stack) { + const s = block.stripe_data; + s.timesteps.push(timestep); + s.offsets.push(s.offsets.at(-1)); + } + } + data.push(summarized_mem); + + return { + max_size, + allocations_over_time: data, + max_at_time, + summarized_mem, + elements_length: elements.length, + context_for_id: id => { + const elem = elements[id]; + let text = `Addr: ${formatAddr(elem)}`; + text = `${text}, Size: ${formatSize(elem.size)} allocation`; + text = `${text}, Total memory used after allocation: ${formatSize( + elem.max_allocated_mem, + )}`; + const context = elem?.compile_context ?? 'None'; + text = `${text}, Compile context: ${context}`; + if (elem.stream !== null) { + text = `${text}, stream ${elem.stream}`; + } + if (elem.segment_pool_id) { + text = `${text}, pool_id (${elem.segment_pool_id[0]}, ${elem.segment_pool_id[1]})`; + } else { + text = `${text}, pool_id unknown`; + } + if (elem.timestamp !== null) { + var d = new Date(elem.time_us / 1000); + text = `${text}, timestamp ${d}`; + } + if (!elem.action.includes('alloc')) { + text = `${text}\nalloc not recorded, stack trace for free:`; + } + if (elem.ghost) { + text = `${text}\n[Ghost block] This block exists in the segment snapshot but has no alloc trace events. ` + + `It was allocated before _record_memory_history() was called, or its alloc event was evicted ` + + `from the trace ring buffer. The block is still active (not freed) at snapshot time.`; + } + const user_metadata_str = format_user_metadata(elem.user_metadata); + if (user_metadata_str) { + text = `${text}\n${user_metadata_str}`; + } + text = `${text}\n${format_frames(elem.frames)}`; + text = `${text}${format_forward_frames(elem.forward_frames)}`; + return text; + }, + }; +} + +export { process_alloc_data, isPrivatePoolId, formatSize, formatAddr, + elideRepeats, frameFilter, format_user_metadata, + format_forward_frames, format_frames }; diff --git a/torch/xpu/__init__.py b/torch/xpu/__init__.py index bc3cf07c88e91..0c77d45201c09 100644 --- a/torch/xpu/__init__.py +++ b/torch/xpu/__init__.py @@ -7,16 +7,22 @@ :func:`is_available()` to determine if your system supports XPU. """ +from __future__ import annotations + import threading import traceback -from collections.abc import Callable from functools import lru_cache -from typing import Any, NewType, Optional +from typing import Any, NewType, TYPE_CHECKING import torch import torch._C from torch._utils import _dummy_type, _LazySeedTracker -from torch.types import Device + + +if TYPE_CHECKING: + from collections.abc import Callable + + from torch.types import Device from ._utils import _get_device_index from .graphs import ( @@ -266,7 +272,7 @@ def get_device_capability(device: Device = None) -> dict[str, Any]: def get_device_properties( device: Device = None, -) -> _XpuDeviceProperties: # pyrefly: ignore # not-a-type +) -> _XpuDeviceProperties: r"""Get the properties of a device. Returns _XpuDeviceProperties containing the following device properties: - ``name`` (str): device name. @@ -279,6 +285,8 @@ def get_device_properties( - ``gpu_eu_count`` (int): number of EUs (Execution Unit). - ``max_work_group_size``: (int): maximum number of work-items permitted in a work-group. - ``max_num_sub_groups`` (int): maximum number of sub-groups supported in a work-group. + - ``memory_clock_rate`` (int) maximum clock rate of device's global memory in MHz. + - ``memory_bus_width`` (int) maximum bus width between device and memory in bits. - ``sub_group_sizes``: (list[int]): a list of supported sub-group sizes. - ``local_mem_size`` (int): device local memory capacity that can be allocated per work-group in bytes. - ``has_fp16`` (bool): whether float16 dtype is supported. @@ -353,9 +361,9 @@ class StreamContext: .. note:: Streams are per-device. """ - cur_stream: Optional["torch.xpu.Stream"] + cur_stream: torch.xpu.Stream | None - def __init__(self, stream: Optional["torch.xpu.Stream"]) -> None: + def __init__(self, stream: torch.xpu.Stream | None) -> None: self.stream = stream self.idx = _get_device_index(None, True) if self.idx is None: @@ -384,7 +392,7 @@ def __exit__(self, type: Any, value: Any, traceback: Any): torch.xpu.set_stream(self.src_prev_stream) -def stream(stream: Optional["torch.xpu.Stream"]) -> StreamContext: +def stream(stream: torch.xpu.Stream | None) -> StreamContext: r"""Wrap around the Context-manager StreamContext that selects a given stream. Arguments: diff --git a/torch/xpu/graphs.py b/torch/xpu/graphs.py index 415408f62407b..51780050f5937 100644 --- a/torch/xpu/graphs.py +++ b/torch/xpu/graphs.py @@ -2,7 +2,7 @@ import typing from collections.abc import Callable -from typing import Optional, overload, TYPE_CHECKING, TypeAlias, Union +from typing import overload, TYPE_CHECKING, TypeAlias from typing_extensions import ParamSpec, Self, TypeVar import torch @@ -34,7 +34,6 @@ "_xpu_isCurrentStreamCapturing" ) -# pyrefly: ignore [missing-module-attribute] from torch._C import _xpu_graph_pool_handle, _xpu_isCurrentStreamCapturing, _XPUGraph @@ -71,7 +70,7 @@ class XPUGraph(_XPUGraph): def __new__(cls, keep_graph: bool = False) -> Self: return super().__new__(cls, keep_graph) - def capture_begin(self, pool: Optional[_POOL_HANDLE] = None) -> None: + def capture_begin(self, pool: _POOL_HANDLE | None = None) -> None: r"""Begin capturing XPU work on the current xpu stream. Typically, you shouldn't call ``capture_begin`` yourself. @@ -165,13 +164,13 @@ class graph: """ # noqa: B950 - default_capture_stream: Optional[torch.xpu.Stream] = None + default_capture_stream: torch.xpu.Stream | None = None def __init__( self, xpu_graph: XPUGraph, - pool: Optional[_POOL_HANDLE] = None, - stream: Optional[torch.xpu.Stream] = None, + pool: _POOL_HANDLE | None = None, + stream: torch.xpu.Stream | None = None, ): # Lazy-init of default_capture_stream helps avoid circular-import errors. # Not thread safe, but graphs already have the general (explicitly documented) @@ -179,9 +178,7 @@ def __init__( if self.__class__.default_capture_stream is None: self.__class__.default_capture_stream = torch.xpu.Stream() - self.pool: Union[tuple[()], tuple[_POOL_HANDLE]] = ( - () if pool is None else (pool,) - ) + self.pool: tuple[()] | tuple[_POOL_HANDLE] = () if pool is None else (pool,) self.capture_stream = ( stream if stream is not None else self.__class__.default_capture_stream ) @@ -204,7 +201,7 @@ def __exit__(self, *args: object) -> None: self.stream_ctx.__exit__(*args) -_ModuleOrCallable: TypeAlias = Union["torch.nn.Module", Callable[..., object]] +_ModuleOrCallable: TypeAlias = torch.nn.Module | Callable[..., object] @overload @@ -213,7 +210,7 @@ def make_graphed_callables( sample_args: tuple[Tensor, ...], num_warmup_iters: int = 3, allow_unused_input: bool = False, - pool: Optional[_POOL_HANDLE] = None, + pool: _POOL_HANDLE | None = None, ) -> _ModuleOrCallable: ... @@ -223,17 +220,17 @@ def make_graphed_callables( sample_args: tuple[tuple[Tensor, ...], ...], num_warmup_iters: int = 3, allow_unused_input: bool = False, - pool: Optional[_POOL_HANDLE] = None, + pool: _POOL_HANDLE | None = None, ) -> tuple[_ModuleOrCallable, ...]: ... def make_graphed_callables( - callables: Union[_ModuleOrCallable, tuple[_ModuleOrCallable, ...]], - sample_args: Union[tuple[Tensor, ...], tuple[tuple[Tensor, ...], ...]], + callables: _ModuleOrCallable | tuple[_ModuleOrCallable, ...], + sample_args: tuple[Tensor, ...] | tuple[tuple[Tensor, ...], ...], num_warmup_iters: int = 3, allow_unused_input: bool = False, - pool: Optional[_POOL_HANDLE] = None, -) -> Union[_ModuleOrCallable, tuple[_ModuleOrCallable, ...]]: + pool: _POOL_HANDLE | None = None, +) -> _ModuleOrCallable | tuple[_ModuleOrCallable, ...]: r"""Accept callables (functions or :class:`nn.Module`\ s) and returns graphed versions. Each graphed callable's forward pass runs its source callable's @@ -445,7 +442,7 @@ def make_graphed_autograd_function( output_unflatten_spec: torch.utils._pytree.TreeSpec, static_input_surface: tuple[Tensor, ...], static_outputs: tuple[Tensor, ...], - static_grad_outputs: tuple[Optional[Tensor], ...], + static_grad_outputs: tuple[Tensor | None, ...], static_grad_inputs: tuple[Tensor, ...], ) -> Callable[..., object]: class Graphed(torch.autograd.Function): @@ -478,9 +475,7 @@ def backward(ctx: object, *grads: Tensor) -> tuple[Tensor, ...]: if not isinstance(static_grad_inputs, tuple): raise RuntimeError("static_grad_inputs must be a tuple") return tuple( - # pyrefly: ignore [bad-argument-type] - b.detach() if b is not None else b - for b in static_grad_inputs + b.detach() if b is not None else b for b in static_grad_inputs ) def functionalized(*user_args: object) -> object: diff --git a/torch/xpu/memory.py b/torch/xpu/memory.py index 9e70381e17b4b..04ca0dd9fc397 100644 --- a/torch/xpu/memory.py +++ b/torch/xpu/memory.py @@ -251,7 +251,6 @@ def set_per_process_memory_fraction(fraction: float, device: Device = None) -> N device = _get_device_index(device, optional=True) if not isinstance(fraction, float): raise TypeError("Invalid type for fraction argument, must be `float`") - # pyrefly: ignore [missing-attribute] torch._C._xpu_setMemoryFraction(fraction, device) @@ -271,7 +270,6 @@ def memory_snapshot( """ if not is_initialized(): return [] - # pyrefly: ignore [missing-attribute] return torch._C._xpu_memorySnapshot(mempool_id)["segments"] @@ -358,7 +356,6 @@ class TraceEntry(TypedDict): Returns: The Snapshot dictionary object """ - # pyrefly: ignore [missing-attribute] s = torch._C._xpu_memorySnapshot(None) if augment_with_fx_traces: s = _augment_memory_snapshot_stack_traces(s) # type: ignore[assignment, arg-type] @@ -467,7 +464,6 @@ def _record_memory_history( Defaults to ``None`` (record all actions). """ - # pyrefly: ignore [missing-attribute] torch._C._xpu_recordMemoryHistory( enabled, context, @@ -481,7 +477,6 @@ def _record_memory_history( class _XPUAllocator: r"""Wrapper over internal XPU memory allocators.""" - # pyrefly: ignore [missing-attribute] def __init__(self, allocator: torch._C._xpu_XPUAllocator): self._allocator = allocator @@ -528,7 +523,6 @@ def __init__(self, path_to_lib_file: str, alloc_fn_name: str, free_fn_name: str) "Failed to load allocator symbols from the shared library." ) - # pyrefly: ignore [missing-attribute] self._allocator = torch._C._xpu_customAllocator(alloc_fn_addr, free_fn_addr) @@ -541,7 +535,6 @@ def change_current_allocator(allocator: _XPUAllocator) -> None: Arguments: allocator (torch.xpu.memory._XPUAllocator): allocator to be set as the active one. """ - # pyrefly: ignore [missing-attribute] torch._C._xpu_changeCurrentAllocator(allocator.allocator()) @@ -551,7 +544,6 @@ def _get_current_allocator() -> _XPUAllocator: Returns: _XPUAllocator: the allocator being currently used. """ - # pyrefly: ignore [missing-attribute] return _XPUAllocator(torch._C._xpu_getAllocator()) diff --git a/torch/xpu/streams.py b/torch/xpu/streams.py index 87c953ce7c323..3e5068ee3b44b 100644 --- a/torch/xpu/streams.py +++ b/torch/xpu/streams.py @@ -121,7 +121,6 @@ class Event(torch._C._XpuEventBase): def __new__(cls, enable_timing=False): return super().__new__(cls, enable_timing=enable_timing) - # pyrefly: ignore [bad-override] def record(self, stream: Stream | torch.Stream | None = None) -> None: r"""Record the event in a given stream. @@ -131,7 +130,6 @@ def record(self, stream: Stream | torch.Stream | None = None) -> None: """ if stream is None: stream = torch.xpu.current_stream() - # pyrefly: ignore [bad-argument-type] super().record(stream) def wait(self, stream: Stream | torch.Stream | None = None) -> None: @@ -142,7 +140,6 @@ def wait(self, stream: Stream | torch.Stream | None = None) -> None: """ if stream is None: stream = torch.xpu.current_stream() - # pyrefly: ignore [bad-argument-type] super().wait(stream) def query(self) -> bool: diff --git a/torchgen/_autoheuristic/benchmark_runner.py b/torchgen/_autoheuristic/benchmark_runner.py index 117058b3373f3..67e24aec82d36 100644 --- a/torchgen/_autoheuristic/benchmark_runner.py +++ b/torchgen/_autoheuristic/benchmark_runner.py @@ -55,12 +55,15 @@ def add_base_arguments(self) -> None: def run(self) -> None: torch.set_default_device("cuda") args = self.parser.parse_args() + # Set environment variables to control autoheuristic behavior + import os + if args.use_heuristic: - torch._inductor.config.autoheuristic_use = self.name - torch._inductor.config.autoheuristic_collect = "" + os.environ["TORCHINDUCTOR_AUTOHEURISTIC_USE"] = self.name + os.environ["TORCHINDUCTOR_AUTOHEURISTIC_COLLECT"] = "" else: - torch._inductor.config.autoheuristic_use = "" - torch._inductor.config.autoheuristic_collect = self.name + os.environ["TORCHINDUCTOR_AUTOHEURISTIC_USE"] = "" + os.environ["TORCHINDUCTOR_AUTOHEURISTIC_COLLECT"] = self.name torch._inductor.config.autoheuristic_log_path = args.o if args.device is not None: torch.cuda.set_device(args.device) diff --git a/torchgen/_autoheuristic/pad_mm/collect_known_mm_shapes.py b/torchgen/_autoheuristic/pad_mm/collect_known_mm_shapes.py new file mode 100644 index 0000000000000..45a5854044b33 --- /dev/null +++ b/torchgen/_autoheuristic/pad_mm/collect_known_mm_shapes.py @@ -0,0 +1,235 @@ +import argparse +import csv +import sys +from pathlib import Path + + +# Add parent directory to path for imports +sys.path.append(str(Path(__file__).absolute().parents[1])) +sys.path.append( + str( + Path(__file__).absolute().parents[3] + / "benchmarks" + / "dynamo" + / "microbenchmarks" + ) +) + +from operator_inp_utils import ( # type: ignore[import-not-found] + deserialize_args, + OperatorInputsLoader, +) + +import torch +from torch._inductor.fx_passes.pad_mm import ( + get_alignment_size_dtype, # type: ignore[import-not-found] +) +from torch._subclasses.fake_tensor import FakeTensorMode + + +def is_aligned(dim: int, align_size: int) -> bool: + """Check if dimension is aligned to the given alignment size.""" + return dim % align_size == 0 + + +def extract_mm_shapes_from_loader( + loader: OperatorInputsLoader, +) -> list[tuple[int, int, int, torch.dtype, torch.dtype]]: + """Extract matrix multiplication shapes from an OperatorInputsLoader using deserialize_args with FakeTensorMode.""" + shapes = [] + + # Matrix multiplication operators to look for + mm_operators = ["aten.mm.default", "aten.addmm.default", "aten.bmm.default"] + + # Use FakeTensorMode to avoid instantiating actual tensors + with FakeTensorMode(): + for op_name in mm_operators: + if op_name not in loader.operator_db: + continue + + # Count shapes extracted from this operator + shape_count = 0 + + # Access the raw string data directly from operator_db and reuse existing parsing + for input_str in loader.operator_db[op_name]: + try: + # Use deserialize_args to parse inputs - will create fake tensors + args, kwargs = deserialize_args(input_str) + + if op_name == "aten.mm.default": + # mm(input, mat2) -> result + if len(args) >= 2: + a, b = args[0], args[1] + if isinstance(a, torch.Tensor) and isinstance( + b, torch.Tensor + ): + a_shape, a_dtype = tuple(a.shape), a.dtype + b_shape, b_dtype = tuple(b.shape), b.dtype + if len(a_shape) == 2 and len(b_shape) == 2: + m, k = a_shape + k2, n = b_shape + if k == k2: # Valid matrix multiplication + shapes.append((m, k, n, a_dtype, b_dtype)) + shape_count += 1 + + elif op_name == "aten.addmm.default": + # addmm(bias, input, mat2) -> result + if len(args) >= 3: + _, a, b = args[0], args[1], args[2] + if isinstance(a, torch.Tensor) and isinstance( + b, torch.Tensor + ): + a_shape, a_dtype = tuple(a.shape), a.dtype + b_shape, b_dtype = tuple(b.shape), b.dtype + if len(a_shape) == 2 and len(b_shape) == 2: + m, k = a_shape + k2, n = b_shape + if k == k2: # Valid matrix multiplication + shapes.append((m, k, n, a_dtype, b_dtype)) + shape_count += 1 + + elif op_name == "aten.bmm.default": + # bmm(input, mat2) -> result (batch matrix multiplication) + if len(args) >= 2: + a, b = args[0], args[1] + if isinstance(a, torch.Tensor) and isinstance( + b, torch.Tensor + ): + a_shape, a_dtype = tuple(a.shape), a.dtype + b_shape, b_dtype = tuple(b.shape), b.dtype + if len(a_shape) == 3 and len(b_shape) == 3: + batch1, m, k = a_shape + batch2, k2, n = b_shape + if ( + batch1 == batch2 and k == k2 + ): # Valid batch matrix multiplication + shapes.append((m, k, n, a_dtype, b_dtype)) + shape_count += 1 + + except Exception: + # Skip invalid inputs + continue + + print(f" Extracted {shape_count} shapes from {op_name}") + + return shapes + + +def filter_unaligned_shapes( + shapes: list[tuple[int, int, int, torch.dtype, torch.dtype]], +) -> list[tuple[int, int, int, torch.dtype, torch.dtype]]: + """Filter shapes to keep only those that are not completely aligned (so padding is relevant).""" + filtered_shapes = [] + + for m, k, n, dtype1, dtype2 in shapes: + # Use the primary dtype for alignment calculation (assume both dtypes are similar for alignment purposes) + dtype = dtype1 + try: + align_size = get_alignment_size_dtype(dtype) + + # Only keep shapes where not all dimensions are aligned + if not all(is_aligned(dim, align_size) for dim in [m, k, n]): + filtered_shapes.append((m, k, n, dtype1, dtype2)) + + except Exception: + # If we can't get alignment size, skip this shape + continue + + return filtered_shapes + + +def collect_known_mm_shapes() -> list[tuple[int, int, int, torch.dtype, torch.dtype]]: + """ + Collect known matrix multiplication shapes from HuggingFace, TIMM, and TorchBench datasets. + + Returns: + List of tuples containing (m, k, n, dtype1, dtype2) for matrix multiplication shapes + that are not completely aligned (so padding is relevant). + """ + all_shapes = [] + + loaders = [] + + # Try to load each dataset + try: + hf_loader = OperatorInputsLoader.get_huggingface_loader() + loaders.append(("HuggingFace", hf_loader)) + except Exception as e: + print(f"Warning: Could not load HuggingFace dataset: {e}") + + try: + timm_loader = OperatorInputsLoader.get_timm_loader() + loaders.append(("TIMM", timm_loader)) + except Exception as e: + print(f"Warning: Could not load TIMM dataset: {e}") + + try: + torchbench_loader = OperatorInputsLoader.get_torchbench_loader() + loaders.append(("TorchBench", torchbench_loader)) + except Exception as e: + print(f"Warning: Could not load TorchBench dataset: {e}") + + # Extract shapes from each loader + for dataset_name, loader in loaders: + print(f"Extracting shapes from {dataset_name}...") + + shapes = extract_mm_shapes_from_loader(loader) + print(f"Found {len(shapes)} matrix multiplication shapes from {dataset_name}") + all_shapes.extend(shapes) + + # Remove duplicates + unique_shapes = list(set(all_shapes)) + print(f"Total unique shapes before filtering: {len(unique_shapes)}") + + # Filter for unaligned shapes only + filtered_shapes = filter_unaligned_shapes(unique_shapes) + print(f"Shapes after filtering for unaligned: {len(filtered_shapes)}") + + return filtered_shapes + + +def main(output_file="mm_shapes.csv"): + shapes = collect_known_mm_shapes() + + print(f"\nCollected {len(shapes)} real-world matrix multiplication shapes") + + # Convert dtype objects to strings and filter for desired dtypes + dtype_map = { + torch.float16: "float16", + torch.bfloat16: "bfloat16", + torch.float32: "float32", + } + + # Convert to desired format and filter dtypes + csv_rows = [] + for m, k, n, dtype1, dtype2 in shapes: + # Use the first dtype and convert to string + if dtype1 in dtype_map: + dtype_str = dtype_map[dtype1] + csv_rows.append([m, k, n, dtype_str]) + + # Save to CSV file + with open(output_file, "w", newline="") as csvfile: + writer = csv.writer(csvfile) + # Write header + writer.writerow(["M", "K", "N", "dtype"]) + # Write data rows + writer.writerows(csv_rows) + + print(f"Saved matrix multiplication shapes to {output_file}") + + +if __name__ == "__main__": + parser = argparse.ArgumentParser( + description="Collect matrix multiplication shapes from real-world datasets and save to CSV" + ) + parser.add_argument( + "--output", + "-o", + type=str, + default="mm_shapes.csv", + help="Output CSV filename (default: mm_shapes.csv)", + ) + + args = parser.parse_args() + main(args.output) diff --git a/torchgen/_autoheuristic/pad_mm/evaluate_pad_mm_heuristics.py b/torchgen/_autoheuristic/pad_mm/evaluate_pad_mm_heuristics.py new file mode 100644 index 0000000000000..42b37d74c3092 --- /dev/null +++ b/torchgen/_autoheuristic/pad_mm/evaluate_pad_mm_heuristics.py @@ -0,0 +1,426 @@ +#!/usr/bin/env python3 + +import argparse +import csv +import functools + +import torch +from torch._inductor.autoheuristic.autoheuristic_utils import AHContext, AHMetadata +from torch._inductor.fx_passes.pad_mm import get_alignment_size_dtype +from torch._inductor.runtime.benchmarking import benchmarker +from torch._inductor.utils import get_gpu_shared_memory + + +def fits_in_memory(dtype, m: int, k: int, n: int) -> bool: + threshold_memory = torch.cuda.get_device_properties(0).total_memory / 4 + return dtype.itemsize * (m * k + k * n + m * n) < threshold_memory + + +def set_precision(dtype, float32_precision: str = "highest") -> None: + precision = float32_precision if dtype == torch.float32 else "high" + torch.set_float32_matmul_precision(precision) + + +def get_heuristic_decision(m: int, k: int, n: int, dtype: torch.dtype) -> str | None: + from torch._inductor.autoheuristic.autoheuristic import AutoHeuristic, LocalFeedback + from torch._inductor.fx_passes.pad_mm import ( + get_alignment_size, + get_context, + get_padded_length, + pad_mm_operations, + pad_mm_precondition, + ) + + torch._inductor.config.autoheuristic_use.pad_mm = True + + if not torch._inductor.config.run_autoheuristic("pad_mm"): + return None + + a = torch.randn(m, k, dtype=dtype, device="cuda") + b = torch.randn(k, n, dtype=dtype, device="cuda") + + m_padded_length = get_padded_length(m, get_alignment_size(a)) + k_padded_length = get_padded_length(k, get_alignment_size(a)) + n_padded_length = get_padded_length(n, get_alignment_size(b)) + + context = get_context( + a, + b, + mat1_pre_padded=False, + mat2_pre_padded=False, + m_padded_length=m_padded_length, + k_padded_length=k_padded_length, + n_padded_length=n_padded_length, + ) + + def dummy_feedback(choice: str) -> float: + return 1.0 + + def fallback() -> str: + return "no_decision" + + autoheuristic = AutoHeuristic( + fallback=fallback, + choices=["orig", "pad"], + feedback=LocalFeedback(dummy_feedback), + context=context, + name="pad_mm", + augment_context=pad_mm_operations(), + precondition=pad_mm_precondition, + ) + + choice = autoheuristic.get_choice() + return choice + + +def benchmark_both_choices( + m: int, + k: int, + n: int, + dtype: torch.dtype, + num_reps: int = 3, + float32_precision: str = "highest", +) -> tuple[float, float]: + set_precision(dtype, float32_precision) + a = torch.randn(m, k, dtype=dtype, device="cuda") + b = torch.randn(k, n, dtype=dtype, device="cuda") + + # Use existing benchmarking infrastructure with proper cache management + # benchmarker returns time in milliseconds, so convert to seconds for consistency + orig_time_ms = benchmarker.benchmark( + torch.mm, fn_args=(a, b), rep=num_reps, is_vetted_benchmarking=True + ) + orig_time = orig_time_ms / 1000.0 # Convert ms to seconds + + from torch._inductor.fx_passes.pad_mm import ( + get_alignment_size, + get_padded_length, + pad_mm, + ) + + m_padded_length = get_padded_length(a.shape[0], get_alignment_size(a)) + k_padded_length = get_padded_length(a.shape[1], get_alignment_size(a)) + n_padded_length = get_padded_length(b.shape[1], get_alignment_size(b)) + + if m_padded_length == 0 and k_padded_length == 0 and n_padded_length == 0: + return orig_time, orig_time + + pad_time_ms = benchmarker.benchmark( + pad_mm, + fn_args=(a, b, m_padded_length, k_padded_length, n_padded_length), + rep=num_reps, + is_vetted_benchmarking=True, + ) + pad_time = pad_time_ms / 1000.0 # Convert ms to seconds + + return orig_time, pad_time + + +def load_shapes_from_csv(csv_file: str) -> list: + shapes = [] + with open(csv_file) as f: + reader = csv.DictReader(f) + for row in reader: + m, k, n = int(row["M"]), int(row["K"]), int(row["N"]) + dtype_str = row["dtype"] + + if dtype_str == "float16": + dtype = torch.float16 + elif dtype_str == "bfloat16": + dtype = torch.bfloat16 + elif dtype_str == "float32": + dtype = torch.float32 + else: + continue + + shapes.append((m, k, n, dtype)) + + print(f"Loaded {len(shapes)} shapes from {csv_file}") + return shapes + + +@functools.cache +def get_shared_mem_size(): + return get_gpu_shared_memory() + + +def check_shape_passes_precondition(m: int, k: int, n: int, dtype: torch.dtype) -> bool: + """ + Check if a shape passes the same precondition used by the actual pad_mm AutoHeuristics. + + This uses the exact same pad_mm_precondition function that the AutoHeuristic system + uses, avoiding hardcoded magic numbers by delegating to the source of truth. + """ + from torch._inductor.autoheuristic.autoheuristic_utils import pad_mm_precondition + + shared_memory = get_shared_mem_size() + device_capa = torch.cuda.get_device_capability() + + # Create the same metadata and context that AutoHeuristics uses + metadata = AHMetadata( + shared_memory=shared_memory, + device_capa=device_capa, + choices=["orig", "pad"], # Required but not used for precondition check + name="pad_mm", # Required but not used for precondition check + ) + + context = AHContext() + context.add_feature("m", m) + context.add_feature("k", k) + context.add_feature("n", n) + + # Use the actual pad_mm_precondition function - no hardcoded values! + return pad_mm_precondition(metadata, context) + + +def filter_shapes(shapes: list) -> list: + filtered = [] + aligned_count = 0 + precondition_failed_count = 0 + memory_count = 0 + + for m, k, n, dtype in shapes: + # Check if already aligned + align_size = get_alignment_size_dtype(dtype) + is_aligned = all((dim % align_size == 0) for dim in [m, k, n]) + + if is_aligned: + aligned_count += 1 + continue + + # Check if passes the actual precondition used by pad_mm AutoHeuristics + if not check_shape_passes_precondition(m, k, n, dtype): + precondition_failed_count += 1 + continue + + # Check if fits in memory + if not fits_in_memory(dtype, m, k, n): + memory_count += 1 + continue + + # This shape is suitable for evaluation + filtered.append((m, k, n, dtype)) + + print("Filtering results:") + print(f" Already aligned (skipped): {aligned_count}") + print(f" Failed pad_mm_precondition (skipped): {precondition_failed_count}") + print(f" Too large for memory (skipped): {memory_count}") + print(f" Suitable for evaluation: {len(filtered)}") + + return filtered + + +def main(): + parser = argparse.ArgumentParser( + description="Evaluate trained AutoHeuristics for pad_mm optimization" + ) + parser.add_argument("csv_file", help="Path to CSV file with M,K,N,dtype columns") + parser.add_argument( + "--num-reps", type=int, default=3, help="Benchmark repetitions (default: 3)" + ) + parser.add_argument( + "--device", type=int, default=None, help="CUDA device (default: current)" + ) + parser.add_argument( + "--max-shapes", + type=int, + default=10000, + help="Max shapes to test (default: 10000)", + ) + parser.add_argument( + "--float32_matmul_precision", + type=str, + choices=["high", "highest"], + default="highest", + help="Matmul precision for float32 (default: highest). Non-fp32 always uses 'high'.", + ) + + args = parser.parse_args() + + torch.set_default_device("cuda") + if args.device is not None: + torch.cuda.set_device(args.device) + + print(f"Using CUDA device: {torch.cuda.current_device()}") + print() + + shapes = load_shapes_from_csv(args.csv_file) + if not shapes: + print("No shapes found!") + return + + shapes = filter_shapes(shapes) + if not shapes: + print("No suitable shapes found!") + return + + if len(shapes) > args.max_shapes: + shapes = shapes[: args.max_shapes] + print(f"Limited to first {args.max_shapes} shapes") + + print(f"Evaluating {len(shapes)} shapes with {args.num_reps} reps each") + print() + + total_decisions = 0 + correct_decisions = 0 + true_positives = 0 # Chose pad, should pad + true_negatives = 0 # Chose orig, should orig + false_positives = 0 # Chose pad, should orig + false_negatives = 0 # Chose orig, should pad + no_decision_shapes = 0 + + tp_speedups = [] # Speed-up percentages for true positives + fp_slowdowns = [] # Speed-down percentages for false positives + + # Track non-confident decisions and confident decisions by dtype + no_decision_shape_list = [] # List of (M, K, N, dtype) where heuristic chose no_decision + confident_by_dtype = {} # Count of confident decisions by dtype + + for i, (m, k, n, dtype) in enumerate(shapes, 1): + print(f"Shape {i}/{len(shapes)}: M={m}, K={k}, N={n}, dtype={dtype}") + + heuristic_choice = get_heuristic_decision(m, k, n, dtype) + print(f" Heuristic: {heuristic_choice}") + + orig_time, pad_time = benchmark_both_choices( + m, k, n, dtype, args.num_reps, args.float32_matmul_precision + ) + ground_truth = "pad" if pad_time < orig_time else "orig" + + print(f" Times: orig={orig_time:.6f}s, pad={pad_time:.6f}s") + print(f" Ground truth: {ground_truth}") + + if heuristic_choice == "no_decision": + # Heuristic punted to benchmarking - this is correct behavior for small/uncertain shapes + no_decision_shapes += 1 + no_decision_shape_list.append((m, k, n, dtype)) + print(" Heuristic chose to benchmark (conservative)") + else: + # Heuristic made a confident decision - evaluate accuracy + total_decisions += 1 + # Track confident decisions by dtype + dtype_str = str(dtype).replace("torch.", "") + confident_by_dtype[dtype_str] = confident_by_dtype.get(dtype_str, 0) + 1 + if heuristic_choice == ground_truth: + correct_decisions += 1 + print(" ✓ CORRECT") + if heuristic_choice == "pad": + true_positives += 1 # Correctly chose pad + # Calculate speed-up: (orig_time - pad_time) / orig_time * 100 + speedup = (orig_time - pad_time) / orig_time * 100 + tp_speedups.append(speedup) + print(f" Speed-up: {speedup:.1f}%") + else: + true_negatives += 1 # Correctly chose orig + else: + print(" ✗ WRONG") + if heuristic_choice == "pad" and ground_truth == "orig": + false_positives += 1 + # Calculate speed-down: (pad_time - orig_time) / orig_time * 100 + slowdown = (pad_time - orig_time) / orig_time * 100 + fp_slowdowns.append(slowdown) + print(f" Speed-down: {slowdown:.1f}%") + elif heuristic_choice == "orig" and ground_truth == "pad": + false_negatives += 1 + + print(f" Confidence Rate: {total_decisions}/{i}") + if total_decisions > 0: + accuracy = correct_decisions / total_decisions * 100 + tp_rate = true_positives / total_decisions * 100 + tn_rate = true_negatives / total_decisions * 100 + fp_rate = false_positives / total_decisions * 100 + fn_rate = false_negatives / total_decisions * 100 + + # Compute average speedup/slowdown + avg_tp_speedup = sum(tp_speedups) / len(tp_speedups) if tp_speedups else 0 + avg_fp_slowdown = ( + sum(fp_slowdowns) / len(fp_slowdowns) if fp_slowdowns else 0 + ) + + print( + f" Accuracy: {correct_decisions}/{total_decisions} ({accuracy:.1f}%) " + f"| TP: {tp_rate:.1f}% (avg speedup: {avg_tp_speedup:.1f}%) " + f"| TN: {tn_rate:.1f}% " + f"| FP: {fp_rate:.1f}% (avg slowdown: {avg_fp_slowdown:.1f}%)" + f"| FN: {fn_rate:.1f}%" + ) + + print() + + print("=== FINAL RESULTS ===") + print(f"Confident decisions: {total_decisions}") + print(f"#Shapes without confident decisions: {no_decision_shapes}") + + if total_decisions > 0: + accuracy = correct_decisions / total_decisions * 100 + tp_rate = true_positives / total_decisions * 100 + tn_rate = true_negatives / total_decisions * 100 + fp_rate = false_positives / total_decisions * 100 + fn_rate = false_negatives / total_decisions * 100 + + avg_tp_speedup = sum(tp_speedups) / len(tp_speedups) if tp_speedups else 0 + avg_fp_slowdown = sum(fp_slowdowns) / len(fp_slowdowns) if fp_slowdowns else 0 + + print( + f"\nConfident decision accuracy: {accuracy:.1f}% ({correct_decisions}/{total_decisions})" + ) + + if tp_speedups: + print( + f"True Positives (chose pad, should pad): {tp_rate:.1f}% ({true_positives}) " + f"| Avg speed-up: {avg_tp_speedup:.1f}%" + ) + else: + print( + f"True Positives (chose pad, should pad): {tp_rate:.1f}% ({true_positives})" + ) + + print( + f"True Negatives (chose orig, should orig): {tn_rate:.1f}% ({true_negatives})" + ) + + if fp_slowdowns: + print( + f"False Positives (chose pad, should orig): {fp_rate:.1f}% ({false_positives}) " + f"| Avg speed-down: {avg_fp_slowdown:.1f}%" + ) + else: + print( + f"False Positives (chose pad, should orig): {fp_rate:.1f}% ({false_positives})" + ) + + print( + f"False Negatives (chose orig, should pad): {fn_rate:.1f}% ({false_negatives})" + ) + else: + print("No confident decisions made!") + + total_evaluated = total_decisions + no_decision_shapes + if total_evaluated > 0: + print( + f"\nConfidence rate: ({total_decisions}/{total_evaluated} made confident decisions)" + ) + + # Print shapes where AutoHeuristics did not make a confident decision + print(f"\n=== NON-CONFIDENT DECISIONS ({len(no_decision_shape_list)}) ===") + if no_decision_shape_list: + print("Shapes where AutoHeuristics chose 'no_decision' (non-confident):") + for m, k, n, dtype in no_decision_shape_list: + dtype_str = str(dtype).replace("torch.", "") + print(f" M={m}, K={k}, N={n}, dtype={dtype_str}") + else: + print("All shapes had confident decisions!") + + # Print confident decisions by dtype + print("\n=== CONFIDENT DECISIONS BY DTYPE ===") + if confident_by_dtype: + print("Number of confident decisions per dtype:") + for dtype_str, count in sorted(confident_by_dtype.items()): + print(f" {dtype_str}: {count} confident decisions") + print(f"Total confident decisions: {sum(confident_by_dtype.values())}") + else: + print("No confident decisions made!") + + +if __name__ == "__main__": + main() diff --git a/torchgen/_autoheuristic/pad_mm/gen_data_pad_mm.py b/torchgen/_autoheuristic/pad_mm/gen_data_pad_mm.py index b476bacfb67db..eba7dd19b53b2 100644 --- a/torchgen/_autoheuristic/pad_mm/gen_data_pad_mm.py +++ b/torchgen/_autoheuristic/pad_mm/gen_data_pad_mm.py @@ -1,5 +1,7 @@ +import csv import random import sys +from collections.abc import Generator from pathlib import Path from typing import Any @@ -10,9 +12,13 @@ from benchmark_utils import ( # type: ignore[import-not-found] fits_in_memory, get_mm_tensors, + get_random_between_pow2, set_precision, transpose_tensors, ) +from collect_known_mm_shapes import ( + collect_known_mm_shapes, # type: ignore[import-not-found] +) import torch from torch._inductor.fx_passes.pad_mm import ( # type: ignore[import-not-found] @@ -29,10 +35,121 @@ class BenchmarkRunnerPadMM(BenchmarkRunner): # type: ignore[misc, no-any-unimpo def __init__(self) -> None: super().__init__("pad_mm") + # Add CLI argument for additional shape CSV files + self.parser.add_argument( + "--additional-shape-csv", + nargs="*", + default=[], + help="List of CSV files containing additional matrix multiplication shapes (M,K,N,dtype format)", + ) + + # Initialize additional_shape_collections + self.additional_shape_collections: list[ + list[tuple[int, int, int, torch.dtype, torch.dtype]] + ] = [] + + # Initialize the shape generator (will be set up after parsing args) + self.shape_generator = None + + def load_shapes_from_csv( + self, csv_file: str + ) -> list[tuple[int, int, int, torch.dtype, torch.dtype]]: + """Load matrix multiplication shapes from a CSV file in M,K,N,dtype format.""" + shapes = [] + dtype_map = { + "float16": torch.float16, + "bfloat16": torch.bfloat16, + "float32": torch.float32, + } + + try: + with open(csv_file) as f: + reader = csv.DictReader(f) + for row in reader: + m = int(row["M"]) + k = int(row["K"]) + n = int(row["N"]) + dtype_str = row["dtype"] + + if dtype_str in dtype_map: + dtype = dtype_map[dtype_str] + # Store as (m, k, n, dtype1, dtype2) with same dtype for both + shapes.append((m, k, n, dtype, dtype)) + else: + print( + f"Warning: Unknown dtype '{dtype_str}' in {csv_file}, skipping row" + ) + + print(f"Loaded {len(shapes)} shapes from {csv_file}") + except Exception as e: + print(f"Error loading shapes from {csv_file}: {e}") + + return shapes + + def setup_shape_collections(self, csv_files: list[str]) -> None: + """Setup additional shape collections from CSV files and built-in collection.""" + self.additional_shape_collections = [] + + # Load shapes from provided CSV files first + for csv_file in csv_files: + shapes = self.load_shapes_from_csv(csv_file) + if shapes: + self.additional_shape_collections.append(shapes) + + self.additional_shape_collections.append(collect_known_mm_shapes()) + self.shape_generator = self.generate_mm_shapes() + + def generate_mm_shapes(self) -> Generator[tuple[int, int, int, Any], None, None]: + """Generator that yields (m, k, n, dtype) tuples for matrix multiplication. + + First exhausts all shapes from additional_shape_collections, then generates random shapes. + Only yields unaligned shapes since external CSV shapes may not be pre-filtered. + """ + # Phase 1: Use all shapes from additional shape collections + for collection in self.additional_shape_collections: + for m, k, n, dtype1, _ in collection: + # Filter for unaligned shapes only (external CSVs may not be pre-filtered) + align_size = get_alignment_size_dtype(dtype1) + if not all(self.is_aligned(dim, align_size) for dim in [m, k, n]): + # Check if it fits in memory + if fits_in_memory(dtype1, m, k, n): + yield (m, k, n, dtype1) + + # Phase 2: Generate infinite random shapes + + while True: + # Generate random dtype + dtype_choices = [torch.float16, torch.bfloat16, torch.float32] + dtype = random.choices(dtype_choices)[0] + + # Generate random shape for this dtype + uniform = random.choices([True, False])[0] + align_size = get_alignment_size_dtype(dtype) + + # Keep trying until we get a valid unaligned shape that fits in memory + while True: + if uniform: + m = random.randint(1, 65536) + k = random.randint(1, 65536) + n = random.randint(1, 65536) + else: + m = self.get_random_dim() + k = self.get_random_dim() + n = self.get_random_dim() + + # Skip if all dimensions are aligned (we need unaligned for padding to be relevant) + if all(self.is_aligned(dim, align_size) for dim in [m, k, n]): + continue + + # Check if it fits in memory + if fits_in_memory(dtype, m, k, n): + yield (m, k, n, dtype) + break + def create_input(self) -> tuple[Any, ...]: - dtype = self.get_dtype() + # Get the next shape from the generator + m, k, n, dtype = next(self.shape_generator) set_precision(dtype) - m, k, n = self.get_m_k_n(dtype) (transpose_left, transpose_right) = transpose_tensors() prepadded_left = self.prepadded() @@ -107,40 +224,45 @@ def get_random_dim( return 2 ** random.randint(min_power2, max_power2) # type: ignore[no-any-return] else: # choose a random number between 2^i and 2^(i+1) - return self.get_random_between_pow2(min_power2, max_power2) # type: ignore[no-any-return] + return get_random_between_pow2(min_power2, max_power2) # type: ignore[no-any-return] def is_aligned(self, dim: int, align_size: int) -> bool: return dim % align_size == 0 - def get_m_k_n(self, dtype: Any) -> tuple[int, int, int]: - uniform = random.choices([True, False])[0] - align_size = get_alignment_size_dtype(dtype) + def prepadded(self, p_prepadded: float = 0.2) -> bool: + # p_prepadded: probability that a tensor is "prepadded", i.e. pad_mm excludes time it takes to pad from benchmarking + return random.choices([True, False], [p_prepadded, 1 - p_prepadded])[0] + + def run(self) -> None: + """Override run to setup shape collections before running.""" + import time - # repeat until tensors fit in memory - while True: - if uniform: - m = random.randint(1, 65536) - k = random.randint(1, 65536) - n = random.randint(1, 65536) - else: - m = self.get_random_dim() - k = self.get_random_dim() - n = self.get_random_dim() + from tqdm import tqdm - if all(self.is_aligned(dim, align_size) for dim in [m, k, n]): - # skip if already aligned - continue + torch.set_default_device("cuda") + args = self.parser.parse_args() - if fits_in_memory(dtype, m, k, n): - return (m, k, n) + # Setup shape collections based on CLI arguments + self.setup_shape_collections(args.additional_shape_csv) - def prepadded(self, p_prepadded: float = 0.2) -> bool: - # p_prepadded: probability that a tensor is "prepadded", i.e. pad_mm excludes time it takes to pad from benchmarking - return random.choices([True, False], [p_prepadded, 1 - p_prepadded])[0] + # Set up torch configuration (copied from parent run method) + + if args.use_heuristic: + torch._inductor.config.autoheuristic_use.pad_mm = True + torch._inductor.config.autoheuristic_collect.pad_mm = False + else: + torch._inductor.config.autoheuristic_use.pad_mm = False + torch._inductor.config.autoheuristic_collect.pad_mm = True + torch._inductor.config.autoheuristic_log_path = args.o + if args.device is not None: + torch.cuda.set_device(args.device) + random.seed(time.time()) - def get_dtype(self) -> Any: - dtype_choices = [torch.float16, torch.bfloat16, torch.float32] - return random.choices(dtype_choices)[0] + # Run the main benchmarking loop + for _ in tqdm(range(args.num_samples)): + input = self.create_input() + for _ in range(args.num_reps): + self.run_benchmark(*input) if __name__ == "__main__": diff --git a/torchgen/aoti/fallback_ops.py b/torchgen/aoti/fallback_ops.py index 93a2b3362efe2..84d413d426d19 100644 --- a/torchgen/aoti/fallback_ops.py +++ b/torchgen/aoti/fallback_ops.py @@ -36,16 +36,18 @@ "aten._fft_c2c.default": {}, "aten._fft_r2c.default": {}, "aten._flash_attention_backward.default": {}, - "aten._flash_attention_forward.default": {}, + "aten._flash_attention_forward.default": {"v2": ["block_table", "num_splits"]}, + "aten._flash_attention_forward_no_dropout_inplace.default": {"v2": ["num_splits"]}, "aten._flash_attention_forward.quantized": {}, "aten._fused_moving_avg_obs_fq_helper_functional.default": {}, "aten._fused_moving_avg_obs_fq_helper.default": {}, "aten._fused_rms_norm.default": {}, + "aten._grouped_mm.default": {}, "aten._histogramdd_from_bin_cts.default": {}, "aten._int_mm.out": {}, "aten._pdist_backward.default": {}, "aten._pdist_forward.default": {}, - "aten._scaled_dot_product_attention_math_for_mps.default": {}, + "aten._scaled_dot_product_attention_math_for_mps.default": {"v2": ["enable_gqa"]}, "aten._scaled_dot_product_cudnn_attention_backward.default": {}, "aten._scaled_dot_product_cudnn_attention.default": {}, "aten._scaled_dot_product_efficient_attention_backward.default": {}, @@ -150,6 +152,12 @@ "aten.randn.default": {}, "aten.randn.generator": {}, "aten.randperm.default": {}, + "aten.rand_like.default": {}, + "aten.rand_like.generator": {}, + "aten.randint_like.default": {}, + "aten.randint_like.low_dtype": {}, + "aten.randn_like.default": {}, + "aten.randn_like.generator": {}, "aten.repeat_interleave.Tensor": {}, "aten.replication_pad1d_backward.default": {}, "aten.replication_pad2d_backward.default": {}, diff --git a/torchgen/api/python.py b/torchgen/api/python.py index ca971e854b234..254d7c1ee9b43 100644 --- a/torchgen/api/python.py +++ b/torchgen/api/python.py @@ -288,7 +288,7 @@ def argument_str_pyi( name += "_" # pyi merges the _out and functional variants into the same signature, with an optional out arg - if name == "out" and type_str == "Tensor" and not deprecated: + if name == "out" and not deprecated: type_str = f"{type_str} | None".replace(" | None | None", " | None") # pyi deprecated signatures don't get defaults for their out arg @@ -975,14 +975,17 @@ def argument_type_str_pyi(t: Type) -> str: if str(t.elem) == "int": ret = "_int | _size" if t.size is not None else "_size" elif t.is_tensor_like(): - # TODO: this doesn't seem right... - # Tensor?[] currently translates to tuple[Tensor, ...] | list[Tensor] | None - # It should probably translate to tuple[Tensor | None, ...] | list[Tensor | None] - add_optional = True + # Tensor?[] translates to tuple[Tensor | None, ...] | list[Tensor | None] | None + # Tensor[] translates to tuple[Tensor, ...] | list[Tensor] + if isinstance(t.elem, OptionalType): + add_optional = True + elem_str = "Tensor | None" + else: + elem_str = "Tensor" ret = ( - "Tensor | tuple[Tensor, ...] | list[Tensor]" + f"Tensor | tuple[{elem_str}, ...] | list[{elem_str}]" if t.size is not None - else "tuple[Tensor, ...] | list[Tensor]" + else f"tuple[{elem_str}, ...] | list[{elem_str}]" ) elif str(t.elem) == "float": ret = "Sequence[_float]" diff --git a/torchgen/api/types/types_base.py b/torchgen/api/types/types_base.py index 08085fa0fa2bf..322ae1c39c1ed 100644 --- a/torchgen/api/types/types_base.py +++ b/torchgen/api/types/types_base.py @@ -18,7 +18,7 @@ from abc import ABC, abstractmethod from dataclasses import dataclass from enum import auto, Enum -from typing import TYPE_CHECKING, Union +from typing import TYPE_CHECKING if TYPE_CHECKING: @@ -35,7 +35,7 @@ class SpecialArgName(Enum): possibly_redundant_memory_format = auto() -ArgName = Union[str, SpecialArgName] +ArgName = str | SpecialArgName # This class shouldn't be created directly; instead, use/create one of the singletons below. diff --git a/torchgen/gen.py b/torchgen/gen.py index e8147a86d7ad5..10f727b9dfba0 100644 --- a/torchgen/gen.py +++ b/torchgen/gen.py @@ -2101,6 +2101,7 @@ def gen_headers( static_dispatch_idx: list[BackendIndex], selector: SelectiveBuilder, backend_indices: dict[DispatchKey, BackendIndex], + headeronly_fm: FileManager, core_fm: FileManager, cpu_fm: FileManager, device_fms: dict[str, FileManager], @@ -2227,7 +2228,7 @@ def gen_aten_interned_strings() -> dict[str, str]: def gen_tags_enum() -> dict[str, str]: return {"enum_of_valid_tags": (",\n".join(sorted(valid_tags)))} - core_fm.write("enum_tag.h", gen_tags_enum) + headeronly_fm.write("enum_tag.h", gen_tags_enum) def gen_source_files( @@ -2782,6 +2783,14 @@ def main() -> None: help="output directory for AOTInductor shim", default="torch/csrc/inductor/aoti_torch/generated", ) + parser.add_argument( + "--headeronly-install-dir", + "--headeronly_install_dir", + help="output directory for header-only generated files (e.g. enum_tag.h). " + "Defaults to `/core` when --install-dir is set, otherwise " + "`build/torch/headeronly/core`.", + default=None, + ) parser.add_argument( "--rocm", action="store_true", @@ -2979,12 +2988,23 @@ def main() -> None: aoti_install_dir = f"{options.aoti_install_dir}" Path(aoti_install_dir).mkdir(parents=True, exist_ok=True) + if options.headeronly_install_dir is not None: + headeronly_install_dir = options.headeronly_install_dir + elif options.install_dir is not None: + headeronly_install_dir = f"{options.install_dir}/core" + else: + headeronly_install_dir = "build/torch/headeronly/core" + Path(headeronly_install_dir).mkdir(parents=True, exist_ok=True) + core_fm = make_file_manager(options=options, install_dir=core_install_dir) cpu_fm = make_file_manager(options=options) cpu_vec_fm = make_file_manager(options=options) cuda_fm = make_file_manager(options=options) ops_fm = make_file_manager(options=options, install_dir=ops_install_dir) aoti_fm = make_file_manager(options=options, install_dir=aoti_install_dir) + headeronly_fm = make_file_manager( + options=options, install_dir=headeronly_install_dir + ) device_fms = {"cuda": cuda_fm} if options.xpu: device_fms["xpu"] = make_file_manager(options=options) @@ -3034,6 +3054,7 @@ def main() -> None: static_dispatch_idx=static_dispatch_idx, selector=selector, backend_indices=backend_indices, + headeronly_fm=headeronly_fm, core_fm=core_fm, cpu_fm=cpu_fm, device_fms=device_fms, diff --git a/torchgen/gen_aoti_c_shim.py b/torchgen/gen_aoti_c_shim.py index 3f626955108f6..dde6a6c8eda98 100644 --- a/torchgen/gen_aoti_c_shim.py +++ b/torchgen/gen_aoti_c_shim.py @@ -214,7 +214,16 @@ def gen_arguments( callsite_exprs: list[str] = [] for arg in flat_arguments: if arg.name in skipped_args: - callsite_exprs.append("std::nullopt") + # Pass the arg's schema default when available (e.g. "false" for + # a bool arg with default=False), so non-optional args with defaults + # can be versioned too. Fall back to std::nullopt for optional args + # with no default (matches historical behavior). + if arg.default is not None: + from torchgen.api.cpp import default_expr + + callsite_exprs.append(default_expr(arg.default, arg.type, symint=False)) + else: + callsite_exprs.append("std::nullopt") continue new_types, names, _, new_callsite_exprs = convert_arg_type_and_name( arg.type, arg.name, arg.is_write diff --git a/torchgen/gen_backend_stubs.py b/torchgen/gen_backend_stubs.py index 6580613a60f4b..efe63a80249eb 100644 --- a/torchgen/gen_backend_stubs.py +++ b/torchgen/gen_backend_stubs.py @@ -549,9 +549,13 @@ def gen_dispatcher_registrations( def run( source_yaml: str, output_dir: str, dry_run: bool, impl_path: str | None = None ) -> None: - # Assumes that this file lives at PYTORCH_ROOT/torchgen/gen_backend_stubs.py - pytorch_root = Path(__file__).absolute().parent.parent - template_dir = os.path.join(pytorch_root, "aten/src/ATen/templates") + # Assumes that this file lives at torchgen/gen_backend_stubs.py + root = Path(__file__).absolute().parent.parent + common_dir = os.path.join(root, "aten/src") # Assumes root is pytorch_root + if not os.path.exists(common_dir): # This file is out-of-tree. + common_dir = os.path.join(root, "torchgen/packaged") + + template_dir = os.path.join(common_dir, "ATen/templates") def make_file_manager(install_dir: str) -> FileManager: return FileManager( @@ -560,10 +564,8 @@ def make_file_manager(install_dir: str) -> FileManager: fm = make_file_manager(output_dir) - native_yaml_path = os.path.join( - pytorch_root, "aten/src/ATen/native/native_functions.yaml" - ) - tags_yaml_path = os.path.join(pytorch_root, "aten/src/ATen/native/tags.yaml") + native_yaml_path = os.path.join(common_dir, "ATen/native/native_functions.yaml") + tags_yaml_path = os.path.join(common_dir, "ATen/native/tags.yaml") parsed_yaml = parse_native_yaml(native_yaml_path, tags_yaml_path) native_functions, backend_indices = ( parsed_yaml.native_functions, diff --git a/torchgen/gen_functionalization_type.py b/torchgen/gen_functionalization_type.py index ffd0fc3d45fad..251ba64248a3c 100644 --- a/torchgen/gen_functionalization_type.py +++ b/torchgen/gen_functionalization_type.py @@ -37,6 +37,7 @@ NativeFunction, NativeFunctionsGroup, NativeFunctionsViewGroup, + OperatorName, Return, SchemaKind, SelfArgument, @@ -71,9 +72,21 @@ "resize_as_", # This function is used as for testing purposes only. "_fill_mem_eff_dropout_mask_", + # Inference-only op called behind a custom op graph break. + "_flash_attention_forward_no_dropout_inplace", ] ) +# Eager cumulative out variants compute in the out dtype when dtype is omitted. +# Functionalization normally lowers mutable ops through their functional variants, +# so these need to thread the out dtype explicitly to preserve eager semantics. +CUMULATIVE_OUT_OPS_PRESERVING_OUT_DTYPE = { + OperatorName.parse("cumsum.out"), + OperatorName.parse("cumprod.out"), + OperatorName.parse("cumsum.dimname_out"), + OperatorName.parse("cumprod.dimname_out"), +} + # This file contains codegen that relates to the functionalization pass. # It includes: # - gen_functionalization_definition @@ -605,6 +618,40 @@ def wrap_propagate_mutations_and_return( {returns_str}""" +def maybe_replace_cumulative_out_dtype_exprs( + f: NativeFunction, + functional_sig: DispatcherSignature, + functional_exprs: list[str], +) -> list[str]: + if ( + f.func.kind() != SchemaKind.out + or f.func.name not in CUMULATIVE_OUT_OPS_PRESERVING_OUT_DTYPE + ): + return functional_exprs + + if len(f.func.arguments.out) != 1: + raise AssertionError( + f"Expected a single out argument for cumulative out op: {f.func.name}" + ) + + dtype_arg_idx = next( + (i for i, arg in enumerate(functional_sig.arguments()) if arg.name == "dtype"), + None, + ) + if dtype_arg_idx is None: + raise AssertionError( + f"Expected dtype argument for cumulative out op: {f.func.name}" + ) + + adjusted_exprs = functional_exprs.copy() + dtype_expr = adjusted_exprs[dtype_arg_idx] + adjusted_exprs[dtype_arg_idx] = ( + f"{dtype_expr}.has_value() ? {dtype_expr} : " + f"std::optional({f.func.arguments.out[0].name}_.scalar_type())" + ) + return adjusted_exprs + + # Generates the Functionalization kernel for: # - mutation ops (inplace and out= ops) @with_native_function_and @@ -676,6 +723,9 @@ def emit_inplace_functionalization_body( e.expr for e in translate(unwrapped_args_ctx, functional_sig.arguments(), method=False) ] + functional_exprs = maybe_replace_cumulative_out_dtype_exprs( + f, functional_sig, functional_exprs + ) meta_conversion_str, meta_call_ctx = convert_to_meta_tensors(dispatcher_sig) # We don't want to run the inplace meta func for ops like .set_(), because: diff --git a/torchgen/static_runtime/generator.py b/torchgen/static_runtime/generator.py index e15b2514830ac..d8aba4d13bcde 100644 --- a/torchgen/static_runtime/generator.py +++ b/torchgen/static_runtime/generator.py @@ -270,7 +270,7 @@ def is_supported(g: NativeFunctionsGroup | NativeFunctionsViewGroup) -> bool: # the string, just test the dang thing directly if "at::Tensor" != cpp.returns_type(func.returns, symint=False).cpp_type(): # Returns a non-Tensor value. - logger.info("NON-TENSOR RET TYPE: %s", str(func)) + logger.info("NON-TENSOR RET TYPE: %s", func) return False return True diff --git a/version.txt b/version.txt index f925b7d0ce58a..d8b698973a491 100644 --- a/version.txt +++ b/version.txt @@ -1 +1 @@ -2.11.0a0 +2.12.0