Skip to content

[SYCL-MLIR] Opaque pointer support Polygeist-to-LLVM #8767

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
sommerlukas merged 6 commits into from
Mar 28, 2023
Merged

[SYCL-MLIR] Opaque pointer support Polygeist-to-LLVM #8767

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
e77750c
[SYCL-MLIR] Opaque pointer Polygeist to LLVM
sommerlukas Mar 17, 2023
9bbb47c
Remove unncessary bitcasts and fix tests
sommerlukas Mar 24, 2023
49f93e1
Version BareMemRefToLLVM patterns
sommerlukas Mar 24, 2023
b1063d4
Add versioning and opaque pointer option
sommerlukas Mar 24, 2023
cb0c9b0
Address PR feedback
sommerlukas Mar 27, 2023
f7823e0
Add typed pointer tests and test for nested pointer
sommerlukas Mar 28, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion polygeist/include/mlir/Conversion/PolygeistPasses.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,10 @@ def ConvertPolygeistToLLVM : Pass<"convert-polygeist-to-llvm", "mlir::ModuleOp">
Option<"dataLayout", "data-layout", "std::string",
/*default=*/"\"\"",
"String description (LLVM format) of the data layout that is "
"expected on the produced module">
"expected on the produced module">,
Option<"useOpaquePointers", "use-opaque-pointers", "bool",
/*default=*/"false", "Generate LLVM IR using opaque pointers "
"instead of typed pointers">,
];
}

Expand Down
3 changes: 2 additions & 1 deletion polygeist/include/mlir/Dialect/Polygeist/Transforms/Passes.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,8 @@ namespace polygeist {
/// MemRef dialect to the LLVM dialect forcing a "bare pointer" calling
/// convention.
void populateBareMemRefToLLVMConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns);
RewritePatternSet &patterns,
bool useOpaquePointers = false);

#define GEN_PASS_DECL
#include "mlir/Dialect/Polygeist/Transforms/Passes.h.inc"
Expand Down
663 changes: 630 additions & 33 deletions polygeist/lib/Conversion/PolygeistToLLVM/PolygeistToLLVM.cpp
View file
Open in desktop

Large diffs are not rendered by default.

377 changes: 347 additions & 30 deletions polygeist/lib/Dialect/Polygeist/Transforms/BareMemRefToLLVM.cpp
View file
Open in desktop

Large diffs are not rendered by default.

573 changes: 573 additions & 0 deletions polygeist/test/polygeist-opt/bareptrlowering-typed-pointer.mlir
View file
Open in desktop

Large diffs are not rendered by default.

279 changes: 135 additions & 144 deletions polygeist/test/polygeist-opt/bareptrlowering.mlir
View file
Open in desktop

Large diffs are not rendered by default.

45 changes: 45 additions & 0 deletions polygeist/test/polygeist-opt/sycl/cast-typed-pointer.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,45 @@
// RUN: polygeist-opt --convert-polygeist-to-llvm='use-opaque-pointers=0' --split-input-file %s | FileCheck %s

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

// CHECK-LABEL: llvm.func @test1(
// 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>)>> {
// 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>)>>
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
// CHECK: }

func.func @test1(%arg0: memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_> {
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_>
func.return %0 : memref<?x!sycl_array_1_>
}

// -----

// CHECK-LABEL: llvm.func @test2(
// 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>)>> {
// 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>)>>
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>>
// CHECK: }

!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
func.func @test2(%arg0: memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_> {
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_>
func.return %0: memref<?x!sycl_array_1_>
}

// -----

// CHECK-LABEL: llvm.func @test_addrspaces(
// 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> {
// 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>
// CHECK: llvm.return %[[VAL_1]] : !llvm.ptr<struct<"class.sycl::_V1::detail::array.1", (array<1 x i64>)>, 4>
// CHECK: }

!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
func.func @test_addrspaces(%arg0: memref<?x!sycl_id_1_, 4>) -> memref<?x!sycl_array_1_, 4> {
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_id_1_, 4>) -> memref<?x!sycl_array_1_, 4>
func.return %0: memref<?x!sycl_array_1_, 4>
}
26 changes: 13 additions & 13 deletions polygeist/test/polygeist-opt/sycl/cast.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,45 +1,45 @@
// RUN: polygeist-opt --convert-polygeist-to-llvm --split-input-file %s | FileCheck %s
// RUN: polygeist-opt --convert-polygeist-to-llvm='use-opaque-pointers=1' --split-input-file %s | FileCheck %s

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

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

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

// -----

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

!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
func.func @test2(%arg0: memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_> {
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_>
%0 = sycl.cast %arg0 : memref<?x!sycl_id_1_> to memref<?x!sycl_array_1_>
func.return %0: memref<?x!sycl_array_1_>
}

// -----

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

!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
func.func @test_addrspaces(%arg0: memref<?x!sycl_id_1_, 4>) -> memref<?x!sycl_array_1_, 4> {
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_id_1_, 4>) -> memref<?x!sycl_array_1_, 4>
%0 = sycl.cast %arg0 : memref<?x!sycl_id_1_, 4> to memref<?x!sycl_array_1_, 4>
func.return %0: memref<?x!sycl_array_1_, 4>
}
68 changes: 68 additions & 0 deletions polygeist/test/polygeist-opt/sycl/subindex-typed-pointer.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,68 @@
// RUN: polygeist-opt --convert-polygeist-to-llvm='use-opaque-pointers=0' --split-input-file %s | FileCheck %s

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

!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
func.func @test_1(%arg0: memref<?x!llvm.struct<(!sycl_id_1_)>>) -> memref<?x!sycl_id_1_> {
%c0 = arith.constant 0 : index
%0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!llvm.struct<(!sycl_id_1_)>>, index) -> memref<?x!sycl_id_1_>
return %0 : memref<?x!sycl_id_1_>
}

// -----

// CHECK-LABEL: @test_2
// CHECK: llvm.return %{{.*}} : !llvm.ptr<struct<"class.sycl::_V1::detail::AccessorImplDevice

!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
!sycl_accessor_impl_device_1_ = !sycl.accessor_impl_device<[1], (!sycl_id_1_, !sycl_range_1_, !sycl_range_1_)>
!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>)>)>

func.func @test_2(%arg0: memref<?x!sycl_accessor_1_>) -> memref<?x!sycl_accessor_impl_device_1_> {
%c0 = arith.constant 0 : index
%0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!sycl_accessor_1_>, index) -> memref<?x!sycl_accessor_impl_device_1_>
return %0 : memref<?x!sycl_accessor_impl_device_1_>
}

// -----

// CHECK: llvm.func @test_3([[A0:.*]]: !llvm.ptr<struct<(i32)>>) -> !llvm.ptr<i32> {
// CHECK: [[ZERO:%.*]] = llvm.mlir.constant(0 : i64) : i64
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO]], 0] : (!llvm.ptr<struct<(i32)>>, i64) -> !llvm.ptr<i32>
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr<i32>

func.func @test_3(%arg0: memref<?x!llvm.struct<(i32)>>) -> memref<?xi32> {
%c0 = arith.constant 0 : index
%0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!llvm.struct<(i32)>>, index) -> memref<?xi32>
return %0 : memref<?xi32>
}

// -----

// 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>)>)>)>> {
// CHECK: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[A5]]] : (!llvm.ptr<struct<([[IDTYPE]])>>, i64) -> !llvm.ptr<struct<([[IDTYPE]])>>
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr<struct<([[IDTYPE]])>>

!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
func.func @test_4(%arg0: memref<1x!llvm.struct<(!sycl_id_1_)>>, %arg1: index) -> memref<?x!llvm.struct<(!sycl_id_1_)>> {
%0 = "polygeist.subindex"(%arg0, %arg1) : (memref<1x!llvm.struct<(!sycl_id_1_)>>, index) -> memref<?x!llvm.struct<(!sycl_id_1_)>>
return %0 : memref<?x!llvm.struct<(!sycl_id_1_)>>
}

// -----

// 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> {
// CHECK-DAG: [[ZERO1:%.*]] = llvm.mlir.constant(0 : index) : i64
// CHECK-DAG: [[ZERO2:%.*]] = llvm.mlir.constant(0 : i64) : i64
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO2]], 0, [[ZERO1]]] : (!llvm.ptr<[[ARRTYPE]], 4>, i64, i64) -> !llvm.ptr<i64, 4>

!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
func.func @test_5(%arg0: memref<?x!sycl.array<[1], (memref<1xi64, 4>)>, 4>) -> memref<1xi64, 4> {
%c0 = arith.constant 0 : index
%0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!sycl.array<[1], (memref<1xi64, 4>)>, 4>, index) -> memref<1xi64, 4>
return %0 : memref<1xi64, 4>
}
24 changes: 13 additions & 11 deletions polygeist/test/polygeist-opt/sycl/subindex.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,8 +1,8 @@
// RUN: polygeist-opt --convert-polygeist-to-llvm --split-input-file %s | FileCheck %s
// RUN: polygeist-opt --convert-polygeist-to-llvm='use-opaque-pointers=1' --split-input-file %s | FileCheck %s

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

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

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

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

// -----

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

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

// -----

// 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>)>)>)>> {
// CHECK: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[A5]]] : (!llvm.ptr<struct<([[IDTYPE]])>>, i64) -> !llvm.ptr<struct<([[IDTYPE]])>>
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr<struct<([[IDTYPE]])>>
// CHECK: llvm.func @test_4([[A0:%.*]]: !llvm.ptr, [[A5:%.*]]: i64) -> !llvm.ptr {
// CHECK: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[A5]]] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(struct<"class.sycl::_V1::id.1", {{.*}})>
// CHECK-NEXT: llvm.return [[GEP]] : !llvm.ptr

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

// -----

// 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> {
// CHECK: llvm.func @test_5([[A0:%.*]]: !llvm.ptr<4>) -> !llvm.ptr<4> {
// CHECK-DAG: [[ZERO1:%.*]] = llvm.mlir.constant(0 : index) : i64
// CHECK-DAG: [[ZERO2:%.*]] = llvm.mlir.constant(0 : i64) : i64
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO2]], 0, [[ZERO1]]] : (!llvm.ptr<[[ARRTYPE]], 4>, i64, i64) -> !llvm.ptr<i64, 4>
// CHECK-NEXT: [[GEP:%.*]] = llvm.getelementptr [[A0]][[[ZERO2]], [[ZERO2]], [[ZERO1]]] : (!llvm.ptr<4>, i64, i64, i64) -> !llvm.ptr<4>, i64

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