Merge remote-tracking branch 'origin/main' into quinnlp/fused-xeforge

Author: quinnlp

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

@@ -1,5 +1,15 @@
 name: Build benchmarks wheel
 
+# Wheels are signed with keyless Sigstore (OIDC). Each *.whl is accompanied by a
+# *.whl.sigstore.json bundle in the benchmarks-wheel-py* artifacts.
+#
+# Verify a downloaded benchmarks wheel + bundle:
+#   cosign verify-blob triton_kernels_benchmark-*.whl \
+#     --bundle triton_kernels_benchmark-*.whl.sigstore.json \
+#     --certificate-oidc-issuer https://token.actions.githubusercontent.com \
+#     --certificate-identity-regexp \
+#       'https://github.com/intel/intel-xpu-backend-for-triton/.github/workflows/build-benchmarks-wheel.yml@.*'
+
 on:
   workflow_dispatch:
   workflow_call:
@@ -42,6 +52,9 @@ jobs:
       fail-fast: false
       max-parallel: 2
     timeout-minutes: 120
+    permissions:
+      id-token: write   # required for keyless Sigstore/OIDC signing
+      contents: read
     defaults:
       run:
         shell: bash -noprofile --norc -eo pipefail -c "source /opt/intel/oneapi/setvars.sh > /dev/null; source {0}"
@@ -66,6 +79,12 @@ jobs:
           cd benchmarks
           python -m build --wheel --no-isolation
 
+      - name: Sign wheels with Sigstore (keyless)
+        # SHA-pinned for supply-chain safety; comment tracks the human-readable version.
+        uses: sigstore/gh-action-sigstore-python@a5caf349bc536fbef3668a10ed7f5cd309a4b53d # v3.2.0
+        with:
+          inputs: benchmarks/dist/*.whl
+
       - name: Install benchmarks wheel
         run: |
           # Uninstall first — see install-benchmarks/action.yml for full rationale.
@@ -81,4 +100,6 @@ jobs:
         uses: actions/upload-artifact@v7
         with:
           name: benchmarks-wheel-py${{ matrix.python }}
-          path: benchmarks/dist/*.whl
+          path: |
+            benchmarks/dist/*.whl
+            benchmarks/dist/*.whl.sigstore.json

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

@@ -389,6 +389,6 @@ jobs:
           use_pyenv: ${{ inputs.use_pyenv_python }}
 
       - name: Ensure job status
-        uses: re-actors/alls-green@05ac9388f0aebcb5727afa17fcccfecd6f8ec5fe
+        uses: re-actors/alls-green@b5b5b37504aa4183270bd3d855c52a67f212be35
         with:
           jobs: ${{ toJSON(needs) }}

.github/workflows/codeql.yml

@@ -104,7 +104,7 @@ jobs:
         /opt/google/chrome/chrome --version
 
     - name: Generate Security Report
-      uses: rsdmike/github-security-report-action@v3.0.4
+      uses: rsdmike/github-security-report-action@1df22b1e0a7e15b32f728ccf7bab259e46c80589 # v4.0.1
       with:
         template: report
         token: ${{ secrets.SECURITY_TOKEN }}

.github/workflows/integration-tests-amd.yml

@@ -47,8 +47,8 @@ jobs:
       TRITON_BUILD_WITH_CLANG_LLD: "TRUE"
       TRITON_USE_ASSERT_ENABLED_LLVM: "TRUE"
       TRITON_DISABLE_LINE_INFO: 1
-      PROTON_SKIP_CUDAGRAPH_TEST: 1
       PROTON_SKIP_PC_SAMPLING_TEST: 1
+      PROTON_SKIP_CUDAGRAPH_TEST: 1
       PYTHON: "python3"
       PYTEST_ADDOPTS: "-p triton._compile_warmup"
       CCACHE_COMPRESS: "true"
@@ -221,7 +221,7 @@ jobs:
       - name: Run Proton tests
         run: |
           unset HIP_VISIBLE_DEVICES
-          make test-proton
+          TRITON_DISABLE_LINE_INFO=0 make test-proton
       - name: Inspect cache directories
         run: |
           mkdir -p ~/.triton
@@ -346,7 +346,7 @@ jobs:
             fi
           done
           export LD_LIBRARY_PATH="$ROCM_SDK_LIB:$LD_LIBRARY_PATH"
-          make test-proton
+          TRITON_DISABLE_LINE_INFO=0 make test-proton
       - name: Inspect cache directories
         run: |
           mkdir -p ~/.triton

.github/workflows/nightly-wheels.yml

@@ -1,5 +1,15 @@
 name: Triton wheels
 
+# Wheels are signed with keyless Sigstore (OIDC). Each *.whl is accompanied by a
+# *.whl.sigstore.json bundle in the wheels-* artifacts.
+#
+# Verify a downloaded wheel + bundle:
+#   cosign verify-blob <wheel_name>.whl \
+#     --bundle <wheel_name>.whl.sigstore.json \
+#     --certificate-oidc-issuer https://token.actions.githubusercontent.com \
+#     --certificate-identity-regexp \
+#       'https://github.com/intel/intel-xpu-backend-for-triton/.github/workflows/nightly-wheels.yml@.*'
+
 on:
   workflow_dispatch:
   schedule:
@@ -41,9 +51,12 @@ jobs:
           - "3.12"
           - "3.13"
           - "3.14"
-          - "3.15.0-beta.2"
+          - "3.15.0-rc.1"
       fail-fast: false
       max-parallel: 2
+    permissions:
+      id-token: write   # required for keyless Sigstore/OIDC signing
+      contents: read
     defaults:
       run:
         shell: bash -noprofile --norc -eo pipefail -c "source /opt/intel/oneapi/setvars.sh > /dev/null; source {0}"
@@ -134,13 +147,26 @@ jobs:
           cp -L timm*/dist/*.whl wheels/ 2>/dev/null || true
           ls -lh wheels/
 
+      - name: Sign PyTorch-related wheels with Sigstore (keyless)
+        # SHA-pinned for supply-chain safety; comment tracks the human-readable version.
+        uses: sigstore/gh-action-sigstore-python@a5caf349bc536fbef3668a10ed7f5cd309a4b53d # v3.2.0
+        with:
+          inputs: wheels/*.whl
+
       - name: Prepare vLLM wheels for upload
         if: contains(fromJSON('["3.10", "3.11", "3.12"]'), matrix.python)
         run: |
           mkdir -p wheels-vllm
           cp -L vllm-xpu-kernels/dist/*.whl wheels-vllm/
           ls -lh wheels-vllm/
 
+      - name: Sign vLLM wheels with Sigstore (keyless)
+        if: contains(fromJSON('["3.10", "3.11", "3.12"]'), matrix.python)
+        # SHA-pinned for supply-chain safety; comment tracks the human-readable version.
+        uses: sigstore/gh-action-sigstore-python@a5caf349bc536fbef3668a10ed7f5cd309a4b53d # v3.2.0
+        with:
+          inputs: wheels-vllm/*.whl
+
       - name: Report environment details
         run: |
           reports_file="wheels/.env"

.github/workflows/vllm-tests-reusable.yml

@@ -144,7 +144,7 @@ jobs:
           - vllm-tdesc
           - vllm-rest
     runs-on: ${{ fromJSON((inputs.runner_label == '' || startsWith(inputs.runner_label, 'max')) && format('["linux", "rolling", "{0}"]', inputs.runner_label || 'max1100') || format('["linux", "{0}"]', inputs.runner_label)) }}
-    timeout-minutes: 180
+    timeout-minutes: 240
     defaults:
       run:
         shell: bash -noprofile --norc -eo pipefail -c "source /opt/intel/oneapi/setvars.sh > /dev/null; source {0}"

bin/RegisterTritonDialects.h

@@ -76,6 +76,7 @@ void registerTestMembarPass();
 void registerTestAMDGPUMembarPass();
 void registerTestTritonAMDGPURangeAnalysis();
 void registerTestLoopPeelingPass();
+void registerTestRegisterPressurePass();
 namespace proton {
 void registerTestScopeIdAllocationPass();
 } // namespace proton
@@ -103,6 +104,7 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
   mlir::test::registerTestLoopPeelingPass();
   mlir::test::registerTestAMDGPUMembarPass();
   mlir::test::registerTestTritonAMDGPURangeAnalysis();
+  mlir::test::registerTestRegisterPressurePass();
   mlir::triton::registerConvertTritonToTritonGPUPass();
   mlir::triton::intel::registerTritonIntelDescriptorVersioning();
   mlir::triton::intel::registerTritonIntelFuseReshape();

include/triton/Analysis/Allocation.h

@@ -9,6 +9,7 @@
 #include "llvm/ADT/SetVector.h"
 
 #include <limits>
+#include <optional>
 
 namespace mlir {
 
@@ -24,6 +25,12 @@ unsigned defaultAllocationAnalysisScratchSizeFn(Operation *op);
 /// Returns whether an operation uses scratch memory across CTAs.
 bool hasCrossCTAScratch(Operation *op);
 
+/// For atomic-result scratch, returns the CTA bits broadcast from
+/// each group leader. Physical scratch owners have these bits clear.
+/// Scalar results are broadcast from CTA0 across all CTAs.
+/// Callers check whether scratch has been allocated.
+std::optional<uint16_t> getAtomicScratchBroadcastMask(Operation *op);
+
 unsigned getNumScratchElemsSwizzledCvt(const LinearLayout &srcLayout,
                                        const LinearLayout &dstLayout,
                                        int bitwidth, int numBanks = 32,

include/triton/Analysis/BufferRegion.h

@@ -2,6 +2,7 @@
 #define TRITON_ANALYSIS_BUFFER_REGION_H
 
 #include <cstdint>
+#include <memory>
 #include <optional>
 #include <set>
 #include <tuple>
@@ -19,6 +20,11 @@
 #include "llvm/ADT/SparseBitVector.h"
 #include "llvm/ADT/UniqueVector.h"
 
+namespace mlir {
+class Allocation;
+class ModuleAllocation;
+} // namespace mlir
+
 namespace mlir::triton::gpu {
 enum class SharedKind : uint32_t;
 }
@@ -121,11 +127,8 @@ struct BufferRegionView {
   uint32_t affineCTAOffset = 0;
   /// Deterministically interned identity of the owning allocation frame.
   uint32_t allocationFrame = 0;
-
-  bool intersects(const BufferRegionView &other) const {
-    return allocationFrame == other.allocationFrame &&
-           region.intersects(other.region);
-  }
+  /// Descriptor allocation supplying these views; null for implicit scratch.
+  Operation *allocation = nullptr;
 
   bool contains(const BufferRegionView &other) const {
     return allocationFrame == other.allocationFrame &&
@@ -138,7 +141,8 @@ struct BufferRegionView {
 private:
   auto key() const {
     return std::tie(allocationFrame, region, storageBase, affineOffset,
-                    affinePartitionOffset, affineCTAOffset, partitionBases);
+                    affinePartitionOffset, affineCTAOffset, partitionBases,
+                    allocation);
   }
 
 public:
@@ -151,9 +155,6 @@ struct BufferRegionView {
   }
 };
 
-/// An exact access view, or no view when the physical region is unknown.
-using BufferRegionAccess = std::optional<BufferRegionView>;
-
 //===----------------------------------------------------------------------===//
 // Buffer state planning
 //===----------------------------------------------------------------------===//
@@ -223,6 +224,20 @@ struct RegionInfo {
   }
 };
 
+/// Complete MAY-views of a descriptor, shared by all of its memory effects.
+/// Runtime indices may select several physical views across loop iterations.
+/// Addresses, rather than runtime descriptor keys, define the geometry; the
+/// memory space distinguishes shared bytes from TMEM's 32-bit storage words.
+struct BufferRegionFootprint {
+  Attribute memorySpace;
+  RegionInfo regionInfo;
+};
+
+/// Prove disjointness only if every candidate pair is disjoint. Missing views
+/// and unnormalized allocation frames may overlap, never denote empty storage.
+bool mayOverlap(const BufferRegionFootprint *lhs,
+                const BufferRegionFootprint *rhs);
+
 enum class RW { Read, Write };
 
 struct MemoryAccess {
@@ -250,6 +265,7 @@ bool hasSharedAccess(Operation *op,
 //
 // Produces a RegionInfo lattice for each MemDesc/ptr-like SSA value,
 // and also collects a global list of all discovered BufferRegions.
+// Requires completed allocation analyses or their materialized offsets.
 //
 class BufferRegionAnalysis : public dataflow::SparseForwardDataFlowAnalysis<
                                  dataflow::Lattice<RegionInfo>> {
@@ -258,23 +274,37 @@ class BufferRegionAnalysis : public dataflow::SparseForwardDataFlowAnalysis<
   using Base =
       dataflow::SparseForwardDataFlowAnalysis<dataflow::Lattice<RegionInfo>>;
   using Base::getLatticeElement;
-  using Base::SparseForwardDataFlowAnalysis;
+  enum class Mode { AllMemory, TensorMemoryOnly };
+
+  explicit BufferRegionAnalysis(DataFlowSolver &solver,
+                                Mode mode = Mode::AllMemory,
+                                ModuleAllocation *allocation = nullptr)
+      : Base(solver), mode(mode), moduleAllocation(allocation) {}
 
   enum RegionType { SHARED_MEMORY, TENSOR_MEMORY, BARRIER, NUM_REGION_TYPES };
 
   const RegionInfo &getRegionInfo(Value value) {
     return getLatticeElement(value)->getValue();
   }
 
-  /// Return every exact view an access may reference. A null view represents
-  /// an unknown region and therefore may alias any other view.
-  llvm::SmallVector<BufferRegionAccess> getAccessRegions(Value value);
+  /// Return an immutable footprint covering every possible view, owned by this
+  /// analysis. Call after solver convergence. Return null for unknown geometry
+  /// or views outside allocationFrame, when set.
+  const BufferRegionFootprint *
+  getFootprint(Value value, FunctionOpInterface allocationFrame = {});
+
+  /// Describe an operation's allocated shared scratch, including cross-CTA
+  /// accesses.
+  const BufferRegionFootprint *getScratchFootprint(Operation *op);
 
-  /// Translate a callee-local view into the caller's allocation frame.
-  BufferRegionAccess translateToCallsite(BufferRegionAccess view,
-                                         CallOpInterface call,
-                                         FunctionOpInterface caller,
-                                         FunctionOpInterface callee) const;
+  /// Translate callee-local views into the caller's allocation frame, caching
+  /// the complete footprint and preserving views already in other frames.
+  const BufferRegionFootprint *
+  translateToCallsite(const BufferRegionFootprint *footprint,
+                      CallOpInterface call, FunctionOpInterface callee);
+
+  /// Shared-memory offset of the callee frame in the caller's allocation.
+  uint32_t getCallOffset(CallOpInterface call) const;
 
   uint32_t getOperationId(Operation *operation) const {
     return operationInterner.idFor(operation);
@@ -313,6 +343,11 @@ class BufferRegionAnalysis : public dataflow::SparseForwardDataFlowAnalysis<
   LogicalResult initialize(Operation *top) override;
 
 private:
+  const Mode mode;
+  ModuleAllocation *moduleAllocation;
+
+  Allocation *getAllocation(Operation *op) const;
+
   BufferRegionView getAllocView(Value allocation, uint32_t storageBase,
                                 llvm::ArrayRef<uint32_t> partitionBases = {});
   BufferRegionView getSubView(Type type, const BufferRegionView &view,
@@ -323,7 +358,15 @@ class BufferRegionAnalysis : public dataflow::SparseForwardDataFlowAnalysis<
   // Global registry of all regions
   std::set<BufferRegion> usedBufferRegions[NUM_REGION_TYPES];
   bool usedUnknownBufferRegions[NUM_REGION_TYPES] = {};
-  llvm::DenseMap<std::pair<Type, uint32_t>, AddressSet> footprintCache;
+  llvm::DenseMap<std::pair<Type, uint32_t>,
+                 llvm::SmallVector<BufferRegion::CTAAddresses, 2>>
+      footprintCache;
+  llvm::DenseMap<Value, std::unique_ptr<BufferRegionFootprint>> valueFootprints;
+  llvm::DenseMap<Operation *, std::unique_ptr<BufferRegionFootprint>>
+      scratchFootprints;
+  llvm::DenseMap<std::pair<const BufferRegionFootprint *, Operation *>,
+                 std::unique_ptr<BufferRegionFootprint>>
+      callsiteFootprints;
   llvm::UniqueVector<Operation *> operationInterner;
 };