Skip to content

Commit 4928589

Browse files
Reland changes that depends on llvm/llvm-project@1cec5fffd8fd (#3959)
This PR relands changes that depends on llvm/llvm-project@1cec5fffd8fd. In order to use the updated LLVM pin, SPIRV translator needs to be also updated. We need to apply temporary changes to SPIRV translator, as the version of SPIRV translator used in IGC is different. --------- Signed-off-by: Whitney Tsang <[email protected]> Co-authored-by: Anatoly Myachev <[email protected]>
2 parents 465b0b6 + 1c364b6 commit 4928589

File tree

15 files changed

+757
-404
lines changed

15 files changed

+757
-404
lines changed

cmake/llvm-hash.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1 @@
1-
adba14acea99cc6a17d837763a3248c9d4a2fadf
1+
1cec5fffd8fddd9d85b516f876093b0e3f0eec5f

test/TritonGPU/amd/amd-range-analysis.mlir

Lines changed: 432 additions & 257 deletions
Large diffs are not rendered by default.

third_party/amd/include/Analysis/RangeAnalysis.h

Lines changed: 14 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,10 @@
66
#include "mlir/Dialect/Arith/IR/Arith.h"
77
#include "mlir/Interfaces/LoopLikeInterface.h"
88

9+
namespace mlir::triton {
10+
class FuncOp;
11+
}
12+
913
namespace mlir::triton::AMD {
1014

1115
/// This struct (analysis) adapt's upstream's IntegerRangeAnalysis (inferring
@@ -35,11 +39,15 @@ struct TritonIntegerRangeAnalysis : dataflow::IntegerRangeAnalysis {
3539

3640
void setToEntryState(dataflow::IntegerValueRangeLattice *lattice) override;
3741

42+
void initializeFuncOp(triton::FuncOp funcOp);
43+
3844
LogicalResult visitOperation(
3945
Operation *op,
4046
ArrayRef<const dataflow::IntegerValueRangeLattice *> operands,
4147
ArrayRef<dataflow::IntegerValueRangeLattice *> resultsLattices) override;
4248

49+
std::optional<int64_t> maybeGetTripCount(LoopLikeOpInterface loop);
50+
4351
/// This method (which overloads
4452
/// AbstractSparseForwardDataFlowAnalysis::visitRegionSuccessors)
4553
/// implements "abstract interpretation" of loops with statically known bounds
@@ -89,6 +97,8 @@ struct TritonIntegerRangeAnalysis : dataflow::IntegerRangeAnalysis {
8997
/// [0, 2147483647] ∩ [-2147483648, 128] = [0, 128]
9098
std::optional<ConstantIntRanges> maybeGetAssumedRange(Value anchor) const;
9199

100+
int64_t getTotalLoopTripCount(LoopLikeOpInterface loop);
101+
92102
/// Trip counts of all loops with static loop bounds contained under the root
93103
/// operation being analyzed. Note, nested loops have trip counts computed as
94104
/// a product of enclosing loops; i.e. for
@@ -117,7 +127,7 @@ struct TritonIntegerRangeAnalysis : dataflow::IntegerRangeAnalysis {
117127
llvm::DenseMap<Value, SetVector<Operation *>> assumptions;
118128
};
119129

120-
std::optional<SmallVector<ConstantIntRanges>>
130+
std::optional<SmallVector<std::optional<ConstantIntRanges>>>
121131
collectRanges(const DataFlowSolver &solver, ValueRange values);
122132

123133
bool cmpIIsStaticallyTrue(const DataFlowSolver &solver, arith::CmpIOp cmpOp);
@@ -127,6 +137,9 @@ bool isEmptyInitializedRange(ConstantIntRanges rv);
127137
void populateFoldTrueCmpIOpPatterns(RewritePatternSet &patterns,
128138
DataFlowSolver *solver);
129139

140+
void initializeFuncOps(Operation *op,
141+
TritonIntegerRangeAnalysis *rangeAnalysis);
142+
130143
} // namespace mlir::triton::AMD
131144

132145
#endif

0 commit comments

Comments
 (0)