Skip to content

Commit 1e90026

Browse files
whitneywhtsangwhitneywhtsang
committed
[SYCL-MLIR] Fix merge
Signed-off-by: Tsang, Whitney <whitney.tsang@intel.com>
1 parent 328f5fc commit 1e90026

File tree

13 files changed

+76
-69
lines changed

13 files changed

+76
-69
lines changed

‎mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOpsTypes.h

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -58,23 +58,23 @@ enum class DecoratedAccess : int { No = 0, Yes = 1, Legacy = 2 };
5858
llvm::StringRef
5959
memoryAccessModeAsString(mlir::sycl::MemoryAccessMode MemAccessMode);
6060
LogicalResult parseMemoryAccessMode(AsmParser &Parser,
61-
FailureOr<MemoryAccessMode> &MemAccessMode);
61+
MemoryAccessMode &MemAccessMode);
6262
void printMemoryAccessMode(AsmPrinter &Printer, MemoryAccessMode MemAccessMode);
6363

6464
llvm::StringRef
6565
memoryTargetModeAsString(mlir::sycl::MemoryTargetMode MemTargetMode);
6666
LogicalResult parseMemoryTargetMode(AsmParser &Parser,
67-
FailureOr<MemoryTargetMode> &MemTargetMode);
67+
MemoryTargetMode &MemTargetMode);
6868
void printMemoryTargetMode(AsmPrinter &Printer, MemoryTargetMode MemTargetMode);
6969

7070
std::string accessAddressSpaceAsString(mlir::sycl::AccessAddrSpace AccAddress);
7171
LogicalResult parseAccessAddrSpace(AsmParser &Parser,
72-
FailureOr<AccessAddrSpace> &AccAddress);
72+
AccessAddrSpace &AccAddress);
7373
void printAccessAddrSpace(AsmPrinter &Printer, AccessAddrSpace AccAddress);
7474

7575
std::string decoratedAccessAsString(mlir::sycl::DecoratedAccess DecAccess);
7676
LogicalResult parseDecoratedAccess(AsmParser &Parser,
77-
FailureOr<DecoratedAccess> &DecAccess);
77+
DecoratedAccess &DecAccess);
7878
void printDecoratedAccess(AsmPrinter &Printer, DecoratedAccess DecAccess);
7979

8080
template <typename Parameter> class SYCLInheritanceTypeTrait {

‎mlir-sycl/lib/Dialect/IR/SYCLOpsTypes.cpp

Lines changed: 27 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -30,26 +30,26 @@ llvm::StringRef mlir::sycl::memoryAccessModeAsString(
3030
llvm_unreachable("Invalid MemoryAccessMode");
3131
}
3232

33-
mlir::LogicalResult mlir::sycl::parseMemoryAccessMode(
34-
mlir::AsmParser &Parser,
35-
mlir::FailureOr<mlir::sycl::MemoryAccessMode> &MemAccessMode) {
33+
mlir::LogicalResult
34+
mlir::sycl::parseMemoryAccessMode(mlir::AsmParser &Parser,
35+
mlir::sycl::MemoryAccessMode &MemAccessMode) {
3636
mlir::StringRef Keyword;
3737
if (Parser.parseKeyword(&Keyword)) {
3838
return mlir::ParseResult::failure();
3939
}
4040

4141
if (Keyword == "read") {
42-
MemAccessMode.emplace(mlir::sycl::MemoryAccessMode::Read);
42+
MemAccessMode = mlir::sycl::MemoryAccessMode::Read;
4343
} else if (Keyword == "write") {
44-
MemAccessMode.emplace(mlir::sycl::MemoryAccessMode::Write);
44+
MemAccessMode = mlir::sycl::MemoryAccessMode::Write;
4545
} else if (Keyword == "read_write") {
46-
MemAccessMode.emplace(mlir::sycl::MemoryAccessMode::ReadWrite);
46+
MemAccessMode = mlir::sycl::MemoryAccessMode::ReadWrite;
4747
} else if (Keyword == "discard_write") {
48-
MemAccessMode.emplace(mlir::sycl::MemoryAccessMode::DiscardWrite);
48+
MemAccessMode = mlir::sycl::MemoryAccessMode::DiscardWrite;
4949
} else if (Keyword == "discard_read_write") {
50-
MemAccessMode.emplace(mlir::sycl::MemoryAccessMode::DiscardReadWrite);
50+
MemAccessMode = mlir::sycl::MemoryAccessMode::DiscardReadWrite;
5151
} else if (Keyword == "atomic") {
52-
MemAccessMode.emplace(mlir::sycl::MemoryAccessMode::Atomic);
52+
MemAccessMode = mlir::sycl::MemoryAccessMode::Atomic;
5353
} else {
5454
return Parser.emitError(Parser.getCurrentLocation(),
5555
"expected valid MemoryAccessMode keyword");
@@ -84,28 +84,28 @@ llvm::StringRef mlir::sycl::memoryTargetModeAsString(
8484
llvm_unreachable("Invalid MemoryTargetMode");
8585
}
8686

87-
mlir::LogicalResult mlir::sycl::parseMemoryTargetMode(
88-
mlir::AsmParser &Parser,
89-
mlir::FailureOr<mlir::sycl::MemoryTargetMode> &MemTargetMode) {
87+
mlir::LogicalResult
88+
mlir::sycl::parseMemoryTargetMode(mlir::AsmParser &Parser,
89+
mlir::sycl::MemoryTargetMode &MemTargetMode) {
9090
mlir::StringRef Keyword;
9191
if (Parser.parseKeyword(&Keyword)) {
9292
return mlir::ParseResult::failure();
9393
}
9494

9595
if (Keyword == "global_buffer") {
96-
MemTargetMode.emplace(mlir::sycl::MemoryTargetMode::GlobalBuffer);
96+
MemTargetMode = mlir::sycl::MemoryTargetMode::GlobalBuffer;
9797
} else if (Keyword == "constant_buffer") {
98-
MemTargetMode.emplace(mlir::sycl::MemoryTargetMode::ConstantBuffer);
98+
MemTargetMode = mlir::sycl::MemoryTargetMode::ConstantBuffer;
9999
} else if (Keyword == "local") {
100-
MemTargetMode.emplace(mlir::sycl::MemoryTargetMode::Local);
100+
MemTargetMode = mlir::sycl::MemoryTargetMode::Local;
101101
} else if (Keyword == "image") {
102-
MemTargetMode.emplace(mlir::sycl::MemoryTargetMode::Image);
102+
MemTargetMode = mlir::sycl::MemoryTargetMode::Image;
103103
} else if (Keyword == "host_buffer") {
104-
MemTargetMode.emplace(mlir::sycl::MemoryTargetMode::HostBuffer);
104+
MemTargetMode = mlir::sycl::MemoryTargetMode::HostBuffer;
105105
} else if (Keyword == "host_image") {
106-
MemTargetMode.emplace(mlir::sycl::MemoryTargetMode::HostImage);
106+
MemTargetMode = mlir::sycl::MemoryTargetMode::HostImage;
107107
} else if (Keyword == "image_array") {
108-
MemTargetMode.emplace(mlir::sycl::MemoryTargetMode::ImageArray);
108+
MemTargetMode = mlir::sycl::MemoryTargetMode::ImageArray;
109109
} else {
110110
return Parser.emitError(Parser.getCurrentLocation(),
111111
"expected valid MemoryTargetMode keyword");
@@ -124,9 +124,9 @@ mlir::sycl::accessAddressSpaceAsString(mlir::sycl::AccessAddrSpace AccAddress) {
124124
return std::to_string(static_cast<int>(AccAddress));
125125
}
126126

127-
mlir::LogicalResult mlir::sycl::parseAccessAddrSpace(
128-
mlir::AsmParser &Parser,
129-
mlir::FailureOr<mlir::sycl::AccessAddrSpace> &AccAddress) {
127+
mlir::LogicalResult
128+
mlir::sycl::parseAccessAddrSpace(mlir::AsmParser &Parser,
129+
mlir::sycl::AccessAddrSpace &AccAddress) {
130130

131131
int AddSpaceInt;
132132
if (Parser.parseInteger<int>(AddSpaceInt)) {
@@ -137,7 +137,7 @@ mlir::LogicalResult mlir::sycl::parseAccessAddrSpace(
137137
assert(0 <= AddSpaceInt && AddSpaceInt <= 6 &&
138138
"Expecting address space value between 0 and 6 (inclusive)");
139139

140-
AccAddress.emplace(static_cast<mlir::sycl::AccessAddrSpace>(AddSpaceInt));
140+
AccAddress = static_cast<mlir::sycl::AccessAddrSpace>(AddSpaceInt);
141141
return mlir::ParseResult::success();
142142
}
143143

@@ -151,9 +151,9 @@ mlir::sycl::decoratedAccessAsString(mlir::sycl::DecoratedAccess DecAccess) {
151151
return std::to_string(static_cast<int>(DecAccess));
152152
}
153153

154-
mlir::LogicalResult mlir::sycl::parseDecoratedAccess(
155-
mlir::AsmParser &Parser,
156-
mlir::FailureOr<mlir::sycl::DecoratedAccess> &DecAccess) {
154+
mlir::LogicalResult
155+
mlir::sycl::parseDecoratedAccess(mlir::AsmParser &Parser,
156+
mlir::sycl::DecoratedAccess &DecAccess) {
157157

158158
int DecAccessInt;
159159
if (Parser.parseInteger<int>(DecAccessInt)) {
@@ -163,7 +163,7 @@ mlir::LogicalResult mlir::sycl::parseDecoratedAccess(
163163
assert(0 <= DecAccessInt && DecAccessInt <= 2 &&
164164
"Expecting Decorated Access value between 0 and 2 (inclusive)");
165165

166-
DecAccess.emplace(static_cast<mlir::sycl::DecoratedAccess>(DecAccessInt));
166+
DecAccess = static_cast<mlir::sycl::DecoratedAccess>(DecAccessInt);
167167
return mlir::ParseResult::success();
168168
}
169169

‎polygeist/include/mlir/Dialect/Polygeist/Utils/Utils.h

Lines changed: 8 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -137,15 +137,19 @@ inline bool canBeLoweredToBarePtr(mlir::MemRefType memRefType) {
137137
inline LLVM::LLVMFuncOp getFreeFn(LLVMTypeConverter &typeConverter,
138138
ModuleOp module) {
139139
return typeConverter.getOptions().useGenericFunctions
140-
? LLVM::lookupOrCreateGenericFreeFn(module)
141-
: LLVM::lookupOrCreateFreeFn(module);
140+
? LLVM::lookupOrCreateGenericFreeFn(
141+
module, typeConverter.useOpaquePointers())
142+
: LLVM::lookupOrCreateFreeFn(module,
143+
typeConverter.useOpaquePointers());
142144
}
143145

144146
inline LLVM::LLVMFuncOp getAllocFn(LLVMTypeConverter &typeConverter,
145147
ModuleOp module, Type indexType) {
146148
return typeConverter.getOptions().useGenericFunctions
147-
? LLVM::lookupOrCreateGenericAllocFn(module, indexType)
148-
: LLVM::lookupOrCreateMallocFn(module, indexType);
149+
? LLVM::lookupOrCreateGenericAllocFn(
150+
module, indexType, typeConverter.useOpaquePointers())
151+
: LLVM::lookupOrCreateMallocFn(module, indexType,
152+
typeConverter.useOpaquePointers());
149153
}
150154

151155
} // namespace mlir

‎polygeist/lib/Conversion/PolygeistToLLVM/PolygeistToLLVM.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -942,7 +942,7 @@ struct ConvertPolygeistToLLVMPass
942942
populatePolygeistToLLVMConversionPatterns(converter, patterns);
943943
populateSCFToControlFlowConversionPatterns(patterns);
944944
cf::populateControlFlowToLLVMConversionPatterns(converter, patterns);
945-
populateMemRefToLLVMConversionPatterns(converter, patterns);
945+
populateFinalizeMemRefToLLVMConversionPatterns(converter, patterns);
946946
populateFuncToLLVMConversionPatterns(converter, patterns);
947947
populateMathToLLVMConversionPatterns(converter, patterns);
948948
populateOpenMPToLLVMConversionPatterns(converter, patterns);

‎polygeist/lib/Dialect/Polygeist/Transforms/LICM.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -493,9 +493,8 @@ Value SCFForGuardBuilder::createGuardExpr() const {
493493
}
494494

495495
scf::IfOp SCFForGuardBuilder::createGuard() const {
496-
TypeRange types(loop->getResults());
497496
return builder.create<scf::IfOp>(
498-
loop.getLoc(), types, createGuardExpr(),
497+
loop.getLoc(), createGuardExpr(),
499498
[&](OpBuilder &b, Location loc) {
500499
b.create<scf::YieldOp>(loc, loop.getResults());
501500
},
@@ -535,9 +534,8 @@ Value SCFParallelGuardBuilder::createGuardExpr() const {
535534
}
536535

537536
scf::IfOp SCFParallelGuardBuilder::createGuard() const {
538-
TypeRange types(loop->getResults());
539537
return builder.create<scf::IfOp>(
540-
loop.getLoc(), types, createGuardExpr(),
538+
loop.getLoc(), createGuardExpr(),
541539
[&](OpBuilder &b, Location loc) {
542540
b.create<scf::YieldOp>(loc, loop.getResults());
543541
},

‎polygeist/test/polygeist-opt/bareptrlowering.mlir

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -30,23 +30,23 @@ func.func private @ptr_args_and_ret(%arg0: memref<1xi64>, %arg1: memref<?xi64>)
3030

3131
// -----
3232

33-
// CHECK-LABEL: llvm.func @ptr_args_and_ret_with_attrs(!llvm.ptr<i64> {llvm.byval = !llvm.ptr<i64>}, !llvm.ptr<i64> {llvm.byval = !llvm.ptr<i64>}) -> !llvm.ptr<i64>
33+
// CHECK-LABEL: llvm.func @ptr_args_and_ret_with_attrs(!llvm.ptr<i64> {llvm.byval = i64}, !llvm.ptr<i64> {llvm.byval = i64}) -> !llvm.ptr<i64>
3434

35-
func.func private @ptr_args_and_ret_with_attrs(%arg0: memref<1xi64> {llvm.byval = memref<1xi64>},
36-
%arg1: memref<?xi64> {llvm.byval = memref<?xi64>}) -> memref<?x4x4xi64>
35+
func.func private @ptr_args_and_ret_with_attrs(%arg0: memref<1xi64> {llvm.byval = i64},
36+
%arg1: memref<?xi64> {llvm.byval = i64}) -> memref<?x4x4xi64>
3737

3838
// -----
3939

4040
gpu.module @kernels {
4141

4242
// CHECK-LABEL: llvm.func @kernel(
43-
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<i64> {llvm.byval = !llvm.ptr<i64>},
44-
// CHECK-SAME: %[[VAL_1:.*]]: !llvm.ptr<i64> {llvm.byval = !llvm.ptr<i64>}) attributes {gpu.kernel, workgroup_attributions = 0 : i64} {
43+
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<i64> {llvm.byval = i64},
44+
// CHECK-SAME: %[[VAL_1:.*]]: !llvm.ptr<i64> {llvm.byval = i64}) attributes {gpu.kernel, workgroup_attributions = 0 : i64} {
4545
// CHECK-NEXT: llvm.return
4646
// CHECK-NEXT: }
4747

48-
gpu.func @kernel(%arg0: memref<1xi64> {llvm.byval = memref<1xi64>},
49-
%arg1: memref<?xi64> {llvm.byval = memref<?xi64>}) kernel {
48+
gpu.func @kernel(%arg0: memref<1xi64> {llvm.byval = i64},
49+
%arg1: memref<?xi64> {llvm.byval = i64}) kernel {
5050
gpu.return
5151
}
5252
}

‎polygeist/tools/cgeist/Lib/CGCall.cc

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -544,12 +544,12 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *Expr) {
544544
assert(0);
545545

546546
if (A0.getType().isa<MemRefType>())
547-
return ValueCategory(Builder.create<memref::AtomicRMWOp>(
548-
Loc, A1.getType(), Op, A1, A0,
549-
std::vector<Value>({getConstantIndex(0)})),
550-
/*isReference*/ false);
547+
return ValueCategory(
548+
Builder.create<memref::AtomicRMWOp>(
549+
Loc, Op, A1, A0, std::vector<Value>({getConstantIndex(0)})),
550+
/*isReference*/ false);
551551
return ValueCategory(
552-
Builder.create<LLVM::AtomicRMWOp>(Loc, A1.getType(), Lop, A0, A1,
552+
Builder.create<LLVM::AtomicRMWOp>(Loc, Lop, A0, A1,
553553
LLVM::AtomicOrdering::acq_rel),
554554
/*isReference*/ false);
555555
}

‎polygeist/tools/cgeist/Lib/CGExpr.cc

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1227,7 +1227,7 @@ ValueCategory MLIRScanner::VisitAtomicExpr(clang::AtomicExpr *BO) {
12271227
Loc, A1.getType(), Op, A1, A0,
12281228
std::vector<mlir::Value>({getConstantIndex(0)}));
12291229
else
1230-
V = Builder.create<LLVM::AtomicRMWOp>(Loc, A1.getType(), Lop, A0, A1,
1230+
V = Builder.create<LLVM::AtomicRMWOp>(Loc, Lop, A0, A1,
12311231
LLVM::AtomicOrdering::acq_rel);
12321232

12331233
if (Ty.isa<mlir::IntegerType>())

‎polygeist/tools/cgeist/Test/Verification/free.c

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2,16 +2,16 @@
22

33
void free(void*);
44

5-
void metafree(void* x, void (*foo)(int), void (*bar)()) {
5+
void metafree(void* x, void (*foo)(int), void (*bar)(void)) {
66
foo(0);
77
bar();
88
free(x);
99
}
1010

11-
// CHECK: func @metafree(%arg0: !llvm.ptr<i8>, %arg1: !llvm.ptr<func<void (i32)>>, %arg2: !llvm.ptr<func<void (...)>>)
11+
// CHECK: func @metafree(%arg0: !llvm.ptr<i8>, %arg1: !llvm.ptr<func<void (i32)>>, %arg2: !llvm.ptr<func<void ()>>)
1212
// CHECK-NEXT: %c0_i32 = arith.constant 0 : i32
13-
// CHECK-NEXT: llvm.call %arg1(%c0_i32) : (i32) -> ()
14-
// CHECK-NEXT: llvm.call %arg2() : () -> ()
13+
// CHECK-NEXT: llvm.call %arg1(%c0_i32) : !llvm.ptr<func<void (i32)>>, (i32) -> ()
14+
// CHECK-NEXT: llvm.call %arg2() : !llvm.ptr<func<void ()>>, () -> ()
1515
// CHECK-NEXT: llvm.call @free(%arg0) : (!llvm.ptr<i8>) -> ()
1616
// CHECK-NEXT: return
1717
// CHECK-NEXT: }

‎polygeist/tools/cgeist/Test/Verification/ptrarith.c

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -67,7 +67,7 @@ void *f4(void *ptr, size_t index) {
6767
// CHECK: %[[VOIDPTR_0:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr<func<i32 ()>> to !llvm.ptr<i8>
6868
// CHECK: %[[PTR:.*]] = llvm.getelementptr %[[VOIDPTR_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr<i8>, i64) -> !llvm.ptr<i8>
6969
// CHECK: %[[FUNCPTR:.*]] = llvm.bitcast %[[PTR]] : !llvm.ptr<i8> to !llvm.ptr<func<i32 ()>>
70-
// CHECK: %[[RES:.*]] = llvm.call %[[FUNCPTR]]() : () -> i32
70+
// CHECK: %[[RES:.*]] = llvm.call %[[FUNCPTR]]() : !llvm.ptr<func<i32 ()>>, () -> i32
7171
// CHECK: return %[[RES]] : i32
7272
// CHECK: }
7373

0 commit comments

Comments
 (0)