Skip to content

Commit b1e03ed

Browse files
sommerlukassommerlukas
authored
[SYCL-MLIR] Opaque pointer support Polygeist-to-LLVM (#8767)
First step to add opaque pointer support for the SYCL-MLIR project: Make the lowering patterns in the Polygeist-to-LLVM lowering compatible with opaque pointers. Partly resolves #8616. --------- Signed-off-by: Lukas Sommer <lukas.sommer@codeplay.com>
1 parent 30fc9c9 commit b1e03ed

File tree

10 files changed

+1830
-233
lines changed

10 files changed

+1830
-233
lines changed

‎polygeist/include/mlir/Conversion/PolygeistPasses.td

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,10 @@ def ConvertPolygeistToLLVM : Pass<"convert-polygeist-to-llvm", "mlir::ModuleOp">
2424
Option<"dataLayout", "data-layout", "std::string",
2525
/*default=*/"\"\"",
2626
"String description (LLVM format) of the data layout that is "
27-
"expected on the produced module">
27+
"expected on the produced module">,
28+
Option<"useOpaquePointers", "use-opaque-pointers", "bool",
29+
/*default=*/"false", "Generate LLVM IR using opaque pointers "
30+
"instead of typed pointers">,
2831
];
2932
}
3033

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,8 @@ namespace polygeist {
2626
/// MemRef dialect to the LLVM dialect forcing a "bare pointer" calling
2727
/// convention.
2828
void populateBareMemRefToLLVMConversionPatterns(LLVMTypeConverter &converter,
29-
RewritePatternSet &patterns);
29+
RewritePatternSet &patterns,
30+
bool useOpaquePointers = false);
3031

3132
#define GEN_PASS_DECL
3233
#include "mlir/Dialect/Polygeist/Transforms/Passes.h.inc"

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

Lines changed: 630 additions & 33 deletions
Large diffs are not rendered by default.

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

Lines changed: 347 additions & 30 deletions
Large diffs are not rendered by default.

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

Lines changed: 573 additions & 0 deletions
Large diffs are not rendered by default.

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

Lines changed: 135 additions & 144 deletions
Large diffs are not rendered by default.

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

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
// RUN: polygeist-opt --convert-polygeist-to-llvm='use-opaque-pointers=0' --split-input-file %s | FileCheck %s
2+
3+
!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
4+
!sycl_range_1_ = !sycl.range<[1], (!sycl_array_1_)>
5+
6+
// CHECK-LABEL: llvm.func @test1(
7+
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<struct<"class.sycl::_V1::range.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>>) -> !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>> {
8+
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr<struct<"class.sycl::_V1::range.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>> to !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
9+
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
10+
// CHECK: }
11+
12+
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_>
14+
func.return %0 : memref<?x!sycl_array_1_>
15+
}
16+
17+
// -----
18+
19+
// CHECK-LABEL: llvm.func @test2(
20+
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>>) -> !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>> {
21+
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr<struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>> to !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
22+
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
23+
// CHECK: }
24+
25+
!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
26+
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
27+
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_>
29+
func.return %0: memref<?x!sycl_array_1_>
30+
}
31+
32+
// -----
33+
34+
// CHECK-LABEL: llvm.func @test_addrspaces(
35+
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>, 4>) -> !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>, 4> {
36+
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr<struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>, 4> to !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>, 4>
37+
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>, 4>
38+
// CHECK: }
39+
40+
!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
41+
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
42+
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>
44+
func.return %0: memref<?x!sycl_array_1_, 4>
45+
}

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

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,45 +1,45 @@
1-
// RUN: polygeist-opt --convert-polygeist-to-llvm --split-input-file %s | FileCheck %s
1+
// RUN: polygeist-opt --convert-polygeist-to-llvm='use-opaque-pointers=1' --split-input-file %s | FileCheck %s
22

33
!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
44
!sycl_range_1_ = !sycl.range<[1], (!sycl_array_1_)>
55

66
// CHECK-LABEL: llvm.func @test1(
7-
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<struct<"class.sycl::_V1::range.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>>) -> !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>> {
8-
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr<struct<"class.sycl::_V1::range.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>> to !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
9-
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
7+
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr) -> !llvm.ptr {
8+
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr to !llvm.ptr
9+
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr
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

1717
// -----
1818

1919
// CHECK-LABEL: llvm.func @test2(
20-
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>>) -> !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>> {
21-
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr<struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>> to !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
22-
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
20+
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr) -> !llvm.ptr {
21+
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr to !llvm.ptr
22+
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr
2323
// CHECK: }
2424

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

3232
// -----
3333

3434
// CHECK-LABEL: llvm.func @test_addrspaces(
35-
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>, 4>) -> !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>, 4> {
36-
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr<struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>, 4> to !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>, 4>
37-
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>, 4>
35+
// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr<4>) -> !llvm.ptr<4> {
36+
// CHECK: %[[VAL_1:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr<4> to !llvm.ptr<4>
37+
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<4>
3838
// CHECK: }
3939

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-typed-pointer.mlir

Lines changed: 68 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,68 @@
1+
// RUN: polygeist-opt --convert-polygeist-to-llvm='use-opaque-pointers=0' --split-input-file %s | FileCheck %s
2+
3+
// CHECK-LABEL: @test_1
4+
// CHECK: [[ZERO:%.*]] = llvm.mlir.constant(0 : i64) : i64
5+
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr %{{.*}}[[[ZERO]], 0] : (!llvm.ptr<struct<([[SYCLIDSTRUCT:struct<"class.sycl::_V1::id.1"]], {{.*}} -> !llvm.ptr<[[SYCLIDSTRUCT]], {{.*}}
6+
// CHECK-NEXT: llvm.return [[GEP]]
7+
8+
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
9+
func.func @test_1(%arg0: memref<?x!llvm.struct<(!sycl_id_1_)>>) -> memref<?x!sycl_id_1_> {
10+
%c0 = arith.constant 0 : index
11+
%0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!llvm.struct<(!sycl_id_1_)>>, index) -> memref<?x!sycl_id_1_>
12+
return %0 : memref<?x!sycl_id_1_>
13+
}
14+
15+
// -----
16+
17+
// CHECK-LABEL: @test_2
18+
// CHECK: llvm.return %{{.*}} : !llvm.ptr<struct<"class.sycl::_V1::detail::AccessorImplDevice
19+
20+
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
21+
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
22+
!sycl_accessor_impl_device_1_ = !sycl.accessor_impl_device<[1], (!sycl_id_1_, !sycl_range_1_, !sycl_range_1_)>
23+
!sycl_accessor_1_ = !sycl.accessor<[1, i32, read_write, global_buffer], (!sycl.accessor_impl_device<[1], (!sycl_id_1_, !sycl_range_1_, !sycl_range_1_)>, !llvm.struct<(ptr<i32, 1>)>)>
24+
25+
func.func @test_2(%arg0: memref<?x!sycl_accessor_1_>) -> memref<?x!sycl_accessor_impl_device_1_> {
26+
%c0 = arith.constant 0 : index
27+
%0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!sycl_accessor_1_>, index) -> memref<?x!sycl_accessor_impl_device_1_>
28+
return %0 : memref<?x!sycl_accessor_impl_device_1_>
29+
}
30+
31+
// -----
32+
33+
// CHECK: llvm.func @test_3([[A0:.*]]: !llvm.ptr<struct<(i32)>>) -> !llvm.ptr<i32> {
34+
// CHECK: [[ZERO:%.*]] = llvm.mlir.constant(0 : i64) : i64
35+
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO]], 0] : (!llvm.ptr<struct<(i32)>>, i64) -> !llvm.ptr<i32>
36+
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr<i32>
37+
38+
func.func @test_3(%arg0: memref<?x!llvm.struct<(i32)>>) -> memref<?xi32> {
39+
%c0 = arith.constant 0 : index
40+
%0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!llvm.struct<(i32)>>, index) -> memref<?xi32>
41+
return %0 : memref<?xi32>
42+
}
43+
44+
// -----
45+
46+
// CHECK: llvm.func @test_4([[A0:%.*]]: !llvm.ptr<struct<([[IDTYPE:struct<"class.sycl::_V1::id.1", \(struct<"class.sycl::_V1::detail::array.1", \(array<1 x i64>\)>\)>]])>>, [[A5:%.*]]: i64) -> !llvm.ptr<struct<(struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>)>> {
47+
// CHECK: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[A5]]] : (!llvm.ptr<struct<([[IDTYPE]])>>, i64) -> !llvm.ptr<struct<([[IDTYPE]])>>
48+
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr<struct<([[IDTYPE]])>>
49+
50+
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
51+
func.func @test_4(%arg0: memref<1x!llvm.struct<(!sycl_id_1_)>>, %arg1: index) -> memref<?x!llvm.struct<(!sycl_id_1_)>> {
52+
%0 = "polygeist.subindex"(%arg0, %arg1) : (memref<1x!llvm.struct<(!sycl_id_1_)>>, index) -> memref<?x!llvm.struct<(!sycl_id_1_)>>
53+
return %0 : memref<?x!llvm.struct<(!sycl_id_1_)>>
54+
}
55+
56+
// -----
57+
58+
// CHECK: llvm.func @test_5([[A0:%.*]]: !llvm.ptr<[[ARRTYPE:struct<"class.sycl::_V1::detail::array.1", \(array<1 x i64>\)>]], 4>) -> !llvm.ptr<i64, 4> {
59+
// CHECK-DAG: [[ZERO1:%.*]] = llvm.mlir.constant(0 : index) : i64
60+
// CHECK-DAG: [[ZERO2:%.*]] = llvm.mlir.constant(0 : i64) : i64
61+
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO2]], 0, [[ZERO1]]] : (!llvm.ptr<[[ARRTYPE]], 4>, i64, i64) -> !llvm.ptr<i64, 4>
62+
63+
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
64+
func.func @test_5(%arg0: memref<?x!sycl.array<[1], (memref<1xi64, 4>)>, 4>) -> memref<1xi64, 4> {
65+
%c0 = arith.constant 0 : index
66+
%0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!sycl.array<[1], (memref<1xi64, 4>)>, 4>, index) -> memref<1xi64, 4>
67+
return %0 : memref<1xi64, 4>
68+
}

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

Lines changed: 13 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,8 @@
1-
// RUN: polygeist-opt --convert-polygeist-to-llvm --split-input-file %s | FileCheck %s
1+
// RUN: polygeist-opt --convert-polygeist-to-llvm='use-opaque-pointers=1' --split-input-file %s | FileCheck %s
22

33
// CHECK-LABEL: @test_1
44
// CHECK: [[ZERO:%.*]] = llvm.mlir.constant(0 : i64) : i64
5-
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr %{{.*}}[[[ZERO]], 0] : (!llvm.ptr<struct<([[SYCLIDSTRUCT:struct<"class.sycl::_V1::id.1"]], {{.*}} -> !llvm.ptr<[[SYCLIDSTRUCT]], {{.*}}
5+
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr %{{.*}}[[[ZERO]], 0] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::id.1", {{.*}}
66
// CHECK-NEXT: llvm.return [[GEP]]
77

88
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
@@ -15,7 +15,8 @@ func.func @test_1(%arg0: memref<?x!llvm.struct<(!sycl_id_1_)>>) -> memref<?x!syc
1515
// -----
1616

1717
// CHECK-LABEL: @test_2
18-
// CHECK: llvm.return %{{.*}} : !llvm.ptr<struct<"class.sycl::_V1::detail::AccessorImplDevice
18+
// CHECK: [[GEP:%.*]] = llvm.getelementptr %{{.*}}[%{{.*}}, {{.*}}] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<"class.sycl::_V1::detail::AccessorImplDevice.1", {{.*}}
19+
// CHECK-NEXT: llvm.return [[GEP]]
1920

2021
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
2122
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
@@ -30,10 +31,11 @@ func.func @test_2(%arg0: memref<?x!sycl_accessor_1_>) -> memref<?x!sycl_accessor
3031

3132
// -----
3233

33-
// CHECK: llvm.func @test_3([[A0:.*]]: !llvm.ptr<struct<(i32)>>) -> !llvm.ptr<i32> {
34+
// CHECK: llvm.func @test_3([[A0:.*]]: !llvm.ptr) -> !llvm.ptr {
35+
// CHECK: [[IDX_ZERO:%.*]] = llvm.mlir.constant(0 : index) : i64
3436
// CHECK: [[ZERO:%.*]] = llvm.mlir.constant(0 : i64) : i64
35-
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO]], 0] : (!llvm.ptr<struct<(i32)>>, i64) -> !llvm.ptr<i32>
36-
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr<i32>
37+
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO]], [[IDX_ZERO]]] : (!llvm.ptr, i64, i64) -> !llvm.ptr, i32
38+
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr
3739

3840
func.func @test_3(%arg0: memref<?x!llvm.struct<(i32)>>) -> memref<?xi32> {
3941
%c0 = arith.constant 0 : index
@@ -43,9 +45,9 @@ func.func @test_3(%arg0: memref<?x!llvm.struct<(i32)>>) -> memref<?xi32> {
4345

4446
// -----
4547

46-
// CHECK: llvm.func @test_4([[A0:%.*]]: !llvm.ptr<struct<([[IDTYPE:struct<"class.sycl::_V1::id.1", \(struct<"class.sycl::_V1::detail::array.1", \(array<1 x i64>\)>\)>]])>>, [[A5:%.*]]: i64) -> !llvm.ptr<struct<(struct<"class.sycl::_V1::id.1", (struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>)>)>> {
47-
// CHECK: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[A5]]] : (!llvm.ptr<struct<([[IDTYPE]])>>, i64) -> !llvm.ptr<struct<([[IDTYPE]])>>
48-
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr<struct<([[IDTYPE]])>>
48+
// CHECK: llvm.func @test_4([[A0:%.*]]: !llvm.ptr, [[A5:%.*]]: i64) -> !llvm.ptr {
49+
// CHECK: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[A5]]] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(struct<"class.sycl::_V1::id.1", {{.*}})>
50+
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr
4951

5052
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
5153
func.func @test_4(%arg0: memref<1x!llvm.struct<(!sycl_id_1_)>>, %arg1: index) -> memref<?x!llvm.struct<(!sycl_id_1_)>> {
@@ -55,10 +57,10 @@ func.func @test_4(%arg0: memref<1x!llvm.struct<(!sycl_id_1_)>>, %arg1: index) ->
5557

5658
// -----
5759

58-
// CHECK: llvm.func @test_5([[A0:%.*]]: !llvm.ptr<[[ARRTYPE:struct<"class.sycl::_V1::detail::array.1", \(array<1 x i64>\)>]], 4>) -> !llvm.ptr<i64, 4> {
60+
// CHECK: llvm.func @test_5([[A0:%.*]]: !llvm.ptr<4>) -> !llvm.ptr<4> {
5961
// CHECK-DAG: [[ZERO1:%.*]] = llvm.mlir.constant(0 : index) : i64
6062
// CHECK-DAG: [[ZERO2:%.*]] = llvm.mlir.constant(0 : i64) : i64
61-
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO2]], 0, [[ZERO1]]] : (!llvm.ptr<[[ARRTYPE]], 4>, i64, i64) -> !llvm.ptr<i64, 4>
63+
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO2]], [[ZERO2]], [[ZERO1]]] : (!llvm.ptr<4>, i64, i64, i64) -> !llvm.ptr<4>, i64
6264

6365
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
6466
func.func @test_5(%arg0: memref<?x!sycl.array<[1], (memref<1xi64, 4>)>, 4>) -> memref<1xi64, 4> {

0 commit comments

Comments
 (0)