Detect write-after-read hazards during stream capture with pinned memory.

This issue was first called out by Natalia here: https://github.com/pytorch/pytorch/pull/166765#issuecomment-3487509660

To be specific, this detects what I am calling
CPU-write-after-GPU-read hazards.

Essentially, during stream capture to a cuda graph, CPU code is eager,
but GPU code is lazy. No GPU kernels will launch until the cuda graph
is replayed. If a CUDA memcpy or kernel node reads from pinned host
memory (which is announce to us via record_event()), we want to throw
an error if the CPU ever writes to that pinned host memory any time
afterward.

In practice, we cannot check if an arbitrary CUDA kernel reads from a
particular host memory buffer, but we can check if a memcpy node
does. Fortunately, the only usage of record_event() in pytorch always
comes after a memcpy node, so I am confident that the current solution
will catch >99% of problems.

The core algorithm is to compute a hash of any read-from host buffer
at the time that record_event() is called. If that host buffer's hash
code is not the same when stream capture ends, we throw an error.

I use SHA-1 to hash because it was already available in pytorch's C++
code. I did not feel confident in my own ability to write a
low-collision hash function and also did not want to add a new
dependency.

This still requires test cases, but the algorithm itself is there.

This also requires thinking about the appropriate behaviors if a
memory pool spans multiple graph captures. Right now, I delete all
sections_read_under_stream_capture_ entries at the end of stream
capture, which I am pretty sure is the appropriate behavior, but am
not sure yet.


ghstack-source-id: 341ca4f6e4f839692269c4af9ebdaaec84962efc
Pull-Request: https://github.com/pytorch/pytorch/pull/167508

.bc-linter.yml

@@ -13,4 +13,3 @@ exclude:
   - "**/benchmarks/**"
   - "**/test_*.py"
   - "**/*_test.py"
-  - "tools/**"

.ci/docker/almalinux/Dockerfile

@@ -7,13 +7,13 @@ ENV LC_ALL en_US.UTF-8
 ENV LANG en_US.UTF-8
 ENV LANGUAGE en_US.UTF-8
 
-ARG DEVTOOLSET_VERSION=13
+ARG DEVTOOLSET_VERSION=11
 
 RUN yum -y update
 RUN yum -y install epel-release
 # install glibc-langpack-en make sure en_US.UTF-8 locale is available
 RUN yum -y install glibc-langpack-en
-RUN 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 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}-toolchain
 # 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 '*'
 ENV PATH=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/bin:$PATH
@@ -41,7 +41,6 @@ RUN bash ./install_conda.sh && rm install_conda.sh
 # Install CUDA
 FROM base as cuda
 ARG CUDA_VERSION=12.6
-ARG DEVTOOLSET_VERSION=13
 RUN rm -rf /usr/local/cuda-*
 ADD ./common/install_cuda.sh install_cuda.sh
 COPY ./common/install_nccl.sh install_nccl.sh
@@ -51,8 +50,7 @@ ENV CUDA_HOME=/usr/local/cuda-${CUDA_VERSION}
 # Preserve CUDA_VERSION for the builds
 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
-
+ENV PATH=/usr/local/cuda-${CUDA_VERSION}/bin:$PATH
 
 FROM cuda as cuda12.6
 RUN bash ./install_cuda.sh 12.6
@@ -70,22 +68,8 @@ FROM cuda as cuda13.0
 RUN bash ./install_cuda.sh 13.0
 ENV DESIRED_CUDA=13.0
 
-FROM ${ROCM_IMAGE} as rocm_base
-ARG DEVTOOLSET_VERSION=13
-ENV LC_ALL en_US.UTF-8
-ENV LANG en_US.UTF-8
-ENV LANGUAGE en_US.UTF-8
-# Install devtoolset on ROCm base image
-RUN yum -y update && \
-    yum -y install epel-release && \
-    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 '*'
-ENV PATH=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/bin:$PATH
-
-FROM rocm_base as rocm
+FROM ${ROCM_IMAGE} as rocm
 ARG PYTORCH_ROCM_ARCH
-ARG DEVTOOLSET_VERSION=13
 ENV PYTORCH_ROCM_ARCH ${PYTORCH_ROCM_ARCH}
 ADD ./common/install_mkl.sh install_mkl.sh
 RUN bash ./install_mkl.sh && rm install_mkl.sh
@@ -104,7 +88,6 @@ COPY --from=cuda13.0  /usr/local/cuda-13.0 /usr/local/cuda-13.0
 
 # 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

.ci/docker/almalinux/build.sh

@@ -63,7 +63,7 @@ docker build \
   --target final \
   --progress plain \
   --build-arg "BASE_TARGET=${BASE_TARGET}" \
-  --build-arg "DEVTOOLSET_VERSION=13" \
+  --build-arg "DEVTOOLSET_VERSION=11" \
   ${EXTRA_BUILD_ARGS} \
   -t ${tmp_tag} \
   $@ \

.ci/docker/build.sh

@@ -261,9 +261,9 @@ case "$tag" in
     PYTHON_VERSION=3.10
     CUDA_VERSION=12.8.1
     ;;
-  pytorch-linux-jammy-aarch64-py3.10-gcc13)
+  pytorch-linux-jammy-aarch64-py3.10-gcc11)
     ANACONDA_PYTHON_VERSION=3.10
-    GCC_VERSION=13
+    GCC_VERSION=11
     ACL=yes
     VISION=yes
     OPENBLAS=yes
@@ -271,19 +271,9 @@ case "$tag" in
     # from pytorch/llvm:9.0.1 is x86 specific
     SKIP_LLVM_SRC_BUILD_INSTALL=yes
     ;;
-  pytorch-linux-jammy-aarch64-py3.10-clang21)
+  pytorch-linux-jammy-aarch64-py3.10-gcc11-inductor-benchmarks)
     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
+    GCC_VERSION=11
     ACL=yes
     VISION=yes
     OPENBLAS=yes

.ci/docker/ci_commit_pins/triton.txt

@@ -1 +1 @@
-bfeb066872bc1e8b2d2bc0a3b295b99dd77206e7
+7416ffcb92cdbe98d9f97e4e6f95247e46dfc9fd

.ci/docker/common/install_acl.sh

@@ -3,7 +3,7 @@
 
 set -eux
 
-ACL_VERSION=${ACL_VERSION:-"v52.6.0"}
+ACL_VERSION=${ACL_VERSION:-"v25.02"}
 ACL_INSTALL_DIR="/acl"
 
 # Clone ACL

.ci/docker/common/install_clang.sh

@@ -8,8 +8,8 @@ if [ -n "$CLANG_VERSION" ]; then
     # work around ubuntu apt-get conflicts
     sudo apt-get -y -f install
     wget --no-check-certificate -O - https://apt.llvm.org/llvm-snapshot.gpg.key | sudo apt-key add -
-    if [[ $CLANG_VERSION -ge 18 ]]; then
-      apt-add-repository "deb http://apt.llvm.org/jammy/ llvm-toolchain-jammy-${CLANG_VERSION} main"
+    if [[ $CLANG_VERSION == 18 ]]; then
+      apt-add-repository "deb http://apt.llvm.org/jammy/ llvm-toolchain-jammy-18 main"
     fi
   fi
 

.ci/docker/common/install_gcc.sh

@@ -7,11 +7,11 @@ if [ -n "$GCC_VERSION" ]; then
   # Need the official toolchain repo to get alternate packages
   add-apt-repository ppa:ubuntu-toolchain-r/test
   apt-get update
-  apt-get install -y g++-$GCC_VERSION gfortran-$GCC_VERSION
+  apt-get install -y g++-$GCC_VERSION
   update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-"$GCC_VERSION" 50
   update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-"$GCC_VERSION" 50
   update-alternatives --install /usr/bin/gcov gcov /usr/bin/gcov-"$GCC_VERSION" 50
-  update-alternatives --install /usr/bin/gfortran gfortran /usr/bin/gfortran-"$GCC_VERSION" 50
+
 
   # Cleanup package manager
   apt-get autoclean && apt-get clean

.ci/docker/common/install_libgomp.sh

@@ -1,56 +0,0 @@
-#!/bin/bash
-# Script used only in CD pipeline
-
-set -ex
-
-# install dependencies
-dnf -y install gmp-devel libmpc-devel texinfo flex bison
-
-cd /usr/local/src
-# fetch source for gcc 13
-git clone --depth 1 --single-branch -b releases/gcc-13.3.0 https://github.com/gcc-mirror/gcc.git gcc-13.3.0
-
-mkdir -p gcc-13.3.0/build-gomp
-cd gcc-13.3.0/build-gomp
-
-# configure gcc build
-# I got these flags by:
-# 1. downloading the source rpm for gcc-11 on AlmaLinux 8 container
-#    dnf install -y dnf-plugins-core rpmdevtools
-#   dnf download --source libgomp
-# 2. extracting the gcc.spec from the source.
-#    rpmdev-extract gcc-xx.src.rpm
-# 3. extracting optflags and ld_flags from gcc.spec:
-#    rpm --eval '%{optflags}'
-#    rpm --eval '%{build_ldflags}'
-#
-# I had to remove the following flags because they didn't compile for this version of libgomp:
-#   -Werror=format-security
-#   -specs=/usr/lib/rpm/redhat/redhat-hardened-cc1
-#   -specs=/usr/lib/rpm/redhat/redhat-annobin-cc1
-#
-# I added -march=armv8-a -mtune=generic to make them explicit. I don't think they're strictly needed.
-
-OPT_FLAGS='-O2 -march=armv8-a -mtune=generic'\
-' -fexceptions -g -grecord-gcc-switches -pipe -Wall'\
-' -Wp,-D_FORTIFY_SOURCE=2 -Wp,-D_GLIBCXX_ASSERTIONS'\
-' -fstack-protector-strong -fasynchronous-unwind-tables'\
-' -fstack-clash-protection'
-
-LDFLAGS='-Wl,-z,relro -Wl,--as-needed -Wl,-z,now'
-
-CFLAGS="$OPT_FLAGS" \
-CXXFLAGS="$OPT_FLAGS" \
-LDFLAGS="$LDFLAGS" \
-../configure \
-  --prefix=/usr \
-  --libdir=/usr/lib64 \
-  --enable-languages=c,c++ \
-  --disable-multilib \
-  --disable-bootstrap \
-  --enable-libgomp
-
-# only build libgomp
-make -j$(nproc) all-target-libgomp
-
-make install-target-libgomp

.ci/docker/common/install_openblas.sh

@@ -10,7 +10,6 @@ git clone https://github.com/OpenMathLib/OpenBLAS.git -b "${OPENBLAS_VERSION}" -
 
 OPENBLAS_CHECKOUT_DIR="OpenBLAS"
 OPENBLAS_BUILD_FLAGS="
-CC=gcc
 NUM_THREADS=128
 USE_OPENMP=1
 NO_SHARED=0

.ci/docker/common/install_rocm_magma.sh

@@ -12,8 +12,8 @@ function do_install() {
 
     rocm_version_nodot=${rocm_version//./}
 
-    # post merge of https://github.com/icl-utk-edu/magma/pull/65
-    MAGMA_VERSION=c0792ae825fb36872784892ea643dd6f3456bc5f
+    # https://github.com/icl-utk-edu/magma/pull/65
+    MAGMA_VERSION=d6e4117bc88e73f06d26c6c2e14f064e8fc3d1ec
     magma_archive="magma-rocm${rocm_version_nodot}-${MAGMA_VERSION}-1.tar.bz2"
 
     rocm_dir="/opt/rocm"

.ci/docker/manywheel/Dockerfile_2_28

@@ -149,7 +149,7 @@ FROM cpu_final as rocm_final
 ARG ROCM_VERSION=6.0
 ARG PYTORCH_ROCM_ARCH
 ENV PYTORCH_ROCM_ARCH ${PYTORCH_ROCM_ARCH}
-ARG DEVTOOLSET_VERSION=13
+ARG DEVTOOLSET_VERSION=11
 ENV LDFLAGS="-Wl,-rpath=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/lib64 -Wl,-rpath=/opt/rh/gcc-toolset-${DEVTOOLSET_VERSION}/root/usr/lib"
 # Somewhere in ROCm stack, we still use non-existing /opt/rocm/hip path,
 # below workaround helps avoid error

.ci/docker/manywheel/Dockerfile_2_28_aarch64

@@ -50,10 +50,6 @@ RUN rm install_ninja.sh
 ENV PATH=/opt/rh/gcc-toolset-${GCCTOOLSET_VERSION}/root/usr/bin:$PATH
 ENV LD_LIBRARY_PATH=/opt/rh/gcc-toolset-${GCCTOOLSET_VERSION}/root/usr/lib64:/opt/rh/gcc-toolset-${GCCTOOLSET_VERSION}/root/usr/lib:$LD_LIBRARY_PATH
 
-# Build a newer version of libgomp than that supported in in Almalinux 8.
-COPY ./common/install_libgomp.sh install_libgomp.sh
-RUN bash ./install_libgomp.sh && rm install_libgomp.sh
-
 # git236+ would refuse to run git commands in repos owned by other users
 # Which causes version check to fail, as pytorch repo is bind-mounted into the image
 # Override this behaviour by treating every folder as safe

.ci/docker/manywheel/build.sh

@@ -97,7 +97,7 @@ case ${image} in
     manylinux2_28-builder:xpu)
         TARGET=xpu_final
         GPU_IMAGE=amd64/almalinux:8
-        DOCKER_GPU_BUILD_ARG=" --build-arg DEVTOOLSET_VERSION=13"
+        DOCKER_GPU_BUILD_ARG=" --build-arg DEVTOOLSET_VERSION=11"
         MANY_LINUX_VERSION="2_28"
         ;;
     *)

.ci/docker/requirements-docs.txt

@@ -1,11 +1,15 @@
-sphinx==7.2.6
+sphinx==5.3.0
 #Description: This is used to generate PyTorch docs
-#Pinned versions: 7.2.6
+#Pinned versions: 5.3.0
 
-pytorch_sphinx_theme2==0.2.0
-#Description: This is needed to generate PyTorch docs
-#Pinned versions: 0.2.0
+standard-imghdr==3.13.0; python_version >= "3.13"
+#Description: This is needed by Sphinx, so it needs to be added here.
+# The reasons are as follows:
+# 1) This module has been removed from the Python standard library since Python 3.13(https://peps.python.org/pep-0594/#imghdr);
+# 2) The current version of Sphinx (5.3.0) is not compatible with Python 3.13.
+# Once Sphinx is upgraded to a version compatible with Python 3.13 or later, we can remove this dependency.
 
+-e git+https://github.com/pytorch/pytorch_sphinx_theme.git@71e55749be14ceb56e7f8211a9fb649866b87ad4#egg=pytorch_sphinx_theme2
 # 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.
@@ -32,17 +36,17 @@ tensorboard==2.18.0 ; python_version >= "3.13"
 #Description: This is used to generate PyTorch docs
 #Pinned versions: 2.13.0
 
-breathe==4.36.0
+breathe==4.34.0
 #Description: This is used to generate PyTorch C++ docs
-#Pinned versions: 4.36.0
+#Pinned versions: 4.34.0
 
-exhale==0.3.7
+exhale==0.2.3
 #Description: This is used to generate PyTorch C++ docs
-#Pinned versions: 0.3.7
+#Pinned versions: 0.2.3
 
-docutils==0.20
+docutils==0.16
 #Description: This is used to generate PyTorch C++ docs
-#Pinned versions: 0.20
+#Pinned versions: 0.16
 
 bs4==0.0.1
 #Description: This is used to generate PyTorch C++ docs
@@ -52,13 +56,13 @@ IPython==8.12.0
 #Description: This is used to generate PyTorch functorch docs
 #Pinned versions: 8.12.0
 
-myst-nb==1.3.0
+myst-nb==0.17.2
 #Description: This is used to generate PyTorch functorch and torch.compile docs.
-#Pinned versions: 1.3.0
+#Pinned versions: 0.17.2
 
 # The following are required to build torch.distributed.elastic.rendezvous.etcd* docs
 python-etcd==0.4.5
 sphinx-copybutton==0.5.0
-sphinx-design==0.6.1
+sphinx-design==0.4.0
 sphinxcontrib-mermaid==1.0.0
-myst-parser==4.0.1
+myst-parser==0.18.1

.ci/docker/triton_version.txt

@@ -1 +1 @@
-3.5.1
+3.5.0

.ci/magma-rocm/Makefile

@@ -1,7 +1,7 @@
 SHELL=/usr/bin/env bash
 
 DOCKER_CMD ?= docker
-DESIRED_ROCM ?= 7.1
+DESIRED_ROCM ?= 7.0
 DESIRED_ROCM_SHORT = $(subst .,,$(DESIRED_ROCM))
 PACKAGE_NAME = magma-rocm
 # inherit this from underlying docker image, do not pass this env var to docker
@@ -16,7 +16,6 @@ DOCKER_RUN = set -eou pipefail; ${DOCKER_CMD} run --rm -i \
 	magma-rocm/build_magma.sh
 
 .PHONY: all
-all: magma-rocm71
 all: magma-rocm70
 all: magma-rocm64
 
@@ -25,11 +24,6 @@ clean:
 	$(RM) -r magma-*
 	$(RM) -r output
 
-.PHONY: magma-rocm71
-magma-rocm71: DESIRED_ROCM := 7.1
-magma-rocm71:
-	$(DOCKER_RUN)
-
 .PHONY: magma-rocm70
 magma-rocm70: DESIRED_ROCM := 7.0
 magma-rocm70:

.ci/pytorch/python_doc_push_script.sh

@@ -89,41 +89,23 @@ if [ "$is_main_doc" = true ]; then
 
   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 |
+  # Count the number of lines in the file and turn that number into a variable
+  # $lines. The `cut -f1 ...` is to only parse the number, not the filename
+  # Skip the report header by subtracting 2: the header will be output even if
+  # there are no undocumented items.
   #
   # 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
+  lines=$(wc -l build/coverage/python.txt 2>/dev/null |cut -f1 -d' ')
+  undocumented=$((lines - 2))
+  if [ $undocumented -lt 0 ]; 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
-    fi
-    echo ""
+  elif [ $undocumented -gt 0 ]; then
+    echo undocumented objects found:
+    cat build/coverage/python.txt
     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
+    echo "You can reproduce locally by running 'cd docs && make coverage && cat build/coverage/python.txt'"
     exit 1
   fi
 else

.ci/pytorch/test.sh

@@ -337,7 +337,7 @@ test_python() {
 
 test_python_smoke() {
   # Smoke tests for H100/B200
-  time python test/run_test.py --include test_matmul_cuda test_scaled_matmul_cuda inductor/test_fp8 inductor/test_max_autotune inductor/test_cutedsl_grouped_mm $PYTHON_TEST_EXTRA_OPTION --upload-artifacts-while-running
+  time python test/run_test.py --include test_matmul_cuda test_scaled_matmul_cuda inductor/test_fp8 inductor/test_max_autotune $PYTHON_TEST_EXTRA_OPTION --upload-artifacts-while-running
   assert_git_not_dirty
 }
 
@@ -1653,7 +1653,7 @@ test_operator_microbenchmark() {
 
   cd "${TEST_DIR}"/benchmarks/operator_benchmark
 
-  for OP_BENCHMARK_TESTS in matmul mm addmm bmm conv; do
+  for OP_BENCHMARK_TESTS in matmul mm addmm bmm; do
     $TASKSET python -m pt.${OP_BENCHMARK_TESTS}_test --tag-filter long \
       --output-json-for-dashboard "${TEST_REPORTS_DIR}/operator_microbenchmark_${OP_BENCHMARK_TESTS}_compile.json" \
       --benchmark-name "PyTorch operator microbenchmark" --use-compile

.ci/pytorch/win-test-helpers/arm64/build_pytorch.ps1

@@ -70,7 +70,7 @@ sccache --zero-stats
 sccache --show-stats
 
 # Build the wheel
-python -m build --wheel --no-isolation
+python -m build --wheel --no-build-isolation
 if ($LASTEXITCODE -ne 0) { exit 1 }
 
 # Install the wheel locally

.clang-tidy

@@ -60,11 +60,9 @@ performance-*,
 readability-container-size-empty,
 readability-delete-null-pointer,
 readability-duplicate-include,
-readability-named-parameter,
 readability-misplaced-array-index,
 readability-redundant*,
 readability-simplify-subscript-expr,
-readability-static-definition-in-anonymous-namespace
 readability-string-compare,
 -readability-redundant-access-specifiers,
 -readability-redundant-control-flow,

.claude/skills/add-uint-support/SKILL.md

@@ -1,319 +0,0 @@
----
-name: add-uint-support
-description: Add unsigned integer (uint) type support to PyTorch operators by updating AT_DISPATCH macros. Use when adding support for uint16, uint32, uint64 types to operators, kernels, or when user mentions enabling unsigned types, barebones unsigned types, or uint support.
----
-
-# Add Unsigned Integer (uint) Support to Operators
-
-This skill helps add support for unsigned integer types (uint16, uint32, uint64) to PyTorch operators by updating their AT_DISPATCH macros.
-
-## When to use this skill
-
-Use this skill when:
-- Adding uint16, uint32, or uint64 support to an operator
-- User mentions "unsigned types", "uint support", "barebones unsigned types"
-- Enabling support for kUInt16, kUInt32, kUInt64 in kernels
-- Working with operator implementations that need expanded type coverage
-
-## Quick reference
-
-**Add unsigned types to existing dispatch:**
-```cpp
-// Before
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES));
-
-// After (method 1: add unsigned types explicitly)
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES));
-
-// After (method 2: use V2 integral types if AT_INTEGRAL_TYPES present)
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_INTEGRAL_TYPES_V2), AT_EXPAND(AT_FLOATING_TYPES));
-```
-
-## Type group reference
-
-**Unsigned type groups:**
-- `AT_BAREBONES_UNSIGNED_TYPES`: kUInt16, kUInt32, kUInt64
-- `AT_INTEGRAL_TYPES_V2`: AT_INTEGRAL_TYPES + AT_BAREBONES_UNSIGNED_TYPES
-
-**Relationship:**
-```cpp
-AT_INTEGRAL_TYPES          // kByte, kChar, kInt, kLong, kShort
-AT_BAREBONES_UNSIGNED_TYPES  // kUInt16, kUInt32, kUInt64
-AT_INTEGRAL_TYPES_V2       // INTEGRAL_TYPES + BAREBONES_UNSIGNED_TYPES
-```
-
-## Instructions
-
-### Step 1: Determine if conversion to V2 is needed
-
-Check if the file uses AT_DISPATCH_V2:
-
-**If using old AT_DISPATCH:**
-- First convert to AT_DISPATCH_V2 using the at-dispatch-v2 skill
-- Then proceed with adding uint support
-
-**If already using AT_DISPATCH_V2:**
-- Proceed directly to Step 2
-
-### Step 2: Analyze the current dispatch macro
-
-Identify what type groups are currently in use:
-
-```cpp
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  // body
-}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBFloat16);
-    ^^^^^^^^^^^^^^^^^^^^^^^^^
-    Current type coverage
-```
-
-Common patterns:
-- `AT_EXPAND(AT_ALL_TYPES)` → includes AT_INTEGRAL_TYPES + AT_FLOATING_TYPES
-- `AT_EXPAND(AT_INTEGRAL_TYPES)` → signed integers only
-- `AT_EXPAND(AT_FLOATING_TYPES)` → floating point types
-
-### Step 3: Choose the uint addition method
-
-Two approaches:
-
-**Method 1: Add AT_BAREBONES_UNSIGNED_TYPES explicitly**
-- Use when: You want to be explicit about adding uint support
-- Add `AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES)` to the type list
-
-**Method 2: Substitute AT_INTEGRAL_TYPES with AT_INTEGRAL_TYPES_V2**
-- Use when: The dispatch already uses `AT_EXPAND(AT_INTEGRAL_TYPES)`
-- More concise: replaces one type group with its superset
-- Only applicable if AT_INTEGRAL_TYPES is present
-
-### Step 4: Apply the transformation
-
-**Method 1 example:**
-```cpp
-// Before
-AT_DISPATCH_V2(
-    dtype,
-    "min_values_cuda",
-    AT_WRAP([&]() {
-      kernel_impl<scalar_t>(iter);
-    }),
-    AT_EXPAND(AT_ALL_TYPES),
-    kBFloat16, kHalf, kBool
-);
-
-// After (add unsigned types)
-AT_DISPATCH_V2(
-    dtype,
-    "min_values_cuda",
-    AT_WRAP([&]() {
-      kernel_impl<scalar_t>(iter);
-    }),
-    AT_EXPAND(AT_ALL_TYPES),
-    AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES),
-    kBFloat16, kHalf, kBool
-);
-```
-
-**Method 2 example:**
-```cpp
-// Before
-AT_DISPATCH_V2(
-    dtype,
-    "integral_op",
-    AT_WRAP([&]() {
-      kernel<scalar_t>();
-    }),
-    AT_EXPAND(AT_INTEGRAL_TYPES)
-);
-
-// After (substitute with V2)
-AT_DISPATCH_V2(
-    dtype,
-    "integral_op",
-    AT_WRAP([&]() {
-      kernel<scalar_t>();
-    }),
-    AT_EXPAND(AT_INTEGRAL_TYPES_V2)
-);
-```
-
-### Step 5: Handle AT_ALL_TYPES vs individual type groups
-
-If the dispatch uses `AT_EXPAND(AT_ALL_TYPES)`:
-- `AT_ALL_TYPES` = `AT_INTEGRAL_TYPES` + `AT_FLOATING_TYPES`
-- To add uint: add `AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES)` to the list
-
-If the dispatch separately lists INTEGRAL and FLOATING:
-```cpp
-// Before
-AT_EXPAND(AT_INTEGRAL_TYPES), AT_EXPAND(AT_FLOATING_TYPES)
-
-// After (Method 2 preferred)
-AT_EXPAND(AT_INTEGRAL_TYPES_V2), AT_EXPAND(AT_FLOATING_TYPES)
-```
-
-### Step 6: Verify all dispatch sites
-
-Check the file for ALL dispatch macros that need uint support:
-- Some operators have multiple dispatch sites (CPU, CUDA, different functions)
-- Apply the transformation consistently across all sites
-- Ensure each gets the same type coverage updates
-
-### Step 7: Validate the changes
-
-Check that:
-- [ ] AT_DISPATCH_V2 format is used (not old AT_DISPATCH)
-- [ ] Unsigned types are added via one of the two methods
-- [ ] All relevant dispatch sites in the file are updated
-- [ ] Type groups use `AT_EXPAND()`
-- [ ] Arguments are properly formatted and comma-separated
-
-## Common patterns
-
-### Pattern 1: AT_ALL_TYPES + extras
-
-```cpp
-// Before
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBFloat16);
-
-// After
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES), kHalf, kBFloat16);
-```
-
-### Pattern 2: Separate INTEGRAL + FLOATING
-
-```cpp
-// Before
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_INTEGRAL_TYPES), AT_EXPAND(AT_FLOATING_TYPES));
-
-// After
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_INTEGRAL_TYPES_V2), AT_EXPAND(AT_FLOATING_TYPES));
-```
-
-### Pattern 3: Old dispatch needs conversion first
-
-```cpp
-// Before (needs v2 conversion first)
-AT_DISPATCH_ALL_TYPES_AND2(kHalf, kBFloat16, dtype, "op", [&]() {
-  kernel<scalar_t>();
-});
-
-// After v2 conversion
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBFloat16);
-
-// After adding uint support
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES), kHalf, kBFloat16);
-```
-
-## Multiple dispatch sites example
-
-For a file with multiple functions:
-
-```cpp
-void min_values_kernel_cuda(TensorIterator& iter) {
-  AT_DISPATCH_V2(iter.dtype(), "min_values_cuda", AT_WRAP([&]() {
-    impl<scalar_t>(iter);
-  }), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES), kBFloat16, kHalf);
-  //                           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-  //                           Added uint support
-}
-
-void min_launch_kernel(TensorIterator &iter) {
-  AT_DISPATCH_V2(iter.input_dtype(), "min_cuda", AT_WRAP([&]() {
-    gpu_reduce_kernel<scalar_t>(iter);
-  }), AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES), kBFloat16, kHalf);
-  //                           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-  //                           Added uint support here too
-}
-```
-
-## Decision tree
-
-Use this decision tree to determine the approach:
-
-```
-Is the file using AT_DISPATCH_V2?
-├─ No → Use at-dispatch-v2 skill first, then continue
-└─ Yes
-   └─ Does it use AT_EXPAND(AT_INTEGRAL_TYPES)?
-      ├─ Yes → Replace with AT_EXPAND(AT_INTEGRAL_TYPES_V2)
-      └─ No → Add AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES) to type list
-```
-
-## Edge cases
-
-### Case 1: Dispatch with only floating types
-
-If the operator only supports floating point types, don't add uint support:
-
-```cpp
-// Leave as-is - floating point only operator
-AT_DISPATCH_V2(dtype, "float_op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_FLOATING_TYPES), kHalf);
-```
-
-### Case 2: Complex types present
-
-Unsigned types work alongside complex types:
-
-```cpp
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES),
-    AT_EXPAND(AT_BAREBONES_UNSIGNED_TYPES),
-    AT_EXPAND(AT_COMPLEX_TYPES),
-    kHalf, kBFloat16);
-```
-
-### Case 3: Already has uint support
-
-Check if uint types are already present:
-- If `AT_INTEGRAL_TYPES_V2` is used → already has uint support
-- If `AT_BAREBONES_UNSIGNED_TYPES` is already in list → already has uint support
-- Skip the file if uint support is already present
-
-## Workflow
-
-When asked to add uint support:
-
-1. Read the target file
-2. Check if using AT_DISPATCH_V2:
-   - If not → use at-dispatch-v2 skill first
-3. Identify all dispatch macro sites
-4. For each dispatch:
-   - Analyze current type groups
-   - Choose method (add BAREBONES_UNSIGNED or upgrade to V2)
-   - Apply transformation with Edit tool
-5. Show the user the changes
-6. Explain what was modified
-
-## Important notes
-
-- Always check if v2 conversion is needed first
-- Apply changes consistently across all dispatch sites in the file
-- Method 2 (AT_INTEGRAL_TYPES_V2) is cleaner when applicable
-- Method 1 (explicit AT_BAREBONES_UNSIGNED_TYPES) is more explicit
-- Unsigned types are: kUInt16, kUInt32, kUInt64 (not kByte which is uint8)
-- Some operators may not semantically support unsigned types - use judgment
-
-## Testing
-
-After adding uint support, the operator should accept uint16, uint32, and uint64 tensors. The user is responsible for functional testing.

.claude/skills/at-dispatch-v2/SKILL.md

@@ -1,305 +0,0 @@
----
-name: at-dispatch-v2
-description: Convert PyTorch AT_DISPATCH macros to AT_DISPATCH_V2 format in ATen C++ code. Use when porting AT_DISPATCH_ALL_TYPES_AND*, AT_DISPATCH_FLOATING_TYPES*, or other dispatch macros to the new v2 API. For ATen kernel files, CUDA kernels, and native operator implementations.
----
-
-# AT_DISPATCH to AT_DISPATCH_V2 Converter
-
-This skill helps convert PyTorch's legacy AT_DISPATCH macros to the new AT_DISPATCH_V2 format, as defined in `aten/src/ATen/Dispatch_v2.h`.
-
-## When to use this skill
-
-Use this skill when:
-- Converting AT_DISPATCH_* macros to AT_DISPATCH_V2
-- Porting ATen kernels to use the new dispatch API
-- Working with files in `aten/src/ATen/native/` that use dispatch macros
-- User mentions "AT_DISPATCH", "dispatch v2", "Dispatch_v2.h", or macro conversion
-
-## Quick reference
-
-**Old format:**
-```cpp
-AT_DISPATCH_ALL_TYPES_AND3(kBFloat16, kHalf, kBool, dtype, "kernel_name", [&]() {
-  // lambda body
-});
-```
-
-**New format:**
-```cpp
-AT_DISPATCH_V2(dtype, "kernel_name", AT_WRAP([&]() {
-  // lambda body
-}), AT_EXPAND(AT_ALL_TYPES), kBFloat16, kHalf, kBool);
-```
-
-## Key transformations
-
-1. **Reorder arguments**: `scalar_type` and `name` come first, then lambda, then types
-2. **Wrap the lambda**: Use `AT_WRAP(lambda)` to handle internal commas
-3. **Expand type groups**: Use `AT_EXPAND(AT_ALL_TYPES)` instead of implicit expansion
-4. **List individual types**: Add extra types (kHalf, kBFloat16, etc.) after expanded groups
-5. **Add include**: `#include <ATen/Dispatch_v2.h>` near other Dispatch includes
-
-## Instructions
-
-### Step 1: Add the Dispatch_v2.h include
-
-Add the v2 header near the existing `#include <ATen/Dispatch.h>`:
-
-```cpp
-#include <ATen/Dispatch.h>
-#include <ATen/Dispatch_v2.h>
-```
-
-Keep the old Dispatch.h include for now (other code may still need it).
-
-### Step 2: Identify the old dispatch pattern
-
-Common patterns to convert:
-
-- `AT_DISPATCH_ALL_TYPES_AND{2,3,4}(type1, type2, ..., scalar_type, name, lambda)`
-- `AT_DISPATCH_FLOATING_TYPES_AND{2,3}(type1, type2, ..., scalar_type, name, lambda)`
-- `AT_DISPATCH_ALL_TYPES_AND_COMPLEX_AND{2,3}(type1, ..., scalar_type, name, lambda)`
-- `AT_DISPATCH_FLOATING_AND_COMPLEX_TYPES_AND{2,3}(type1, ..., scalar_type, name, lambda)`
-
-### Step 3: Map the old macro to type groups
-
-Identify which type group macro corresponds to the base types:
-
-| Old macro base | AT_DISPATCH_V2 type group |
-|----------------|---------------------------|
-| `ALL_TYPES` | `AT_EXPAND(AT_ALL_TYPES)` |
-| `FLOATING_TYPES` | `AT_EXPAND(AT_FLOATING_TYPES)` |
-| `INTEGRAL_TYPES` | `AT_EXPAND(AT_INTEGRAL_TYPES)` |
-| `COMPLEX_TYPES` | `AT_EXPAND(AT_COMPLEX_TYPES)` |
-| `ALL_TYPES_AND_COMPLEX` | `AT_EXPAND(AT_ALL_TYPES_AND_COMPLEX)` |
-
-For combined patterns, use multiple `AT_EXPAND()` entries:
-```cpp
-// Old: AT_DISPATCH_ALL_TYPES_AND_COMPLEX_AND2(...)
-// New: AT_EXPAND(AT_ALL_TYPES), AT_EXPAND(AT_COMPLEX_TYPES), type1, type2
-```
-
-### Step 4: Extract the individual types
-
-From `AT_DISPATCH_*_AND2(type1, type2, ...)` or `AT_DISPATCH_*_AND3(type1, type2, type3, ...)`, extract the individual types (type1, type2, etc.).
-
-These become the trailing arguments after the type group:
-```cpp
-AT_DISPATCH_V2(..., AT_EXPAND(AT_ALL_TYPES), kBFloat16, kHalf, kBool)
-                                             ^^^^^^^^^^^^^^^^^^^^^^^^
-                                             Individual types from AND3
-```
-
-### Step 5: Transform to AT_DISPATCH_V2
-
-Apply the transformation:
-
-**Pattern:**
-```cpp
-AT_DISPATCH_V2(
-  scalar_type,           // 1st: The dtype expression
-  "name",                // 2nd: The debug string
-  AT_WRAP(lambda),       // 3rd: The lambda wrapped in AT_WRAP
-  type_groups,           // 4th+: Type groups with AT_EXPAND()
-  individual_types       // Last: Individual types
-)
-```
-
-**Example transformation:**
-```cpp
-// BEFORE
-AT_DISPATCH_ALL_TYPES_AND3(
-    kBFloat16, kHalf, kBool,
-    iter.dtype(),
-    "min_values_cuda",
-    [&]() {
-      min_values_kernel_cuda_impl<scalar_t>(iter);
-    }
-);
-
-// AFTER
-AT_DISPATCH_V2(
-    iter.dtype(),
-    "min_values_cuda",
-    AT_WRAP([&]() {
-      min_values_kernel_cuda_impl<scalar_t>(iter);
-    }),
-    AT_EXPAND(AT_ALL_TYPES),
-    kBFloat16, kHalf, kBool
-);
-```
-
-### Step 6: Handle multi-line lambdas
-
-For lambdas with internal commas or complex expressions, AT_WRAP is essential:
-
-```cpp
-AT_DISPATCH_V2(
-    dtype,
-    "complex_kernel",
-    AT_WRAP([&]() {
-      gpu_reduce_kernel<scalar_t, scalar_t>(
-        iter,
-        MinOps<scalar_t>{},
-        thrust::pair<scalar_t, int64_t>(upper_bound(), 0)  // Commas inside!
-      );
-    }),
-    AT_EXPAND(AT_ALL_TYPES)
-);
-```
-
-### Step 7: Verify the conversion
-
-Check that:
-- [ ] `AT_WRAP()` wraps the entire lambda
-- [ ] Type groups use `AT_EXPAND()`
-- [ ] Individual types don't have `AT_EXPAND()` (just `kBFloat16`, not `AT_EXPAND(kBFloat16)`)
-- [ ] Argument order is: scalar_type, name, lambda, types
-- [ ] Include added: `#include <ATen/Dispatch_v2.h>`
-
-## Type group reference
-
-Available type group macros (use with `AT_EXPAND()`):
-
-```cpp
-AT_INTEGRAL_TYPES      // kByte, kChar, kInt, kLong, kShort
-AT_FLOATING_TYPES      // kDouble, kFloat
-AT_COMPLEX_TYPES       // kComplexDouble, kComplexFloat
-AT_QINT_TYPES         // kQInt8, kQUInt8, kQInt32
-AT_ALL_TYPES          // INTEGRAL_TYPES + FLOATING_TYPES
-AT_ALL_TYPES_AND_COMPLEX  // ALL_TYPES + COMPLEX_TYPES
-AT_INTEGRAL_TYPES_V2  // INTEGRAL_TYPES + unsigned types
-AT_BAREBONES_UNSIGNED_TYPES  // kUInt16, kUInt32, kUInt64
-AT_FLOAT8_TYPES       // Float8 variants
-```
-
-## Common patterns
-
-### Pattern: AT_DISPATCH_ALL_TYPES_AND2
-
-```cpp
-// Before
-AT_DISPATCH_ALL_TYPES_AND2(kHalf, kBFloat16, dtype, "op", [&]() {
-  kernel<scalar_t>(data);
-});
-
-// After
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>(data);
-}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBFloat16);
-```
-
-### Pattern: AT_DISPATCH_FLOATING_TYPES_AND3
-
-```cpp
-// Before
-AT_DISPATCH_FLOATING_TYPES_AND3(kHalf, kBFloat16, kFloat8_e4m3fn,
-    tensor.scalar_type(), "float_op", [&] {
-  process<scalar_t>(tensor);
-});
-
-// After
-AT_DISPATCH_V2(tensor.scalar_type(), "float_op", AT_WRAP([&] {
-  process<scalar_t>(tensor);
-}), AT_EXPAND(AT_FLOATING_TYPES), kHalf, kBFloat16, kFloat8_e4m3fn);
-```
-
-### Pattern: AT_DISPATCH_ALL_TYPES_AND_COMPLEX_AND2
-
-```cpp
-// Before
-AT_DISPATCH_ALL_TYPES_AND_COMPLEX_AND2(
-    kComplexHalf, kHalf,
-    self.scalar_type(),
-    "complex_op",
-    [&] {
-      result = compute<scalar_t>(self);
-    }
-);
-
-// After
-AT_DISPATCH_V2(
-    self.scalar_type(),
-    "complex_op",
-    AT_WRAP([&] {
-      result = compute<scalar_t>(self);
-    }),
-    AT_EXPAND(AT_ALL_TYPES),
-    AT_EXPAND(AT_COMPLEX_TYPES),
-    kComplexHalf,
-    kHalf
-);
-```
-
-## Edge cases
-
-### Case 1: No extra types (rare)
-
-```cpp
-// Before
-AT_DISPATCH_ALL_TYPES(dtype, "op", [&]() { kernel<scalar_t>(); });
-
-// After
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES));
-```
-
-### Case 2: Many individual types (AND4, AND5, etc.)
-
-```cpp
-// Before
-AT_DISPATCH_FLOATING_TYPES_AND4(kHalf, kBFloat16, kFloat8_e4m3fn, kFloat8_e5m2,
-    dtype, "float8_op", [&]() { kernel<scalar_t>(); });
-
-// After
-AT_DISPATCH_V2(dtype, "float8_op", AT_WRAP([&]() {
-  kernel<scalar_t>();
-}), AT_EXPAND(AT_FLOATING_TYPES), kHalf, kBFloat16, kFloat8_e4m3fn, kFloat8_e5m2);
-```
-
-### Case 3: Lambda with no captures
-
-```cpp
-// Before
-AT_DISPATCH_ALL_TYPES_AND2(kHalf, kBool, dtype, "op", []() {
-  static_kernel<scalar_t>();
-});
-
-// After
-AT_DISPATCH_V2(dtype, "op", AT_WRAP([]() {
-  static_kernel<scalar_t>();
-}), AT_EXPAND(AT_ALL_TYPES), kHalf, kBool);
-```
-
-## Benefits of AT_DISPATCH_V2
-
-1. **No arity in macro name**: Don't need different macros for AND2, AND3, AND4
-2. **Composable type sets**: Mix and match type groups with `AT_EXPAND()`
-3. **Extensible**: Easy to add more types without hitting macro limits
-4. **Clearer**: Type groups are explicit, not implicit in macro name
-
-## Important notes
-
-- Keep `#include <ATen/Dispatch.h>` - other code may need it
-- The `AT_WRAP()` is mandatory - prevents comma parsing issues in the lambda
-- Type groups need `AT_EXPAND()`, individual types don't
-- The v2 API is in `aten/src/ATen/Dispatch_v2.h` - refer to it for full docs
-- See the header file for the Python script to regenerate the macro implementation
-
-## Workflow
-
-When asked to convert AT_DISPATCH macros:
-
-1. Read the file to identify all AT_DISPATCH uses
-2. Add `#include <ATen/Dispatch_v2.h>` if not present
-3. For each dispatch macro:
-   - Identify the pattern and extract components
-   - Map the base type group
-   - Extract individual types
-   - Construct the AT_DISPATCH_V2 call
-   - Apply with Edit tool
-4. Show the user the complete converted file
-5. Explain what was changed
-
-Do NOT compile or test the code - focus on accurate conversion only.

.github/ISSUE_TEMPLATE/release-feature-request.yml

@@ -1,11 +1,11 @@
-name: 🚀 New Feature for Release
+name: 🚀 Release highlight for proposed Feature
 description: Submit a Release highlight for proposed Feature
 labels: ["release-feature-request"]
 
 body:
 - type: textarea
   attributes:
-    label: New Feature for Release
+    label: Release highlight for proposed Feature
     description: >
       Example: “A torch.special module, analogous to SciPy's special module.”
 - type: input

.github/actions/pytest-cache-download/action.yml

@@ -38,9 +38,9 @@ runs:
       run: |
         python3 .github/scripts/pytest_cache.py \
           --download \
-          --cache_dir "$GITHUB_WORKSPACE/$CACHE_DIR" \
-          --pr_identifier "$GITHUB_REF" \
-          --job_identifier "$JOB_IDENTIFIER" \
-          --temp_dir "$RUNNER_TEMP" \
-          --repo "$REPO" \
-          --bucket "$BUCKET" \
+          --cache_dir $GITHUB_WORKSPACE/$CACHE_DIR \
+          --pr_identifier $GITHUB_REF \
+          --job_identifier $JOB_IDENTIFIER \
+          --temp_dir $RUNNER_TEMP \
+          --repo $REPO \
+          --bucket $BUCKET \

.github/actions/pytest-cache-upload/action.yml

@@ -47,11 +47,11 @@ runs:
       run: |
         python3 .github/scripts/pytest_cache.py \
           --upload \
-          --cache_dir "$GITHUB_WORKSPACE/$CACHE_DIR" \
-          --pr_identifier "$GITHUB_REF" \
-          --job_identifier "$JOB_IDENTIFIER" \
-          --sha "$SHA" \
-          --test_config "$TEST_CONFIG" \
-          --shard "$SHARD" \
-          --repo "$REPO" \
-          --temp_dir "$RUNNER_TEMP" \
+          --cache_dir $GITHUB_WORKSPACE/$CACHE_DIR \
+          --pr_identifier $GITHUB_REF \
+          --job_identifier $JOB_IDENTIFIER \
+          --sha $SHA \
+          --test_config $TEST_CONFIG \
+          --shard $SHARD \
+          --repo $REPO \
+          --temp_dir $RUNNER_TEMP \

.github/ci_commit_pins/audio.txt

@@ -1 +1 @@
-ad5816f0eee1c873df1b7d371c69f1f811a89387
+3b0e7a6f192ca2715e7e6cbe5db007aea7165fe2

.github/ci_commit_pins/vision.txt

@@ -1 +1 @@
-cfbc5c2f1c798991715a6b06bb3ce46478c4487c
+218d2ab791d437309f91e0486eb9fa7f00badc17

.github/ci_commit_pins/xla.txt

@@ -1 +1 @@
-c8b09f5f77d6bf6fb7ed7a9aa83e5d8156b3a5e9
+df6798dfb931ce7c7fe5bed2447cd1092a5981af

.github/copilot-instructions.md

@@ -1,125 +0,0 @@
-# PyTorch Copilot Instructions
-
-This is the PyTorch machine learning framework codebase. These instructions help AI agents navigate and contribute effectively.
-
-## Architecture Overview
-
-### Core Components
-
-- **c10/** - Core library (C++-10 compatible) for essential, binary-size-conscious functionality
-- **aten/** - ATen tensor library (C++), PyTorch's foundation without autograd
-  - `aten/src/ATen/native/` - Modern operator implementations (CPU/CUDA/MPS/sparse)
-  - `aten/src/ATen/native/native_functions.yaml` - **Critical**: Declarative operator registry
-- **torch/** - Python bindings and public API
-  - `torch/csrc/` - C++ Python bindings (hand-written and generated)
-  - `torch/csrc/autograd/` - Reverse-mode automatic differentiation
-  - `torch/csrc/jit/` - TorchScript JIT compiler
-- **torchgen/** - Code generation tooling that reads `native_functions.yaml`
-- **tools/** - Build scripts, autograd derivatives, code generation
-
-### The Code Generation Workflow
-
-**Most operator changes require editing `native_functions.yaml`**, not direct C++ files. This YAML file:
-1. Declares operator signatures, variants (function/method), and dispatch behavior
-2. Gets processed by `torchgen/` to generate C++/Python bindings
-3. Produces headers in `build/aten/src/ATen/` during compilation
-
-Example entry structure:
-```yaml
-- func: my_op(Tensor self, Scalar alpha=1) -> Tensor
-  variants: function, method
-  dispatch:
-    CPU: my_op_cpu
-    CUDA: my_op_cuda
-```
-
-After editing `native_functions.yaml`, implement kernels in `aten/src/ATen/native/` (see `aten/src/ATen/native/README.md`).
-
-## Development Workflows
-
-### Building from Source
-
-**Never run `setup.py` directly** - use pip with editable install:
-```bash
-python -m pip install --no-build-isolation -v -e .
-```
-
-Speed up builds:
-- `DEBUG=1` - Debug symbols with `-g -O0`
-- `USE_CUDA=0` - Skip CUDA compilation
-- `BUILD_TEST=0` - Skip C++ test binaries
-- Install `ninja` (`pip install ninja`) for faster builds
-- Use `ccache` for incremental compilation caching
-
-Rebuild specific targets: `(cd build && ninja <target>)`
-
-### Testing
-
-**Critical**: DO NOT run entire test suites. Run specific tests only:
-```bash
-python test/test_torch.py TestTorch.test_specific_case
-```
-
-**Test structure**: All tests use `torch.testing._internal.common_utils`:
-```python
-from torch.testing._internal.common_utils import run_tests, TestCase
-
-class TestFeature(TestCase):
-    def test_something(self):
-        # Use self.assertEqual for tensor comparisons
-        pass
-
-if __name__ == "__main__":
-    run_tests()
-```
-
-**For bug fixes**: Create a standalone reproduction script first, verify it fails, then fix and add to appropriate test file.
-
-### Linting
-
-Run linter (not pre-commit): `lintrunner -a` (auto-applies fixes)
-
-## Project-Specific Conventions
-
-### Memory and Storage
-- **Storage is never nullptr** (but `StorageImpl.data` may be nullptr for unallocated outputs)
-- CUDA device info lives in storage objects
-
-### Python-C++ Integration (`torch/csrc/`)
-- Always include `Python.h` **first** to avoid `_XOPEN_SOURCE` redefinition errors
-- Use `pybind11::gil_scoped_acquire` before calling Python API or using `THPObjectPtr`
-- Wrap entry points with `HANDLE_TH_ERRORS` / `END_HANDLE_TH_ERRORS` for exception conversion
-
-### Dispatch System
-- PyTorch uses operator dispatch to route calls to backend-specific kernels
-- Prefer `CompositeExplicitAutograd` dispatch when writing device-agnostic compound ops
-- See `aten/src/ATen/native/README.md` for dispatch keyword guidance
-
-## Git Workflow (AI Agent Specific)
-
-When preparing PRs from this environment:
-```bash
-git stash -u
-git reset --hard $(cat /tmp/orig_work.txt)  # Reset to LOCAL branch
-git stash pop
-# Resolve conflicts if necessary
-```
-
-## Common Gotchas
-
-1. **Editing generated files** - If it's in `build/`, don't edit it. Edit the source template or `native_functions.yaml`
-2. **NVCC template compilation** - NVCC is stricter about C++ than gcc/clang; code working on Linux may fail Windows CI
-3. **Windows symbol visibility** - Use `TORCH_API` macros for exported symbols (required on Windows, optional on Linux)
-4. **No internet access** - DO NOT attempt to install dependencies during development
-
-## Key Files Reference
-
-- `AGENTS.md` - Instructions specific to AI coding agents
-- `CONTRIBUTING.md` - Comprehensive human contributor guide
-- `GLOSSARY.md` - Terminology (ATen, kernels, operations, JIT, TorchScript)
-- `aten/src/ATen/native/README.md` - Operator implementation guide
-- `tools/autograd/derivatives.yaml` - Gradient definitions for autograd
-
-## Performance Debugging
-
-Use `TORCH_SHOW_CPP_STACKTRACES=1` for C++ traces in Python errors. For profiling, prefer `py-spy` over manual instrumentation.

.github/scripts/generate_binary_build_matrix.py

@@ -11,24 +11,18 @@ architectures:
     * Latest XPU
 """
 
-import json
 import os
-import re
-from pathlib import Path
 from typing import Optional
 
 
-SCRIPT_DIR = Path(__file__).absolute().parent
-REPO_ROOT = SCRIPT_DIR.parent.parent
-
-
+# NOTE: Please also update the CUDA sources in `PIP_SOURCES` in tools/nightly.py when changing this
 CUDA_ARCHES = ["12.6", "12.8", "12.9", "13.0"]
 CUDA_STABLE = "12.8"
 CUDA_ARCHES_FULL_VERSION = {
     "12.6": "12.6.3",
     "12.8": "12.8.1",
     "12.9": "12.9.1",
-    "13.0": "13.0.0",
+    "13.0": "13.0.2",
 }
 CUDA_ARCHES_CUDNN_VERSION = {
     "12.6": "9",
@@ -37,7 +31,8 @@ CUDA_ARCHES_CUDNN_VERSION = {
     "13.0": "9",
 }
 
-ROCM_ARCHES = ["7.0", "7.1"]
+# NOTE: Please also update the ROCm sources in `PIP_SOURCES` in tools/nightly.py when changing this
+ROCM_ARCHES = ["6.4", "7.0"]
 
 XPU_ARCHES = ["xpu"]
 
@@ -142,48 +137,9 @@ PYTORCH_EXTRA_INSTALL_REQUIREMENTS = {
 }
 
 
-# Used by tools/nightly.py
-PYTORCH_NIGHTLY_PIP_INDEX_URL = "https://download.pytorch.org/whl/nightly"
-NIGHTLY_SOURCE_MATRIX = {
-    "cpu": dict(
-        name="cpu",
-        index_url=f"{PYTORCH_NIGHTLY_PIP_INDEX_URL}/cpu",
-        supported_platforms=["Linux", "macOS", "Windows"],
-        accelerator="cpu",
-    )
-}
-CUDA_NIGHTLY_SOURCE_MATRIX = {
-    f"cuda-{major}.{minor}": dict(
-        name=f"cuda-{major}.{minor}",
-        index_url=f"{PYTORCH_NIGHTLY_PIP_INDEX_URL}/cu{major}{minor}",
-        supported_platforms=["Linux", "Windows"],
-        accelerator="cuda",
-    )
-    for major, minor in (map(int, version.split(".")) for version in CUDA_ARCHES)
-}
-ROCM_NIGHTLY_SOURCE_MATRIX = {
-    f"rocm-{major}.{minor}": dict(
-        name=f"rocm-{major}.{minor}",
-        index_url=f"{PYTORCH_NIGHTLY_PIP_INDEX_URL}/rocm{major}.{minor}",
-        supported_platforms=["Linux"],
-        accelerator="rocm",
-    )
-    for major, minor in (map(int, version.split(".")) for version in ROCM_ARCHES)
-}
-XPU_NIGHTLY_SOURCE_MATRIX = {
-    "xpu": dict(
-        name="xpu",
-        index_url=f"{PYTORCH_NIGHTLY_PIP_INDEX_URL}/xpu",
-        supported_platforms=["Linux"],
-        accelerator="xpu",
-    )
-}
-NIGHTLY_SOURCE_MATRIX.update(CUDA_NIGHTLY_SOURCE_MATRIX)
-NIGHTLY_SOURCE_MATRIX.update(ROCM_NIGHTLY_SOURCE_MATRIX)
-NIGHTLY_SOURCE_MATRIX.update(XPU_NIGHTLY_SOURCE_MATRIX)
-
-
 def get_nccl_wheel_version(arch_version: str) -> str:
+    import re
+
     requirements = map(
         str.strip, re.split("[;|]", PYTORCH_EXTRA_INSTALL_REQUIREMENTS[arch_version])
     )
@@ -191,14 +147,17 @@ def get_nccl_wheel_version(arch_version: str) -> str:
 
 
 def read_nccl_pin(arch_version: str) -> str:
-    nccl_pin_path = (
-        REPO_ROOT
-        / ".ci"
-        / "docker"
-        / "ci_commit_pins"
-        / f"nccl-cu{arch_version[:2]}.txt"
+    from pathlib import Path
+
+    nccl_pin_path = os.path.join(
+        Path(__file__).absolute().parents[2],
+        ".ci",
+        "docker",
+        "ci_commit_pins",
+        f"nccl-cu{arch_version[:2]}.txt",
     )
-    return nccl_pin_path.read_text().strip()
+    with open(nccl_pin_path) as f:
+        return f.read().strip()
 
 
 def validate_nccl_dep_consistency(arch_version: str) -> None:
@@ -206,8 +165,7 @@ def validate_nccl_dep_consistency(arch_version: str) -> None:
     wheel_ver = get_nccl_wheel_version(arch_version)
     if not nccl_release_tag.startswith(f"v{wheel_ver}"):
         raise RuntimeError(
-            f"{arch_version} NCCL release tag version {nccl_release_tag} "
-            f"does not correspond to wheel version {wheel_ver}"
+            f"{arch_version} NCCL release tag version {nccl_release_tag} does not correspond to wheel version {wheel_ver}"
         )
 
 
@@ -454,14 +412,7 @@ def generate_wheels_matrix(
     return ret
 
 
-arch_version = ""
-for arch_version in CUDA_ARCHES:
-    validate_nccl_dep_consistency(arch_version)
-del arch_version
-
-
-if __name__ == "__main__":
-    # Used by tools/nightly.py
-    (SCRIPT_DIR / "nightly_source_matrix.json").write_text(
-        json.dumps(NIGHTLY_SOURCE_MATRIX, indent=4) + "\n"
-    )
+validate_nccl_dep_consistency("13.0")
+validate_nccl_dep_consistency("12.9")
+validate_nccl_dep_consistency("12.8")
+validate_nccl_dep_consistency("12.6")

.github/workflows/_rocm-test.yml

@@ -97,8 +97,8 @@ jobs:
         shell: bash
         run: |
           ngpu=$(rocminfo | grep -c -E 'Name:.*\sgfx')
-          if [[ $ngpu -lt 2 ]]; then #We are temporarily reducing this down to 2 from 4 so that we can run tests on nodes with less gpus.
-            echo "Error: only $ngpu GPU(s) detected, at least 2 GPUs are needed for distributed jobs"
+          if [[ $ngpu -lt 4 ]]; then
+            echo "Error: only $ngpu GPU(s) detected, at least 4 GPUs are needed for distributed jobs"
             exit 1
           fi
 

.github/workflows/_xpu-test.yml

@@ -344,21 +344,5 @@ jobs:
           if-no-files-found: ignore
           path: ./**/core.[1-9]*
 
-      - name: Authenticate with AWS
-        uses: aws-actions/configure-aws-credentials@ececac1a45f3b08a01d2dd070d28d111c5fe6722 # v4.1.0
-        with:
-          role-to-assume: arn:aws:iam::308535385114:role/gha_workflow_upload-benchmark-results
-          # The max duration enforced by the server side
-          role-duration-seconds: 18000
-          aws-region: us-east-1
-
-      - name: Upload the benchmark results
-        uses: pytorch/test-infra/.github/actions/upload-benchmark-results@main
-        with:
-          benchmark-results-dir: test/test-reports
-          dry-run: false
-          schema-version: v3
-          github-token: ${{ secrets.GITHUB_TOKEN }}
-
       - name: Teardown XPU
         uses: ./.github/actions/teardown-xpu

.github/workflows/build-almalinux-images.yml

@@ -36,7 +36,7 @@ jobs:
     runs-on: linux.9xlarge.ephemeral
     strategy:
       matrix:
-        tag: ["cuda12.6", "cuda12.8", "cuda12.9", "cuda13.0", "rocm7.0", "rocm7.1", "cpu"]
+        tag: ["cuda12.6", "cuda12.8", "cuda12.9", "cuda13.0", "rocm6.4", "rocm7.0", "rocm7.1", "cpu"]
     steps:
       - name: Build docker image
         uses: pytorch/pytorch/.github/actions/binary-docker-build@main

.github/workflows/build-libtorch-images.yml

@@ -52,6 +52,7 @@ jobs:
           { tag: "cuda12.9" },
           { tag: "cuda12.8" },
           { tag: "cuda12.6" },
+          { tag: "rocm6.4"  },
           { tag: "rocm7.0"  },
           { tag: "rocm7.1"  },
           { tag: "cpu"      },

.github/workflows/build-magma-rocm-linux.yml

@@ -34,7 +34,7 @@ jobs:
       id-token: write
     strategy:
       matrix:
-        rocm_version: ["71", "70"]
+        rocm_version: ["70", "64"]
     steps:
       - name: Checkout PyTorch
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2

.github/workflows/build-manywheel-images.yml

@@ -54,6 +54,7 @@ jobs:
           { name: "manylinuxaarch64-builder",       tag: "cuda12.9",          runner: "linux.arm64.2xlarge.ephemeral" },
           { name: "manylinuxaarch64-builder",       tag: "cuda12.8",          runner: "linux.arm64.2xlarge.ephemeral" },
           { name: "manylinuxaarch64-builder",       tag: "cuda12.6",          runner: "linux.arm64.2xlarge.ephemeral" },
+          { name: "manylinux2_28-builder",          tag: "rocm6.4",           runner: "linux.9xlarge.ephemeral" },
           { name: "manylinux2_28-builder",          tag: "rocm7.0",           runner: "linux.9xlarge.ephemeral" },
           { name: "manylinux2_28-builder",          tag: "rocm7.1",           runner: "linux.9xlarge.ephemeral" },
           { name: "manylinux2_28-builder",          tag: "cpu",               runner: "linux.9xlarge.ephemeral" },

.github/workflows/build-triton-wheel.yml

@@ -55,7 +55,7 @@ jobs:
         docker-image: ["pytorch/manylinux2_28-builder:cpu"]
         include:
           - device: "rocm"
-            rocm_version: "7.1"
+            rocm_version: "7.0"
             runs_on: "${{ needs.get-label-type.outputs.label-type }}linux.4xlarge"
           - device: "cuda"
             rocm_version: ""
@@ -159,7 +159,12 @@ jobs:
             WITH_CLANG_LDD="--with-clang-ldd"
           fi
 
-          docker exec -t "${container_name}" bash -c "${PYTHON_EXECUTABLE} /pytorch/.github/scripts/build_triton_wheel.py --device=$BUILD_DEVICE $RELEASE $WITH_CLANG_LDD"
+          if [[ "${BUILD_DEVICE}" == xpu ]]; then
+            docker exec -t "${container_name}" bash -c "dnf install -y gcc-toolset-13-gcc-c++"
+            docker exec -t "${container_name}" bash -c "source /opt/rh/gcc-toolset-13/enable && ${PYTHON_EXECUTABLE} /pytorch/.github/scripts/build_triton_wheel.py --device=$BUILD_DEVICE $RELEASE"
+          else
+            docker exec -t "${container_name}" bash -c "${PYTHON_EXECUTABLE} /pytorch/.github/scripts/build_triton_wheel.py --device=$BUILD_DEVICE $RELEASE $WITH_CLANG_LDD"
+          fi
 
           if [[ ("${{ matrix.device }}" == "cuda" || "${{ matrix.device }}" == "xpu") ]]; then
             docker exec -t "${container_name}"  bash -c "auditwheel repair --plat ${PLATFORM} //artifacts/*.whl"

.github/workflows/docker-builds.yml

@@ -77,11 +77,9 @@ jobs:
           pytorch-linux-noble-riscv64-py3.12-gcc14
         ]
         include:
-          - docker-image-name: pytorch-linux-jammy-aarch64-py3.10-gcc13
+          - docker-image-name: pytorch-linux-jammy-aarch64-py3.10-gcc11
             runner: linux.arm64.m7g.4xlarge
-          - docker-image-name: pytorch-linux-jammy-aarch64-py3.10-clang21
-            runner: linux.arm64.m7g.4xlarge
-          - docker-image-name: pytorch-linux-jammy-aarch64-py3.10-gcc13-inductor-benchmarks
+          - docker-image-name: pytorch-linux-jammy-aarch64-py3.10-gcc11-inductor-benchmarks
             runner: linux.arm64.m7g.4xlarge
             timeout-minutes: 600
     # Docker uploads fail from LF runners, see https://github.com/pytorch/pytorch/pull/137358

.github/workflows/docker-release.yml

@@ -8,7 +8,6 @@ on:
       - docker.Makefile
       - .github/workflows/docker-release.yml
       - .github/scripts/generate_docker_release_matrix.py
-      - .github/scripts/generate_binary_build_matrix.py
   push:
     branches:
       - nightly

.github/workflows/generated-linux-binary-libtorch-nightly.yml

@@ -384,6 +384,124 @@ jobs:
       github-token: ${{ secrets.GITHUB_TOKEN }}
     uses: ./.github/workflows/_binary-upload.yml
 
+  libtorch-rocm6_4-shared-with-deps-release-build:
+    if: ${{ github.repository_owner == 'pytorch' }}
+    uses: ./.github/workflows/_binary-build-linux.yml
+    needs: get-label-type
+    with:
+      PYTORCH_ROOT: /pytorch
+      PACKAGE_TYPE: libtorch
+      # TODO: This is a legacy variable that we eventually want to get rid of in
+      #       favor of GPU_ARCH_VERSION
+      DESIRED_CUDA: rocm6.4
+      GPU_ARCH_VERSION: "6.4"
+      GPU_ARCH_TYPE: rocm
+      DOCKER_IMAGE: libtorch-cxx11-builder
+      DOCKER_IMAGE_TAG_PREFIX: rocm6.4
+      LIBTORCH_CONFIG: release
+      LIBTORCH_VARIANT: shared-with-deps
+      runner_prefix: "${{ needs.get-label-type.outputs.label-type }}"
+      timeout-minutes: 300
+      build_name: libtorch-rocm6_4-shared-with-deps-release
+      build_environment: linux-binary-libtorch
+    secrets:
+      github-token: ${{ secrets.GITHUB_TOKEN }}
+  libtorch-rocm6_4-shared-with-deps-release-test:  # Testing
+    if: ${{ github.repository_owner == 'pytorch' }}
+    needs:
+      - libtorch-rocm6_4-shared-with-deps-release-build
+      - get-label-type
+    runs-on: linux.rocm.gpu.mi250
+    timeout-minutes: 240
+    env:
+      PYTORCH_ROOT: /pytorch
+      PACKAGE_TYPE: libtorch
+      # TODO: This is a legacy variable that we eventually want to get rid of in
+      #       favor of GPU_ARCH_VERSION
+      DESIRED_CUDA: rocm6.4
+      GPU_ARCH_VERSION: "6.4"
+      GPU_ARCH_TYPE: rocm
+      SKIP_ALL_TESTS: 1
+      DOCKER_IMAGE: libtorch-cxx11-builder
+      DOCKER_IMAGE_TAG_PREFIX: rocm6.4
+      LIBTORCH_CONFIG: release
+      LIBTORCH_VARIANT: shared-with-deps
+    permissions:
+      id-token: write
+      contents: read
+    steps:
+      - name: Setup ROCm
+        uses: ./.github/actions/setup-rocm
+      - uses: actions/download-artifact@v4.1.7
+        name: Download Build Artifacts
+        with:
+          name: libtorch-rocm6_4-shared-with-deps-release
+          path: "${{ runner.temp }}/artifacts/"
+      - name: Checkout PyTorch
+        uses: actions/checkout@v4
+        with:
+          ref: ${{ github.event_name == 'pull_request' && github.event.pull_request.head.sha || github.sha }}
+          submodules: recursive
+          path: pytorch
+          show-progress: false
+      - name: Clean PyTorch checkout
+        run: |
+          # Remove any artifacts from the previous checkouts
+          git clean -fxd
+        working-directory: pytorch
+      - name: ROCm set GPU_FLAG
+        run: |
+          echo "GPU_FLAG=--device=/dev/mem --device=/dev/kfd --device=/dev/dri --group-add video --group-add daemon" >> "${GITHUB_ENV}"
+      - name: configure aws credentials
+        id: aws_creds
+        if: ${{ startsWith(github.event.ref, 'refs/tags/ciflow/') }}
+        uses: aws-actions/configure-aws-credentials@v4
+        with:
+          role-to-assume: arn:aws:iam::308535385114:role/gha_workflow_s3_and_ecr_read_only
+          aws-region: us-east-1
+          role-duration-seconds: 18000
+      - name: Calculate docker image
+        id: calculate-docker-image
+        uses: pytorch/test-infra/.github/actions/calculate-docker-image@main
+        with:
+          docker-registry: ${{ startsWith(github.event.ref, 'refs/tags/ciflow/') && '308535385114.dkr.ecr.us-east-1.amazonaws.com' || 'docker.io' }}
+          docker-image-name: libtorch-cxx11-builder
+          custom-tag-prefix: rocm6.4
+          docker-build-dir: .ci/docker
+          working-directory: pytorch
+      - name: Pull Docker image
+        uses: pytorch/test-infra/.github/actions/pull-docker-image@main
+        with:
+          docker-image: ${{ steps.calculate-docker-image.outputs.docker-image }}
+      - name: Test Pytorch binary
+        uses: ./pytorch/.github/actions/test-pytorch-binary
+        env:
+          DOCKER_IMAGE: ${{ steps.calculate-docker-image.outputs.docker-image }}
+      - name: Teardown ROCm
+        uses: ./.github/actions/teardown-rocm
+  libtorch-rocm6_4-shared-with-deps-release-upload:  # Uploading
+    if: ${{ github.repository_owner == 'pytorch' }}
+    permissions:
+      id-token: write
+      contents: read
+    needs: libtorch-rocm6_4-shared-with-deps-release-test
+    with:
+      PYTORCH_ROOT: /pytorch
+      PACKAGE_TYPE: libtorch
+      # TODO: This is a legacy variable that we eventually want to get rid of in
+      #       favor of GPU_ARCH_VERSION
+      DESIRED_CUDA: rocm6.4
+      GPU_ARCH_VERSION: "6.4"
+      GPU_ARCH_TYPE: rocm
+      DOCKER_IMAGE: libtorch-cxx11-builder
+      DOCKER_IMAGE_TAG_PREFIX: rocm6.4
+      LIBTORCH_CONFIG: release
+      LIBTORCH_VARIANT: shared-with-deps
+      build_name: libtorch-rocm6_4-shared-with-deps-release
+    secrets:
+      github-token: ${{ secrets.GITHUB_TOKEN }}
+    uses: ./.github/workflows/_binary-upload.yml
+
   libtorch-rocm7_0-shared-with-deps-release-build:
     if: ${{ github.repository_owner == 'pytorch' }}
     uses: ./.github/workflows/_binary-build-linux.yml
@@ -501,121 +619,3 @@ jobs:
     secrets:
       github-token: ${{ secrets.GITHUB_TOKEN }}
     uses: ./.github/workflows/_binary-upload.yml
-
-  libtorch-rocm7_1-shared-with-deps-release-build:
-    if: ${{ github.repository_owner == 'pytorch' }}
-    uses: ./.github/workflows/_binary-build-linux.yml
-    needs: get-label-type
-    with:
-      PYTORCH_ROOT: /pytorch
-      PACKAGE_TYPE: libtorch
-      # TODO: This is a legacy variable that we eventually want to get rid of in
-      #       favor of GPU_ARCH_VERSION
-      DESIRED_CUDA: rocm7.1
-      GPU_ARCH_VERSION: "7.1"
-      GPU_ARCH_TYPE: rocm
-      DOCKER_IMAGE: libtorch-cxx11-builder
-      DOCKER_IMAGE_TAG_PREFIX: rocm7.1
-      LIBTORCH_CONFIG: release
-      LIBTORCH_VARIANT: shared-with-deps
-      runner_prefix: "${{ needs.get-label-type.outputs.label-type }}"
-      timeout-minutes: 300
-      build_name: libtorch-rocm7_1-shared-with-deps-release
-      build_environment: linux-binary-libtorch
-    secrets:
-      github-token: ${{ secrets.GITHUB_TOKEN }}
-  libtorch-rocm7_1-shared-with-deps-release-test:  # Testing
-    if: ${{ github.repository_owner == 'pytorch' }}
-    needs:
-      - libtorch-rocm7_1-shared-with-deps-release-build
-      - get-label-type
-    runs-on: linux.rocm.gpu.mi250
-    timeout-minutes: 240
-    env:
-      PYTORCH_ROOT: /pytorch
-      PACKAGE_TYPE: libtorch
-      # TODO: This is a legacy variable that we eventually want to get rid of in
-      #       favor of GPU_ARCH_VERSION
-      DESIRED_CUDA: rocm7.1
-      GPU_ARCH_VERSION: "7.1"
-      GPU_ARCH_TYPE: rocm
-      SKIP_ALL_TESTS: 1
-      DOCKER_IMAGE: libtorch-cxx11-builder
-      DOCKER_IMAGE_TAG_PREFIX: rocm7.1
-      LIBTORCH_CONFIG: release
-      LIBTORCH_VARIANT: shared-with-deps
-    permissions:
-      id-token: write
-      contents: read
-    steps:
-      - name: Setup ROCm
-        uses: ./.github/actions/setup-rocm
-      - uses: actions/download-artifact@v4.1.7
-        name: Download Build Artifacts
-        with:
-          name: libtorch-rocm7_1-shared-with-deps-release
-          path: "${{ runner.temp }}/artifacts/"
-      - name: Checkout PyTorch
-        uses: actions/checkout@v4
-        with:
-          ref: ${{ github.event_name == 'pull_request' && github.event.pull_request.head.sha || github.sha }}
-          submodules: recursive
-          path: pytorch
-          show-progress: false
-      - name: Clean PyTorch checkout
-        run: |
-          # Remove any artifacts from the previous checkouts
-          git clean -fxd
-        working-directory: pytorch
-      - name: ROCm set GPU_FLAG
-        run: |
-          echo "GPU_FLAG=--device=/dev/mem --device=/dev/kfd --device=/dev/dri --group-add video --group-add daemon" >> "${GITHUB_ENV}"
-      - name: configure aws credentials
-        id: aws_creds
-        if: ${{ startsWith(github.event.ref, 'refs/tags/ciflow/') }}
-        uses: aws-actions/configure-aws-credentials@v4
-        with:
-          role-to-assume: arn:aws:iam::308535385114:role/gha_workflow_s3_and_ecr_read_only
-          aws-region: us-east-1
-          role-duration-seconds: 18000
-      - name: Calculate docker image
-        id: calculate-docker-image
-        uses: pytorch/test-infra/.github/actions/calculate-docker-image@main
-        with:
-          docker-registry: ${{ startsWith(github.event.ref, 'refs/tags/ciflow/') && '308535385114.dkr.ecr.us-east-1.amazonaws.com' || 'docker.io' }}
-          docker-image-name: libtorch-cxx11-builder
-          custom-tag-prefix: rocm7.1
-          docker-build-dir: .ci/docker
-          working-directory: pytorch
-      - name: Pull Docker image
-        uses: pytorch/test-infra/.github/actions/pull-docker-image@main
-        with:
-          docker-image: ${{ steps.calculate-docker-image.outputs.docker-image }}
-      - name: Test Pytorch binary
-        uses: ./pytorch/.github/actions/test-pytorch-binary
-        env:
-          DOCKER_IMAGE: ${{ steps.calculate-docker-image.outputs.docker-image }}
-      - name: Teardown ROCm
-        uses: ./.github/actions/teardown-rocm
-  libtorch-rocm7_1-shared-with-deps-release-upload:  # Uploading
-    if: ${{ github.repository_owner == 'pytorch' }}
-    permissions:
-      id-token: write
-      contents: read
-    needs: libtorch-rocm7_1-shared-with-deps-release-test
-    with:
-      PYTORCH_ROOT: /pytorch
-      PACKAGE_TYPE: libtorch
-      # TODO: This is a legacy variable that we eventually want to get rid of in
-      #       favor of GPU_ARCH_VERSION
-      DESIRED_CUDA: rocm7.1
-      GPU_ARCH_VERSION: "7.1"
-      GPU_ARCH_TYPE: rocm
-      DOCKER_IMAGE: libtorch-cxx11-builder
-      DOCKER_IMAGE_TAG_PREFIX: rocm7.1
-      LIBTORCH_CONFIG: release
-      LIBTORCH_VARIANT: shared-with-deps
-      build_name: libtorch-rocm7_1-shared-with-deps-release
-    secrets:
-      github-token: ${{ secrets.GITHUB_TOKEN }}
-    uses: ./.github/workflows/_binary-upload.yml

.github/workflows/inductor-perf-test-nightly-aarch64.yml

@@ -72,7 +72,7 @@ jobs:
       runner_prefix: "${{ needs.get-label-type.outputs.label-type }}"
       runner: linux.arm64.m7g.4xlarge
       build-environment: linux-jammy-aarch64-py3.10
-      docker-image-name: ci-image:pytorch-linux-jammy-aarch64-py3.10-gcc13-inductor-benchmarks
+      docker-image-name: ci-image:pytorch-linux-jammy-aarch64-py3.10-gcc11-inductor-benchmarks
       test-matrix: |
         { include: [
           { config: "inductor_huggingface_perf_cpu_aarch64", shard: 1, num_shards: 9, runner: "linux.arm64.m7g.metal" },

.github/workflows/inductor-rocm.yml

@@ -1,10 +1,9 @@
 name: inductor-rocm
 
 on:
-  schedule:
-    - cron: 0 */3 * * *
   push:
     branches:
+      - main
       - release/*
     tags:
       - ciflow/inductor-rocm/*

.github/workflows/inductor-unittest.yml

@@ -115,10 +115,10 @@ jobs:
       runner_prefix: "${{ needs.get-label-type.outputs.label-type }}"
       test-matrix: |
         { include: [
-          { config: "inductor_amx", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
-          { config: "inductor_amx", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
-          { config: "inductor_avx2", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.avx2" },
-          { config: "inductor_avx2", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.avx2" },
+          { config: "inductor_amx", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
+          { config: "inductor_amx", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
+          { config: "inductor_avx2", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.10xlarge.avx2" },
+          { config: "inductor_avx2", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.10xlarge.avx2" },
         ]}
     secrets: inherit
 

.github/workflows/inductor.yml

@@ -84,13 +84,13 @@ jobs:
       runner_prefix: "${{ needs.get-label-type.outputs.label-type }}"
       test-matrix: |
         { include: [
-          { config: "cpu_inductor_torchbench", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
-          { config: "cpu_inductor_torchbench", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
-          { config: "dynamic_cpu_inductor_huggingface", shard: 1, num_shards: 1, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
-          { config: "dynamic_cpu_inductor_timm", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
-          { config: "dynamic_cpu_inductor_timm", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
-          { config: "dynamic_cpu_inductor_torchbench", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
-          { config: "dynamic_cpu_inductor_torchbench", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge.amx" },
+          { config: "cpu_inductor_torchbench", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
+          { config: "cpu_inductor_torchbench", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
+          { config: "dynamic_cpu_inductor_huggingface", shard: 1, num_shards: 1, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
+          { config: "dynamic_cpu_inductor_timm", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
+          { config: "dynamic_cpu_inductor_timm", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
+          { config: "dynamic_cpu_inductor_torchbench", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
+          { config: "dynamic_cpu_inductor_torchbench", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.8xlarge.amx" },
           { config: "inductor_torchbench_cpu_smoketest_perf", shard: 1, num_shards: 1, runner: "${{ needs.get-label-type.outputs.label-type }}linux.24xl.spr-metal" },
         ]}
       build-additional-packages: "vision audio torchao"

.github/workflows/lint.yml

@@ -76,12 +76,11 @@ jobs:
 
   # NOTE: mypy needs its own job because it depends on --all-files, without assessing all files it sometimes
   #       fails to find types when it should
-  # NOTE: We should be able to disable this and consolidate with Pyrefly
-  lintrunner-pyrefly:
+  lintrunner-mypy:
     uses: pytorch/test-infra/.github/workflows/linux_job_v2.yml@main
-    name: lintrunner-pyrefly-${{ needs.get-changed-files.outputs.changed-files == '*' && 'all' || 'partial' }}
+    name: lintrunner-mypy-${{ needs.get-changed-files.outputs.changed-files == '*' && 'all' || 'partial' }}
     needs: [get-label-type, get-changed-files]
-    # Only run if there are changed files relevant to pyrefly
+    # Only run if there are changed files relevant to mypy
     if: |
       github.repository_owner == 'pytorch' && (
         needs.get-changed-files.outputs.changed-files == '*' ||
@@ -99,8 +98,8 @@ jobs:
       ref: ${{ github.event_name == 'pull_request' && github.event.pull_request.head.sha || github.sha }}
       script: |
         CHANGED_FILES="${{ needs.get-changed-files.outputs.changed-files }}"
-        echo "Running pyrefly"
-        ADDITIONAL_LINTRUNNER_ARGS="--take PYREFLY --all-files" .github/scripts/lintrunner.sh
+        echo "Running mypy"
+        ADDITIONAL_LINTRUNNER_ARGS="--take MYPY,MYPYSTRICT --all-files" .github/scripts/lintrunner.sh
 
   lintrunner-noclang:
     uses: pytorch/test-infra/.github/workflows/linux_job_v2.yml@main
@@ -119,9 +118,9 @@ jobs:
         CHANGED_FILES="${{ needs.get-changed-files.outputs.changed-files }}"
         echo "Running all other linters"
         if [ "$CHANGED_FILES" = '*' ]; then
-          ADDITIONAL_LINTRUNNER_ARGS="--skip CLANGTIDY,CLANGFORMAT,PYREFLY --all-files" .github/scripts/lintrunner.sh
+          ADDITIONAL_LINTRUNNER_ARGS="--skip CLANGTIDY,CLANGFORMAT,MYPY,MYPYSTRICT,PYREFLY --all-files" .github/scripts/lintrunner.sh
         else
-          ADDITIONAL_LINTRUNNER_ARGS="--skip CLANGTIDY,CLANGFORMAT,PYREFLY ${CHANGED_FILES}" .github/scripts/lintrunner.sh
+          ADDITIONAL_LINTRUNNER_ARGS="--skip CLANGTIDY,CLANGFORMAT,MYPY,MYPYSTRICT,PYREFLY ${CHANGED_FILES}" .github/scripts/lintrunner.sh
         fi
 
   quick-checks:

.github/workflows/linux-aarch64.yml

@@ -33,7 +33,7 @@ jobs:
     with:
       runner_prefix: ${{ needs.get-label-type.outputs.label-type }}
       build-environment: linux-jammy-aarch64-py3.10
-      docker-image-name: ci-image:pytorch-linux-jammy-aarch64-py3.10-gcc13
+      docker-image-name: ci-image:pytorch-linux-jammy-aarch64-py3.10-gcc11
       runner: linux.arm64.m7g.4xlarge
       test-matrix: |
         { include: [

.github/workflows/nightly.yml

@@ -41,7 +41,7 @@ jobs:
     uses: ./.github/workflows/_linux-build.yml
     needs: get-label-type
     with:
-      runner_prefix: "${{ needs.get-label-type.outputs.label-type }}"
+      runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge"
       build-environment: linux-jammy-py3.10-gcc11
       docker-image-name: ci-image:pytorch-linux-jammy-py3.10-gcc11
     secrets: inherit

.github/workflows/operator_benchmark.yml

@@ -60,7 +60,7 @@ jobs:
     with:
       build-environment: linux-jammy-aarch64-py3.10
       runner: linux.arm64.m7g.4xlarge
-      docker-image-name: ci-image:pytorch-linux-jammy-aarch64-py3.10-gcc13
+      docker-image-name: ci-image:pytorch-linux-jammy-aarch64-py3.10-gcc11
       test-matrix: |
         { include: [
           { config: "cpu_operator_benchmark_short", shard: 1, num_shards: 1, runner: "linux.arm64.m8g.4xlarge" },

.github/workflows/pull.yml

@@ -66,10 +66,10 @@ jobs:
           { config: "default", shard: 5, num_shards: 5, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
           { config: "docs_test", shard: 1, num_shards: 1,  runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
           { config: "jit_legacy", shard: 1, num_shards: 1, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
-          { config: "backwards_compat", shard: 1, num_shards: 1, runner: "${{ needs.get-label-type.outputs.label-type }}linux.c7i.2xlarge" },
+          { config: "backwards_compat", shard: 1, num_shards: 1, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
           { config: "distributed", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
           { config: "distributed", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
-          { config: "numpy_2_x", shard: 1, num_shards: 1, runner: "${{ needs.get-label-type.outputs.label-type }}linux.c7i.2xlarge" },
+          { config: "numpy_2_x", shard: 1, num_shards: 1, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
         ]}
     secrets: inherit
 
@@ -167,8 +167,8 @@ jobs:
       docker-image-name: ci-image:pytorch-linux-jammy-py3-clang12-onnx
       test-matrix: |
         { include: [
-          { config: "default", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.c7i.2xlarge" },
-          { config: "default", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.c7i.2xlarge" },
+          { config: "default", shard: 1, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
+          { config: "default", shard: 2, num_shards: 2, runner: "${{ needs.get-label-type.outputs.label-type }}linux.2xlarge" },
         ]}
     secrets: inherit
 

.github/workflows/rocm.yml

@@ -3,14 +3,13 @@ name: rocm
 on:
   push:
     branches:
+      - main
       - release/*
     tags:
       - ciflow/rocm/*
   workflow_dispatch:
   schedule:
     - cron: 29 8 * * *  # about 1:29am PDT
-    - cron: 0 */3 * * *
-
 
 concurrency:
   group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref_name }}-${{ github.ref_type == 'branch' && github.sha }}-${{ github.event_name == 'workflow_dispatch' }}-${{ github.event_name == 'schedule' }}

.github/workflows/trunk.yml

@@ -204,7 +204,6 @@ jobs:
         { include: [
           { config: "default", shard: 1, num_shards: 2, runner: "linux.rocm.gpu.gfx942.1" },
           { config: "default", shard: 2, num_shards: 2, runner: "linux.rocm.gpu.gfx942.1" },
-          { config: "distributed", shard: 1, num_shards: 1, runner: "linux.rocm.gpu.gfx942.4" },
         ]}
     secrets: inherit
 
@@ -222,7 +221,7 @@ jobs:
       build-environment: linux-jammy-rocm-py3.10
       docker-image: ${{ needs.linux-jammy-rocm-py3_10-build.outputs.docker-image }}
       test-matrix: ${{ needs.linux-jammy-rocm-py3_10-build.outputs.test-matrix }}
-      tests-to-include: "test_nn test_torch test_cuda test_ops test_unary_ufuncs test_binary_ufuncs test_autograd inductor/test_torchinductor distributed/test_c10d_common distributed/test_c10d_nccl"
+      tests-to-include: "test_nn test_torch test_cuda test_ops test_unary_ufuncs test_binary_ufuncs test_autograd inductor/test_torchinductor"
     secrets: inherit
 
   inductor-build:

.gitignore

@@ -127,7 +127,6 @@ torch/test/
 torch/utils/benchmark/utils/valgrind_wrapper/callgrind.h
 torch/utils/benchmark/utils/valgrind_wrapper/valgrind.h
 torch/version.py
-torch/_inductor/kernel/vendored_templates/*
 minifier_launcher.py
 aten/src/ATen/native/transformers/hip/flash_attn/ck/fmha_fwd_d*
 aten/src/ATen/native/transformers/hip/flash_attn/ck/fmha_bwd_d*
@@ -144,7 +143,6 @@ scripts/release_notes/*.json
 sccache-stats*.json
 lint.json
 merge_record.json
-.github/scripts/nightly_source_matrix.json
 
 # These files get copied over on invoking setup.py
 torchgen/packaged/*
@@ -399,4 +397,3 @@ CLAUDE.local.md
 /test_*.py
 /debug_*.py
 CLAUDE_CONTEXT/
-/.claude/settings.local.json

.lintrunner.toml

@@ -121,6 +121,94 @@ command = [
 ]
 is_formatter = true
 
+[[linter]]
+code = 'MYPY'
+include_patterns = [
+    'setup.py',
+    'functorch/dim/**/*.py',
+    'torch/**/*.py',
+    'torch/**/*.pyi',
+    'caffe2/**/*.py',
+    'caffe2/**/*.pyi',
+    'test/test_bundled_images.py',
+    'test/test_bundled_inputs.py',
+    'test/test_complex.py',
+    'test/test_datapipe.py',
+    'test/test_futures.py',
+    'test/test_numpy_interop.py',
+    'test/test_torch.py',
+    'test/test_type_hints.py',
+    'test/test_type_info.py',
+    'test/test_utils.py',
+]
+exclude_patterns = [
+    '**/fb/**',
+]
+command = [
+    'python3',
+    'tools/linter/adapters/mypy_linter.py',
+    '--config=mypy.ini',
+    '--',
+    '@{{PATHSFILE}}'
+]
+init_command = [
+    'python3',
+    'tools/linter/adapters/pip_init.py',
+    '--dry-run={{DRYRUN}}',
+    'numpy==1.26.4 ; python_version >= "3.10" and python_version <= "3.11"',
+    'numpy==2.1.0 ; python_version >= "3.12"',
+    'expecttest==0.3.0',
+    'mypy==1.16.0',
+    'sympy==1.13.3',
+    'types-requests==2.27.25',
+    'types-pyyaml==6.0.2',
+    'types-tabulate==0.8.8',
+    'types-protobuf==5.29.1.20250403',
+    'types-setuptools==79.0.0.20250422',
+    'types-jinja2==2.11.9',
+    'types-colorama==0.4.6',
+    'filelock==3.18.0',
+    'junitparser==2.1.1',
+    'rich==14.1.0',
+    'pyyaml==6.0.2',
+    'optree==0.13.0',
+    'dataclasses-json==0.6.7',
+    'pandas==2.2.3',
+]
+
+[[linter]]
+code = 'MYPYSTRICT'
+include_patterns = [
+    '.github/**/*.py',
+    'benchmarks/instruction_counts/**/*.py',
+    'tools/**/*.py',
+    'torchgen/**/*.py',
+    'torch/utils/_pytree.py',
+    'torch/utils/_cxx_pytree.py',
+    'torch/utils/benchmark/utils/common.py',
+    'torch/utils/benchmark/utils/timer.py',
+    'torch/utils/benchmark/utils/valgrind_wrapper/**/*.py',
+]
+exclude_patterns = [
+    # (linbinyu) copied from internal repo
+    '**/fb/**',
+    'tools/code_analyzer/gen_operators_yaml.py',
+    'tools/dynamo/verify_dynamo.py',
+    'tools/gen_vulkan_spv.py',
+    'tools/test/gen_operators_yaml_test.py',
+    'tools/test/gen_oplist_test.py',
+    'tools/test/test_selective_build.py',
+    'tools/experimental/torchfuzz/**',
+]
+command = [
+    'python3',
+    'tools/linter/adapters/mypy_linter.py',
+    '--config=mypy-strict.ini',
+    '--code=MYPYSTRICT',
+    '--',
+    '@{{PATHSFILE}}'
+]
+
 
 [[linter]]
 code = 'PYREFLY'
@@ -142,7 +230,6 @@ init_command = [
     'python3',
     'tools/linter/adapters/pip_init.py',
     '--dry-run={{DRYRUN}}',
-    'numpy==1.26.4 ; python_version >= "3.10" and python_version <= "3.11"',
     'numpy==2.1.0 ; python_version >= "3.12"',
     'expecttest==0.3.0',
     'pyrefly==0.36.2',
@@ -211,6 +298,7 @@ exclude_patterns = [
     '**/*pb.h',
     '**/*inl.h',
     'aten/src/ATen/cpu/FlushDenormal.cpp',
+    'aten/src/ATen/cpu/Utils.cpp',
     'aten/src/ATen/cpu/vml.h',
     'aten/src/ATen/CPUFixedAllocator.h',
     'aten/src/ATen/Parallel*.h',
@@ -229,6 +317,8 @@ exclude_patterns = [
     'c10/util/win32-headers.h',
     'c10/test/**/*.h',
     'third_party/**/*',
+    'torch/csrc/api/include/torch/nn/modules/common.h',
+    'torch/csrc/api/include/torch/linalg.h',
     'torch/csrc/autograd/generated/**',
     'torch/csrc/distributed/**/*.cu',
     'torch/csrc/distributed/c10d/WinSockUtils.hpp',
@@ -240,6 +330,7 @@ exclude_patterns = [
     'torch/csrc/utils/generated_serialization_types.h',
     'torch/csrc/utils/pythoncapi_compat.h',
     'torch/csrc/inductor/aoti_runtime/sycl_runtime_wrappers.h',
+    'aten/src/ATen/ExpandBase.h',
 ]
 init_command = [
     'python3',

CMakeLists.txt

@@ -234,17 +234,7 @@ option(USE_COLORIZE_OUTPUT "Colorize output during compilation" ON)
 option(USE_ASAN "Use Address+Undefined Sanitizers" OFF)
 option(USE_LSAN "Use Leak Sanitizer" OFF)
 option(USE_TSAN "Use Thread Sanitizer" OFF)
-
-# Track whether USE_CUDA was explicitly set by the user (before option() is called)
-# If USE_CUDA is already defined in cache, it means user explicitly set it
-if(DEFINED CACHE{USE_CUDA})
-  set(_USE_CUDA_EXPLICITLY_SET TRUE)
-else()
-  set(_USE_CUDA_EXPLICITLY_SET FALSE)
-endif()
-
 option(USE_CUDA "Use CUDA" ON)
-
 option(USE_XPU "Use XPU" ON)
 cmake_dependent_option(
   BUILD_LAZY_CUDA_LINALG "Build cuda linalg ops as separate library" ON

CONTRIBUTING.md

@@ -11,6 +11,7 @@ aspects of contributing to PyTorch.
 <!-- toc -->
 
 - [Developing PyTorch](#developing-pytorch)
+  - [Setup the development environment](#setup-the-development-environment)
   - [Tips and Debugging](#tips-and-debugging)
 - [Nightly Checkout & Pull](#nightly-checkout--pull)
 - [Codebase structure](#codebase-structure)
@@ -18,7 +19,7 @@ aspects of contributing to PyTorch.
   - [Python Unit Testing](#python-unit-testing)
   - [Better local unit tests with `pytest`](#better-local-unit-tests-with-pytest)
   - [Local linting](#local-linting)
-    - [Running `pyrefly`](#running-pyrefly)
+    - [Running `mypy`](#running-mypy)
   - [C++ Unit Testing](#c-unit-testing)
   - [Run Specific CI Jobs](#run-specific-ci-jobs)
 - [Merging your Change](#merging-your-change)
@@ -66,6 +67,23 @@ aspects of contributing to PyTorch.
 
 Follow the instructions for [installing PyTorch from source](https://github.com/pytorch/pytorch#from-source). If you get stuck when developing PyTorch on your machine, check out the [tips and debugging](#tips-and-debugging) section below for common solutions.
 
+### Setup the development environment
+
+First, you need to [fork the PyTorch project on GitHub](https://github.com/pytorch/pytorch/fork) and follow the instructions at [Connecting to GitHub with SSH](https://docs.github.com/en/authentication/connecting-to-github-with-ssh) to setup your SSH authentication credentials.
+
+Then clone the PyTorch project and setup the development environment:
+
+```bash
+git clone git@github.com:<USERNAME>/pytorch.git
+cd pytorch
+git remote add upstream git@github.com:pytorch/pytorch.git
+
+make setup-env
+# Or run `make setup-env-cuda` for pre-built CUDA binaries
+# Or run `make setup-env-rocm` for pre-built ROCm binaries
+source venv/bin/activate  # or `. .\venv\Scripts\activate` on Windows
+```
+
 ### Tips and Debugging
 
 * If you want to have no-op incremental rebuilds (which are fast), see [Make no-op build fast](#make-no-op-build-fast) below.
@@ -281,7 +299,7 @@ dependencies as well as the nightly binaries into the repo directory.
 **Prerequisites**:
 The following packages should be installed with `pip`:
 - `expecttest` and `hypothesis` - required to run tests
-- `pyrefly` - recommended for type checking. [Pyrefly](https://pyrefly.org/)
+- `mypy` - recommended for linting
 - `pytest` - recommended to run tests more selectively
 Running
 ```
@@ -350,32 +368,15 @@ make lint
 
 Learn more about the linter on the [lintrunner wiki page](https://github.com/pytorch/pytorch/wiki/lintrunner)
 
-#### Running `pyrefly`
+#### Running `mypy`
 
-[Pyrefly](https://pyrefly.org/) is a high-performance static type checker for Python. It provides fast type checking along with IDE features like autocomplete and instant error feedback.
-
-PyTorch uses Pyrefly for type checking across the codebase. The configuration is managed in `pyrefly.toml` at the root of the repository.
-
-**Getting Started with Pyrefly:**
-
-To run type checking on the PyTorch codebase:
-```bash
-pyrefly check
-```
-
-For more detailed error information with summaries:
-```bash
-pyrefly check --summarize-errors
-```
-
-**Learn More:**
-- [Pyrefly Configuration](https://pyrefly.org/en/docs/configuration/) - Detailed configuration options
-- [Pyrefly IDE Features](https://pyrefly.org/en/docs/IDE-features/) - Set up Pyrefly in your editor for real-time type checking
-- [Python Typing Tutorial](https://pyrefly.org/en/docs/typing-for-python-developers/) - Learn about Python type annotations
+`mypy` is an optional static type checker for Python. We have multiple `mypy`
+configs for the PyTorch codebase that are automatically validated against whenever the linter is run.
 
 See [Guide for adding type annotations to
 PyTorch](https://github.com/pytorch/pytorch/wiki/Guide-for-adding-type-annotations-to-PyTorch)
-for PyTorch-specific guidance on how to set up `pyrefly` and tackle type annotation tasks in this codebase.
+for more information on how to set up `mypy` and tackle type annotation
+tasks.
 
 ### C++ Unit Testing
 

SECURITY.md

@@ -1,7 +1,7 @@
 # Security Policy
 
  - [**Reporting a Vulnerability**](#reporting-a-vulnerability)
- - [**Using PyTorch Securely**](#using-pytorch-securely)
+ - [**Using Pytorch Securely**](#using-pytorch-securely)
    - [Untrusted models](#untrusted-models)
    - [TorchScript models](#torchscript-models)
    - [Untrusted inputs](#untrusted-inputs)
@@ -10,28 +10,28 @@
 - [**CI/CD security principles**](#cicd-security-principles)
 ## Reporting Security Issues
 
-Beware that none of the topics under [Using PyTorch Securely](#using-pytorch-securely) are considered vulnerabilities of PyTorch.
+Beware that none of the topics under [Using Pytorch Securely](#using-pytorch-securely) are considered vulnerabilities of Pytorch.
 
 However, if you believe you have found a security vulnerability in PyTorch, we encourage you to let us know right away. We will investigate all legitimate reports and do our best to quickly fix the problem.
 
 Please report security issues using https://github.com/pytorch/pytorch/security/advisories/new
 
-All reports submitted through the security advisories mechanism would **either be made public or dismissed by the team within 90 days of the submission**. If advisory has been closed on the grounds that it is not a security issue, please do not hesitate to create an [new issue](https://github.com/pytorch/pytorch/issues/new?template=bug-report.yml) as it is still likely a valid issue within the framework.
+All reports submitted thru the security advisories mechanism would **either be made public or dismissed by the team within 90 days of the submission**. If advisory has been closed on the grounds that it is not a security issue, please do not hesitate to create an [new issue](https://github.com/pytorch/pytorch/issues/new?template=bug-report.yml) as it is still likely a valid issue within the framework.
 
 Please refer to the following page for our responsible disclosure policy, reward guidelines, and those things that should not be reported:
 
 https://www.facebook.com/whitehat
 
 
-## Using PyTorch Securely
-**PyTorch models are programs**, so treat its security seriously -- running untrusted models is equivalent to running untrusted code. In general we recommend that model weights and the python code for the model are distributed independently. That said, be careful about where you get the python code from and who wrote it (preferentially check for a provenance or checksums, do not run any pip installed package).
+## Using Pytorch Securely
+**Pytorch models are programs**, so treat its security seriously -- running untrusted models is equivalent to running untrusted code. In general we recommend that model weights and the python code for the model are distributed independently. That said, be careful about where you get the python code from and who wrote it (preferentially check for a provenance or checksums, do not run any pip installed package).
 
 ### Untrusted models
 Be careful when running untrusted models. This classification includes models created by unknown developers or utilizing data obtained from unknown sources[^data-poisoning-sources].
 
 **Prefer to execute untrusted models within a secure, isolated environment such as a sandbox** (e.g., containers, virtual machines). This helps protect your system from potentially malicious code. You can find further details and instructions in [this page](https://developers.google.com/code-sandboxing).
 
-**Be mindful of risky model formats**. Give preference to share and load weights with the appropriate format for your use case. [Safetensors](https://huggingface.co/docs/safetensors/en/index) gives the most safety but is the most restricted in what it supports. [`torch.load`](https://pytorch.org/docs/stable/generated/torch.load.html#torch.load) has a significantly larger surface of attack but is more flexible in what it can serialize. See the documentation for more details.
+**Be mindful of risky model formats**. Give preference to share and load weights with the appropriate format for your use case. [safetensors](https://huggingface.co/docs/safetensors/en/index) gives the most safety but is the most restricted in what it supports. [`torch.load`](https://pytorch.org/docs/stable/generated/torch.load.html#torch.load) has a significantly larger surface of attack but is more flexible in what it can serialize. See the documentation for more details.
 
 Even for more secure serialization formats, unexpected inputs to the downstream system can cause diverse security threats (e.g. denial of service, out of bound reads/writes) and thus we recommend extensive validation of any untrusted inputs.
 
@@ -43,7 +43,7 @@ Important Note: The trustworthiness of a model is not binary. You must always de
 
 ### TorchScript models
 
-TorchScript models should be treated the same way as locally executable code from an unknown source. Only run TorchScript models if you trust the provider. Please note, that tools for introspecting TorchScript models (such as `torch.utils.model_dump`) may also execute partial or full code stored in those models, therefore they should be used only if you trust the provider of the binary you are about to load.
+TorchScript models should treated the same way as locally executable code from an unknown source. Only run TorchScript models if you trust the provider. Please note, that tools for introspecting TorchScript models (such as `torch.utils.model_dump`) may also execute partial or full code stored in those models, therefore they should be used only if you trust the provider of the binary you are about to load.
 
 ### Untrusted inputs during training and prediction
 
@@ -59,9 +59,9 @@ If applicable, prepare your model against bad inputs and prompt injections. Some
 
 ### Data privacy
 
-**Take special security measures if you train your models with sensitive data**. Prioritize [sandboxing](https://developers.google.com/code-sandboxing) your models and:
-- Do not feed sensitive data to an untrusted model (even if runs in a sandboxed environment)
-- If you consider publishing a model that was partially trained with sensitive data, be aware that data can potentially be recovered from the trained weights (especially if the model overfits).
+**Take special security measures if your model if you train models with sensitive data**. Prioritize [sandboxing](https://developers.google.com/code-sandboxing) your models and:
+- Do not feed sensitive data to untrusted model (even if runs in a sandboxed environment)
+- If you consider publishing a model that was partially trained with sensitive data, be aware that data can potentially be recovered from the trained weights (especially if model overfits).
 
 ### Using distributed features
 

aten/src/ATen/CPUGeneratorImpl.cpp

@@ -181,7 +181,7 @@ c10::intrusive_ptr<c10::TensorImpl> CPUGeneratorImpl::get_state() const {
   static const size_t size = sizeof(CPUGeneratorImplState);
   static_assert(std::is_standard_layout_v<CPUGeneratorImplState>, "CPUGeneratorImplState is not a PODType");
 
-  auto state_tensor = at::detail::empty_cpu({static_cast<int64_t>(size)}, ScalarType::Byte, std::nullopt, std::nullopt, std::nullopt, std::nullopt);
+  auto state_tensor = at::detail::empty_cpu({(int64_t)size}, ScalarType::Byte, std::nullopt, std::nullopt, std::nullopt, std::nullopt);
   auto rng_state = state_tensor.data_ptr();
 
   // accumulate generator data to be copied into byte tensor

aten/src/ATen/Context.cpp

@@ -23,6 +23,8 @@ C10_DIAGNOSTIC_POP()
 #endif
 namespace at {
 
+namespace {
+
 /*
   These const variables defined the fp32 precisions for different backend
   We have "generic", "cuda", "mkldnn" backend now and we can choose fp32
@@ -39,6 +41,16 @@ namespace at {
                 ->rnn
 */
 
+  C10_ALWAYS_INLINE void warn_deprecated_fp32_precision_api(){
+    TORCH_WARN_ONCE(
+      "Please use the new API settings to control TF32 behavior, such as torch.backends.cudnn.conv.fp32_precision = 'tf32' "
+      "or torch.backends.cuda.matmul.fp32_precision = 'ieee'. Old settings, e.g, torch.backends.cuda.matmul.allow_tf32 = True, "
+      "torch.backends.cudnn.allow_tf32 = True, allowTF32CuDNN() and allowTF32CuBLAS() will be deprecated after Pytorch 2.9. Please see "
+      "https://pytorch.org/docs/main/notes/cuda.html#tensorfloat-32-tf32-on-ampere-and-later-devices"
+    );
+  }
+} // namespace
+
 Float32Backend str2backend(const std::string& name) {
   if (name == "generic")
     return Float32Backend::GENERIC;
@@ -194,6 +206,7 @@ bool Context::allowTF32CuDNN(std::optional<Float32Op> op) const {
   } else {
     return float32Precision(Float32Backend::CUDA, op.value()) == Float32Precision::TF32;
   }
+  warn_deprecated_fp32_precision_api();
   return allow_tf32_cudnn;
 }
 
@@ -201,6 +214,7 @@ void Context::setAllowTF32CuDNN(bool b) {
   setFloat32Precision(Float32Backend::CUDA, Float32Op::RNN, b ? Float32Precision::TF32 : Float32Precision::NONE);
   setFloat32Precision(Float32Backend::CUDA, Float32Op::CONV, b ? Float32Precision::TF32 : Float32Precision::NONE);
   allow_tf32_cudnn = b;
+  warn_deprecated_fp32_precision_api();
 }
 
 void Context::setSDPPriorityOrder(const std::vector<int64_t>& order) {
@@ -209,7 +223,7 @@ void Context::setSDPPriorityOrder(const std::vector<int64_t>& order) {
     "setSDPPriority order expected ", sdp_priority_order.size() - 1, " but got ",
     at::num_sdp_backends, " unique backends specified in priority order.");
   for (uint32_t i = 0; i < order.size(); i++) {
-    sdp_priority_order[i] = static_cast<at::SDPBackend>(order[i]);
+    sdp_priority_order[i] = (at::SDPBackend) order[i];
   }
 }
 
@@ -311,6 +325,7 @@ bool Context::allowTF32CuBLAS() const {
       "Current status indicate that you have used mix of the legacy and new APIs to set the TF32 status for cublas matmul. ",
       "We suggest only using the new API to set the TF32 flag. See also: ",
       "https://pytorch.org/docs/main/notes/cuda.html#tensorfloat-32-tf32-on-ampere-and-later-devices");
+  warn_deprecated_fp32_precision_api();
   return allow_tf32_new;
 }
 
@@ -334,6 +349,7 @@ Float32MatmulPrecision Context::float32MatmulPrecision() const {
       "Current status indicate that you have used mix of the legacy and new APIs to set the matmul precision. ",
       "We suggest only using the new API for matmul precision. See also: ",
       "https://pytorch.org/docs/main/notes/cuda.html#tensorfloat-32-tf32-on-ampere-and-later-devices");
+  warn_deprecated_fp32_precision_api();
   return float32_matmul_precision;
 }
 
@@ -361,6 +377,7 @@ Float32Precision Context::float32Precision(Float32Backend backend, Float32Op op)
 
 void Context::setFloat32MatmulPrecision(const std::string &s) {
   auto match = [this](const std::string & s_) {
+    warn_deprecated_fp32_precision_api();
     // TODO: consider if CuDNN field needs to also be set for potential future CuDNN ops like multi-headed attention
     if (s_ == "highest") {
       float32_matmul_precision = at::Float32MatmulPrecision::HIGHEST;
@@ -808,14 +825,6 @@ void Context::setDisplayVmapFallbackWarnings(bool enabled) {
   display_vmap_fallback_warnings_ = enabled;
 }
 
-bool Context::warnOnAccumulateGradStreamMismatch() const {
-  return warn_on_accumulate_grad_stream_mismatch_;
-}
-
-void Context::setWarnOnAccumulateGradStreamMismatch(bool enabled) {
-  warn_on_accumulate_grad_stream_mismatch_ = enabled;
-}
-
 bool Context::isDefaultMobileCPUAllocatorSet() {
   return prev_allocator_ptr_ != nullptr;
 }

aten/src/ATen/Context.h

@@ -404,9 +404,6 @@ class TORCH_API Context {
   void setDisplayVmapFallbackWarnings(bool enabled);
   bool areVmapFallbackWarningsEnabled() const;
 
-  void setWarnOnAccumulateGradStreamMismatch(bool enabled);
-  bool warnOnAccumulateGradStreamMismatch() const;
-
   bool isDefaultMobileCPUAllocatorSet();
   void setDefaultMobileCPUAllocator();
   void unsetDefaultMobileCPUAllocator();
@@ -497,7 +494,6 @@ class TORCH_API Context {
   bool release_original_weights = false;
 #endif
   bool display_vmap_fallback_warnings_ = false;
-  bool warn_on_accumulate_grad_stream_mismatch_ = true;
   std::atomic<at::QEngine> quantized_engine = at::QEngine::NoQEngine;
   bool enable_sparse_tensor_invariant_checks = false;
   bool allow_fp16_reduction_cpu = false;

aten/src/ATen/Dispatch.h

@@ -197,7 +197,6 @@ inline at::ScalarType scalar_type(at::ScalarType s) {
     /* don't use TYPE again in case it is an expensive or side-effect op */ \
     at::ScalarType _st = ::detail::scalar_type(the_type);                   \
     RECORD_KERNEL_FUNCTION_DTYPE(at_dispatch_name, _st);                    \
-    C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum")             \
     switch (_st) {                                                          \
       __VA_ARGS__                                                           \
       default:                                                              \
@@ -209,7 +208,6 @@ inline at::ScalarType scalar_type(at::ScalarType s) {
             toString(_st),                                                  \
             "'");                                                           \
     }                                                                       \
-    C10_DIAGNOSTIC_POP()                                                    \
   }()
 
 #define AT_DISPATCH_CASE_FLOATING_TYPES(...)            \

aten/src/ATen/MapAllocator.cpp

@@ -252,13 +252,13 @@ MapAllocator::MapAllocator(WithFd /*unused*/, std::string_view filename, int fd,
     if (!(flags_ & ALLOCATOR_MAPPED_FROMFD)) {
       if (flags_ & ALLOCATOR_MAPPED_SHARED) {
         // NOLINTNEXTLINE(bugprone-assignment-in-if-condition)
-        if ((fd = open(filename_.c_str(), flags, static_cast<mode_t>(0600))) == -1) {
+        if ((fd = open(filename_.c_str(), flags, (mode_t)0600)) == -1) {
           TORCH_CHECK(false, "unable to open file <", filename_, "> in read-write mode: ", c10::utils::str_error(errno), " (", errno, ")");
         }
       } else if (flags_ & ALLOCATOR_MAPPED_SHAREDMEM) {
 #ifdef HAVE_SHM_OPEN
         // NOLINTNEXTLINE(bugprone-assignment-in-if-condition)
-        if((fd = shm_open(filename_.c_str(), flags, static_cast<mode_t>(0600))) == -1) {
+        if((fd = shm_open(filename_.c_str(), flags, (mode_t)0600)) == -1) {
           TORCH_CHECK(false, "unable to open shared memory object <", filename_, "> in read-write mode: ", c10::utils::str_error(errno), " (", errno, ")");
         }
 #else
@@ -503,7 +503,7 @@ RefcountedMapAllocator::RefcountedMapAllocator(WithFd /*unused*/, const char *fi
 
 void RefcountedMapAllocator::initializeAlloc() {
   TORCH_CHECK(base_ptr_, "base_ptr_ is null");
-  MapInfo *map_info = static_cast<MapInfo*>(base_ptr_);
+  MapInfo *map_info = (MapInfo*)base_ptr_;
 
 #ifdef _WIN32
   ReleaseContext* r_ctx = new ReleaseContext;
@@ -539,7 +539,7 @@ void RefcountedMapAllocator::close() {
   }
 #else /* _WIN32 */
 
-  MapInfo *info = static_cast<MapInfo*>(data);
+  MapInfo *info = (MapInfo*)(data);
   if (--info->refcount == 0) {
 #ifdef HAVE_SHM_UNLINK
     if (shm_unlink(filename_.c_str()) == -1) {

aten/src/ATen/TensorIterator.cpp

@@ -862,7 +862,7 @@ void TensorIteratorBase::narrow(int dim, int64_t start, int64_t size) {
   shape_[dim] = size;
   view_offsets_[dim] += start;
   for (auto& op : operands_) {
-    op.data = (static_cast<char*>(op.data)) + op.stride_bytes[dim] * start;
+    op.data = ((char*)op.data) + op.stride_bytes[dim] * start;
   }
   if (size == 1 && !is_reduction_) {
     coalesce_dimensions();
@@ -873,7 +873,7 @@ void TensorIteratorBase::select_all_keeping_dim(int start_dim, IntArrayRef indic
   TORCH_INTERNAL_ASSERT(start_dim <= ndim());
   for (const auto i : c10::irange(start_dim, ndim())) {
     for (auto& op : operands_) {
-      op.data = (static_cast<char*>(op.data)) + op.stride_bytes[i] * indices[i - start_dim];
+      op.data = ((char*)op.data) + op.stride_bytes[i] * indices[i - start_dim];
     }
     shape_[i] = 1;
   }

aten/src/ATen/TensorIteratorInternal.h

@@ -41,7 +41,7 @@ inline void serial_for_each(
     IntArrayRef strides,
     char** base_ptrs,
     size_t ntensors,
-    TensorIteratorBase::loop2d_t loop,
+    typename TensorIteratorBase::loop2d_t loop,
     Range range) {
   const auto ndim = shape.size();
   TORCH_INTERNAL_ASSERT_DEBUG_ONLY(

aten/src/ATen/VmapModeRegistrations.cpp

@@ -72,16 +72,10 @@ TORCH_LIBRARY_IMPL(aten, VmapMode, m) {
   m.impl("random_", unsupportedRandomOp_<Tensor&, std::optional<Generator>>);
 
   m.impl("rand_like", unsupportedRandomOp<const Tensor&, TENSOROPTIONS, std::optional<MemoryFormat>>);
-  m.impl("rand_like.generator", unsupportedRandomOp<const Tensor&, std::optional<Generator>, TENSOROPTIONS, std::optional<MemoryFormat>>);
   m.impl("randn_like", unsupportedRandomOp<const Tensor&, TENSOROPTIONS, std::optional<MemoryFormat>>);
-  m.impl("randn_like.generator", unsupportedRandomOp<const Tensor&, std::optional<Generator>, TENSOROPTIONS, std::optional<MemoryFormat>>);
 
   m.impl("randint_like", unsupportedRandomOp<const Tensor&, int64_t, TENSOROPTIONS, std::optional<MemoryFormat>>);
-  m.impl("randint_like.Tensor", unsupportedRandomOp<const Tensor&, const Tensor&, TENSOROPTIONS, std::optional<MemoryFormat>>);
   m.impl("randint_like.low_dtype", unsupportedRandomOp<const Tensor&, int64_t, int64_t, TENSOROPTIONS, std::optional<MemoryFormat>>);
-  m.impl("randint_like.generator", unsupportedRandomOp<const Tensor&, int64_t, std::optional<Generator>, TENSOROPTIONS, std::optional<MemoryFormat>>);
-  m.impl("randint_like.Tensor_generator", unsupportedRandomOp<const Tensor&, const Tensor&, std::optional<Generator>, TENSOROPTIONS, std::optional<MemoryFormat>>);
-  m.impl("randint_like.low_generator_dtype", unsupportedRandomOp<const Tensor&, int64_t, int64_t, std::optional<Generator>, TENSOROPTIONS, std::optional<MemoryFormat>>);
 
   m.impl("rand", unsupportedRandomOp<IntArrayRef, TENSOROPTIONS>);
   m.impl("rand.generator", unsupportedRandomOp<IntArrayRef, std::optional<Generator>, TENSOROPTIONS>);

aten/src/ATen/core/CachingHostAllocator.h

@@ -10,7 +10,9 @@
 #include <optional>
 
 #include <deque>
+#include <vector>
 #include <mutex>
+#include <shared_mutex>
 
 C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wunused-parameter")
 namespace at {
@@ -36,6 +38,9 @@ struct HostBlock {
   bool allocated_{false}; // in-use flag
   size_t event_count_{0}; // number of related events
   ska::flat_hash_set<S> streams_; // streams on which the block was used
+  c10::MempoolId_t owning_pool_{0,0};
+  bool was_allocated_during_stream_capture_;
+  std::vector<std::tuple<void *, size_t, std::string>> sections_read_under_stream_capture_;
 };
 
 template <typename B>
@@ -220,11 +225,52 @@ struct alignas(hardware_destructive_interference_size) HostStatsStaged {
  * we hold the same locks as empty_cache, to ensure the fidelity of the stats.
  */
 
+// Generic per-pool structures for the host caching allocator. These are
+// templated so they can reference the allocator's template parameters B and E.
+template <typename S_, typename E_, typename B_>
+struct HostBlockPool {
+  HostBlockPool() = default;
+  HostBlockPool(const HostBlockPool&) = delete;
+  HostBlockPool& operator=(const HostBlockPool&) = delete;
+
+  alignas(hardware_destructive_interference_size) std::mutex blocks_mutex_;
+  ska::flat_hash_set<B_*> blocks_; // all blocks in this pool
+  ska::flat_hash_map<void*, B_*> ptr_to_block_;
+
+  // Per-size free lists guarded by their own mutexes.
+  alignas(hardware_destructive_interference_size) std::vector<FreeBlockList<B_>> free_list_ =
+      std::vector<FreeBlockList<B_>>(MAX_SIZE_INDEX);
+
+  // Events pending for blocks in this pool.
+  alignas(hardware_destructive_interference_size) std::mutex events_mutex_;
+  std::deque<std::pair<E_, B_*>> events_;
+};
+
+// The HostBlockPool owned by one HostPrivatePool cannot share blocks
+// with any other HostPrivatePool. This invariant is necessary to
+// implement cuda graph capture correctly, which depends upon virtual
+// addresses staying alive for the duration of a graph's existence.
+template <typename S_, typename E_, typename B_>
+struct HostPrivatePool {
+  explicit HostPrivatePool(c10::MempoolId_t id_) : id(id_) {}
+  HostPrivatePool(const HostPrivatePool&) = delete;
+  HostPrivatePool& operator=(const HostPrivatePool&) = delete;
+
+  c10::MempoolId_t id;
+  int use_count{1};
+  HostBlockPool<S_, E_, B_> blocks;
+};
+
+
 template <
     typename S,
     typename E,
     typename B = HostBlock<S>>
 struct CachingHostAllocatorImpl {
+
+  using BlockPool = HostBlockPool<S, E, B>;
+  using PrivatePool = HostPrivatePool<S, E, B>;
+
   virtual ~CachingHostAllocatorImpl() {
     active_ = false;
     if (pinned_use_background_threads()) {
@@ -232,37 +278,47 @@ struct CachingHostAllocatorImpl {
     }
   }
 
- public:
   // return data_ptr and block pair.
   virtual std::pair<void*, void*> allocate(size_t size) {
     if (size == 0) {
       return {nullptr, nullptr};
     }
 
+    auto&& [mempool_id, pool] = get_allocation_pool(get_current_stream());
+
     // If we are using background threads, we can process events in the
     // background.
     if (!pinned_use_background_threads()) {
-      process_events();
+      process_events(pool);
     }
 
     // Round up the allocation to the nearest power of two to improve reuse.
     // These power of two sizes are also used to index into the free list.
     size_t roundSize = c10::llvm::PowerOf2Ceil(size);
 
-    // First, try to allocate from the free list
-    auto* block = get_free_block(roundSize);
+    // First, try to allocate from the free list of the chosen pool
+    auto* block = get_free_block(roundSize, pool);
     if (block) {
+      block->was_allocated_during_stream_capture_ = stream_is_capturing(get_current_stream());
       return {block->ptr_, reinterpret_cast<void*>(block)};
     }
 
     // Check in the recently freed blocks with pending events to see if we
     // can reuse them. Call get_free_block again after processing events
     if (pinned_use_background_threads()) {
+      process_events_for_specific_size(roundSize, pool);
+      block = get_free_block(roundSize, pool);
+      if (block) {
+        block->was_allocated_during_stream_capture_ = stream_is_capturing(get_current_stream());
+        return {block->ptr_, reinterpret_cast<void*>(block)};
+      }
+
       // Launch the background thread and process events in a loop.
       static bool background_thread_flag [[maybe_unused]] = [this] {
         getBackgroundThreadPool()->run([&]() {
           while (active_) {
-            process_events();
+            // Background thread conservatively processes default pool events.
+            process_events(default_pool_);
             std::this_thread::sleep_for(std::chrono::microseconds(100));
           }
         });
@@ -278,8 +334,9 @@ struct CachingHostAllocatorImpl {
     // Then, create a new block.
     block = new B(roundSize, ptr);
     block->allocated_ = true;
-
-    add_allocated_block(block);
+    block->owning_pool_ = mempool_id;
+    block->was_allocated_during_stream_capture_ = stream_is_capturing(get_current_stream());
+    add_allocated_block(block, pool);
     return {block->ptr_, reinterpret_cast<void*>(block)};
   }
 
@@ -309,75 +366,65 @@ struct CachingHostAllocatorImpl {
       }
     }
 
-    // Event recording must be done outside the mutex to avoid potential
-    // deadlocks (e.g., when Python GIL is involved)
-    for (auto stream : streams) {
-      record_stream(events, stream);
+    // cudaEventRecord() does not work for indicating when a block can
+    // be reused while stream capture is happening, since no actual
+    // GPU work happens during stream capture.
+    if (!block->was_allocated_during_stream_capture_) {
+      // Event recording must be done outside the mutex to avoid potential
+      // deadlocks (e.g., when Python GIL is involved)
+      for (auto stream : streams) {
+        record_stream(events, stream);
+      }
     }
 
     if (!events) {
+      auto& pool = pool_from_block(block);
       auto index = size_index(block->size_);
-      std::lock_guard<std::mutex> g(free_list_[index].mutex_);
-      free_list_[index].list_.push_back(block);
+      std::lock_guard<std::mutex> g(pool.free_list_[index].mutex_);
+      pool.free_list_[index].list_.push_back(block);
     } else {
+      TORCH_INTERNAL_ASSERT(!block->was_allocated_during_stream_capture_ || events->empty());
       // restore these events that record by used streams.
-      std::lock_guard<std::mutex> g(events_mutex_);
+      auto& pool = pool_from_block(block);
+      std::lock_guard<std::mutex> g(pool.events_mutex_);
       for (auto&& event : *events) {
-        events_.emplace_front(std::move(event), block);
+        pool.events_.emplace_front(std::move(event), block);
       }
     }
   }
 
   virtual bool record_event(void* ptr, void* ctx, c10::Stream s) {
-    S stream = S(s);
-    auto* block = reinterpret_cast<B*>(ctx);
-
-    // Note: we need to check if the passed-in `ctx` is valid. This is because
-    // `record_event` (via `CachingHostAllocator_recordEvent`) can be invoked on
-    // an arbitrary tensor, and is not guaranteed to correspond to a pinned
-    // memory allocation. Therefore, we need to check that `ctx` is valid before
-    // proceeding.
-    {
-      std::lock_guard<std::mutex> g(blocks_mutex_);
-      if (blocks_.find(block) != blocks_.end()) {
-        // Now we know this object is safe to access.
-        std::lock_guard<std::mutex> gb(block->mutex_);
-        TORCH_INTERNAL_ASSERT(block->allocated_);
-        block->streams_.insert(stream);
-        return true;
-      }
-      auto it = ptr_to_block_.find(ptr);
-      if (it != ptr_to_block_.end()) {
-        block = it->second;
-        std::lock_guard<std::mutex> g(block->mutex_);
-        TORCH_INTERNAL_ASSERT(block->allocated_);
-        block->streams_.insert(stream);
-        return true;
-      }
+    B *block = nullptr;
+    if (block_exists(ctx)) {
+      block = reinterpret_cast<B*>(ctx);
+    } else {
+      block = get_block_from_ptr(ptr);
     }
-
-    return false;
+    if (block == nullptr) {
+      return false;
+    }
+    S stream = S(s);
+    std::lock_guard<std::mutex> gb(block->mutex_);
+    TORCH_INTERNAL_ASSERT(block->allocated_);
+    block->streams_.insert(stream);
+    return true;
   }
 
-  virtual void empty_cache() {
-    // Flush any available blocks into the free_list.
-    process_events();
+  void free_from_pool(BlockPool &pool) {
+    for (size_t i = 0; i < pool.free_list_.size(); ++i) {
+      std::lock(pool.free_list_[i].mutex_, pool.blocks_mutex_);
+      std::lock_guard<std::mutex> gf(pool.free_list_[i].mutex_, std::adopt_lock);
+      std::lock_guard<std::mutex> gb(pool.blocks_mutex_, std::adopt_lock);
 
-    // Remove all elements from the free list, remove them from the blocks
-    // list, and free the associated pinned memory allocation. This requires
-    // concurrently holding both the free list mutexes and the blocks mutex, and
-    // is the only function that concurrently holds multiple mutexes.
-    for (size_t i = 0; i < free_list_.size(); ++i) {
-      std::lock(free_list_[i].mutex_, blocks_mutex_);
-      std::lock_guard<std::mutex> gf(free_list_[i].mutex_, std::adopt_lock);
-      std::lock_guard<std::mutex> gb(blocks_mutex_, std::adopt_lock);
-
-      std::vector<B*> blocks_to_remove(free_list_[i].list_.begin(), free_list_[i].list_.end());
-      free_list_[i].list_.clear();
+      std::vector<B*> blocks_to_remove(
+                                       pool.free_list_[i].list_.begin(), pool.free_list_[i].list_.end());
+      pool.free_list_[i].list_.clear();
 
       for (auto* block : blocks_to_remove) {
-        blocks_.erase(block);
-        ptr_to_block_.erase(block->ptr_);
+        // this suggests that you can be part of the free_list_
+        // while still being in blocks_. Hmm...
+        pool.blocks_.erase(block);
+        pool.ptr_to_block_.erase(block->ptr_);
         auto index = size_index(block->size_);
         free_block(block);
         stats_.allocations.decrease(1);
@@ -389,6 +436,30 @@ struct CachingHostAllocatorImpl {
     }
   }
 
+  // TODO: Rethink how this is implemented. Should it take a pool id
+  // like in CUDACachingAllocator?
+  virtual void empty_cache() {
+    // Flush available blocks in all pools into their free_lists.
+    process_events(default_pool_);
+
+    free_from_pool(default_pool_);
+    // Also flush and free from private pools that are marked freeable
+    {
+      std::shared_lock<std::shared_mutex> lg(instance_mutex_);
+      for (auto it = graph_pools_freeable_.begin(); it != graph_pools_freeable_.end();) {
+        process_events(it->second->blocks);
+        free_from_pool(it->second->blocks);
+        if (it->second->blocks.blocks_.empty()) {
+          auto erase_count = graph_pools_.erase(it->first);
+          TORCH_INTERNAL_ASSERT(erase_count == 1);
+          it = graph_pools_freeable_.erase(it);
+        } else {
+          ++it;
+        }
+      }
+    }
+  }
+
   inline size_t size_index(size_t size) {
     return c10::llvm::Log2_64_Ceil(size);
   }
@@ -421,10 +492,13 @@ struct CachingHostAllocatorImpl {
     // Accurate reading of memory stats requires concurrently holding both the
     // free list mutexes and the blocks mutex. Previously, this was only done in
     // empty_cache function.
-    for (size_t i = 0; i < free_list_.size(); ++i) {
-      std::lock(free_list_[i].mutex_, blocks_mutex_);
-      std::lock_guard<std::mutex> gf(free_list_[i].mutex_, std::adopt_lock);
-      std::lock_guard<std::mutex> gb(blocks_mutex_, std::adopt_lock);
+    for (size_t i = 0; i < default_pool_.free_list_.size(); ++i) {
+      std::lock(
+          default_pool_.free_list_[i].mutex_, default_pool_.blocks_mutex_);
+      std::lock_guard<std::mutex> gf(
+          default_pool_.free_list_[i].mutex_, std::adopt_lock);
+      std::lock_guard<std::mutex> gb(
+          default_pool_.blocks_mutex_, std::adopt_lock);
 
       // We collect the slow-path stats only once, since they are not collected
       // per bucket (we pick index 0 arbitrarily). These are also all the host
@@ -458,10 +532,13 @@ struct CachingHostAllocatorImpl {
     // Resetting accumulated memory stats requires concurrently holding both the
     // free list mutexes and the blocks mutex. Previously, this was only done in
     // empty_cache function.
-    for (size_t i = 0; i < free_list_.size(); ++i) {
-      std::lock(free_list_[i].mutex_, blocks_mutex_);
-      std::lock_guard<std::mutex> gf(free_list_[i].mutex_, std::adopt_lock);
-      std::lock_guard<std::mutex> gb(blocks_mutex_, std::adopt_lock);
+    for (size_t i = 0; i < default_pool_.free_list_.size(); ++i) {
+      std::lock(
+          default_pool_.free_list_[i].mutex_, default_pool_.blocks_mutex_);
+      std::lock_guard<std::mutex> gf(
+          default_pool_.free_list_[i].mutex_, std::adopt_lock);
+      std::lock_guard<std::mutex> gb(
+          default_pool_.blocks_mutex_, std::adopt_lock);
 
       if (i == 0) {
         stats_.allocations.reset_accumulated();
@@ -485,10 +562,13 @@ struct CachingHostAllocatorImpl {
     // Resetting peak memory stats requires concurrently holding both the
     // free list mutexes and the blocks mutex. Previously, this was only done in
     // empty_cache function.
-    for (size_t i = 0; i < free_list_.size(); ++i) {
-      std::lock(free_list_[i].mutex_, blocks_mutex_);
-      std::lock_guard<std::mutex> gf(free_list_[i].mutex_, std::adopt_lock);
-      std::lock_guard<std::mutex> gb(blocks_mutex_, std::adopt_lock);
+    for (size_t i = 0; i < default_pool_.free_list_.size(); ++i) {
+      std::lock(
+          default_pool_.free_list_[i].mutex_, default_pool_.blocks_mutex_);
+      std::lock_guard<std::mutex> gf(
+          default_pool_.free_list_[i].mutex_, std::adopt_lock);
+      std::lock_guard<std::mutex> gb(
+          default_pool_.blocks_mutex_, std::adopt_lock);
 
       if (i == 0) {
         stats_.allocations.reset_peak();
@@ -508,13 +588,12 @@ struct CachingHostAllocatorImpl {
     }
   }
 
- private:
-  virtual void add_allocated_block(B* block) {
-    std::lock_guard<std::mutex> g(blocks_mutex_);
-    blocks_.insert(block);
+  void add_allocated_block(B* block, BlockPool& pool) {
+    std::lock_guard<std::mutex> g(pool.blocks_mutex_);
+    pool.blocks_.insert(block);
     stats_.allocations.increase(1);
     stats_.allocated_bytes.increase(block->size_);
-    ptr_to_block_.insert({block->ptr_, block});
+    pool.ptr_to_block_.insert({block->ptr_, block});
 
     // Unfortunately, we have to, on the slow path, quickly
     // lock the bucket to record the allocation. This should
@@ -522,7 +601,7 @@ struct CachingHostAllocatorImpl {
     auto size = block->size_;
     auto index = size_index(size);
     {
-      std::lock_guard<std::mutex> g(free_list_[index].mutex_);
+      std::lock_guard<std::mutex> gf(pool.free_list_[index].mutex_);
       stats_.allocation_bucket_stats[index].increase(1);
       stats_.allocated_bytes_bucket_stats[index].increase(size);
       stats_.active_bucket_stats[index].increase(1);
@@ -530,46 +609,45 @@ struct CachingHostAllocatorImpl {
     }
   }
 
-  virtual B* get_free_block(size_t size) {
+  B* get_free_block(size_t size, BlockPool& pool) {
     auto index = size_index(size);
-    std::lock_guard<std::mutex> g(free_list_[index].mutex_);
-    if (!free_list_[index].list_.empty()) {
-      B* block = free_list_[index].list_.back();
-      free_list_[index].list_.pop_back();
+    std::lock_guard<std::mutex> g(pool.free_list_[index].mutex_);
+    if (!pool.free_list_[index].list_.empty()) {
+      B* block = pool.free_list_[index].list_.back();
+      pool.free_list_[index].list_.pop_back();
       block->allocated_ = true;
-      stats_.active_bucket_stats[index].increase(1);
-      stats_.active_bytes_bucket_stats[index].increase(size);
+      stats_.allocation_bucket_stats[index].increase(1);
+      stats_.allocated_bytes_bucket_stats[index].increase(size);
       return block;
     }
     return nullptr;
   }
 
-  virtual void process_events() {
-    // process all events until the last unready event, not for specific size.
-    process_events_for_specific_size(-1);
+ private:
+
+  void process_events(BlockPool& pool) {
+    // process all events in the given pool until the last unready event.
+    process_events_for_specific_size(-1, pool);
   }
 
   // If size is -1, process all events from backwards until the last unready
   // event. Otherwise, process events for a specific size and on first ready block
   // is found, add it to the free list and return.
-  virtual void process_events_for_specific_size(int64_t size) {
+  virtual void process_events_for_specific_size(int64_t size, BlockPool& pool) {
     size_t event_count = 0;
     size_t max_events = 0;
     {
-      std::lock_guard<std::mutex> g(events_mutex_);
-      max_events = events_.size();
+      std::lock_guard<std::mutex> g(pool.events_mutex_);
+      max_events = pool.events_.size();
     }
 
     while (true) {
-      // Avoid calling cudaEventDestroy while holding a mutex, so move
-      // intermediate events out of the lock into this object.
-      // process the last event
       std::optional<std::pair<E, B*>> processed;
       {
-        std::lock_guard<std::mutex> g(events_mutex_);
-        if (!events_.empty()) {
-          processed = std::move(events_.back());
-          events_.pop_back();
+        std::lock_guard<std::mutex> g(pool.events_mutex_);
+        if (!pool.events_.empty()) {
+          processed = std::move(pool.events_.back());
+          pool.events_.pop_back();
         }
       }
 
@@ -580,8 +658,8 @@ struct CachingHostAllocatorImpl {
       if (size != -1) {
         if (event_count++ > max_events) {
           {
-            std::lock_guard<std::mutex> g(events_mutex_);
-            events_.push_front(std::move(*processed));
+            std::lock_guard<std::mutex> g(pool.events_mutex_);
+            pool.events_.push_front(std::move(*processed));
           }
           return;
         }
@@ -589,8 +667,8 @@ struct CachingHostAllocatorImpl {
           // if we are processing a specific size, and the size of the block
           // doesn't match, we can't use it.
           {
-            std::lock_guard<std::mutex> g(events_mutex_);
-            events_.push_front(std::move(*processed));
+            std::lock_guard<std::mutex> g(pool.events_mutex_);
+            pool.events_.push_front(std::move(*processed));
           }
           continue;
         }
@@ -603,25 +681,24 @@ struct CachingHostAllocatorImpl {
         if (!query_event(event)) {
           // push the event onto the back if it's not ready.
           {
-            std::lock_guard<std::mutex> g(events_mutex_);
+            std::lock_guard<std::mutex> g(pool.events_mutex_);
             if (size == -1) {
-              events_.push_back(std::move(*processed));
+              pool.events_.push_back(std::move(*processed));
               return;
             } else {
-              events_.push_front(std::move(*processed));
+              pool.events_.push_front(std::move(*processed));
               continue;
             }
           }
         }
       }
 
-      // Process the events.
       TORCH_INTERNAL_ASSERT(processed);
       auto* block = processed->second;
       bool available = false;
       {
         std::lock_guard<std::mutex> g(block->mutex_);
-        TORCH_INTERNAL_ASSERT(!block->allocated_)
+        TORCH_INTERNAL_ASSERT(!block->allocated_);
         block->event_count_--;
         if (block->event_count_ == 0) {
           available = true;
@@ -630,8 +707,8 @@ struct CachingHostAllocatorImpl {
 
       if (available) {
         auto index = size_index(block->size_);
-        std::lock_guard<std::mutex> g(free_list_[index].mutex_);
-        free_list_[index].list_.push_back(block);
+        std::lock_guard<std::mutex> g(pool.free_list_[index].mutex_);
+        pool.free_list_[index].list_.push_back(block);
         stats_.active_bucket_stats[index].decrease(1);
         stats_.active_bytes_bucket_stats[index].decrease(size);
         if (size != -1) {
@@ -646,6 +723,130 @@ struct CachingHostAllocatorImpl {
     return pool;
   }
 
+ public:
+  void begin_allocate_to_pool(
+      c10::MempoolId_t pool_id,
+      std::function<bool(c10::Stream)> filter) {
+    std::unique_lock<std::shared_mutex> lg(instance_mutex_);
+    create_or_incref_pool(pool_id);
+    for (auto it2 = captures_underway_.begin(); it2 != captures_underway_.end();
+         ++it2) {
+      TORCH_CHECK(
+          it2->first != pool_id,
+          "beginAllocateToPool: already recording to mempool_id");
+    }
+    captures_underway_.emplace_back(pool_id, std::move(filter));
+  }
+
+  void end_allocate_to_pool(c10::MempoolId_t pool_id) {
+    std::unique_lock<std::shared_mutex> lg(instance_mutex_);
+    for (auto it = captures_underway_.begin(); it != captures_underway_.end();
+         ++it) {
+      if (it->first == pool_id) {
+        captures_underway_.erase(it);
+        return;
+      }
+    }
+    TORCH_CHECK(false, "endAllocatePool: not currently recording to mempool_id");
+  }
+
+  void create_or_incref_pool(c10::MempoolId_t pool_id) {
+    auto it = graph_pools_.find(pool_id);
+    if (it == graph_pools_.end()) {
+      graph_pools_.emplace(pool_id, std::make_unique<PrivatePool>(pool_id));
+    } else {
+      TORCH_INTERNAL_ASSERT(it->second->use_count > 0);
+      it->second->use_count++;
+    }
+  }
+
+  void release_pool(c10::MempoolId_t pool_id) {
+    std::unique_lock<std::shared_mutex> lg(instance_mutex_);
+    auto* pp = graph_pools_.at(pool_id).get();
+    TORCH_INTERNAL_ASSERT(pp != nullptr);
+    auto uc = --(pp->use_count);
+    TORCH_INTERNAL_ASSERT(uc >= 0);
+    if (uc == 0) {
+      bool inserted = graph_pools_freeable_.insert({pool_id, pp}).second;
+      TORCH_INTERNAL_ASSERT(inserted);
+    }
+  }
+
+ private:
+  // Helper: returns the correct pool for a newly allocated block.
+  std::tuple<c10::MempoolId_t, BlockPool&> get_allocation_pool(c10::Stream stream) {
+    std::shared_lock<std::shared_mutex> lg(instance_mutex_);
+    if (C10_UNLIKELY(!captures_underway_.empty())) {
+      for (auto& entry : captures_underway_) {
+        if (entry.second(stream)) {
+          auto it = graph_pools_.find(entry.first);
+          TORCH_INTERNAL_ASSERT(it != graph_pools_.end());
+          return {entry.first, it->second->blocks};
+        }
+      }
+    }
+    return {c10::MempoolId_t{0, 0}, default_pool_};
+  }
+
+  // Helper: return the pool containing a block, based on its owning_pool_.
+  BlockPool& pool_from_block(B* block) {
+    auto id = block->owning_pool_;
+    if (id == c10::MempoolId_t{0, 0}) {
+      return default_pool_;
+    }
+    std::shared_lock<std::shared_mutex> lg(instance_mutex_);
+    auto it = graph_pools_.find(id);
+    TORCH_INTERNAL_ASSERT(it != graph_pools_.end());
+    return it->second->blocks;
+  }
+
+protected:
+  B* get_block_from_ptr(void *ptr) {
+    std::shared_lock<std::shared_mutex> lk(instance_mutex_);
+    {
+      std::lock_guard<std::mutex> lk(default_pool_.blocks_mutex_);
+      if (default_pool_.ptr_to_block_.count(ptr)) {
+        return default_pool_.ptr_to_block_.at(ptr);
+      }
+    }
+    if (C10_LIKELY(graph_pools_.empty())) {
+      return nullptr;
+    } else {
+      for (auto &&[_, private_pool]: graph_pools_) {
+        BlockPool& pool = private_pool->blocks;
+        std::lock_guard<std::mutex> lk(pool.blocks_mutex_);
+        if (pool.ptr_to_block_.count(ptr)) {
+          return pool.ptr_to_block_.at(ptr);
+        }
+      }
+      return nullptr;
+    }
+  }
+
+  bool block_exists(void *block_) {
+    B *block = reinterpret_cast<B*>(block_);
+    std::shared_lock<std::shared_mutex> lk(instance_mutex_);
+    {
+      std::lock_guard<std::mutex> lk(default_pool_.blocks_mutex_);
+      if (default_pool_.blocks_.count(block)) {
+        return true;
+      }
+    }
+    if (C10_LIKELY(graph_pools_.empty())) {
+      return false;
+    } else {
+      for (auto &&[_, private_pool]: graph_pools_) {
+        BlockPool& pool = private_pool->blocks;
+        std::lock_guard<std::mutex> lk(pool.blocks_mutex_);
+        if (pool.blocks_.count(block)) {
+          return true;
+        }
+      }
+      return false;
+    }
+  }
+
+private:
   /* These following functions are runtime-related. */
 
   // Allocate page-locked memory on the host.
@@ -669,24 +870,36 @@ struct CachingHostAllocatorImpl {
     TORCH_CHECK_NOT_IMPLEMENTED(false, "Not implemented for query_event");
   }
 
-  alignas(hardware_destructive_interference_size) std::mutex blocks_mutex_;
-  ska::flat_hash_set<B*> blocks_; // block list
-  ska::flat_hash_map<void*, B*> ptr_to_block_;
+  virtual S get_current_stream() const {
+    TORCH_CHECK_NOT_IMPLEMENTED(false, "Not implemented for get_current_stream");
+  }
 
-  // We keep free list as a vector of free lists, one for each power of two
-  // size. This allows us to quickly find a free block of the right size.
-  // We use deque to store per size free list and guard the list with its own
-  // mutex.
-  alignas(hardware_destructive_interference_size) std::vector<FreeBlockList<B>>
-      free_list_{MAX_SIZE_INDEX};
+  virtual bool stream_is_capturing(S s) const {
+    TORCH_CHECK_NOT_IMPLEMENTED(false, "Not implemented for stream_is_capturing");
+  }
 
-  alignas(hardware_destructive_interference_size) std::mutex events_mutex_;
-  std::deque<std::pair<E, B*>> events_; // event queue paired with block
+  // instance variables
+protected:
+  alignas(hardware_destructive_interference_size) std::shared_mutex instance_mutex_;
+
+  // corresponds to c10::MempoolId_t{0,0}
+  BlockPool default_pool_;
+
+  // Private pools for captures
+  ska::flat_hash_map<c10::MempoolId_t, std::unique_ptr<PrivatePool>, c10::MempoolIdHash>
+      graph_pools_;
+
+  ska::flat_hash_map<c10::MempoolId_t, PrivatePool*, c10::MempoolIdHash>
+      graph_pools_freeable_;
+
+  // Track active capture contexts requesting allocations to specific pools
+  std::vector<
+      std::pair<c10::MempoolId_t, std::function<bool(c10::Stream)>>> captures_underway_;
 
   // Indicates whether the object is active.
   // Set to false in the destructor to signal background threads to stop.
   std::atomic<bool> active_{true};
-protected:
+
   alignas(hardware_destructive_interference_size) HostStatsStaged stats_;
 };
 
@@ -709,6 +922,25 @@ struct TORCH_API HostAllocator : public at::Allocator {
 
   // Resets the peak memory usage metrics
   virtual void reset_peak_stats() = 0;
+
+  virtual void begin_allocate_to_pool(
+      c10::MempoolId_t pool_id,
+      std::function<bool(c10::Stream)> filter) {
+    TORCH_CHECK_NOT_IMPLEMENTED(false, "Not implemented for begin_allocate_to_pool");
+  }
+
+  virtual void end_allocate_to_pool(c10::MempoolId_t pool_id) {
+    TORCH_CHECK_NOT_IMPLEMENTED(false, "Not implemented for end_allocate_to_pool");
+  }
+
+  virtual void create_or_incref_pool(
+      c10::MempoolId_t pool_id) {
+    TORCH_CHECK_NOT_IMPLEMENTED(false, "Not implemented for create_or_incref_pool");
+  }
+
+  virtual void release_pool(c10::MempoolId_t pool_id) {
+    TORCH_CHECK_NOT_IMPLEMENTED(false, "Not implemented for release_pool");
+  }
 };
 
 template <typename T, c10::DeleterFnPtr deleteFunc>
@@ -753,6 +985,25 @@ struct CachingHostAllocatorInterface : public HostAllocator {
     impl_->resetPeakStats();
   }
 
+  void begin_allocate_to_pool(
+      c10::MempoolId_t pool_id,
+      std::function<bool(c10::Stream)> filter) override {
+    impl_->begin_allocate_to_pool(pool_id, std::move(filter));
+  }
+
+  void end_allocate_to_pool(c10::MempoolId_t pool_id) override {
+    impl_->end_allocate_to_pool(pool_id);
+  }
+
+  void create_or_incref_pool(
+      c10::MempoolId_t pool_id) override {
+    impl_->create_or_incref_pool(pool_id);
+  }
+
+  void release_pool(c10::MempoolId_t pool_id) override {
+    impl_->release_pool(pool_id);
+  }
+
   std::unique_ptr<T> impl_;
 };
 

aten/src/ATen/core/IListRef.h

@@ -190,14 +190,12 @@ class IListRef;
  * it to a function (e.g. `ImplT::<dispatch-function>(this_)`).
  */
 #define TORCH_ILISTREF_UNWRAP(TAG, BODY)                         \
-  C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum")    \
   switch (TAG) {                                                 \
     TORCH_ILISTREF_FORALL_TAGS(TORCH_ILISTREF_UNWRAP_CASE, BODY) \
     break;                                                       \
     default:                                                     \
       TORCH_INTERNAL_ASSERT(false, "invalid IListRef tag.");     \
-  } \
-  C10_DIAGNOSTIC_POP()
+  }
 
 enum class IListRefTag {
 #define DEFINE_TAG(tag, ...) tag,

aten/src/ATen/core/TransformationHelper.h

@@ -56,7 +56,7 @@ C10_HOST_DEVICE inline T uniform_int_full_range(V val) {
  * in this overloaded version
  */
 template <typename T, typename V>
-C10_HOST_DEVICE inline std::enable_if_t<!std::is_floating_point_v<T>, T>uniform_int(V val) {
+C10_HOST_DEVICE inline std::enable_if_t<!(std::is_floating_point_v<T>), T>uniform_int(V val) {
   if constexpr (std::is_same_v<T, bool>) {
     return static_cast<bool>(val & 1);
   } else if constexpr (std::is_same_v<T, int64_t>) {

aten/src/ATen/core/boxing/KernelFunction_impl.h

@@ -114,25 +114,25 @@ inline typename remove_symint<T>::type unpackSymInt(T x) {
 }
 
 template <>
-inline remove_symint<c10::SymInt>::type unpackSymInt(c10::SymInt x) {
+inline typename remove_symint<c10::SymInt>::type unpackSymInt(c10::SymInt x) {
   return x.guard_int(__FILE__, __LINE__);
 }
 
 template <>
-inline remove_symint<c10::SymIntArrayRef>::type unpackSymInt(
+inline typename remove_symint<c10::SymIntArrayRef>::type unpackSymInt(
     c10::SymIntArrayRef x) {
   return C10_AS_INTARRAYREF_SLOW(x);
 }
 
 template <>
-inline remove_symint<std::optional<c10::SymInt>>::type unpackSymInt(
+inline typename remove_symint<std::optional<c10::SymInt>>::type unpackSymInt(
     std::optional<c10::SymInt> x) {
   return x.has_value() ? std::make_optional(x->guard_int(__FILE__, __LINE__))
                        : std::nullopt;
 }
 
 template <>
-inline remove_symint<at::OptionalSymIntArrayRef>::type unpackSymInt(
+inline typename remove_symint<at::OptionalSymIntArrayRef>::type unpackSymInt(
     at::OptionalSymIntArrayRef x) {
   return x.has_value() ? std::make_optional(C10_AS_INTARRAYREF_SLOW(*x))
                        : std::nullopt;

aten/src/ATen/core/boxing/impl/make_boxed_from_unboxed_functor.h

@@ -631,8 +631,8 @@ call_functor_with_args_from_stack_(
     Stack* stack,
     std::index_sequence<ivalue_arg_indices...> /*unused*/,
     guts::typelist::typelist<ArgTypes...>* /*unused*/) {
-  (void)stack; // when sizeof...(ivalue_arg_indices) == 0, this argument would
-               // be unused and we have to silence the compiler warning.
+  (void)(stack); // when sizeof...(ivalue_arg_indices) == 0, this argument would
+                 // be unused and we have to silence the compiler warning.
 
   // We're explicitly filtering out DispatchKeySet from the argument list.
   // Some kernels take a DispatchKeySet as their first argument in order to

aten/src/ATen/core/enum_type.h

@@ -18,7 +18,6 @@ struct TORCH_API EnumType : public NamedType {
       TypePtr value,
       std::vector<EnumNameValue> enum_names_values,
       std::weak_ptr<::torch::jit::CompilationUnit> cu) {
-    C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum")
     switch (value->kind()) {
       case TypeKind::IntType:
       case TypeKind::FloatType:
@@ -35,7 +34,6 @@ struct TORCH_API EnumType : public NamedType {
             value->str(),
             "', only int, float and string are supported");
     }
-    C10_DIAGNOSTIC_POP()
   }
 
   std::string str() const override {

aten/src/ATen/core/ivalue.cpp

@@ -601,8 +601,8 @@ std::ostream& IValue::repr(
       double d = v.toDouble();
       int c = std::fpclassify(d);
       if ((c == FP_NORMAL || c == FP_ZERO ) && std::abs(d) < 1e10) {
-        int64_t i = static_cast<int64_t>(d);
-        if (static_cast<double>(i) == d) {
+        int64_t i = int64_t(d);
+        if (double(i) == d) {
           // -0.0 (signed zero) needs to be parsed as -0.
           if (i == 0 && std::signbit(d)) {
             return out << "-" << i << ".";
@@ -799,8 +799,8 @@ std::ostream& operator<<(std::ostream & out, const IValue & v) {
       double d = v.toDouble();
       int c = std::fpclassify(d);
       if (c == FP_NORMAL || c == FP_ZERO) {
-        int64_t i = static_cast<int64_t>(d);
-        if (static_cast<double>(i) == d) {
+        int64_t i = int64_t(d);
+        if (double(i) == d) {
           return out << i << ".";
         }
       }

aten/src/ATen/core/jit_type.h

@@ -41,7 +41,7 @@ void standardizeVectorForUnion(std::vector<TypePtr>* to_flatten);
 inline bool is_contiguous_strides(
     const IntArrayRef sizes,
     const IntArrayRef strides) {
-  size_t n_dim = sizes.size();
+  int n_dim = static_cast<int>(sizes.size());
   if (n_dim == 0) {
     return true;
   }
@@ -50,7 +50,7 @@ inline bool is_contiguous_strides(
     return false;
   }
 
-  for (int i = static_cast<int>(n_dim) - 2; i >= 0; i--) {
+  for (int i = n_dim - 2; i >= 0; i--) {
     if (strides[i] != strides[i + 1] * sizes[i + 1]) {
       return false;
     }
@@ -922,7 +922,6 @@ struct TORCH_API DictType : public SharedType {
     if (auto dyn = key->castRaw<DynamicType>()) {
       kind = dyn->dynamicKind();
     }
-    C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wswitch-enum")
     switch (kind) {
       case TypeKind::AnyType:
       case TypeKind::IntType:
@@ -939,7 +938,6 @@ struct TORCH_API DictType : public SharedType {
             key->str(),
             "', only int, float, complex, Tensor, device and string keys are supported");
     }
-    C10_DIAGNOSTIC_POP()
   }
 
   // aligned with the format in FunctionSchema
@@ -2373,7 +2371,7 @@ private:
 };
 
 template<>
-inline detail::CastReturnType<NamedType>::type Type::cast<NamedType>() {
+inline typename detail::CastReturnType<NamedType>::type Type::cast<NamedType>() {
   if (kind() == TypeKind::TupleType || kind() == TypeKind::FunctionType ||
       kind() == TypeKind::ClassType || kind() == TypeKind::InterfaceType) {
     return std::static_pointer_cast<NamedType>(static_cast<NamedType *>(this)->shared_from_this());
@@ -2382,7 +2380,7 @@ inline detail::CastReturnType<NamedType>::type Type::cast<NamedType>() {
 }
 
 template<>
-inline detail::CastConstReturnType<NamedType>::type Type::cast<NamedType>() const {
+inline typename detail::CastConstReturnType<NamedType>::type Type::cast<NamedType>() const {
   if (kind() == TypeKind::TupleType || kind() == TypeKind::FunctionType ||
       kind() == TypeKind::ClassType || kind() == TypeKind::InterfaceType) {
     return std::static_pointer_cast<const NamedType>(static_cast<const NamedType *>(this)->shared_from_this());

aten/src/ATen/cpu/vec/sve/vec_bfloat16.h

@@ -191,7 +191,7 @@ class Vectorized<BFloat16> {
     auto vals = svreinterpret_u16_bf16(values);
     vals = sveor_u16_x(ptrue, vals, mask);
     return svreinterpret_bf16_u16(vals);
-  }
+  };
   Vectorized<BFloat16> round() const;
   Vectorized<BFloat16> tan() const;
   Vectorized<BFloat16> tanh() const;
@@ -349,47 +349,47 @@ Vectorized<BFloat16> inline Vectorized<BFloat16>::frac() const {
     return convert_float_bfloat16(v1, v2);                     \
   }
 
-DEFINE_BF16_FUNC_VIA_FLOAT(isnan)
-DEFINE_BF16_FUNC_VIA_FLOAT(angle)
-DEFINE_BF16_FUNC_VIA_FLOAT(acos)
-DEFINE_BF16_FUNC_VIA_FLOAT(acosh)
-DEFINE_BF16_FUNC_VIA_FLOAT(asin)
-DEFINE_BF16_FUNC_VIA_FLOAT(atan)
-DEFINE_BF16_FUNC_VIA_FLOAT(atanh)
-DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(atan2)
-DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(copysign)
-DEFINE_BF16_FUNC_VIA_FLOAT(erf)
-DEFINE_BF16_FUNC_VIA_FLOAT(erfc)
-DEFINE_BF16_FUNC_VIA_FLOAT(exp)
-DEFINE_BF16_FUNC_VIA_FLOAT(exp2)
-DEFINE_BF16_FUNC_VIA_FLOAT(expm1)
-DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(fmod)
-DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(hypot)
-DEFINE_BF16_FUNC_VIA_FLOAT(i0)
-DEFINE_BF16_FUNC_VIA_FLOAT(i0e)
-DEFINE_BF16_FUNC_VIA_FLOAT(digamma)
-DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(igamma)
-DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(igammac)
-DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(nextafter)
-DEFINE_BF16_FUNC_VIA_FLOAT(log)
-DEFINE_BF16_FUNC_VIA_FLOAT(log2)
-DEFINE_BF16_FUNC_VIA_FLOAT(log10)
-DEFINE_BF16_FUNC_VIA_FLOAT(log1p)
-DEFINE_BF16_FUNC_VIA_FLOAT(sin)
-DEFINE_BF16_FUNC_VIA_FLOAT(sinh)
-DEFINE_BF16_FUNC_VIA_FLOAT(cos)
-DEFINE_BF16_FUNC_VIA_FLOAT(cosh)
-DEFINE_BF16_FUNC_VIA_FLOAT(ceil)
-DEFINE_BF16_FUNC_VIA_FLOAT(floor)
-DEFINE_BF16_FUNC_VIA_FLOAT(round)
-DEFINE_BF16_FUNC_VIA_FLOAT(tan)
-DEFINE_BF16_FUNC_VIA_FLOAT(tanh)
-DEFINE_BF16_FUNC_VIA_FLOAT(trunc)
-DEFINE_BF16_FUNC_VIA_FLOAT(lgamma)
-DEFINE_BF16_FUNC_VIA_FLOAT(sqrt)
-DEFINE_BF16_FUNC_VIA_FLOAT(reciprocal)
-DEFINE_BF16_FUNC_VIA_FLOAT(rsqrt)
-DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(pow)
+DEFINE_BF16_FUNC_VIA_FLOAT(isnan);
+DEFINE_BF16_FUNC_VIA_FLOAT(angle);
+DEFINE_BF16_FUNC_VIA_FLOAT(acos);
+DEFINE_BF16_FUNC_VIA_FLOAT(acosh);
+DEFINE_BF16_FUNC_VIA_FLOAT(asin);
+DEFINE_BF16_FUNC_VIA_FLOAT(atan);
+DEFINE_BF16_FUNC_VIA_FLOAT(atanh);
+DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(atan2);
+DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(copysign);
+DEFINE_BF16_FUNC_VIA_FLOAT(erf);
+DEFINE_BF16_FUNC_VIA_FLOAT(erfc);
+DEFINE_BF16_FUNC_VIA_FLOAT(exp);
+DEFINE_BF16_FUNC_VIA_FLOAT(exp2);
+DEFINE_BF16_FUNC_VIA_FLOAT(expm1);
+DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(fmod);
+DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(hypot);
+DEFINE_BF16_FUNC_VIA_FLOAT(i0);
+DEFINE_BF16_FUNC_VIA_FLOAT(i0e);
+DEFINE_BF16_FUNC_VIA_FLOAT(digamma);
+DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(igamma);
+DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(igammac);
+DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(nextafter);
+DEFINE_BF16_FUNC_VIA_FLOAT(log);
+DEFINE_BF16_FUNC_VIA_FLOAT(log2);
+DEFINE_BF16_FUNC_VIA_FLOAT(log10);
+DEFINE_BF16_FUNC_VIA_FLOAT(log1p);
+DEFINE_BF16_FUNC_VIA_FLOAT(sin);
+DEFINE_BF16_FUNC_VIA_FLOAT(sinh);
+DEFINE_BF16_FUNC_VIA_FLOAT(cos);
+DEFINE_BF16_FUNC_VIA_FLOAT(cosh);
+DEFINE_BF16_FUNC_VIA_FLOAT(ceil);
+DEFINE_BF16_FUNC_VIA_FLOAT(floor);
+DEFINE_BF16_FUNC_VIA_FLOAT(round);
+DEFINE_BF16_FUNC_VIA_FLOAT(tan);
+DEFINE_BF16_FUNC_VIA_FLOAT(tanh);
+DEFINE_BF16_FUNC_VIA_FLOAT(trunc);
+DEFINE_BF16_FUNC_VIA_FLOAT(lgamma);
+DEFINE_BF16_FUNC_VIA_FLOAT(sqrt);
+DEFINE_BF16_FUNC_VIA_FLOAT(reciprocal);
+DEFINE_BF16_FUNC_VIA_FLOAT(rsqrt);
+DEFINE_BF16_FUNC_VIA_FLOAT_W_ARG(pow);
 
 Vectorized<BFloat16> inline Vectorized<BFloat16>::operator==(
     const Vectorized<BFloat16>& other) const {

aten/src/ATen/cpu/vec/vec128/vec128_bfloat16_neon.h

@@ -19,13 +19,6 @@ inline namespace CPU_CAPABILITY {
 #error "Big endian is not supported."
 #endif
 
-// GCC does not properly optimize bf16 operators
-#if defined(__ARM_FEATURE_BF16) && (__clang_major__ >= 19)
-#define BF16_ARITHMETIC_SUPPORTED() 1
-#else
-#define BF16_ARITHMETIC_SUPPORTED() 0
-#endif
-
 // Unlike the float16_t family of types, bfloat16_t is not available
 // when we're not targeting bfloat16 hardware support on some
 // platforms (but not Mac, so we have to be careful not to shadow the
@@ -359,72 +352,18 @@ class Vectorized<c10::BFloat16> : public Vectorized16<
         other, &Vectorized<float>::name);                        \
   }
 
+  DEFINE_UNARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD(abs)
   Vectorized frac() const;
+  DEFINE_UNARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD(neg)
   DEFINE_UNARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD(trunc)
   DEFINE_UNARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD(sqrt)
-
-#ifdef __ARM_FEATURE_BF16
-  // Flip sign bit
-  Vectorized<c10::BFloat16> neg() const {
-    return vreinterpretq_bf16_s16(vreinterpretq_s16_bf16(values) ^ (-32768));
-  }
-  // Fast reciprocal is fine because we are truncating results
-  Vectorized<c10::BFloat16> reciprocal() const {
-    auto x = vcvtq_low_f32_bf16(values);
-    auto y = vcvtq_high_f32_bf16(values);
-    x = vrecpeq_f32(x);
-    y = vrecpeq_f32(y);
-    return vcvtq_high_bf16_f32(vcvtq_low_bf16_f32(x), y);
-  }
-  // Clearing the sign bit
-  Vectorized<c10::BFloat16> abs() const {
-    return vreinterpretq_bf16_u16(vreinterpretq_u16_bf16(values) & 0x7FFF);
-  }
-#else
-  DEFINE_UNARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD(abs)
-  DEFINE_UNARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD(neg)
   DEFINE_UNARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD(reciprocal)
-#endif
-
-// These functions are optimized on clang-21+
-#if BF16_ARITHMETIC_SUPPORTED() && (__clang_major__ >= 21)
-  Vectorized<c10::BFloat16> operator==(
-      const Vectorized<c10::BFloat16>& other) const {
-    return values == other.values;
-  }
-
-  Vectorized<c10::BFloat16> operator!=(
-      const Vectorized<c10::BFloat16>& other) const {
-    return values != other.values;
-  }
-
-  Vectorized<c10::BFloat16> operator<(
-      const Vectorized<c10::BFloat16>& other) const {
-    return values < other.values;
-  }
-
-  Vectorized<c10::BFloat16> operator<=(
-      const Vectorized<c10::BFloat16>& other) const {
-    return values <= other.values;
-  }
-
-  Vectorized<c10::BFloat16> operator>(
-      const Vectorized<c10::BFloat16>& other) const {
-    return values > other.values;
-  }
-
-  Vectorized<c10::BFloat16> operator>=(
-      const Vectorized<c10::BFloat16>& other) const {
-    return values >= other.values;
-  }
-#else
   DEFINE_BINARY_COMPARISON_OPERATOR_VIA_FLOAT_METHOD(operator==)
   DEFINE_BINARY_COMPARISON_OPERATOR_VIA_FLOAT_METHOD(operator!=)
   DEFINE_BINARY_COMPARISON_OPERATOR_VIA_FLOAT_METHOD(operator<)
   DEFINE_BINARY_COMPARISON_OPERATOR_VIA_FLOAT_METHOD(operator<=)
   DEFINE_BINARY_COMPARISON_OPERATOR_VIA_FLOAT_METHOD(operator>)
   DEFINE_BINARY_COMPARISON_OPERATOR_VIA_FLOAT_METHOD(operator>=)
-#endif
 
 #undef DEFINE_UNARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD
 #undef DEFINE_BINARY_ELEMENTWISE_FUNC_VIA_FLOAT_METHOD
@@ -473,52 +412,28 @@ template <>
 Vectorized<c10::BFloat16> inline operator+(
     const Vectorized<c10::BFloat16>& a,
     const Vectorized<c10::BFloat16>& b) {
-#if BF16_ARITHMETIC_SUPPORTED()
-  bfloat16x8_t x = a;
-  bfloat16x8_t y = b;
-  return x + y;
-#else
   return binary_operator_via_float(std::plus<Vectorized<float>>(), a, b);
-#endif
 }
 
 template <>
 Vectorized<c10::BFloat16> inline operator-(
     const Vectorized<c10::BFloat16>& a,
     const Vectorized<c10::BFloat16>& b) {
-#if BF16_ARITHMETIC_SUPPORTED()
-  bfloat16x8_t x = a;
-  bfloat16x8_t y = b;
-  return x - y;
-#else
   return binary_operator_via_float(std::minus<Vectorized<float>>(), a, b);
-#endif
 }
 
 template <>
 Vectorized<c10::BFloat16> inline operator*(
     const Vectorized<c10::BFloat16>& a,
     const Vectorized<c10::BFloat16>& b) {
-#if BF16_ARITHMETIC_SUPPORTED()
-  bfloat16x8_t x = a;
-  bfloat16x8_t y = b;
-  return x * y;
-#else
   return binary_operator_via_float(std::multiplies<Vectorized<float>>(), a, b);
-#endif
 }
 
 template <>
 Vectorized<c10::BFloat16> inline operator/(
     const Vectorized<c10::BFloat16>& a,
     const Vectorized<c10::BFloat16>& b) {
-#if BF16_ARITHMETIC_SUPPORTED()
-  bfloat16x8_t x = a;
-  bfloat16x8_t y = b;
-  return x / y;
-#else
   return binary_operator_via_float(std::divides<Vectorized<float>>(), a, b);
-#endif
 }
 
 // frac. Implement this here so we can use subtraction
@@ -629,19 +544,12 @@ Vectorized<c10::BFloat16> inline fmadd(
     const Vectorized<c10::BFloat16>& a,
     const Vectorized<c10::BFloat16>& b,
     const Vectorized<c10::BFloat16>& c) {
-#if BF16_ARITHMETIC_SUPPORTED()
-  bfloat16x8_t x = a;
-  bfloat16x8_t y = b;
-  bfloat16x8_t z = c;
-  return x * y + z;
-#else
   // NOTE [BF16 FMA]: There isn't an FMA that accumulates into BF16!  Also,
   // vbfmlalbq_f32 and vbfmlaltq_f32 take the even and odd-numbered
   // elements, not the bottom and top half, so they don't seem
   // particularly useful here. Ideally we would include dot product in
   // the Vectorized interface...
   return a * b + c;
-#endif
 }
 
 template <>
@@ -649,15 +557,8 @@ Vectorized<c10::BFloat16> inline fnmadd(
     const Vectorized<c10::BFloat16>& a,
     const Vectorized<c10::BFloat16>& b,
     const Vectorized<c10::BFloat16>& c) {
-#if BF16_ARITHMETIC_SUPPORTED()
-  bfloat16x8_t x = a;
-  bfloat16x8_t y = b;
-  bfloat16x8_t z = c;
-  return (-x) * y + z;
-#else
   // See NOTE [BF16 FMA] above.
   return -a * b + c;
-#endif
 }
 
 template <>
@@ -665,15 +566,8 @@ Vectorized<c10::BFloat16> inline fmsub(
     const Vectorized<c10::BFloat16>& a,
     const Vectorized<c10::BFloat16>& b,
     const Vectorized<c10::BFloat16>& c) {
-#if BF16_ARITHMETIC_SUPPORTED()
-  bfloat16x8_t x = a;
-  bfloat16x8_t y = b;
-  bfloat16x8_t z = c;
-  return x * y - z;
-#else
   // See NOTE [BF16 FMA] above.
   return a * b - c;
-#endif
 }
 
 template <>
@@ -681,15 +575,8 @@ Vectorized<c10::BFloat16> inline fnmsub(
     const Vectorized<c10::BFloat16>& a,
     const Vectorized<c10::BFloat16>& b,
     const Vectorized<c10::BFloat16>& c) {
-#if BF16_ARITHMETIC_SUPPORTED()
-  bfloat16x8_t x = a;
-  bfloat16x8_t y = b;
-  bfloat16x8_t z = c;
-  return (-x) * y - z;
-#else
   // See NOTE [BF16 FMA] above.
   return -a * b - c;
-#endif
 }
 
 #endif // !defined(C10_MOBILE) && defined(__aarch64__)

aten/src/ATen/cpu/vec/vec128/vec128_convert.h

@@ -6,9 +6,9 @@ namespace at::vec {
 inline namespace CPU_CAPABILITY {
 #if (defined(__aarch64__) && !defined(CPU_CAPABILITY_SVE256))
 
-// Enable auto-vectorization for clang-17+
+// Enable auto-vectorization for GCC-13+ and clang-17+
 // GCC-12 has a bug: gcc.gnu.org/bugzilla/show_bug.cgi?id=117001
-#if defined(__clang__) && (__clang_major__ >= 17)
+#if __GNUC__ > 12 || (defined(__clang__) && (__clang_major__ >= 17))
 
 template <typename from_type, typename to_type>
 inline void convertImpl(
@@ -191,37 +191,22 @@ inline void convert(const at::Half* src, bool* dst, int64_t n) {
 }
 
 #endif
-
-template <typename to_type>
-inline void convertFromBf16Impl(
-    const c10::BFloat16* __restrict src,
-    to_type* __restrict dst,
-    int64_t n) {
-  const uint16_t* srcPtr = reinterpret_cast<const uint16_t*>(src);
-  uint64_t len = static_cast<uint64_t>(n);
-  for (uint64_t i = 0; i < len; i++) {
-    uint32_t tmp = static_cast<uint32_t>(srcPtr[i]) << 16;
-    float tmpF;
-    __builtin_memcpy(&tmpF, &tmp, sizeof(float));
-    dst[i] = static_cast<to_type>(tmpF);
-  }
-}
-#define CONVERT_FROM_BF16_TEMPLATE(to_type)                                \
-  template <>                                                              \
-  inline void convert(const c10::BFloat16* src, to_type* dst, int64_t n) { \
-    return convertFromBf16Impl<to_type>(src, dst, n);                      \
-  }
-
-CONVERT_FROM_BF16_TEMPLATE(uint8_t)
-CONVERT_FROM_BF16_TEMPLATE(int8_t)
-CONVERT_FROM_BF16_TEMPLATE(int16_t)
-CONVERT_FROM_BF16_TEMPLATE(int32_t)
-CONVERT_FROM_BF16_TEMPLATE(int64_t)
-CONVERT_FROM_BF16_TEMPLATE(float)
-CONVERT_FROM_BF16_TEMPLATE(double)
-#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
-CONVERT_FROM_BF16_TEMPLATE(float16_t)
-#endif
+#ifdef __ARM_FEATURE_BF16
+CONVERT_TEMPLATE(bfloat16_t, uint8_t)
+CONVERT_TEMPLATE(bfloat16_t, int8_t)
+CONVERT_TEMPLATE(bfloat16_t, int16_t)
+CONVERT_TEMPLATE(bfloat16_t, int32_t)
+CONVERT_TEMPLATE(bfloat16_t, int64_t)
+CONVERT_TEMPLATE(bfloat16_t, bfloat16_t)
+CONVERT_TEMPLATE(bfloat16_t, float)
+CONVERT_TEMPLATE(bfloat16_t, double)
+CONVERT_TEMPLATE(uint8_t, bfloat16_t)
+CONVERT_TEMPLATE(int8_t, bfloat16_t)
+CONVERT_TEMPLATE(int16_t, bfloat16_t)
+CONVERT_TEMPLATE(int32_t, bfloat16_t)
+CONVERT_TEMPLATE(int64_t, bfloat16_t)
+CONVERT_TEMPLATE(float, bfloat16_t)
+CONVERT_TEMPLATE(double, bfloat16_t)
 
 inline void convertBoolToBfloat16Impl(
     const bool* __restrict src,
@@ -262,6 +247,8 @@ inline void convert(const c10::BFloat16* src, bool* dst, int64_t n) {
 
 #endif
 
+#endif
+
 template <typename src_t>
 struct VecConvert<
     float,

aten/src/ATen/cpu/vec/vec256/vec256_qint.h

@@ -514,7 +514,7 @@ struct Vectorized<c10::qint8> : public Vectorizedqi {
 
   using float_vec_return_type = std::array<Vectorized<float>, kFloatNumVecs>;
   using int_vec_return_type = std::array<Vectorized<c10::qint32>, kIntNumVecs>;
-  using value_type = c10::qint8::underlying;
+  using value_type = typename c10::qint8::underlying;
 
  public:
   using Vectorizedqi::Vectorizedqi;
@@ -727,7 +727,7 @@ struct Vectorized<c10::quint8> : public Vectorizedqi {
 
   using float_vec_return_type = std::array<Vectorized<float>, kFloatNumVecs>;
   using int_vec_return_type = std::array<Vectorized<c10::qint32>, kIntNumVecs>;
-  using value_type = c10::quint8::underlying;
+  using value_type = typename c10::quint8::underlying;
 
  public:
   using Vectorizedqi::Vectorizedqi;

aten/src/ATen/cpu/vec/vec512/vec512_qint.h

@@ -567,7 +567,7 @@ struct Vectorized<c10::qint8> : public Vectorizedqi {
 
   using float_vec_return_type = std::array<Vectorized<float>, 4>;
   using int_vec_return_type = std::array<Vectorized<c10::qint32>, 4>;
-  using value_type = c10::qint8::underlying;
+  using value_type = typename c10::qint8::underlying;
 
  public:
   using Vectorizedqi::Vectorizedqi;
@@ -804,7 +804,7 @@ struct Vectorized<c10::quint8> : public Vectorizedqi {
 
   using float_vec_return_type = std::array<Vectorized<float>, 4>;
   using int_vec_return_type = std::array<Vectorized<c10::qint32>, 4>;
-  using value_type = c10::quint8::underlying;
+  using value_type = typename c10::quint8::underlying;
 
  public:
   using Vectorizedqi::Vectorizedqi;

aten/src/ATen/cpu/vec/vec_base.h

@@ -672,7 +672,7 @@ struct Vectorized {
     return map(std::sqrt);
   }
   Vectorized<T> reciprocal() const {
-    return map([](T x) { return (T)1 / x; });
+    return map([](T x) { return (T)(1) / x; });
   }
   Vectorized<T> rsqrt() const {
     return map([](T x) { return (T)1 / std::sqrt(x); });

aten/src/ATen/cpu/vml.h

@@ -46,7 +46,7 @@ inline void vrsqrt(scalar_t* out, scalar_t* in, int64_t size) {
   parallel_for(0, size, 2048, [out, in](int64_t begin, int64_t end) {
     map(
         [](const Vectorized<scalar_t>& x) {
-          return Vectorized<scalar_t>((scalar_t)1) / x.sqrt();
+          return Vectorized<scalar_t>((scalar_t)(1)) / x.sqrt();
         },
         out + begin,
         in + begin,

aten/src/ATen/cuda/CUDABlas.cpp

@@ -388,7 +388,6 @@ static inline bool bgemm_internal_cublaslt(CUDABLAS_BGEMM_ARGTYPES_AND_C_DTYPE(D
 #ifndef USE_ROCM
   at::Half halpha;
   at::Half hbeta;
-  uint32_t mask = -1;
 #endif
   void * alpha_ptr = α
   void * beta_ptr = β
@@ -428,7 +427,7 @@ static inline bool bgemm_internal_cublaslt(CUDABLAS_BGEMM_ARGTYPES_AND_C_DTYPE(D
     auto fp16_reduction = at::globalContext().allowFP16ReductionCuBLAS();
     if (fp16_reduction !=
         at::CuBLASReductionOption::AllowReducedPrecisionWithSplitK) {
-      mask =
+      uint32_t mask =
           fp16_reduction ==
                   at::CuBLASReductionOption::DisallowReducedPrecisionAllowSplitK
               ? (CUBLASLT_REDUCTION_SCHEME_COMPUTE_TYPE |
@@ -445,7 +444,7 @@ static inline bool bgemm_internal_cublaslt(CUDABLAS_BGEMM_ARGTYPES_AND_C_DTYPE(D
     auto bf16_reduction = at::globalContext().allowBF16ReductionCuBLAS();
     if (bf16_reduction !=
         at::CuBLASReductionOption::AllowReducedPrecisionWithSplitK) {
-      mask =
+      uint32_t mask =
           bf16_reduction ==
                   at::CuBLASReductionOption::DisallowReducedPrecisionAllowSplitK
               ? (CUBLASLT_REDUCTION_SCHEME_COMPUTE_TYPE |
@@ -512,41 +511,17 @@ static inline bool bgemm_internal_cublaslt(CUDABLAS_BGEMM_ARGTYPES_AND_C_DTYPE(D
   cublasStatus_t cublasStatus = CUBLAS_STATUS_SUCCESS;
   cublasLtMatmulHeuristicResult_t heuristicResult = {};
   int returnedResult = 0;
-  // on Blackwell+, we fake a n > 1 matmul when querying heuristics
-  // to prevent cuBLASLt from dispatching to a GEMV kernel for batch-invariance
-#ifndef USE_ROCM
-  const bool lie_to_cublaslt = mask == CUBLASLT_REDUCTION_SCHEME_NONE && n == 1 && at::cuda::getCurrentDeviceProperties()->major >= 10;
-#else
-  const bool lie_to_cublaslt = false;
-#endif
-  if (lie_to_cublaslt) {
-     CuBlasLtMatrixLayout FakeBdesc(abType, k, 2, ldb, opb == CUBLAS_OP_T);
-     CuBlasLtMatrixLayout FakeCdesc(cType, m, 2, ldc);
-
-     TORCH_CUDABLAS_CHECK(cublasLtMatmulAlgoGetHeuristic(
-        ltHandle,
-        computeDesc.descriptor(),
-        Adesc.descriptor(),
-        FakeBdesc.descriptor(),
-        FakeCdesc.descriptor(),
-        FakeCdesc.descriptor(),
-        preference.descriptor(),
-        1,
-        &heuristicResult,
-        &returnedResult));
-  } else {
-    TORCH_CUDABLAS_CHECK(cublasLtMatmulAlgoGetHeuristic(
-        ltHandle,
-        computeDesc.descriptor(),
-        Adesc.descriptor(),
-        Bdesc.descriptor(),
-        Cdesc.descriptor(),
-        Cdesc.descriptor(),
-        preference.descriptor(),
-        1,
-        &heuristicResult,
-        &returnedResult));
-  }
+  TORCH_CUDABLAS_CHECK(cublasLtMatmulAlgoGetHeuristic(
+      ltHandle,
+      computeDesc.descriptor(),
+      Adesc.descriptor(),
+      Bdesc.descriptor(),
+      Cdesc.descriptor(),
+      Cdesc.descriptor(),
+      preference.descriptor(),
+      1,
+      &heuristicResult,
+      &returnedResult));
   if (returnedResult == 0) {
     cublasStatus = CUBLAS_STATUS_NOT_SUPPORTED;
   }
@@ -1597,7 +1572,7 @@ bool gemm_and_bias(
   }
 
   using opmath_t = at::opmath_type<Dtype>;
-  opmath_t beta_val = bias ? 0 : 1; // bias is added in epilogue unless nullptr
+  opmath_t beta_val = 0; // bias is added in epilogue
 
   cudaDataType_t abType = CUDA_R_32F;
   cudaDataType_t cType = CUDA_R_32F;
@@ -1686,22 +1661,15 @@ bool gemm_and_bias(
     _syncCurrentWithCarveoutStream(stream, true);
   }
 #endif
-  const auto epilogue = [&]() -> cublasLtEpilogue_t {
-    // The cuBLAS documentation indicates that
-    // *_<ACTIVATION>_BIAS = *_<ACTIVATION>,
-    // but we keep it verbose here for clarity.
-    switch (activation) {
-      case GEMMAndBiasActivationEpilogue::RELU:
-        return bias ? CUBLASLT_EPILOGUE_RELU_BIAS : CUBLASLT_EPILOGUE_RELU;
-      case GEMMAndBiasActivationEpilogue::GELU:
-        return bias ? CUBLASLT_EPILOGUE_GELU_BIAS : CUBLASLT_EPILOGUE_GELU;
-      default:
-        return bias ? CUBLASLT_EPILOGUE_BIAS : CUBLASLT_EPILOGUE_DEFAULT;
-    }
-  }();
-  computeDesc.setAttribute(CUBLASLT_MATMUL_DESC_EPILOGUE, epilogue);
+  cublasLtEpilogue_t epilogue = CUBLASLT_EPILOGUE_BIAS;
+  if (activation == GEMMAndBiasActivationEpilogue::RELU) {
+    epilogue = CUBLASLT_EPILOGUE_RELU_BIAS;
+  } else if (activation == GEMMAndBiasActivationEpilogue::GELU) {
+    epilogue = CUBLASLT_EPILOGUE_GELU_BIAS;
+  }
 
-  if (bias) {
+  if (bias != nullptr) {
+    computeDesc.setAttribute(CUBLASLT_MATMUL_DESC_EPILOGUE, epilogue);
     computeDesc.setAttribute(CUBLASLT_MATMUL_DESC_BIAS_POINTER, bias);
   }
 

aten/src/ATen/cuda/CUDAGeneratorImpl.cpp

@@ -194,8 +194,8 @@ void CUDAGeneratorState::unregister_graph(cuda::CUDAGraph* graph) {
 void CUDAGeneratorState::capture_prologue() {
   capturing_ = true;
   offset_intragraph_ = 0;
-  seed_extragraph_.fill_(static_cast<int64_t>(seed_));
-  offset_extragraph_.fill_(0);
+  seed_extragraph_.fill_(int64_t(seed_));
+  offset_extragraph_.fill_(int64_t(0));
 }
 
 /**
@@ -216,8 +216,8 @@ void CUDAGeneratorState::replay_prologue(uint64_t wholegraph_increment) {
   at::cuda::assertNotCapturing(
       "Cannot prepare for replay during capturing stage.");
   if (wholegraph_increment) {
-      seed_extragraph_.fill_(static_cast<int64_t>(seed_));
-      offset_extragraph_.fill_(static_cast<int64_t>(philox_offset_per_thread_));
+      seed_extragraph_.fill_(int64_t(seed_));
+      offset_extragraph_.fill_(int64_t(philox_offset_per_thread_));
       // Applies the total increment achieved during previous captures to update the
       // offset.
       increase(wholegraph_increment);
@@ -329,7 +329,7 @@ c10::intrusive_ptr<c10::TensorImpl> CUDAGeneratorImpl::get_state() const {
   constexpr size_t offset_size = sizeof(int64_t);
   constexpr size_t total_size = seed_size + offset_size;
 
-  auto state_tensor = at::detail::empty_cpu({static_cast<int64_t>(total_size)}, ScalarType::Byte, std::nullopt, std::nullopt, std::nullopt, std::nullopt);
+  auto state_tensor = at::detail::empty_cpu({(int64_t)total_size}, ScalarType::Byte, std::nullopt, std::nullopt, std::nullopt, std::nullopt);
   auto rng_state = state_tensor.data_ptr<uint8_t>();
   auto current_seed = this->current_seed();
   auto offset = static_cast<int64_t>(this->philox_offset_per_thread()); // Note that old THCGeneratorState had offset as std::atomic<int64_t>

aten/src/ATen/cuda/CUDAGraph.cpp

@@ -1,6 +1,8 @@
+#include <ATen/core/CachingHostAllocator.h>
 #include <ATen/cuda/CUDAGeneratorImpl.h>
 #include <ATen/cuda/CUDAGraph.h>
 #include <ATen/cuda/Exceptions.h>
+#include <ATen/cuda/MemPool.h>
 #include <ATen/Functions.h>
 #include <c10/cuda/CUDAFunctions.h>
 
@@ -13,7 +15,7 @@ static bool _cuda_graphs_debug = false;
 MempoolId_t graph_pool_handle() {
   // Sets just the second value, to distinguish it from MempoolId_ts created from
   // cudaStreamGetCaptureInfo id_s in capture_begin.
-  return c10::cuda::MemPool::graph_pool_handle();
+  return at::cuda::MemPool::graph_pool_handle();
 }
 
 /**
@@ -90,7 +92,7 @@ void CUDAGraph::capture_begin(MempoolId_t pool/*=0*/, cudaStreamCaptureMode capt
   } else {
     // User did not ask us to share a mempool. Create graph pool handle using is_user_created=false.
     // Sets just the first value, to distinguish it from MempoolId_ts created by graph_pool_handle().
-    mempool_id_ = c10::cuda::MemPool::graph_pool_handle(false);
+    mempool_id_ = at::cuda::MemPool::graph_pool_handle(false);
     TORCH_INTERNAL_ASSERT(mempool_id_.first > 0);
   }
 
@@ -104,6 +106,13 @@ void CUDAGraph::capture_begin(MempoolId_t pool/*=0*/, cudaStreamCaptureMode capt
       return status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive && stream_capture_id == capture_id_;
   });
 
+  at::getHostAllocator(at::kCUDA)->begin_allocate_to_pool(mempool_id_, [this](c10::Stream stream) {
+    cudaStreamCaptureStatus status{};
+    CaptureId_t stream_capture_id = 0;
+    AT_CUDA_CHECK(cudaStreamGetCaptureInfo(CUDAStream(CUDAStream::UNCHECKED, stream), &status, &stream_capture_id));
+    return status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive && stream_capture_id == capture_id_;
+  });
+
   // cudaStreamCaptureModeGlobal is the most conservative option to
   // prevent potentially unsafe CUDA API calls during capture.  See
   // https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html#group__CUDART__STREAM_1g9d0535d93a214cbf126835257b16ba85
@@ -124,6 +133,7 @@ void CUDAGraph::capture_end() {
   AT_CUDA_CHECK(cudaStreamEndCapture(capture_stream_, &graph_));
 
   c10::cuda::CUDACachingAllocator::endAllocateToPool(capture_dev_, mempool_id_);
+  at::getHostAllocator(at::kCUDA)->end_allocate_to_pool(mempool_id_);
 
   TORCH_CHECK(graph_ != nullptr, "Invalid capture.");
 
@@ -278,6 +288,7 @@ void CUDAGraph::reset() {
   if (capture_ended_) {
     // notifyCaptureDestroy may throw. How should we handle this?
     c10::cuda::CUDACachingAllocator::releasePool(capture_dev_, mempool_id_);
+    at::getHostAllocator(at::kCUDA)->release_pool(mempool_id_);
     capture_ended_ = false;
   }
   if (has_graph_) {

aten/src/ATen/cuda/CUDAGraphsUtils.cuh

@@ -14,7 +14,7 @@
 
 namespace at::cuda {
 
-using CaptureId_t = c10::cuda::CaptureId_t;
+using CaptureId_t = c10::CaptureId_t;
 using CaptureStatus = c10::cuda::CaptureStatus;
 
 // Use this version where you don't want to create a CUDA context if none exists.

aten/src/ATen/cuda/CUDAGreenContext.cpp

@@ -1,90 +1,78 @@
 #include <ATen/cuda/CUDAGreenContext.h>
 
-#if defined(CUDA_VERSION) && (CUDA_VERSION >= 12030) && !defined(USE_ROCM) && defined(PYTORCH_C10_DRIVER_API_SUPPORTED)
-#include <c10/cuda/driver_api.h>
-#include <stdexcept>
-#include <vector>
-#define HAS_CUDA_GREEN_CONTEXT() 1
-#else
-#define HAS_CUDA_GREEN_CONTEXT() 0
-// Suppress unsued private field warnings as this class is not supposed to be called
-C10_DIAGNOSTIC_PUSH_AND_IGNORED_IF_DEFINED("-Wunused-private-field")
-#endif
-
 namespace at::cuda {
+  GreenContext::GreenContext(uint32_t device_id, uint32_t num_sms) {
+#if CUDA_HAS_GREEN_CONTEXT
+    int driver_version;
+    C10_CUDA_CHECK(cudaDriverGetVersion(&driver_version));
+    TORCH_CHECK(
+        driver_version >= 12080, "cuda driver too old to use green context!");
+    CUcontext pctx = nullptr;
+    C10_CUDA_DRIVER_CHECK(c10::cuda::DriverAPI::get()->cuCtxGetCurrent_(&pctx));
+    if (C10_UNLIKELY(!pctx)) {
+      TORCH_WARN(
+          "Attempted to create a green context but"
+          " there was no primary context! Creating a primary context...");
 
-GreenContext::GreenContext(uint32_t device_id, uint32_t num_sms) {
-#if HAS_CUDA_GREEN_CONTEXT()
-  int driver_version;
-  C10_CUDA_CHECK(cudaDriverGetVersion(&driver_version));
-  TORCH_CHECK(
-      driver_version >= 12080, "cuda driver too old to use green context!");
-  CUcontext pctx = nullptr;
-  C10_CUDA_DRIVER_CHECK(c10::cuda::DriverAPI::get()->cuCtxGetCurrent_(&pctx));
-  if (C10_UNLIKELY(!pctx)) {
-    TORCH_WARN(
-        "Attempted to create a green context but"
-        " there was no primary context! Creating a primary context...");
+      cudaFree(0);
+    }
 
-    cudaFree(0);
-  }
+    CUdevice device;
+    device_id_ = device_id;
+    C10_CUDA_DRIVER_CHECK(
+        c10::cuda::DriverAPI::get()->cuDeviceGet_(&device, device_id));
 
-   CUdevice device;
-  device_id_ = device_id;
-  C10_CUDA_DRIVER_CHECK(
-      c10::cuda::DriverAPI::get()->cuDeviceGet_(&device, device_id));
+    // Get device resources
+    CUdevResource device_resource;
+    C10_CUDA_DRIVER_CHECK(c10::cuda::DriverAPI::get()->cuDeviceGetDevResource_(
+        device, &device_resource, CU_DEV_RESOURCE_TYPE_SM));
 
-  // Get device resources
-  CUdevResource device_resource;
-  C10_CUDA_DRIVER_CHECK(c10::cuda::DriverAPI::get()->cuDeviceGetDevResource_(
-      device, &device_resource, CU_DEV_RESOURCE_TYPE_SM));
+    // Split resources
+    std::vector<CUdevResource> result(1);
+    auto result_data = result.data();
+    unsigned int nb_groups = 1;
+    CUdevResource remaining;
 
-  // Split resources
-  std::vector<CUdevResource> result(1);
-  auto result_data = result.data();
-  unsigned int nb_groups = 1;
-  CUdevResource remaining;
+    C10_CUDA_DRIVER_CHECK(
+        c10::cuda::DriverAPI::get()->cuDevSmResourceSplitByCount_(
+            result_data,
+            &nb_groups,
+            &device_resource,
+            &remaining,
+            0, // default flags
+            num_sms));
 
-  C10_CUDA_DRIVER_CHECK(
-      c10::cuda::DriverAPI::get()->cuDevSmResourceSplitByCount_(
-          result_data,
-          &nb_groups,
-          &device_resource,
-          &remaining,
-          0, // default flags
-          num_sms));
+    TORCH_CHECK(nb_groups == 1, "Failed to create single resource group");
 
-  TORCH_CHECK(nb_groups == 1, "Failed to create single resource group");
+    // Generate resource descriptor
+    CUdevResourceDesc desc;
+    C10_CUDA_DRIVER_CHECK(
+        c10::cuda::DriverAPI::get()->cuDevResourceGenerateDesc_(
+            &desc, result_data, 1));
 
-  // Generate resource descriptor
-  CUdevResourceDesc desc;
-  C10_CUDA_DRIVER_CHECK(
-      c10::cuda::DriverAPI::get()->cuDevResourceGenerateDesc_(
-          &desc, result_data, 1));
+    // Create green context
+    // CU_GREEN_CTX_DEFAULT_STREAM is required per docs:
+    // https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__GREEN__CONTEXTS.html
+    C10_CUDA_DRIVER_CHECK(c10::cuda::DriverAPI::get()->cuGreenCtxCreate_(
+        &green_ctx_, desc, device, CU_GREEN_CTX_DEFAULT_STREAM));
 
-  // Create green context
-  // CU_GREEN_CTX_DEFAULT_STREAM is required per docs:
-  // https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__GREEN__CONTEXTS.html
-  C10_CUDA_DRIVER_CHECK(c10::cuda::DriverAPI::get()->cuGreenCtxCreate_(
-      &green_ctx_, desc, device, CU_GREEN_CTX_DEFAULT_STREAM));
-
-  // Convert to regular context
-  C10_CUDA_DRIVER_CHECK(
-      c10::cuda::DriverAPI::get()->cuCtxFromGreenCtx_(&context_, green_ctx_));
-  TORCH_CHECK(context_, "Green ctx conversion to regular ctx failed!");
+    // Convert to regular context
+    C10_CUDA_DRIVER_CHECK(
+        c10::cuda::DriverAPI::get()->cuCtxFromGreenCtx_(&context_, green_ctx_));
+    TORCH_CHECK(context_, "Green ctx conversion to regular ctx failed!");
 #else
-  TORCH_CHECK(false, "Green Context is only supported on CUDA 12.8+!");
+    TORCH_CHECK(false, "Green Context is only supported on CUDA 12.8+!");
 #endif
   }
 
   std::unique_ptr<GreenContext> GreenContext::create(
       uint32_t num_sms,
       std::optional<uint32_t> device_id) {
-#if HAS_CUDA_GREEN_CONTEXT()
+#if CUDA_HAS_GREEN_CONTEXT
     if (!device_id.has_value()) {
       device_id = at::cuda::current_device();
     }
-    return std::unique_ptr<GreenContext>(new GreenContext(device_id.value(), num_sms));
+    return std::make_unique<GreenContext>(device_id.value(), num_sms);
 #else
     TORCH_CHECK(false, "Green Context is only supported on CUDA 12.8+!");
 #endif
@@ -92,7 +80,7 @@ GreenContext::GreenContext(uint32_t device_id, uint32_t num_sms) {
 
   // Implement move operations
   GreenContext::GreenContext(GreenContext&& other) noexcept{
-#if HAS_CUDA_GREEN_CONTEXT()
+#if CUDA_HAS_GREEN_CONTEXT
     device_id_ = std::exchange(other.device_id_, -1);
     green_ctx_ = std::exchange(other.green_ctx_, nullptr);
     context_ = std::exchange(other.context_, nullptr);
@@ -103,7 +91,7 @@ GreenContext::GreenContext(uint32_t device_id, uint32_t num_sms) {
   }
 
   GreenContext& GreenContext::operator=(GreenContext&& other) noexcept{
-#if HAS_CUDA_GREEN_CONTEXT()
+#if CUDA_HAS_GREEN_CONTEXT
     if (this != &other) {
       // Clean up current resources
       if (green_ctx_) {
@@ -132,7 +120,7 @@ GreenContext::GreenContext(uint32_t device_id, uint32_t num_sms) {
   }
 
   GreenContext::~GreenContext() noexcept{
-#if HAS_CUDA_GREEN_CONTEXT()
+#if CUDA_HAS_GREEN_CONTEXT
     C10_CUDA_DRIVER_CHECK(
         c10::cuda::DriverAPI::get()->cuGreenCtxDestroy_(green_ctx_));
 #else
@@ -140,9 +128,25 @@ GreenContext::GreenContext(uint32_t device_id, uint32_t num_sms) {
 #endif
   }
 
+  // Get the underlying CUDA context
+  CUcontext GreenContext::getContext() const {
+#if CUDA_HAS_GREEN_CONTEXT
+    return context_;
+#else
+    TORCH_CHECK(false, "Green Context is only supported on CUDA 12.8+!");
+#endif
+  }
+
+  // Get the underlying green context
+#if CUDA_HAS_GREEN_CONTEXT
+  CUgreenCtx GreenContext::getGreenContext() const {
+    return green_ctx_;
+  }
+#endif
+
   // Make this context current
   void GreenContext::setContext() {
-#if HAS_CUDA_GREEN_CONTEXT()
+#if CUDA_HAS_GREEN_CONTEXT
     auto current_stream = c10::cuda::getCurrentCUDAStream();
     parent_stream_ = current_stream.stream();
 
@@ -171,7 +175,7 @@ GreenContext::GreenContext(uint32_t device_id, uint32_t num_sms) {
   }
 
   void GreenContext::popContext() {
-#if HAS_CUDA_GREEN_CONTEXT()
+#if CUDA_HAS_GREEN_CONTEXT
     // see above note about stream being hardcoded to the default stream
     at::cuda::CUDAEvent ev;
     ev.record(c10::cuda::getCurrentCUDAStream());

aten/src/ATen/cuda/CUDAGreenContext.h

@@ -1,38 +1,53 @@
 #pragma once
 #include <ATen/cuda/CUDAEvent.h>
-#include <cuda.h>
 
-// Forward declare green context as opaque ptr
-typedef struct CUgreenCtx_st* CUgreenCtx;
+#if defined(CUDA_VERSION) && !defined(USE_ROCM) && defined(PYTORCH_C10_DRIVER_API_SUPPORTED)
+#include <c10/cuda/driver_api.h>
+#include <cuda.h>
+#include <memory>
+#include <stdexcept>
+#include <vector>
+#define CUDA_HAS_GREEN_CONTEXT 1
+#else
+#define CUDA_HAS_GREEN_CONTEXT 0
+#endif
 
 namespace at::cuda {
 
 class TORCH_CUDA_CPP_API GreenContext {
  public:
-  // Green context creation
-  static std::unique_ptr<GreenContext> create(
-      uint32_t num_sms,
-      std::optional<uint32_t> device_id);
-  ~GreenContext() noexcept;
+  GreenContext(uint32_t device_id, uint32_t num_sms);
+
+  static std::unique_ptr<GreenContext> create(uint32_t num_sms, std::optional<uint32_t> device_id);
 
   // Delete copy constructor and assignment
   GreenContext(const GreenContext&) = delete;
   GreenContext& operator=(const GreenContext&) = delete;
 
+  // Implement move operations
+  GreenContext(GreenContext&& other) noexcept;
+  GreenContext& operator=(GreenContext&& other) noexcept;
+  ~GreenContext() noexcept;
+
+  // Get the underlying CUDA context
+  CUcontext getContext() const;
+
+  // Get the underlying green context
+#if CUDA_HAS_GREEN_CONTEXT
+  CUgreenCtx getGreenContext() const;
+#endif
+
   // Make this context current
   void setContext();
 
   void popContext();
 
  private:
-  GreenContext(uint32_t device_id, uint32_t num_sms);
-  // Implement move operations
-  GreenContext(GreenContext&& other) noexcept;
-  GreenContext& operator=(GreenContext&& other) noexcept;
-
+#if CUDA_HAS_GREEN_CONTEXT
   int32_t device_id_ = -1;
   CUgreenCtx green_ctx_ = nullptr;
   CUcontext context_ = nullptr;
   cudaStream_t parent_stream_ = nullptr;
+#endif
 };
 } // namespace at::cuda

aten/src/ATen/cuda/CachingHostAllocator.cpp

@@ -6,6 +6,8 @@
 #include <ATen/detail/CUDAHooksInterface.h>
 #include <c10/core/thread_pool.h>
 #include <c10/cuda/CUDAAllocatorConfig.h>
+#include <c10/cuda/CUDAGraphsC10Utils.h>
+#include <c10/util/hash.h>
 
 #include <cuda_runtime_api.h>
 #include <future>
@@ -72,6 +74,173 @@ struct CUDACachingHostAllocatorImpl
     : public CachingHostAllocatorImpl<CUDAStream, EventPool::Event> {
  private:
   ska::flat_hash_map<void*, bool> use_host_register;
+public:
+  bool record_event(void* ptr, void* ctx, c10::Stream s) override {
+    if (!CachingHostAllocatorImpl<CUDAStream, EventPool::Event>::record_event(ptr, ctx, s)) {
+      return false;
+    }
+
+    // Inspect the preceding captured node on stream s. If it's a
+    // memcpy from pinned host memory that reads from this block, hash
+    // the copied section and record it to detect accidental mutations
+    // during capture. The purpose is to identify cases where the CPU
+    // writes to host pinned memory after the GPU is scheduled to read
+    // from it during cuda graph replay. I am calling this a
+    // CPU-Write-After-GPU-Read hazard. Such a situation is almost
+    // certainly a bug, bcause it doesn't match the semantics of eager
+    // launch. That check happens in end_allocate_to_pool.
+
+    Block *block = nullptr;
+    if (block_exists(ctx)) {
+      block = reinterpret_cast<Block*>(ctx);
+    } else {
+      block = get_block_from_ptr(ptr);
+    }
+
+    assert(block != nullptr);
+
+    CUDAStream stream(CUDAStream::UNCHECKED, s);
+
+    std::optional<std::tuple<const cudaGraphNode_t*, size_t>> terminals_opt = c10::cuda::streamGetTerminalNodes(stream);
+
+    if (!terminals_opt.has_value()) {
+      // no stream capture is happening right now
+      return true;
+    }
+
+    auto&& [terminals, num_terminals] = *terminals_opt;
+
+    if (num_terminals == 0) {
+      TORCH_WARN(
+          "CachingHostAllocator: no nodes precede record_event() during stream capture; "
+          "this is probably a bug.");
+      return true;
+    }
+
+    if (num_terminals > 1) {
+      TORCH_WARN(
+          "CachingHostAllocator: multiple preceding nodes during capture; "
+          "skipping host read verification.");
+      return true;
+    }
+
+    cudaGraphNode_t node = terminals[0];
+    cudaGraphNodeType type{cudaGraphNodeTypeCount};
+    C10_CUDA_CHECK(cudaGraphNodeGetType(node, &type));
+    if (type != cudaGraphNodeTypeMemcpy) {
+      TORCH_WARN(
+          "CachingHostAllocator: preceding captured node is not a memcpy; "
+          "can only check for CPU-write-after-GPU-read errors for memcpy nodes.");
+      return true;
+    }
+
+    cudaMemcpy3DParms params{};
+    C10_CUDA_CHECK(cudaGraphMemcpyNodeGetParams(node, &params));
+
+    // Only handle Host->Device copies (source is host memory)
+    if (params.kind != cudaMemcpyHostToDevice && params.kind != cudaMemcpyDefault) {
+      // don't warn, since Device->Host copy is a valid node to precede record_event()
+      return true;
+    }
+
+    cudaPointerAttributes attr{};
+    cudaError_t attr_err = cudaPointerGetAttributes(&attr, params.srcPtr.ptr);
+    // TODO: Think about what to do if the memory type is
+    // cudaMemoryTypeUnregistered, which I believe refers to pageable
+    // host memory. (Meanwhile, cudaMemoryTypeHost represents pinned
+    // host memory.)
+    if (attr_err != cudaSuccess || attr.type != cudaMemoryTypeHost) {
+      // Clear error if any and return (not pinned host memory)
+      (void)cudaGetLastError();
+      TORCH_WARN("CachingHostAllocator: source of memcpy is not pinned host memory");
+      return true;
+    }
+
+    // Compute the start pointer and a conservative contiguous span size.
+    // For typical 1D host->device copies, height=depth=1 and srcPos.y/z=0.
+    if (params.extent.height > 1 || params.extent.depth > 1) {
+      TORCH_WARN("CachingHostAllocator: non-1D memcopies are not supported for CPU-write-after-GPU-read detection. File an issue if this is important to you.");
+      return true;
+    }
+    char* start_ptr = static_cast<char*>(params.srcPtr.ptr) + params.srcPos.x;
+    size_t n = params.extent.width;
+
+    if (n == 0) {
+      TORCH_WARN("CachingHostAllocator: cudaMemcpyAsync node is copying 0 bytes");
+      return true;
+    }
+
+    // Verify the memcpy source buffer is within this block's bounds
+    char *block_start, *block_end;
+    {
+      std::lock_guard<std::mutex> lg(block->mutex_);
+      block_start = static_cast<char*>(block->ptr_);
+      block_end = block_start + block->size_;
+    }
+    char* src_end = start_ptr + n;
+    if (start_ptr < block_start || src_end > block_end) {
+      TORCH_WARN("CachingHostAllocator: cudaMemcpyAsync node is copying memory outside this block's memory region, which shouldn't be possible.");
+      return true;
+    }
+
+    // Hash the source buffer and record the tuple (ptr, size, hash)
+    c10::sha1 hasher(start_ptr, n);
+    std::string hash_string = hasher.str();
+    {
+      std::lock_guard<std::mutex> lg(block->mutex_);
+      block->sections_read_under_stream_capture_.emplace_back(
+          static_cast<void*>(start_ptr), n, std::move(hash_string));
+    }
+    return true;
+  }
+
+  // See comment in the header: we need to validate that any pinned host
+  // memory regions read by captured memcpy nodes have not been modified during
+  // capture. Perform the validation at the end of allocation-to-pool.
+  void end_allocate_to_pool(c10::MempoolId_t pool_id) {
+    // Look up the private pool for this capture.
+    {
+      std::shared_lock<std::shared_mutex> lg(instance_mutex_);
+      auto it = graph_pools_.find(pool_id);
+      if (it != graph_pools_.end()) {
+        auto& pool = it->second->blocks;
+        // Lock blocks set while we iterate
+        std::lock_guard<std::mutex> gb(pool.blocks_mutex_);
+        for (auto* block : pool.blocks_) {
+          // Copy sections under block lock and clear them to avoid rechecking
+          std::vector<std::tuple<void*, size_t, std::string>> sections;
+          {
+            std::lock_guard<std::mutex> bl(block->mutex_);
+            sections = block->sections_read_under_stream_capture_;
+            block->sections_read_under_stream_capture_.clear();
+          }
+          for (const auto& tup : sections) {
+            void* start = nullptr;
+            size_t len = 0;
+            const std::string* expected_hash = nullptr;
+            // tie cannot bind to temporary from get<>, extract manually
+            start = std::get<0>(tup);
+            len = std::get<1>(tup);
+            expected_hash = &std::get<2>(tup);
+            c10::sha1 hasher(start, len);
+            std::string actual = hasher.str();
+            TORCH_CHECK(
+                actual == *expected_hash,
+                "CUDA graph replay will not match eager execution: a pinned host "
+                "memory buffer read by a memcpy during capture was modified after "
+                "it was recorded. To fix, never write to pinned host memory after "
+                "stream capture records a memcpy that reads from it.");
+          }
+        }
+      }
+    }
+
+    // Defer to base implementation to end allocation to this pool.
+    CachingHostAllocatorImpl<CUDAStream, EventPool::Event>::end_allocate_to_pool(
+        pool_id);
+  }
+
+ private:
 
   void allocate_host_memory(size_t size, void** ptr) override {
     // try allocating from reserve segment first before calling into expensive APIs
@@ -107,6 +276,7 @@ struct CUDACachingHostAllocatorImpl
       allocWithCudaHostRegister(ptr, size);
     } else {
       // Use cudaHostAlloc for allocating pinned memory (global lock in driver)
+      at::cuda::CUDAStreamCaptureModeGuard g{cudaStreamCaptureModeRelaxed};
       C10_CUDA_CHECK(cudaHostAlloc(ptr, size, cudaHostAllocDefault));
     }
 
@@ -173,6 +343,14 @@ struct CUDACachingHostAllocatorImpl
   }
 
   bool query_event(EventPool::Event& event) override {
+    // It is rare, but query_event() can be called during stream
+    // capture. This happens when stream capture is immediately
+    // preceded by allocating to this pool via
+    // _use_cuda_memory_pool_manager. Since capture_begin() is
+    // preceded by torch.cuda.synchronize() in torch/cuda/graphs.py,
+    // we can be certain that cudaEventQuery always returns
+    // cudaSuccess in this rare situation.
+    at::cuda::CUDAStreamCaptureModeGuard g{cudaStreamCaptureModeRelaxed};
     cudaError_t err = cudaEventQuery(*event);
     if (err == cudaErrorNotReady) {
       (void)cudaGetLastError(); // clear CUDA error
@@ -280,9 +458,20 @@ struct CUDACachingHostAllocatorImpl
     }
 
     // Register the mapped pages using cudaHostRegister
+    at::cuda::CUDAStreamCaptureModeGuard g{cudaStreamCaptureModeRelaxed};
     AT_CUDA_CHECK(
         cudaHostRegister(*ptr, roundSize, cudaHostRegisterDefault));
   }
+
+  CUDAStream get_current_stream() const override {
+    return at::cuda::getCurrentCUDAStream();
+  }
+
+  bool stream_is_capturing(CUDAStream s) const override {
+    cudaStreamCaptureStatus status{cudaStreamCaptureStatusNone};
+    C10_CUDA_CHECK(cudaStreamIsCapturing(s, &status));
+    return status != cudaStreamCaptureStatusNone;
+  }
 };
 
 DECLARE_HOST_ALLOCATOR(

aten/src/ATen/cuda/CublasHandlePool.cpp

@@ -155,8 +155,8 @@ size_t parseChosenWorkspaceSize() {
     while (next != end) {
       std::smatch match = *next;
       TORCH_CHECK(match.size() == 3, "Expected CUBLAS_WORKSPACE_SPACE_CONFIG match of size 3 (Format :SIZE:COUNT)");
-      size_t curr_size = std::stoull(match.str(1));
-      size_t count = std::stoull(match.str(2));
+      size_t curr_size = (size_t) std::stoi(match.str(1));
+      size_t count = (size_t) std::stoi(match.str(2));
       total_size += curr_size * 1024 * count;
       next++;
     }

aten/src/ATen/cuda/MemPool.cpp

@@ -0,0 +1,72 @@
+#include <ATen/core/CachingHostAllocator.h>
+#include <ATen/cuda/MemPool.h>
+
+namespace at::cuda {
+
+// uid_ is incremented when a user creates a MemPool,
+// for example: using graph_pool_handle() or c10::cuda::MemPool().
+//
+// uuid_ is incremented when CUDAGraph creates a MemPool
+// as a result of a user not providing a pool.
+//
+// MempoolId_t of {0, 0} is used to denote when no MemPool has been
+// passed to a function, either by user or CUDAGraphs. For example,
+// default value of MempoolId_t for capture_begin function is {0, 0}.
+// That's why uid_ and uuid_ start at 1.
+std::atomic<CaptureId_t> MemPool::uid_{1};
+std::atomic<CaptureId_t> MemPool::uuid_{1};
+
+MemPool::MemPool(
+    CUDACachingAllocator::CUDAAllocator* allocator,
+    bool is_user_created,
+    bool use_on_oom)
+    : allocator_(allocator), is_user_created_(is_user_created) {
+  if (is_user_created_) {
+    id_ = {0, uid_++};
+  } else {
+    id_ = {uuid_++, 0};
+  }
+  device_ = c10::cuda::current_device();
+  CUDACachingAllocator::createOrIncrefPool(device_, id_, allocator);
+  if (use_on_oom) {
+    CUDACachingAllocator::setUseOnOOM(device_, id_);
+  }
+  at::getHostAllocator(at::kCUDA)->create_or_incref_pool(id_);
+}
+
+MemPool::~MemPool() {
+  // TORCH_INTERNAL_ASSERT(use_count() == 1);
+  // We used to assert that TORCH_INTERNAL_ASSERT(use_count() == 1);
+  // However, this assertion is not true if a memory pool is shared
+  // with a cuda graph. That CUDAGraph will increase the use count
+  // until it is reset.
+  CUDACachingAllocator::releasePool(device_, id_);
+  c10::cuda::CUDACachingAllocator::emptyCache(id_);
+  at::getHostAllocator(at::kCUDA)->release_pool(id_);
+  // TODO: at::getHostAllocator(at::kCUDA)->empty_cache(id_)
+}
+
+MempoolId_t MemPool::id() {
+  return id_;
+}
+
+CUDACachingAllocator::CUDAAllocator* MemPool::allocator() {
+  return allocator_;
+}
+
+int MemPool::use_count() {
+  return CUDACachingAllocator::getPoolUseCount(device_, id_);
+}
+
+c10::DeviceIndex MemPool::device() {
+  return device_;
+}
+
+MempoolId_t MemPool::graph_pool_handle(bool is_user_created) {
+  if (is_user_created) {
+    return {0, uid_++};
+  }
+  return {uuid_++, 0};
+}
+
+} // namespace at::cuda

aten/src/ATen/cuda/MemPool.h

@@ -0,0 +1,44 @@
+#pragma once
+
+#include <c10/core/Allocator.h>
+#include <c10/cuda/CUDACachingAllocator.h>
+
+namespace at::cuda {
+
+// Keep BC only
+using c10::CaptureId_t;
+using c10::MempoolId_t;
+
+// MemPool represents a pool of memory in a caching allocator. Currently,
+// it's just the ID of the pool object maintained in the CUDACachingAllocator.
+//
+// An allocator pointer can be passed to the MemPool to define how the
+// allocations should be done in the pool. For example: using a different
+// system allocator such as ncclMemAlloc.
+struct TORCH_CUDA_CPP_API MemPool {
+  MemPool(
+      c10::cuda::CUDACachingAllocator::CUDAAllocator* allocator = nullptr,
+      bool is_user_created = true,
+      bool use_on_oom = false);
+  MemPool(const MemPool&) = delete;
+  MemPool(MemPool&&) = default;
+  MemPool& operator=(const MemPool&) = delete;
+  MemPool& operator=(MemPool&&) = default;
+  ~MemPool();
+
+  MempoolId_t id();
+  c10::cuda::CUDACachingAllocator::CUDAAllocator* allocator();
+  int use_count();
+  c10::DeviceIndex device();
+  static MempoolId_t graph_pool_handle(bool is_user_created = true);
+
+ private:
+  static std::atomic<CaptureId_t> uid_;
+  static std::atomic<CaptureId_t> uuid_;
+  c10::cuda::CUDACachingAllocator::CUDAAllocator* allocator_;
+  bool is_user_created_;
+  MempoolId_t id_;
+  c10::DeviceIndex device_;
+};
+
+} // namespace at::cuda

aten/src/ATen/cuda/NumericLimits.cuh

@@ -55,14 +55,6 @@ struct numeric_limits<int8_t> {
   static inline __host__ __device__ int8_t upper_bound() { return INT8_MAX; }
 };
 
-template <>
-struct numeric_limits<uint16_t> {
-  static inline __host__ __device__ uint16_t lowest() { return 0; }
-  static inline __host__ __device__ uint16_t max() { return UINT16_MAX; }
-  static inline __host__ __device__ uint16_t lower_bound() { return 0; }
-  static inline __host__ __device__ uint16_t upper_bound() { return UINT16_MAX; }
-};
-
 template <>
 struct numeric_limits<int16_t> {
   static inline __host__ __device__ int16_t lowest() { return INT16_MIN; }
@@ -71,14 +63,6 @@ struct numeric_limits<int16_t> {
   static inline __host__ __device__ int16_t upper_bound() { return INT16_MAX; }
 };
 
-template <>
-struct numeric_limits<uint32_t> {
-  static inline __host__ __device__ uint32_t lowest() { return 0; }
-  static inline __host__ __device__ uint32_t max() { return UINT32_MAX; }
-  static inline __host__ __device__ uint32_t lower_bound() { return 0; }
-  static inline __host__ __device__ uint32_t upper_bound() { return UINT32_MAX; }
-};
-
 template <>
 struct numeric_limits<int32_t> {
   static inline __host__ __device__ int32_t lowest() { return INT32_MIN; }
@@ -87,21 +71,6 @@ struct numeric_limits<int32_t> {
   static inline __host__ __device__ int32_t upper_bound() { return INT32_MAX; }
 };
 
-template <>
-struct numeric_limits<uint64_t> {
-#ifdef _MSC_VER
-  static inline __host__ __device__ uint64_t lowest() { return 0; }
-  static inline __host__ __device__ uint64_t max() { return _UI64_MAX; }
-  static inline __host__ __device__ uint64_t lower_bound() { return 0; }
-  static inline __host__ __device__ uint64_t upper_bound() { return _UI64_MAX; }
-#else
-  static inline __host__ __device__ uint64_t lowest() { return 0; }
-  static inline __host__ __device__ uint64_t max() { return UINT64_MAX; }
-  static inline __host__ __device__ uint64_t lower_bound() { return 0; }
-  static inline __host__ __device__ uint64_t upper_bound() { return UINT64_MAX; }
-#endif
-};
-
 template <>
 struct numeric_limits<int64_t> {
 #ifdef _MSC_VER

aten/src/ATen/cuda/cub.h

@@ -24,13 +24,7 @@ namespace detail {
 // radix_sort_pairs doesn't interact with value_t other than to copy
 // the data, so we can save template instantiations by reinterpreting
 // it as an opaque type.
-// We use native integer types for 1/2/4/8-byte values to reduce
-// register usage in CUDA kernels. For sizes > 8 fall back to char array.
 template <int N> struct alignas(N) OpaqueType { char data[N]; };
-template <> struct alignas(1) OpaqueType<1> { uint8_t data; };
-template <> struct alignas(2) OpaqueType<2> { uint16_t data; };
-template <> struct alignas(4) OpaqueType<4> { uint32_t data; };
-template <> struct alignas(8) OpaqueType<8> { uint64_t data; };
 
 template<typename key_t, int value_size>
 void radix_sort_pairs_impl(

aten/src/ATen/cuda/detail/BLASConstants.cu

@@ -2,6 +2,8 @@
 #include <ATen/Tensor.h>
 #include <ATen/cuda/Exceptions.h>
 
+#include <mutex>
+
 namespace at {
 namespace cuda {
 namespace detail {
@@ -10,36 +12,39 @@ __device__ __constant__ float cublas_one_device;
 __device__ __constant__ float cublas_zero_device;
 
 float *get_cublas_device_one() {
-  static float *ptr = nullptr;
-  static auto init_flag = [&]() {
+  static c10::once_flag init_flag;
+
+  c10::call_once(init_flag, []() {
     const float one = 1.f;
     AT_CUDA_CHECK(cudaMemcpyToSymbol(cublas_one_device, &one, sizeof(float)));
-    AT_CUDA_CHECK(cudaGetSymbolAddress(reinterpret_cast<void**>(&ptr), cublas_one_device));
-    return true;
-  }();
+  });
 
+  float *ptr;
+  AT_CUDA_CHECK(cudaGetSymbolAddress(reinterpret_cast<void**>(&ptr), cublas_one_device));
   return ptr;
 }
 
 float *get_cublas_device_zero() {
-  static float *ptr = nullptr;
-  static auto init_flag = [&]() {
+  static c10::once_flag init_flag;
+
+  c10::call_once(init_flag, []() {
     const float zero = 0.f;
     AT_CUDA_CHECK(cudaMemcpyToSymbol(cublas_zero_device, &zero, sizeof(float)));
-    AT_CUDA_CHECK(cudaGetSymbolAddress(reinterpret_cast<void**>(&ptr), cublas_zero_device));
-    return true;
-  }();
+  });
 
+  float *ptr;
+  AT_CUDA_CHECK(cudaGetSymbolAddress(reinterpret_cast<void**>(&ptr), cublas_zero_device));
   return ptr;
 }
 
 float *get_user_alpha_ptr() {
   static float *alpha_ptr;
 
-  static bool init_flag [[maybe_unused]] = []() {
+  static c10::once_flag init_flag;
+
+  c10::call_once(init_flag, []() {
     AT_CUDA_CHECK(cudaMalloc(&alpha_ptr, sizeof(float)));
-    return true;
-  }();
+  });
 
   return alpha_ptr;
 }

aten/src/ATen/cudnn/Descriptors.cpp

@@ -3,7 +3,6 @@
 #include <ATen/ATen.h>
 #include <c10/util/irange.h>
 
-#include <array>
 #include <iostream>
 #include <sstream>
 
@@ -137,9 +136,9 @@ void FilterDescriptor::set(const at::Tensor &t, const at::MemoryFormat memory_fo
     "Weight strides: ", t.strides(), "\n",
     "cuDNN suggested memory_format: ", memory_format);
 
-  std::array<int, CUDNN_DIM_MAX> size;
+  int size[CUDNN_DIM_MAX];
   for (const auto i : c10::irange(dim)) {
-    size[i] = static_cast<int>(t.size(i));
+    size[i] = (int) t.size(i);
   }
   for (const auto i : c10::irange(dim, pad)) {
     size[i] = 1;
@@ -157,7 +156,7 @@ void FilterDescriptor::set(const at::Tensor &t, const at::MemoryFormat memory_fo
     default:
       TORCH_INTERNAL_ASSERT(false, "unsupported memory_format for cuDNN filters");
   }
-  set(getDataType(t), static_cast<int>(dim), size.data(), filter_format);
+  set(getDataType(t), static_cast<int>(dim), size, filter_format);
 }
 
 std::string cudnnMemoryFormatToString(cudnnTensorFormat_t tformat) {

aten/src/ATen/detail/MTIAHooksInterface.h

@@ -9,8 +9,8 @@
 
 #include <c10/core/Allocator.h>
 
-#include <ATen/detail/AcceleratorHooksInterface.h>
 #include <c10/util/python_stub.h>
+#include <ATen/detail/AcceleratorHooksInterface.h>
 
 #include <string>
 namespace at {
@@ -26,7 +26,8 @@ constexpr const char* MTIA_HELP =
 struct TORCH_API MTIAHooksInterface : AcceleratorHooksInterface {
 // this fails the implementation if MTIAHooks functions are called, but
 // MTIA backend is not present.
-#define FAIL_MTIAHOOKS_FUNC(func) TORCH_CHECK(false, "Cannot execute ", func, "() without MTIA backend.");
+#define FAIL_MTIAHOOKS_FUNC(func) \
+  TORCH_CHECK(false, "Cannot execute ", func, "() without MTIA backend.");
 
   ~MTIAHooksInterface() override = default;
 
@@ -91,7 +92,7 @@ struct TORCH_API MTIAHooksInterface : AcceleratorHooksInterface {
     return c10::Stream::unpack3(-1, 0, c10::DeviceType::MTIA);
   }
 
-  virtual void setCurrentStream(const c10::Stream& /*stream*/) const {
+  virtual void setCurrentStream(const c10::Stream& /*stream*/ ) const {
     FAIL_MTIAHOOKS_FUNC(__func__);
   }
 
@@ -123,9 +124,11 @@ struct TORCH_API MTIAHooksInterface : AcceleratorHooksInterface {
     FAIL_MTIAHOOKS_FUNC(__func__);
   }
 
-  virtual void recordMemoryHistory(const std::optional<std::string>& /*enabled*/,
-                                   const std::string& /*stacks*/,
-                                   size_t /*max_entries*/) const {
+
+  virtual void recordMemoryHistory(
+    const std::optional<std::string>& /*enabled*/,
+    const std::string& /*stacks*/,
+    size_t /*max_entries*/) const {
     FAIL_MTIAHOOKS_FUNC(__func__);
   }
 
@@ -156,10 +159,6 @@ struct TORCH_API MTIAHooksInterface : AcceleratorHooksInterface {
     return -1;
   }
 
-  virtual void mtiagraphDestroy(int64_t handle) const {
-    FAIL_MTIAHOOKS_FUNC(__func__);
-  }
-
   virtual void mtiagraphCaptureBegin(int64_t handle, MempoolId_t pool) const {
     FAIL_MTIAHOOKS_FUNC(__func__);
   }
@@ -188,7 +187,8 @@ struct TORCH_API MTIAHooksInterface : AcceleratorHooksInterface {
 struct TORCH_API MTIAHooksArgs {};
 
 TORCH_DECLARE_REGISTRY(MTIAHooksRegistry, MTIAHooksInterface, MTIAHooksArgs);
-#define REGISTER_MTIA_HOOKS(clsname) C10_REGISTER_CLASS(MTIAHooksRegistry, clsname, clsname)
+#define REGISTER_MTIA_HOOKS(clsname) \
+  C10_REGISTER_CLASS(MTIAHooksRegistry, clsname, clsname)
 
 namespace detail {
 TORCH_API const MTIAHooksInterface& getMTIAHooks();

aten/src/ATen/functorch/ADInterpreters.cpp

@@ -198,7 +198,7 @@ static void autogradBasedTransformSendToNext(
   }
 
   // Step 6
-  stack->erase(stack->end() - static_cast<std::ptrdiff_t>(args_size + ret_size), stack->end() - static_cast<std::ptrdiff_t>(ret_size));
+  stack->erase(stack->end() - std::ptrdiff_t(args_size + ret_size), stack->end() - std::ptrdiff_t(ret_size));
 }
 
 void GradInterpreterPtr::processImpl(

aten/src/ATen/functorch/BatchRulesNorm.cpp

@@ -443,14 +443,14 @@ static bool has_same_shape(
   if (!tensor.defined()) {
     return true;
   }
-  if (rankWithoutBatchDim(tensor, tensor_bdim) != static_cast<int64_t>(normalized_shape.size())) {
+  if (rankWithoutBatchDim(tensor, tensor_bdim) != (int64_t) normalized_shape.size()) {
     return false;
   }
   const auto tensor_shape = tensor.sizes();
   for (const auto i : c10::irange(normalized_shape.size())) {
     auto j = i;
     // (0, 1, 2), 1 -> (0, 2, 3)
-    if (tensor_bdim.has_value() && static_cast<int64_t>(i) >= tensor_bdim.value()) {
+    if (tensor_bdim.has_value() && (int64_t)i >= tensor_bdim.value()) {
       j = j + 1;
     }
     if (normalized_shape[i] != tensor_shape[j]) {

aten/src/ATen/functorch/BatchRulesReduceOps.cpp

@@ -135,7 +135,7 @@ static void boxed_reduction_batch_rule(const c10::OperatorHandle& op, torch::jit
     reduction_case = ReductionCase::DimArray;
     dims = arguments[dim_arg_pos].toIntList().vec();
     if (dims.empty()) {
-      auto all_dims = range(0, std::max(static_cast<int64_t>(1), logical_dim));
+      auto all_dims = range(0, std::max((int64_t)1, logical_dim));
       dims = std::vector<int64_t>(all_dims.begin(), all_dims.end());
     }
   } else if (arguments[dim_arg_pos].isInt()) {

aten/src/ATen/functorch/BatchRulesScatterOps.cpp

@@ -432,7 +432,7 @@ namespace {
     // Eg. Given `indexed_shape.size()` is 5 and
     // shape of `values` is (N, 2, 3), then following block
     // will reshape `values` to (N, 1, 1, 2, 3).
-    if ( static_cast<int64_t>(indexed_shape.size()) > values_.dim()) {
+    if ( (int64_t) indexed_shape.size() > values_.dim()) {
       auto values_sizes = values_.sym_sizes();
 
       // number of unit dims (for broadcasting value to indexed_shape)