Skip to content

Commit cb0c9b0

Browse files
sommerlukassommerlukas
committed
Address PR feedback
Signed-off-by: Lukas Sommer <lukas.sommer@codeplay.com>
1 parent b1063d4 commit cb0c9b0

File tree

6 files changed

+32
-35
lines changed

6 files changed

+32
-35
lines changed

‎polygeist/include/mlir/Dialect/Polygeist/Transforms/Passes.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ namespace polygeist {
2727
/// convention.
2828
void populateBareMemRefToLLVMConversionPatterns(LLVMTypeConverter &converter,
2929
RewritePatternSet &patterns,
30-
bool useOpaquePointer = false);
30+
bool useOpaquePointers = false);
3131

3232
#define GEN_PASS_DECL
3333
#include "mlir/Dialect/Polygeist/Transforms/Passes.h.inc"

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

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -462,7 +462,7 @@ struct BaseSubIndexOpLowering : public ConvertOpToLLVMPattern<SubIndexOp> {
462462
.Case<LLVM::LLVMArrayType>(
463463
[&](auto t) { currType = t.getElementType(); })
464464
.Case<LLVM::LLVMPointerType>(
465-
[&](auto t) { assert(false && "Pointer type not allowed"); })
465+
[&](auto t) { llvm_unreachable("Pointer type not allowed"); })
466466
.Default([&](Type t) {
467467
currType = t;
468468
assert(currType == resElemType &&
@@ -532,7 +532,7 @@ struct SubIndexOpLowering : public BaseSubIndexOpLowering {
532532

533533
// Handle the general (non-SYCL) case first.
534534
if (convViewElemType ==
535-
transformed.getSource().getType().cast<MemRefType>().getElementType()) {
535+
cast<MemRefType>(transformed.getSource().getType()).getElementType()) {
536536
auto memRefDesc = createMemRefDescriptor(
537537
loc, viewMemRefType, targetMemRef.allocatedPtr(rewriter, loc),
538538
rewriter.create<LLVM::GEPOp>(loc, prev.getType(), convViewElemType,
@@ -663,10 +663,9 @@ struct Memref2PointerOpLowering
663663
Value baseOffset = targetMemRef.offset(rewriter, loc);
664664
Value ptr = targetMemRef.alignedPtr(rewriter, loc);
665665
Value idxs[] = {baseOffset};
666-
ptr = rewriter.create<LLVM::GEPOp>(
667-
loc, ptr.getType(),
668-
transformed.getSource().getType().cast<MemRefType>().getElementType(),
669-
ptr, idxs);
666+
auto elemType = getTypeConverter()->convertType(
667+
op.getSource().getType().getElementType());
668+
ptr = rewriter.create<LLVM::GEPOp>(loc, ptr.getType(), elemType, ptr, idxs);
670669
assert(ptr.getType().cast<LLVM::LLVMPointerType>().getAddressSpace() ==
671670
op.getType().getAddressSpace() &&
672671
"Expecting Memref2PointerOp source and result types to have the "

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -633,10 +633,10 @@ struct StoreMemRefOpLoweringOld : public MemAccessLowering {
633633

634634
void mlir::polygeist::populateBareMemRefToLLVMConversionPatterns(
635635
mlir::LLVMTypeConverter &converter, RewritePatternSet &patterns,
636-
bool useOpaquePointer) {
636+
bool useOpaquePointers) {
637637
assert(converter.getOptions().useBarePtrCallConv &&
638638
"Expecting \"bare pointer\" calling convention");
639-
if (useOpaquePointer) {
639+
if (useOpaquePointers) {
640640
patterns.add<GetGlobalMemrefOpLowering, ReshapeMemrefOpLowering,
641641
AllocMemrefOpLowering, AllocaMemrefOpLowering,
642642
CastMemrefOpLowering, DeallocOpLowering, LoadMemRefOpLowering,
@@ -653,11 +653,11 @@ void mlir::polygeist::populateBareMemRefToLLVMConversionPatterns(
653653
// Patterns are tried in reverse add order, so this is tried before the
654654
// one added by default.
655655
converter.addConversion(
656-
[&, useOpaquePointer](MemRefType type) -> Optional<Type> {
656+
[&, useOpaquePointers](MemRefType type) -> Optional<Type> {
657657
if (!canBeLoweredToBarePtr(type))
658658
return std::nullopt;
659659

660-
if (useOpaquePointer) {
660+
if (useOpaquePointers) {
661661
return LLVM::LLVMPointerType::get(type.getContext(),
662662
type.getMemorySpaceAsInt());
663663
}

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

Lines changed: 18 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -59,7 +59,7 @@ memref.global @global : memref<3xi64>
5959

6060
// CHECK-LABEL: llvm.func @get_global() -> !llvm.ptr
6161
// CHECK-NEXT: %[[VAL_0:.*]] = llvm.mlir.addressof @global : !llvm.ptr
62-
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.getelementptr inbounds %[[VAL_0]][0, 0] : (!llvm.ptr) -> !llvm.ptr
62+
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.getelementptr inbounds %[[VAL_0]][0, 0] : (!llvm.ptr) -> !llvm.ptr, i64
6363
// CHECK-NEXT: llvm.return %[[VAL_1]] : !llvm.ptr
6464
// CHECK-NEXT: }
6565

@@ -91,8 +91,7 @@ memref.global "private" constant @shape : memref<2xi64> = dense<[2, 2]>
9191

9292
// CHECK-LABEL: llvm.func @reshape(
9393
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr) -> !llvm.ptr
94-
// CHECK: %[[VAL_2:.*]] = llvm.getelementptr inbounds %{{.*}}[0, 0] : (!llvm.ptr) -> !llvm.ptr
95-
// CHECK-NEXT: llvm.return %[[VAL_0]] : !llvm.ptr
94+
// CHECK: llvm.return %[[VAL_0]] : !llvm.ptr
9695
// CHECK-NEXT: }
9796

9897
func.func private @reshape(%arg0: memref<4xi32>) -> memref<2x2xi32> {
@@ -107,8 +106,7 @@ memref.global "private" constant @shape : memref<1xindex>
107106

108107
// CHECK-LABEL: llvm.func @reshape_dyn(
109108
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr) -> !llvm.ptr
110-
// CHECK: %[[VAL_2:.*]] = llvm.getelementptr inbounds %{{.*}}[0, 0] : (!llvm.ptr) -> !llvm.ptr
111-
// CHECK-NEXT: llvm.return %[[VAL_0]] : !llvm.ptr
109+
// CHECK: llvm.return %[[VAL_0]] : !llvm.ptr
112110
// CHECK-NEXT: }
113111

114112
func.func private @reshape_dyn(%arg0: memref<4xi32>) -> memref<?xi32> {
@@ -122,7 +120,7 @@ func.func private @reshape_dyn(%arg0: memref<4xi32>) -> memref<?xi32> {
122120
// CHECK-LABEL: llvm.func @alloca()
123121
// CHECK-NEXT: %[[VAL_0:.*]] = llvm.mlir.null : !llvm.ptr
124122
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.mlir.constant(2 : index) : i64
125-
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr
123+
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
126124
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.ptrtoint %[[VAL_2]] : !llvm.ptr to i64
127125
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.alloca %[[VAL_3]] x i32 : (i64) -> !llvm.ptr
128126
// CHECK-NEXT: llvm.return
@@ -138,7 +136,7 @@ func.func private @alloca() {
138136
// CHECK-LABEL: llvm.func @alloca_nd()
139137
// CHECK-NEXT: %[[VAL_0:.*]] = llvm.mlir.null : !llvm.ptr
140138
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.mlir.constant(60 : index) : i64
141-
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr
139+
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
142140
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.ptrtoint %[[VAL_2]] : !llvm.ptr to i64
143141
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.alloca %[[VAL_3]] x i32 : (i64) -> !llvm.ptr
144142
// CHECK-NEXT: llvm.return
@@ -154,7 +152,7 @@ func.func private @alloca_nd() {
154152
// CHECK-LABEL: llvm.func @alloca_aligned()
155153
// CHECK-NEXT: %[[VAL_0:.*]] = llvm.mlir.null : !llvm.ptr
156154
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.mlir.constant(2 : index) : i64
157-
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr
155+
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
158156
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.ptrtoint %[[VAL_2]] : !llvm.ptr to i64
159157
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.alloca %[[VAL_3]] x i32 {alignment = 8 : i64} : (i64) -> !llvm.ptr
160158
// CHECK-NEXT: llvm.return
@@ -170,7 +168,7 @@ func.func private @alloca_aligned() {
170168
// CHECK-LABEL: llvm.func @alloca_nd_aligned()
171169
// CHECK-NEXT: %[[VAL_0:.*]] = llvm.mlir.null : !llvm.ptr
172170
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.mlir.constant(60 : index) : i64
173-
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr
171+
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
174172
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.ptrtoint %[[VAL_2]] : !llvm.ptr to i64
175173
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.alloca %[[VAL_3]] x i32 {alignment = 8 : i64} : (i64) -> !llvm.ptr
176174
// CHECK-NEXT: llvm.return
@@ -189,7 +187,7 @@ func.func private @alloca_nd_aligned() {
189187
// CHECK-NEXT: %[[VAL_0:.*]] = llvm.mlir.constant(2 : index) : i64
190188
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.mlir.constant(1 : index) : i64
191189
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.mlir.null : !llvm.ptr
192-
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.getelementptr %[[VAL_2]]{{\[}}%[[VAL_0]]] : (!llvm.ptr, i64) -> !llvm.ptr
190+
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.getelementptr %[[VAL_2]]{{\[}}%[[VAL_0]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
193191
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.ptrtoint %[[VAL_3]] : !llvm.ptr to i64
194192
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.call @malloc(%[[VAL_4]]) : (i64) -> !llvm.ptr
195193
// CHECK-NEXT: llvm.return
@@ -210,7 +208,7 @@ func.func private @alloc() {
210208
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.mlir.constant(20 : index) : i64
211209
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.mlir.constant(60 : index) : i64
212210
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.mlir.null : !llvm.ptr
213-
// CHECK-NEXT: %[[VAL_7:.*]] = llvm.getelementptr %[[VAL_6]]{{\[}}%[[VAL_5]]] : (!llvm.ptr, i64) -> !llvm.ptr
211+
// CHECK-NEXT: %[[VAL_7:.*]] = llvm.getelementptr %[[VAL_6]]{{\[}}%[[VAL_5]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
214212
// CHECK-NEXT: %[[VAL_8:.*]] = llvm.ptrtoint %[[VAL_7]] : !llvm.ptr to i64
215213
// CHECK-NEXT: %[[VAL_9:.*]] = llvm.call @malloc(%[[VAL_8]]) : (i64) -> !llvm.ptr
216214
// CHECK-NEXT: llvm.return
@@ -227,7 +225,7 @@ func.func private @alloc_nd() {
227225
// CHECK-NEXT: %[[VAL_0:.*]] = llvm.mlir.constant(2 : index) : i64
228226
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.mlir.constant(1 : index) : i64
229227
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.mlir.null : !llvm.ptr
230-
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.getelementptr %[[VAL_2]]{{\[}}%[[VAL_0]]] : (!llvm.ptr, i64) -> !llvm.ptr
228+
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.getelementptr %[[VAL_2]]{{\[}}%[[VAL_0]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
231229
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.ptrtoint %[[VAL_3]] : !llvm.ptr to i64
232230
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.mlir.constant(8 : index) : i64
233231
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.add %[[VAL_4]], %[[VAL_5]] : i64
@@ -257,7 +255,7 @@ func.func private @alloc_aligned() {
257255
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.mlir.constant(20 : index) : i64
258256
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.mlir.constant(60 : index) : i64
259257
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.mlir.null : !llvm.ptr
260-
// CHECK-NEXT: %[[VAL_7:.*]] = llvm.getelementptr %[[VAL_6]]{{\[}}%[[VAL_5]]] : (!llvm.ptr, i64) -> !llvm.ptr
258+
// CHECK-NEXT: %[[VAL_7:.*]] = llvm.getelementptr %[[VAL_6]]{{\[}}%[[VAL_5]]] : (!llvm.ptr, i64) -> !llvm.ptr, i32
261259
// CHECK-NEXT: %[[VAL_8:.*]] = llvm.ptrtoint %[[VAL_7]] : !llvm.ptr to i64
262260
// CHECK-NEXT: %[[VAL_9:.*]] = llvm.mlir.constant(8 : index) : i64
263261
// CHECK-NEXT: %[[VAL_10:.*]] = llvm.add %[[VAL_8]], %[[VAL_9]] : i64
@@ -307,7 +305,7 @@ func.func private @cast(%arg0: memref<2xi32>) -> memref<?xi32> {
307305
// CHECK-LABEL: llvm.func @load(
308306
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr,
309307
// CHECK-SAME: %[[VAL_1:.*]]: i64) -> f32
310-
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr
308+
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
311309
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.load %[[VAL_2]] : !llvm.ptr
312310
// CHECK-NEXT: llvm.return %[[VAL_3]] : f32
313311
// CHECK-NEXT: }
@@ -326,7 +324,7 @@ func.func private @load(%arg0: memref<100xf32>, %index: index) -> f32 {
326324
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.mlir.constant(100 : index) : i64
327325
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.mul %[[VAL_1]], %[[VAL_3]] : i64
328326
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.add %[[VAL_4]], %[[VAL_2]] : i64
329-
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_5]]] : (!llvm.ptr, i64) -> !llvm.ptr
327+
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_5]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
330328
// CHECK-NEXT: %[[VAL_7:.*]] = llvm.load %[[VAL_6]] : !llvm.ptr
331329
// CHECK-NEXT: llvm.return %[[VAL_7]] : f32
332330
// CHECK-NEXT: }
@@ -345,7 +343,7 @@ func.func private @load_nd(%arg0: memref<100x100xf32>, %index0: index, %index1:
345343
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.mlir.constant(100 : index) : i64
346344
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.mul %[[VAL_1]], %[[VAL_3]] : i64
347345
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.add %[[VAL_4]], %[[VAL_2]] : i64
348-
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_5]]] : (!llvm.ptr, i64) -> !llvm.ptr
346+
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_5]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
349347
// CHECK-NEXT: %[[VAL_7:.*]] = llvm.load %[[VAL_6]] : !llvm.ptr
350348
// CHECK-NEXT: llvm.return %[[VAL_7]] : f32
351349
// CHECK-NEXT: }
@@ -434,7 +432,7 @@ func.func private @call(%arg0: memref<?xf32>, %arg1: index) -> memref<?xf32> {
434432
// CHECK-SAME: %[[VAL_1:.*]]: i64) -> !llvm.ptr
435433
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.mlir.constant(4 : i64) : i64
436434
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.mul %[[VAL_1]], %[[VAL_2]] : i64
437-
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_3]]] : (!llvm.ptr, i64) -> !llvm.ptr
435+
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_3]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
438436
// CHECK-NEXT: llvm.return %[[VAL_4]] : !llvm.ptr
439437
// CHECK-NEXT: }
440438

@@ -448,7 +446,7 @@ func.func private @subindexop_memref(%arg0: memref<4x4xf32>, %arg1: index) -> me
448446
// CHECK-LABEL: llvm.func @subindexop_memref_same_dim(
449447
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr,
450448
// CHECK-SAME: %[[VAL_1:.*]]: i64) -> !llvm.ptr
451-
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr
449+
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr, f32
452450
// CHECK-NEXT: llvm.return %[[VAL_2]] : !llvm.ptr
453451
// CHECK-NEXT: }
454452

@@ -479,7 +477,7 @@ func.func private @subindexop_memref_struct(%arg0: memref<4x!llvm.struct<(f32)>>
479477
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr) -> !llvm.ptr
480478
// CHECK-NEXT: %[[VAL_1:.*]] = llvm.mlir.constant(0 : index) : i64
481479
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.mlir.constant(0 : i64) : i64
482-
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_2]], %[[VAL_2]], %[[VAL_1]]] : (!llvm.ptr, i64, i64, i64) -> !llvm.ptr
480+
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_2]], %[[VAL_2]], %[[VAL_1]]] : (!llvm.ptr, i64, i64, i64) -> !llvm.ptr, f32
483481
// CHECK-NEXT: llvm.return %[[VAL_3]] : !llvm.ptr
484482
// CHECK-NEXT: }
485483

@@ -539,7 +537,7 @@ func.func private @ptr2memref(%arg0: !llvm.ptr) -> memref<?xf32> {
539537
// CHECK-NEXT: %[[VAL_2:.*]] = llvm.extractvalue %[[VAL_0]][1] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
540538
// CHECK-NEXT: %[[VAL_3:.*]] = llvm.mlir.constant(-1 : index) : i64
541539
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.add %[[VAL_3]], %[[VAL_1]] : i64
542-
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.getelementptr %[[VAL_2]][%[[VAL_4]]] : (!llvm.ptr, i64) -> !llvm.ptr
540+
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.getelementptr %[[VAL_2]][%[[VAL_4]]] : (!llvm.ptr, i64) -> !llvm.ptr, i64
543541
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.load %[[VAL_5]] : !llvm.ptr
544542
// CHECK-NEXT: llvm.return %[[VAL_6]] : i64
545543
// CHECK-NEXT: }

‎polygeist/test/polygeist-opt/sycl/cast.mlir‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@
1010
// CHECK: }
1111

1212
func.func @test1(%arg0: memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_> {
13-
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_>
13+
%0 = sycl.cast %arg0 : memref<?x!sycl_range_1_> to memref<?x!sycl_array_1_>
1414
func.return %0 : memref<?x!sycl_array_1_>
1515
}
1616

@@ -25,7 +25,7 @@ func.func @test1(%arg0: memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_> {
2525
!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
2626
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
2727
func.func @test2(%arg0: memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_> {
28-
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_>
28+
%0 = sycl.cast %arg0 : memref<?x!sycl_id_1_> to memref<?x!sycl_array_1_>
2929
func.return %0: memref<?x!sycl_array_1_>
3030
}
3131

@@ -40,6 +40,6 @@ func.func @test2(%arg0: memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_> {
4040
!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
4141
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
4242
func.func @test_addrspaces(%arg0: memref<?x!sycl_id_1_, 4>) -> memref<?x!sycl_array_1_, 4> {
43-
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_id_1_, 4>) -> memref<?x!sycl_array_1_, 4>
43+
%0 = sycl.cast %arg0 : memref<?x!sycl_id_1_, 4> to memref<?x!sycl_array_1_, 4>
4444
func.return %0: memref<?x!sycl_array_1_, 4>
4545
}

‎polygeist/test/polygeist-opt/sycl/subindex.mlir‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,7 @@ func.func @test_3(%arg0: memref<?x!llvm.struct<(i32)>>) -> memref<?xi32> {
4646
// -----
4747

4848
// CHECK: llvm.func @test_4([[A0:%.*]]: !llvm.ptr, [[A5:%.*]]: i64) -> !llvm.ptr {
49-
// CHECK: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[A5]]] : (!llvm.ptr, i64) -> !llvm.ptr
49+
// CHECK: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[A5]]] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(struct<"class.sycl::_V1::id.1", {{.*}})>
5050
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr
5151

5252
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>

0 commit comments

Comments
 (0)