Merge branch 'main' into tkuczynski/enable_test_device_assert

Author: dev-tomek

.github/workflows/build-test-reusable.yml

@@ -127,6 +127,10 @@ jobs:
             DEBUG=1
             python -m build --wheel --no-isolation && pip install dist/*.whl
 
+      - name: Build PTI
+        run: |
+          ./scripts/install-pti.sh --build-level-zero
+
       - name: Set test-triton command line
         id: test-triton
         run: |
@@ -175,6 +179,12 @@ jobs:
           name: triton-${{ inputs.python_version }}-${{ inputs.runner_label || inputs.driver_version }}
           path: dist/*.whl
 
+      - name: Upload PTI wheels
+        uses: actions/upload-artifact@v4
+        with:
+          name: pti-${{ inputs.python_version }}-${{ inputs.runner_label || inputs.driver_version }}
+          path: .scripts_cache/pti/dist/*.whl
+
       - name: Upload test reports
         uses: actions/upload-artifact@v5
         with:
@@ -242,11 +252,23 @@ jobs:
         with:
           name: triton-${{ inputs.python_version }}-${{ inputs.runner_label || inputs.driver_version }}
 
+      - name: Download PTI wheels
+        uses: actions/download-artifact@v5
+        with:
+          name: pti-${{ inputs.python_version }}-${{ inputs.runner_label || inputs.driver_version }}
+
       - name: Install Triton
         run: |
           pip install triton-*.whl
           python -c 'import triton; print(triton.__version__)'
 
+      - name: Install PTI
+        run: |
+          pip install intel_pti-*.whl
+          PTI_LIBS_DIR=$(python ./scripts/pti_lib.py)
+          ls $PTI_LIBS_DIR
+          echo "PTI_LIBS_DIR=$PTI_LIBS_DIR" | tee -a $GITHUB_ENV
+
       - name: Report environment details
         run: |
           if [ "${{ matrix.suite }}" == "minicore" ]; then
@@ -283,21 +305,7 @@ jobs:
       - name: Build PTI && Run Proton tests
         if: matrix.suite == 'rest'
         run: |
-          # `intel-pti` can be installed in "Setup PyTorch" step with `pytorch_mode==wheels`
-          pip uninstall intel-pti -y
-          PTI_COMMIT_ID="$(<.github/pins/pti.txt)"
-          git clone https://github.com/intel/pti-gpu.git
-          cd pti-gpu
-          git checkout $PTI_COMMIT_ID
-          cd sdk
-          cmake --preset linux-icpx-release
-          BUILD_TESTING=1 PTI_BUILD_SAMPLES=1 cmake --build --preset linux-icpx-release
-
-          PTI_LIBS_DIR="$(pwd)/build-linux-icpx-release/lib/"
-          cd ../..
-
-          export LD_LIBRARY_PATH=$PTI_LIBS_DIR:$LD_LIBRARY_PATH
-          export TRITON_XPUPTI_LIB_PATH=$PTI_LIBS_DIR
+          export LD_LIBRARY_PATH=${{ env.PTI_LIBS_DIR }}:$LD_LIBRARY_PATH
           cd third_party/proton/test
           pytest test_api.py test_cmd.py test_lib.py test_profile.py test_viewer.py --device xpu -s -v
           cd ..

.github/workflows/e2e-reusable.yml

@@ -140,7 +140,7 @@ jobs:
 
       - name: Install python test dependencies
         run: |
-          pip install pyyaml pandas scipy 'numpy==1.26.4' psutil pyre_extensions torchrec
+          pip install pyyaml pandas scipy psutil
 
       - name: Install transformers package
         if: ${{ inputs.suite == 'huggingface' }}
@@ -189,11 +189,19 @@ jobs:
         if: ${{ inputs.suite == 'torchbench' }}
         run: |
           cd benchmark
+          sed -i 's/^ *pynvml.*//' requirements.txt
+          pip install -r requirements.txt
+          # for dlrm
+          pip install pyre-extensions
+          curl -fsSL https://raw.githubusercontent.com/facebookresearch/dlrm/refs/heads/torchrec-dlrm/requirements.txt |xargs pip install
+          # for soft_actor_critic, temp fix
+          pip install git+https://github.com/nocoding03/gym@fix-np
+
           if [[ "${{ inputs.only_one_model }}" ]]; then
             python install.py "${{ inputs.only_one_model }}"
           else
             # install all models
-            python install.py
+            python install.py --continue_on_fail
           fi
           pip install -e .
 
@@ -203,6 +211,9 @@ jobs:
         run: |
           cd pytorch
 
+          # Some models are still not compatible with numpy>2.0
+          pip install 'numpy==1.26.4'
+
           export WORKSPACE=$GITHUB_WORKSPACE
 
           if [[ "${{ inputs.TORCH_COMPILE_DEBUG }}" = "1" ]] ; then

python/test/unit/language/test_core.py

@@ -3272,8 +3272,6 @@ def test_dot(M, N, K, num_warps, col_a, col_b, epilogue, input_precision, in_dty
             pytest.xfail(f"input_precision {input_precision} is not supported in the interpreter")
     else:
         if is_xpu():
-            if input_precision in ("bf16x3", "bf16x6"):
-                pytest.skip(f"input_precision {input_precision} is not supported")
             if (M < 8 or N < 16 or (K < 16 and in_dtype == 'float16') or (K < 8 and in_dtype == 'float32')):
                 pytest.xfail("XPU: small dots are not supported")
         elif not is_hip() and K < 16:

scripts/build_pti_data/CMakeLists.txt

@@ -0,0 +1,49 @@
+cmake_minimum_required(VERSION 3.20)
+
+project(
+  pti-python-binaries
+  VERSION 0.0.1
+  LANGUAGES C CXX)
+
+include(FetchContent)
+
+if(DEFINED ENV{PTI_PINNED_COMMIT})
+  set(PTI_PINNED_COMMIT $ENV{PTI_PINNED_COMMIT})
+else()
+  message(FATAL_ERROR "PTI_PINNED_COMMIT env var not defined")
+endif()
+
+if(DEFINED ENV{LEVELZERO_INCLUDE_DIR})
+  set(LevelZero_INCLUDE_DIR "$ENV{LEVELZERO_INCLUDE_DIR}")
+else()
+  message(STATUS "LEVELZERO_INCLUDE_DIR env var not defined; try to use system version")
+endif()
+
+if(DEFINED ENV{LEVELZERO_LIBRARY})
+  set(LevelZero_LIBRARY "$ENV{LEVELZERO_LIBRARY}")
+else()
+  message(STATUS "LEVELZERO_LIBRARY env var not defined; try to use system version")
+endif()
+
+FetchContent_Declare(
+  pti-lib
+  GIT_REPOSITORY https://github.com/intel/pti-gpu.git
+  GIT_TAG        ${PTI_PINNED_COMMIT}
+  SOURCE_SUBDIR sdk
+)
+
+if(NOT APPLE)
+  list(APPEND CMAKE_INSTALL_RPATH $ORIGIN)
+endif()
+
+# Sets the installation directories to be inside the root of the virtual
+# environment, .e.g., .venv/lib/libpti_view.so (note: this is non-standard).
+# However, this is what other oneAPI components and PyTorch XPU (Intel backend)
+# expects.
+set(CMAKE_INSTALL_LIBDIR "${SKBUILD_DATA_DIR}/lib")
+set(CMAKE_INSTALL_INCLUDEDIR "${SKBUILD_DATA_DIR}/include")
+set(CMAKE_INSTALL_BINDIR "${SKBUILD_DATA_DIR}/bin")
+set(CMAKE_INSTALL_DOCDIR "${SKBUILD_DATA_DIR}/share")
+
+
+FetchContent_MakeAvailable(pti-lib)

scripts/build_pti_data/README.md

@@ -0,0 +1,49 @@
+[project]
+name = "intel-pti"
+version = "0.14.0.dev1"
+description = "Add your description here"
+readme = "README.md"
+license = "MIT"
+
+requires-python = ">=3.9"
+dependencies = []
+classifiers = [
+  "Development Status :: 5 - Production/Stable",
+
+  # Indicate who your project is intended for
+  "Intended Audience :: Developers",
+  "Intended Audience :: Education",
+  "Intended Audience :: Science/Research",
+  "Intended Audience :: System Administrators",
+  "Intended Audience :: Other Audience",
+
+  "Topic :: Software Development :: Libraries",
+
+  "Operating System :: Microsoft :: Windows",
+  "Operating System :: POSIX :: Linux",
+
+  "Environment :: GPU",
+]
+
+[tool.scikit-build]
+build.verbose = true
+logging.level = "INFO"
+cmake.version = ">=3.20.0"
+cmake.build-type = "Release"
+cmake.args = ["-GNinja"]
+ninja.version = ">=1.13"
+minimum-version = "build-system.requires"
+sdist.exclude = ["levelzero", "level-zero", "src"]
+
+[tool.uv]
+cache-keys = [{ file = "pyproject.toml" }, { file = "CMakeLists.txt" }]
+
+[build-system]
+requires = ["scikit-build-core>=0.10"]
+build-backend = "scikit_build_core.build"
+
+[project.urls]
+Homepage = "https://github.com/intel/pti-gpu/tree/master/sdk"
+Documentation = "https://github.com/intel/pti-gpu/blob/master/sdk/README.md"
+Repository = "https://github.com/intel/pti-gpu"
+Issues = "https://github.com/intel/pti-gpu/issues"

scripts/build_pti_data/pyproject.toml

@@ -0,0 +1,87 @@
+#!/usr/bin/env bash
+
+set -euo pipefail
+
+# Select what to build.
+BUILD_LEVEL_ZERO=false
+for arg in "$@"; do
+  case $arg in
+    --build-level-zero)
+      BUILD_LEVEL_ZERO=true
+      shift
+      ;;
+    --help)
+      echo "Example usage: ./install-pti.sh [--build-level-zero]"
+      exit 1
+      ;;
+    *)
+      echo "Unknown argument: $arg."
+      exit 1
+      ;;
+  esac
+done
+
+
+# Configure, build and install PyTorch from source.
+
+# intel-xpu-backend-for-triton project root
+ROOT=$(cd "$(dirname "$0")/.." && pwd)
+
+SCRIPTS_DIR=$ROOT/scripts
+PTI_PROJ=$ROOT/.scripts_cache/pti
+LEVEL_ZERO_PROJ=$ROOT/.scripts_cache/level_zero_for_pti
+BASE=$(dirname "$PTI_PROJ")
+
+echo "**** BASE is set to $BASE ****"
+echo "**** PTI_PROJ is set to $PTI_PROJ ****"
+mkdir -p $BASE
+
+function build_level_zero {
+  rm -rf "$LEVEL_ZERO_PROJ"
+  mkdir -p "$LEVEL_ZERO_PROJ"
+  cd "$LEVEL_ZERO_PROJ"
+  LEVEL_ZERO_VERSION=1.24.2
+  LEVEL_ZERO_SHA256=b77e6e28623134ee4e99e2321c127b554bdd5bfa3e80064922eba293041c6c52
+
+  wget --progress=dot:giga -e use_proxy=yes "https://github.com/oneapi-src/level-zero/archive/refs/tags/v${LEVEL_ZERO_VERSION}.tar.gz"
+  echo "${LEVEL_ZERO_SHA256}  v${LEVEL_ZERO_VERSION}.tar.gz" > "v${LEVEL_ZERO_VERSION}.tar.gz.sha256"
+  sha256sum -c "v${LEVEL_ZERO_VERSION}.tar.gz.sha256"
+  tar -xf "v${LEVEL_ZERO_VERSION}.tar.gz"
+  cd "level-zero-${LEVEL_ZERO_VERSION}"
+  echo "${LEVEL_ZERO_VERSION}" | awk -F. '{print $3}' > VERSION_PATCH
+  mkdir build
+  cd build
+  cmake .. -DCMAKE_BUILD_TYPE=Release
+  cmake --build . --config Release --parallel "$(nproc)"
+  # cmake --build . --config Release --target install
+  export LEVELZERO_INCLUDE_DIR="$LEVEL_ZERO_PROJ/level-zero-${LEVEL_ZERO_VERSION}"
+  export LEVELZERO_LIBRARY="$LEVEL_ZERO_PROJ/level-zero-${LEVEL_ZERO_VERSION}/build/lib/libze_loader.so"
+}
+
+function build_pti {
+  rm -rf "$PTI_PROJ"
+  mkdir -p "$PTI_PROJ"
+
+  echo "****** Building $PTI_PROJ ******"
+  cd "$PTI_PROJ"
+  cp "$SCRIPTS_DIR"/build_pti_data/* .
+  pip install uv
+
+  export PTI_PINNED_COMMIT="$(<$ROOT/.github/pins/pti.txt)"
+
+  uv version 0.14.0.dev1
+  uv build
+}
+
+function install_pti {
+  echo "****** Installing PTI ******"
+  cd "$PTI_PROJ"
+  pip install dist/*.whl
+}
+
+if [ "$BUILD_LEVEL_ZERO" = true ]; then
+  build_level_zero
+fi
+
+build_pti
+install_pti

scripts/install-pti.sh

@@ -0,0 +1,22 @@
+"""Prints a lib directory for pti."""
+
+import importlib.metadata
+import pathlib
+
+
+def get_pti_lib_path() -> pathlib.Path:
+    """Returns library path for pti.
+
+    Raises:
+        importlib.metadata.PackageNotFoundError: if 'intel-pti' not installed.
+        AssertionError: if libpti_view.so not found.
+    """
+    files = importlib.metadata.files('intel-pti') or []
+    for f in files:
+        if f.name == 'libpti_view.so':
+            return pathlib.Path(f.locate()).parent.resolve()
+    raise AssertionError('libpti_view.so not found')
+
+
+if __name__ == '__main__':
+    print(get_pti_lib_path())

scripts/pti_lib.py

@@ -30,7 +30,7 @@ class XPUOptions:
     supported_fp8_dtypes: Tuple[str] = ("fp8e5", "fp8e4nv", "fp8e4b15")
     deprecated_fp8_dot_operand_dtypes: Tuple[str] = ()
     default_dot_input_precision: str = "tf32"
-    allowed_dot_input_precisions: Tuple[str] = ("tf32", "tf32x3", "ieee")
+    allowed_dot_input_precisions: Tuple[str] = ("tf32", "tf32x3", "ieee", 'bf16x3', 'bf16x6')
     allow_fp8e4nv: bool = False
     allow_fp8e4b15: bool = True
     grf_mode: tuple = ('small', 'large', 'auto', 'default')

third_party/intel/backend/compiler.py

@@ -1,3 +1,7 @@
+from typing import Optional, Union
+import importlib.metadata
+import pathlib
+import os
 import functools
 import triton
 
@@ -6,7 +10,6 @@
 from .flags import flags
 from .hooks import HookManager, LaunchHook, InstrumentationHook
 from .mode import BaseMode
-from typing import Optional, Union
 
 DEFAULT_PROFILE_NAME = "proton"
 UTILS_CACHE_PATH = None
@@ -21,6 +24,12 @@ def _select_backend() -> str:
     elif backend == "xpu":
         global UTILS_CACHE_PATH
         UTILS_CACHE_PATH = triton.runtime.driver.active.build_proton_help_lib()
+        files = importlib.metadata.files('intel-pti')
+        if files is not None:
+            for f in files:
+                if f.name == 'libpti_view.so':
+                    os.environ["TRITON_XPUPTI_LIB_PATH"] = str(pathlib.Path(f.locate()).parent.resolve())
+                    break
         return "xpupti"
     else:
         raise ValueError("No backend is available for the current target.")