Skip to content

Commit 3509f46

Browse files
Hardcode84Hardcode84
authored
Update to llvm main (#169)
1 parent 8668364 commit 3509f46

File tree

21 files changed

+185
-173
lines changed

21 files changed

+185
-173
lines changed

‎llvm-sha.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1 @@
1-
5489969550a28a5ef63b6e242469b6eb96b2fbbd
1+
85daf6973d2bd3216b9d131be5be625b1227016c

‎mlir/CMakeLists.txt

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -134,9 +134,10 @@ target_link_libraries(${MLIR_EXTENSIONS_LIB} PRIVATE
134134
MLIRIR
135135
MLIRLLVMIR
136136
MLIRTransforms
137-
MLIRStandardOpsTransforms
137+
MLIRFuncTransforms
138138
MLIRLinalgTransforms
139139
MLIRTensorTransforms
140+
MLIRMathToSPIRV
140141
)
141142

142143
target_include_directories(${MLIR_EXTENSIONS_LIB} SYSTEM PRIVATE

‎mlir/include/mlir-extensions/Conversion/gpu_to_gpu_runtime.hpp

Lines changed: 4 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -14,39 +14,11 @@
1414

1515
#pragma once
1616

17-
#include "mlir-extensions/dialect/gpu_runtime/IR/gpu_runtime_ops.hpp"
17+
#include <memory>
1818

19-
#include "mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.h"
20-
#include <llvm/ADT/SmallBitVector.h>
21-
#include <mlir/Analysis/BufferViewFlowAnalysis.h>
22-
#include <mlir/Conversion/AffineToStandard/AffineToStandard.h>
23-
#include <mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h>
24-
#include <mlir/Conversion/GPUCommon/GPUCommonPass.h>
25-
#include <mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h>
26-
#include <mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h>
27-
#include <mlir/Conversion/MathToSPIRV/MathToSPIRV.h>
28-
#include <mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h>
29-
#include <mlir/Conversion/StandardToSPIRV/StandardToSPIRV.h>
30-
#include <mlir/Dialect/Affine/IR/AffineOps.h>
31-
#include <mlir/Dialect/Arithmetic/Transforms/Passes.h>
32-
#include <mlir/Dialect/GPU/ParallelLoopMapper.h>
33-
#include <mlir/Dialect/GPU/Passes.h>
34-
#include <mlir/Dialect/GPU/Utils.h>
35-
#include <mlir/Dialect/LLVMIR/LLVMDialect.h>
36-
#include <mlir/Dialect/Math/IR/Math.h>
37-
#include <mlir/Dialect/MemRef/IR/MemRef.h>
38-
#include <mlir/Dialect/SCF/SCF.h>
39-
#include <mlir/Dialect/SPIRV/IR/SPIRVDialect.h>
40-
#include <mlir/Dialect/SPIRV/IR/SPIRVOps.h>
41-
#include <mlir/Dialect/SPIRV/IR/TargetAndABI.h>
42-
#include <mlir/Dialect/SPIRV/Transforms/Passes.h>
43-
#include <mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h>
44-
#include <mlir/Dialect/StandardOps/IR/Ops.h>
45-
#include <mlir/Pass/PassManager.h>
46-
#include <mlir/Target/SPIRV/Serialization.h>
47-
#include <mlir/Transforms/DialectConversion.h>
48-
#include <mlir/Transforms/GreedyPatternRewriteDriver.h>
49-
#include <mlir/Transforms/Passes.h>
19+
namespace mlir {
20+
class Pass;
21+
}
5022

5123
namespace gpu_runtime {
5224

‎mlir/lib/Conversion/SCFToAffine/SCFToAffine.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -94,7 +94,6 @@ struct SCFToAffinePass
9494
: public mlir::PassWrapper<SCFToAffinePass, mlir::OperationPass<void>> {
9595
virtual void
9696
getDependentDialects(mlir::DialectRegistry &registry) const override {
97-
registry.insert<mlir::StandardOpsDialect>();
9897
registry.insert<mlir::AffineDialect>();
9998
registry.insert<mlir::scf::SCFDialect>();
10099
}

‎mlir/lib/Conversion/gpu_to_gpu_runtime.cpp

Lines changed: 43 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,41 @@
1414

1515
#include "mlir-extensions/Conversion/gpu_to_gpu_runtime.hpp"
1616

17+
#include "mlir-extensions/dialect/gpu_runtime/IR/gpu_runtime_ops.hpp"
18+
19+
#include <mlir/Analysis/BufferViewFlowAnalysis.h>
20+
#include <mlir/Conversion/AffineToStandard/AffineToStandard.h>
21+
#include <mlir/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.h>
22+
#include <mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.h>
23+
#include <mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h>
24+
#include <mlir/Conversion/GPUCommon/GPUCommonPass.h>
25+
#include <mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h>
26+
#include <mlir/Conversion/GPUToSPIRV/GPUToSPIRVPass.h>
27+
#include <mlir/Conversion/MathToSPIRV/MathToSPIRV.h>
28+
#include <mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h>
29+
#include <mlir/Dialect/Affine/IR/AffineOps.h>
30+
#include <mlir/Dialect/Arithmetic/Transforms/Passes.h>
31+
#include <mlir/Dialect/Func/IR/FuncOps.h>
32+
#include <mlir/Dialect/GPU/ParallelLoopMapper.h>
33+
#include <mlir/Dialect/GPU/Passes.h>
34+
#include <mlir/Dialect/GPU/Utils.h>
35+
#include <mlir/Dialect/LLVMIR/LLVMDialect.h>
36+
#include <mlir/Dialect/Math/IR/Math.h>
37+
#include <mlir/Dialect/MemRef/IR/MemRef.h>
38+
#include <mlir/Dialect/SCF/SCF.h>
39+
#include <mlir/Dialect/SPIRV/IR/SPIRVDialect.h>
40+
#include <mlir/Dialect/SPIRV/IR/SPIRVOps.h>
41+
#include <mlir/Dialect/SPIRV/IR/TargetAndABI.h>
42+
#include <mlir/Dialect/SPIRV/Transforms/Passes.h>
43+
#include <mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h>
44+
#include <mlir/Pass/PassManager.h>
45+
#include <mlir/Target/SPIRV/Serialization.h>
46+
#include <mlir/Transforms/DialectConversion.h>
47+
#include <mlir/Transforms/GreedyPatternRewriteDriver.h>
48+
#include <mlir/Transforms/Passes.h>
49+
50+
#include <llvm/ADT/SmallBitVector.h>
51+
1752
static const char *kGpuAllocShared = "gpu.alloc_shared";
1853

1954
struct ParallelLoopGPUMappingPass
@@ -109,7 +144,7 @@ struct InsertGPUAllocs
109144
return {{load.memref()}};
110145
} else if (auto store = mlir::dyn_cast<mlir::memref::StoreOp>(op)) {
111146
return {{store.memref()}};
112-
} else if (auto call = mlir::dyn_cast<mlir::CallOp>(op)) {
147+
} else if (auto call = mlir::dyn_cast<mlir::func::CallOp>(op)) {
113148
mlir::SmallVector<mlir::Value, 4> ret;
114149
for (auto arg : call.operands()) {
115150
if (arg.getType().isa<mlir::MemRefType>())
@@ -131,7 +166,7 @@ struct InsertGPUAllocs
131166
memInterface.hasEffect<mlir::MemoryEffects::Write>())
132167
return true;
133168
}
134-
if (auto call = mlir::dyn_cast<mlir::CallOp>(op)) {
169+
if (auto call = mlir::dyn_cast<mlir::func::CallOp>(op)) {
135170
for (auto arg : call.operands()) {
136171
if (arg.getType().isa<mlir::MemRefType>())
137172
return true;
@@ -193,7 +228,7 @@ struct InsertGPUAllocs
193228
AccessType ret;
194229
for (auto mem : aliases.resolve(memref)) {
195230
for (auto user : mem.getUsers()) {
196-
if (mlir::isa<mlir::ReturnOp>(user)) {
231+
if (mlir::isa<mlir::func::ReturnOp>(user)) {
197232
ret.hostRead = true;
198233
ret.hostWrite = true;
199234
continue;
@@ -220,7 +255,7 @@ struct InsertGPUAllocs
220255

221256
continue;
222257
}
223-
if (mlir::isa<mlir::CallOp>(user)) {
258+
if (mlir::isa<mlir::func::CallOp>(user)) {
224259
bool onDevice = user->getParentOfType<mlir::gpu::LaunchOp>();
225260
(onDevice ? ret.deviceRead : ret.hostRead) = true;
226261
(onDevice ? ret.deviceWrite : ret.hostWrite) = true;
@@ -768,12 +803,12 @@ static mlir::Value lowerFloatSubAtomic(mlir::OpBuilder &builder,
768803
mlir::spirv::MemorySemantics::None, neg);
769804
}
770805

771-
class ConvertAtomicOps : public mlir::OpConversionPattern<mlir::CallOp> {
806+
class ConvertAtomicOps : public mlir::OpConversionPattern<mlir::func::CallOp> {
772807
public:
773-
using mlir::OpConversionPattern<mlir::CallOp>::OpConversionPattern;
808+
using mlir::OpConversionPattern<mlir::func::CallOp>::OpConversionPattern;
774809

775810
mlir::LogicalResult
776-
matchAndRewrite(mlir::CallOp op, mlir::CallOp::Adaptor adaptor,
811+
matchAndRewrite(mlir::func::CallOp op, mlir::func::CallOp::Adaptor adaptor,
777812
mlir::ConversionPatternRewriter &rewriter) const override {
778813
auto operands = adaptor.operands();
779814
if (operands.size() != 2)
@@ -899,7 +934,7 @@ struct GPUToSpirvPass
899934
mlir::ScfToSPIRVContext scfToSpirvCtx;
900935
mlir::populateSCFToSPIRVPatterns(typeConverter, scfToSpirvCtx, patterns);
901936
mlir::populateGPUToSPIRVPatterns(typeConverter, patterns);
902-
mlir::populateStandardToSPIRVPatterns(typeConverter, patterns);
937+
mlir::populateFuncToSPIRVPatterns(typeConverter, patterns);
903938
mlir::cf::populateControlFlowToSPIRVPatterns(typeConverter, patterns);
904939
mlir::arith::populateArithmeticToSPIRVPatterns(typeConverter, patterns);
905940
mlir::populateMathToSPIRVPatterns(typeConverter, patterns);

‎mlir/lib/transforms/common_opts.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,6 @@
2525
#include <mlir/Dialect/Math/IR/Math.h>
2626
#include <mlir/Dialect/MemRef/IR/MemRef.h>
2727
#include <mlir/Dialect/SCF/SCF.h>
28-
#include <mlir/Dialect/StandardOps/IR/Ops.h>
2928
#include <mlir/IR/BuiltinOps.h>
3029
#include <mlir/IR/Matchers.h>
3130
#include <mlir/IR/PatternMatch.h>

‎mlir/lib/transforms/inline_utils.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -16,25 +16,25 @@
1616

1717
#include "mlir-extensions/dialect/plier/dialect.hpp"
1818

19+
#include <mlir/Dialect/Func/IR/FuncOps.h>
1920
#include <mlir/Dialect/SCF/SCF.h>
20-
#include <mlir/Dialect/StandardOps/IR/Ops.h>
2121
#include <mlir/IR/PatternMatch.h>
2222
#include <mlir/Pass/Pass.h>
2323
#include <mlir/Transforms/GreedyPatternRewriteDriver.h>
2424
#include <mlir/Transforms/InliningUtils.h>
2525

2626
namespace {
27-
static bool mustInline(mlir::CallOp call, mlir::FuncOp func) {
27+
static bool mustInline(mlir::func::CallOp call, mlir::FuncOp func) {
2828
auto attr = mlir::StringAttr::get(call.getContext(),
2929
plier::attributes::getForceInlineName());
3030
return call->hasAttr(attr) || func->hasAttr(attr);
3131
}
3232

33-
struct ForceInline : public mlir::OpRewritePattern<mlir::CallOp> {
34-
using mlir::OpRewritePattern<mlir::CallOp>::OpRewritePattern;
33+
struct ForceInline : public mlir::OpRewritePattern<mlir::func::CallOp> {
34+
using OpRewritePattern::OpRewritePattern;
3535

3636
mlir::LogicalResult
37-
matchAndRewrite(mlir::CallOp op,
37+
matchAndRewrite(mlir::func::CallOp op,
3838
mlir::PatternRewriter &rewriter) const override {
3939
auto mod = op->getParentOfType<mlir::ModuleOp>();
4040
assert(mod);
@@ -81,7 +81,7 @@ struct ForceInlinePass
8181
mlir::OperationPass<mlir::ModuleOp>> {
8282
virtual void
8383
getDependentDialects(mlir::DialectRegistry &registry) const override {
84-
registry.insert<mlir::StandardOpsDialect>();
84+
registry.insert<mlir::func::FuncDialect>();
8585
registry.insert<mlir::scf::SCFDialect>();
8686
}
8787

@@ -96,7 +96,7 @@ struct ForceInlinePass
9696
auto mod = getOperation();
9797
(void)mlir::applyPatternsAndFoldGreedily(mod, patterns);
9898

99-
mod->walk([&](mlir::CallOp call) {
99+
mod->walk([&](mlir::func::CallOp call) {
100100
auto func = mod.lookupSymbol<mlir::FuncOp>(call.getCallee());
101101
if (func && mustInline(call, func)) {
102102
call.emitError("Couldn't inline force-inline call");

‎mlir/lib/transforms/type_conversion.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -15,10 +15,10 @@
1515
#include "mlir-extensions/transforms/type_conversion.hpp"
1616

1717
#include <mlir/Dialect/Arithmetic/IR/Arithmetic.h>
18+
#include <mlir/Dialect/Func/IR/FuncOps.h>
19+
#include <mlir/Dialect/Func/Transforms/FuncConversions.h>
1820
#include <mlir/Dialect/SCF/SCF.h>
1921
#include <mlir/Dialect/SCF/Transforms.h>
20-
#include <mlir/Dialect/StandardOps/IR/Ops.h>
21-
#include <mlir/Dialect/StandardOps/Transforms/FuncConversions.h>
2222
#include <mlir/Transforms/DialectConversion.h>
2323

2424
#include "mlir-extensions/dialect/plier/dialect.hpp"
@@ -50,8 +50,8 @@ void plier::populateControlFlowTypeConversionRewritesAndTarget(
5050
typeConverter.isLegal(&op.getBody());
5151
});
5252
mlir::populateCallOpTypeConversionPattern(patterns, typeConverter);
53-
target.addDynamicallyLegalOp<mlir::CallOp>(
54-
[&](mlir::CallOp op) { return typeConverter.isLegal(op); });
53+
target.addDynamicallyLegalOp<mlir::func::CallOp>(
54+
[&](mlir::func::CallOp op) { return typeConverter.isLegal(op); });
5555

5656
mlir::populateBranchOpInterfaceTypeConversionPattern(patterns, typeConverter);
5757
mlir::populateReturnOpTypeConversionPattern(patterns, typeConverter);

‎mlir/lib/transforms/uplift_math.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,8 @@
1818
#include "mlir-extensions/transforms/rewrite_wrapper.hpp"
1919

2020
#include <mlir/Dialect/Arithmetic/IR/Arithmetic.h>
21+
#include <mlir/Dialect/Func/IR/FuncOps.h>
2122
#include <mlir/Dialect/Math/IR/Math.h>
22-
#include <mlir/Dialect/StandardOps/IR/Ops.h>
2323
#include <mlir/IR/PatternMatch.h>
2424

2525
template <typename Op>
@@ -41,11 +41,11 @@ static mlir::Operation *replaceOp2(mlir::OpBuilder &builder, mlir::Location loc,
4141
}
4242

4343
namespace {
44-
struct UpliftMathCalls : public mlir::OpRewritePattern<mlir::CallOp> {
44+
struct UpliftMathCalls : public mlir::OpRewritePattern<mlir::func::CallOp> {
4545
using OpRewritePattern::OpRewritePattern;
4646

4747
mlir::LogicalResult
48-
matchAndRewrite(mlir::CallOp op,
48+
matchAndRewrite(mlir::func::CallOp op,
4949
mlir::PatternRewriter &rewriter) const override {
5050
auto funcName = op.getCallee();
5151
if (funcName.empty())
@@ -120,7 +120,7 @@ struct UpliftFma : public mlir::OpRewritePattern<mlir::arith::AddFOp> {
120120
struct UpliftMathPass
121121
: public plier::RewriteWrapperPass<
122122
UpliftMathPass, void,
123-
plier::DependentDialectsList<mlir::StandardOpsDialect,
123+
plier::DependentDialectsList<mlir::func::FuncDialect,
124124
mlir::arith::ArithmeticDialect,
125125
mlir::math::MathDialect>,
126126
UpliftMathCalls, UpliftFma> {};

‎mlir/tools/level_zero_runner/CMakeLists.txt

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -59,11 +59,9 @@ target_link_libraries(${PROJECT_NAME} PRIVATE
5959
MLIRParser
6060
MLIRSPIRV
6161
MLIRSPIRVTransforms
62-
MLIRStandard
6362
MLIRSupport
6463
MLIRTargetLLVMIRExport
6564
MLIRTransforms
66-
MLIRTranslation
6765
)
6866

6967
target_include_directories(${PROJECT_NAME} PRIVATE

0 commit comments

Comments
 (0)