Merge commit '52e7a4b1b170fed27be8a5851a47671bcda70a98'

Author: whitneywhtsang

.github/workflows/create_release.yml

@@ -0,0 +1,76 @@
+name: Create Release
+
+on:
+  push:
+    branches:
+      - main
+      - release/*
+    tags:
+      # Final Release tags look like: v1.11.0
+      - v[0-9]+.[0-9]+.[0-9]+
+      # Release candidate tags look like: v1.11.0-rc1
+      - v[0-9]+.[0-9]+.[0-9]+-rc[0-9]+
+  release:
+    types: [published]
+  pull_request:
+    paths: [.github/workflows/create_release.yml]
+
+jobs:
+
+  release:
+    if: ${{ github.repository == 'triton-lang/triton' }}
+    name: Create Release
+    runs-on: ubuntu-latest
+    permissions:
+      contents: write
+    outputs:
+      release_name: "${{ steps.release_name.outputs.name }}"
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          show-progress: false
+          submodules: 'recursive'
+          ref: ${{ github.event_name == 'pull_request' && github.event.pull_request.head.sha || github.sha }}
+      - name: Fake name for PRs
+        if: ${{ github.event_name == 'pull_request' }}
+        run: echo "PT_GITHUB_REF=refs/tags/pr-tag" >> "$GITHUB_ENV"
+      - name: Real name for non-PRs
+        if: ${{ github.event_name != 'pull_request' }}
+        run: echo "PT_GITHUB_REF=$GITHUB_REF" >> "$GITHUB_ENV"
+      - name: Set filenames
+        run: |
+          tag_or_branch="${PT_GITHUB_REF#refs/tags/}"
+          tag_or_branch="${tag_or_branch#refs/heads/}"
+          # replace directory separators with _ in branch name
+          tag_or_branch="${tag_or_branch//\//_}"
+          echo "RELEASE_NAME=triton-$tag_or_branch" >> "$GITHUB_ENV"
+          echo "RELEASE_FILE=triton-$tag_or_branch.tar.gz" >> "$GITHUB_ENV"
+      - name: Create source distribution
+        run: |
+            # Create new folder with specified name so extracting the archive yields that
+            rm -rf "/tmp/$RELEASE_NAME"
+            cp -r "$PWD" "/tmp/$RELEASE_NAME"
+            mv "/tmp/$RELEASE_NAME" .
+            # Cleanup
+            find "$RELEASE_NAME" -name '.git*' -exec rm -rv {} \; || true
+            # Create archive
+            tar -czf "$RELEASE_FILE" "$RELEASE_NAME"
+            echo "Created source archive $RELEASE_FILE with content: $(ls -a "$RELEASE_NAME")"
+      - name: Upload source distribution for release
+        if: ${{ github.event_name == 'release' }}
+        uses: softprops/action-gh-release@v2
+        with:
+          files: ${{env.RELEASE_FILE}}
+      - name: Upload source distribution to GHA artifacts for release tags
+        if: ${{ github.event_name == 'push' && startsWith(github.ref, 'refs/tags/v') && contains(github.ref, 'rc') }}
+        uses: actions/[email protected]
+        with:
+          name: ${{ env.RELEASE_FILE }}
+          path: ${{ env.RELEASE_FILE }}
+      - name: Set output
+        id: release_name
+        run: echo "name=release_name::${{ env.RELEASE_NAME }}.tar.gz" >> "${GITHUB_OUTPUT}"
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.sha }}-${{ github.event_name }}
+  cancel-in-progress: true

cmake/llvm-hash.txt

@@ -1 +1 @@
-71a977d0d611f3e9f6137a6b8a26b730b2886ce9
+1d4801f22ab1fd6205b1cf625b690aefc554cd4c

lib/Dialect/TritonGPU/Transforms/WarpSpecialization/AutomaticWarpSpecialization.cpp

@@ -35,9 +35,9 @@ void AutomaticWarpSpecialization::runOnOperation() {
   OpPassManager pm;
   pm.addPass(createTritonGPULoadMMASpecialization({numStages}));
   pm.addPass(createTritonGPURewritePartitionDependencies());
-  // `int-range-optimizations` and SCCP are good at cleaning up loop arithmetic.
-  // FIXME: Re-enable integer range analysis once it is fixed.
-  // pm.addPass(arith::createIntRangeOptimizationsPass());
+  // `int-range-optimizations` is good at cleaning up loop arithmetic involving
+  // circular buffers.
+  pm.addPass(arith::createIntRangeOptimizationsPass());
   pm.addPass(createSCCPPass());
   pm.addPass(createCSEPass());
   pm.addPass(createTritonGPUPartitionLoops());

third_party/amd/include/Analysis/RangeAnalysis.h

@@ -119,8 +119,11 @@ struct TritonIntegerRangeAnalysis : dataflow::IntegerRangeAnalysis {
 
 std::optional<SmallVector<ConstantIntRanges>>
 collectRanges(const DataFlowSolver &solver, ValueRange values);
+
 bool cmpIIsStaticallyTrue(const DataFlowSolver &solver, arith::CmpIOp cmpOp);
 
+bool isEmptyInitializedRange(ConstantIntRanges rv);
+
 } // namespace mlir::triton::AMD
 
 #endif

third_party/amd/lib/Analysis/RangeAnalysis.cpp

@@ -186,6 +186,13 @@ std::optional<ConstantIntRanges> maybeGetAssumedRange(Operation *assumption,
 
 namespace mlir::triton::AMD {
 
+bool isEmptyInitializedRange(ConstantIntRanges rv) {
+  if (!rv.umin().getBitWidth() || !rv.umax().getBitWidth() ||
+      !rv.smin().getBitWidth() || !rv.smax().getBitWidth())
+    return true;
+  return false;
+}
+
 std::optional<SmallVector<ConstantIntRanges>>
 collectRanges(const DataFlowSolver &solver, ValueRange values) {
   SmallVector<ConstantIntRanges> ranges;
@@ -196,6 +203,8 @@ collectRanges(const DataFlowSolver &solver, ValueRange values) {
       return {};
     const ConstantIntRanges &inferredRange =
         maybeInferredRange->getValue().getValue();
+    if (isEmptyInitializedRange(inferredRange))
+      return {};
     ranges.push_back(inferredRange);
   }
   return ranges;

third_party/amd/test/lib/Analysis/TestAMDRangeAnalysis.cpp

@@ -34,6 +34,13 @@ struct TestAMDRangeAnalysisPass
       return signalPassFailure();
 
     auto nonNegativePred = [&solver](Value v) -> bool {
+      if (const auto *r =
+              solver->lookupState<dataflow::IntegerValueRangeLattice>(v)) {
+        if (r->getValue().isUninitialized())
+          return false;
+        if (AMD::isEmptyInitializedRange(r->getValue().getValue()))
+          return false;
+      }
       return succeeded(dataflow::staticallyNonNegative(*solver, v));
     };