Skip to content

Commit 001a03c

Browse files
sommerlukassommerlukas
authored
[SYCL-MLIR] Opaque pointer support in SYCL-to-LLVM conversion (#8944)
Second step to add opaque pointer support for the SYCL-MLIR project: Make the lowering patterns in the SYCL-to-LLVM lowering compatible with opaque pointers. Emitting typed or opaque pointers is controlled by the `use-opaque-pointers` pass option, that was added. Partly resolves #8616. --------- Signed-off-by: Lukas Sommer <lukas.sommer@codeplay.com>
1 parent 00c786d commit 001a03c

22 files changed

+4198
-1161
lines changed

‎mlir-sycl/include/mlir/Conversion/SYCLPasses.td

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,11 @@ def ConvertSYCLToLLVM : Pass<"convert-sycl-to-llvm", "ModuleOp"> {
3030
];
3131
let options = [
3232
Option<"indexBitwidth", "index-bitwidth", "unsigned",
33-
/*default=kDeriveIndexBitwidthFromDataLayout*/"0",
34-
"Bitwidth of the index type, 0 to use size of machine word">
33+
/*default=kDeriveIndexBitwidthFromDataLayout*/"0",
34+
"Bitwidth of the index type, 0 to use size of machine word">,
35+
Option<"useOpaquePointers", "use-opaque-pointers", "bool",
36+
/*default=*/"false", "Generate LLVM IR using opaque pointers "
37+
"instead of typed pointers">,
3538
];
3639
}
3740

‎mlir-sycl/include/mlir/Conversion/SYCLToLLVM/DialectBuilder.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,8 @@ class LLVMBuilder : public DialectBuilder {
7272
public:
7373
LLVMBuilder(OpBuilder &b, Location loc) : DialectBuilder(b, loc) {}
7474

75-
LLVM::AllocaOp genAlloca(Type type, Value size, int64_t align) const;
75+
LLVM::AllocaOp genAlloca(Type type, Type elemType, Value size,
76+
int64_t align) const;
7677
LLVM::BitcastOp genBitcast(Type type, Value val) const;
7778
LLVM::ExtractValueOp genExtractValue(Type type, Value container,
7879
ArrayRef<int64_t> pos) const;

‎mlir-sycl/lib/Conversion/SYCLToLLVM/DialectBuilder.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -103,9 +103,9 @@ func::CallOp FuncBuilder::genCall(StringRef funcName, TypeRange resTypes,
103103
// LLVMBuilder
104104
//===----------------------------------------------------------------------===//
105105

106-
LLVM::AllocaOp LLVMBuilder::genAlloca(Type type, Value size,
106+
LLVM::AllocaOp LLVMBuilder::genAlloca(Type type, Type elemType, Value size,
107107
int64_t align) const {
108-
return create<LLVM::AllocaOp>(type, size, align);
108+
return create<LLVM::AllocaOp>(type, elemType, size, align);
109109
}
110110

111111
LLVM::BitcastOp LLVMBuilder::genBitcast(Type type, Value val) const {

‎mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLToLLVM.cpp

Lines changed: 266 additions & 138 deletions
Large diffs are not rendered by default.

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -146,6 +146,10 @@ LogicalResult SYCLAccessorSubscriptOp::verify() {
146146
.Case<MemRefType>(
147147
[&](auto Ty) { return VerifyElemType(Ty.getElementType()); })
148148
.Case<LLVM::LLVMPointerType>([&](auto Ty) {
149+
if (!Ty.getElementType()) {
150+
// With opaque pointers, there is no element type to inspect.
151+
return success();
152+
}
149153
const Type ElemType = Ty.getElementType();
150154
return (!isa<LLVM::LLVMStructType>(ElemType))
151155
? emitOpError(

‎mlir-sycl/test/Conversion/SYCLToLLVM/sycl-addrspacecast-to-llvm-typed-pointer.mlir

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=0' -verify-diagnostics %s | FileCheck %s
2+
3+
// CHECK-LABEL: llvm.func @PtrCastToGeneric(%arg0: !llvm.ptr<i32>) -> !llvm.ptr<i32, 4> {
4+
// CHECK-NEXT: %0 = llvm.addrspacecast %arg0 : !llvm.ptr<i32> to !llvm.ptr<i32, 4>
5+
// CHECK-NEXT: llvm.return %0 : !llvm.ptr<i32, 4>
6+
// CHECK-NEXT: }
7+
8+
func.func @PtrCastToGeneric(%arg0: memref<?xi32>) -> memref<?xi32, 4> {
9+
%0 = sycl.addrspacecast %arg0 : memref<?xi32> to memref<?xi32, 4>
10+
return %0 : memref<?xi32, 4>
11+
}
12+
13+
// -----
14+
15+
// CHECK-LABEL: llvm.func @GenericCastToPtr(%arg0: !llvm.ptr<i32, 4>) -> !llvm.ptr<i32> {
16+
// CHECK-NEXT: %0 = llvm.addrspacecast %arg0 : !llvm.ptr<i32, 4> to !llvm.ptr<i32>
17+
// CHECK-NEXT: llvm.return %0 : !llvm.ptr<i32>
18+
// CHECK-NEXT: }
19+
20+
func.func @GenericCastToPtr(%arg0: memref<?xi32, 4>) -> memref<?xi32> {
21+
%0 = sycl.addrspacecast %arg0 : memref<?xi32, 4> to memref<?xi32>
22+
return %0 : memref<?xi32>
23+
}

‎mlir-sycl/test/Conversion/SYCLToLLVM/sycl-addrspacecast-to-llvm.mlir

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

3-
// CHECK-LABEL: llvm.func @PtrCastToGeneric(%arg0: !llvm.ptr<i32>) -> !llvm.ptr<i32, 4> {
4-
// CHECK-NEXT: %0 = llvm.addrspacecast %arg0 : !llvm.ptr<i32> to !llvm.ptr<i32, 4>
5-
// CHECK-NEXT: llvm.return %0 : !llvm.ptr<i32, 4>
3+
// CHECK-LABEL: llvm.func @PtrCastToGeneric(%arg0: !llvm.ptr) -> !llvm.ptr<4> {
4+
// CHECK-NEXT: %0 = llvm.addrspacecast %arg0 : !llvm.ptr to !llvm.ptr<4>
5+
// CHECK-NEXT: llvm.return %0 : !llvm.ptr<4>
66
// CHECK-NEXT: }
77

88
func.func @PtrCastToGeneric(%arg0: memref<?xi32>) -> memref<?xi32, 4> {
@@ -12,9 +12,9 @@ func.func @PtrCastToGeneric(%arg0: memref<?xi32>) -> memref<?xi32, 4> {
1212

1313
// -----
1414

15-
// CHECK-LABEL: llvm.func @GenericCastToPtr(%arg0: !llvm.ptr<i32, 4>) -> !llvm.ptr<i32> {
16-
// CHECK-NEXT: %0 = llvm.addrspacecast %arg0 : !llvm.ptr<i32, 4> to !llvm.ptr<i32>
17-
// CHECK-NEXT: llvm.return %0 : !llvm.ptr<i32>
15+
// CHECK-LABEL: llvm.func @GenericCastToPtr(%arg0: !llvm.ptr<4>) -> !llvm.ptr {
16+
// CHECK-NEXT: %0 = llvm.addrspacecast %arg0 : !llvm.ptr<4> to !llvm.ptr
17+
// CHECK-NEXT: llvm.return %0 : !llvm.ptr
1818
// CHECK-NEXT: }
1919

2020
func.func @GenericCastToPtr(%arg0: memref<?xi32, 4>) -> memref<?xi32> {

‎mlir-sycl/test/Conversion/SYCLToLLVM/sycl-call-to-llvm-typed-pointer.mlir

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=0' -verify-diagnostics %s | FileCheck %s
2+
3+
//===-------------------------------------------------------------------------------------------------===//
4+
// sycl.call with non void return type
5+
//===-------------------------------------------------------------------------------------------------===//
6+
7+
// CHECK: llvm.func @foo() -> [[RET_TYPE:i32]]
8+
// CHECK: llvm.func @test() -> [[RET_TYPE]] {
9+
// CHECK-NEXT: %0 = llvm.call @foo() : () -> [[RET_TYPE]]
10+
// CHECK-NEXT: llvm.return %0 : [[RET_TYPE]]
11+
12+
func.func private @foo() -> (i32)
13+
14+
func.func @test() -> (i32) {
15+
%0 = sycl.call @foo() {MangledFunctionName = @foo, TypeName = @accessor} : () -> i32
16+
return %0 : i32
17+
}
18+
19+
// -----
20+
21+
//===-------------------------------------------------------------------------------------------------===//
22+
// Member functions for sycl::accessor
23+
//===-------------------------------------------------------------------------------------------------===//
24+
25+
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
26+
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
27+
!sycl_accessor_1_i32_rw_gb = !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>)>)>
28+
29+
// CHECK: llvm.func @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE([[ARG_TYPES:!llvm.ptr<struct<"class.sycl::_V1::accessor.1",.*]])
30+
func.func private @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE(memref<?x!sycl_accessor_1_i32_rw_gb>, memref<?xi32>, !sycl_range_1_, !sycl_range_1_, !sycl_id_1_)
31+
32+
func.func @accessorInit1(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>, %arg1: memref<?xi32>, %arg2: !sycl_range_1_, %arg3: !sycl_range_1_, %arg4: !sycl_id_1_) {
33+
// CHECK: llvm.call @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE({{.*}}) : ([[ARG_TYPES]]) -> ()
34+
sycl.call @__init(%arg0, %arg1, %arg2, %arg3, %arg4) {MangledFunctionName = @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE, TypeName = @accessor} : (memref<?x!sycl_accessor_1_i32_rw_gb>, memref<?xi32>, !sycl_range_1_, !sycl_range_1_, !sycl_id_1_) -> ()
35+
return
36+
}
37+
38+
// -----

‎mlir-sycl/test/Conversion/SYCLToLLVM/sycl-call-to-llvm.mlir

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm -verify-diagnostics %s | FileCheck %s
1+
// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=1' -verify-diagnostics %s | FileCheck %s
22

33
//===-------------------------------------------------------------------------------------------------===//
44
// sycl.call with non void return type
@@ -24,9 +24,9 @@ func.func @test() -> (i32) {
2424

2525
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
2626
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
27-
!sycl_accessor_1_i32_rw_gb = !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>)>)>
27+
!sycl_accessor_1_i32_rw_gb = !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<1>)>)>
2828

29-
// CHECK: llvm.func @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE([[ARG_TYPES:!llvm.ptr<struct<"class.sycl::_V1::accessor.1",.*]])
29+
// CHECK: llvm.func @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE([[ARG_TYPES:!llvm.ptr,.*]])
3030
func.func private @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE(memref<?x!sycl_accessor_1_i32_rw_gb>, memref<?xi32>, !sycl_range_1_, !sycl_range_1_, !sycl_id_1_)
3131

3232
func.func @accessorInit1(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>, %arg1: memref<?xi32>, %arg2: !sycl_range_1_, %arg3: !sycl_range_1_, %arg4: !sycl_id_1_) {

‎mlir-sycl/test/Conversion/SYCLToLLVM/sycl-cast-to-llvm-typed-pointer.mlir

Lines changed: 92 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,92 @@
1+
// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm='use-opaque-pointers=0' -verify-diagnostics %s | FileCheck %s
2+
3+
!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
4+
!sycl_range_1_ = !sycl.range<[1], (!sycl_array_1_)>
5+
func.func @cast_sycl_range_to_array(%arg0: memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_> {
6+
// CHECK-LABEL: llvm.func @cast_sycl_range_to_array(
7+
// CHECK-SAME: [[SRC:%.*]]: !llvm.ptr<[[RANGE1:.*]]>) -> !llvm.ptr<[[ARRAY1:.*]]>
8+
// CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[RANGE1]]> to !llvm.ptr<[[ARRAY1]]>
9+
// CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[ARRAY1]]>
10+
11+
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_range_1_>) -> memref<?x!sycl_array_1_>
12+
func.return %0 : memref<?x!sycl_array_1_>
13+
}
14+
15+
// -----
16+
17+
!sycl_array_1_ = !sycl.array<[1], (memref<1xi64>)>
18+
!sycl_id_1_ = !sycl.id<[1], (!sycl_array_1_)>
19+
func.func @cast_sycl_id_to_array(%arg0: memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_> {
20+
// CHECK-LABEL: llvm.func @cast_sycl_id_to_array(
21+
// CHECK-SAME: [[SRC:%.*]]: !llvm.ptr<[[ID1:.*]]>) -> !llvm.ptr<[[ARRAY1]]>
22+
// CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[ID1]]> to !llvm.ptr<[[ARRAY1]]>
23+
// CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[ARRAY1]]>
24+
25+
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_id_1_>) -> memref<?x!sycl_array_1_>
26+
func.return %0: memref<?x!sycl_array_1_>
27+
}
28+
29+
// -----
30+
31+
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
32+
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
33+
!sycl_accessor_1_i32_rw_gb = !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>)>)>
34+
func.func @cast_sycl_accessor_to_accessor_common(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl.accessor_common> {
35+
// CHECK-LABEL: llvm.func @cast_sycl_accessor_to_accessor_common(
36+
// CHECK-SAME: [[SRC:%.*]]: !llvm.ptr<[[ACC1:.*]]>) -> !llvm.ptr<[[COMMON:.*]]>
37+
// CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[ACC1]]> to !llvm.ptr<[[COMMON]]>
38+
// CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[COMMON]]>
39+
40+
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl.accessor_common>
41+
func.return %0: memref<?x!sycl.accessor_common>
42+
}
43+
44+
!sycl_LocalAccessorBaseDevice_1_ = !sycl.LocalAccessorBaseDevice<[1], (!sycl_range_1_, !sycl_range_1_, !sycl_id_1_)>
45+
!sycl_local_accessor_base_1_i32_rw = !sycl.local_accessor_base<[1, i32, read_write], (!sycl_LocalAccessorBaseDevice_1_, memref<?xi32, 3>)>
46+
func.func @cast_sycl_accessor_to_local_accessor_base(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl_local_accessor_base_1_i32_rw> {
47+
// CHECK-LABEL: llvm.func @cast_sycl_accessor_to_local_accessor_base(
48+
// CHECK-SAME: [[SRC:%.*]]: !llvm.ptr<[[ACC1]]>) -> !llvm.ptr<[[LOCALBASE:.*]]>
49+
// CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[ACC1]]> to !llvm.ptr<[[LOCALBASE]]>
50+
// CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[LOCALBASE]]>
51+
52+
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl_local_accessor_base_1_i32_rw>
53+
func.return %0: memref<?x!sycl_local_accessor_base_1_i32_rw>
54+
}
55+
56+
func.func @cast_sycl_accessor_to_owner_less_base(%arg0: memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl.owner_less_base> {
57+
// CHECK-LABEL: llvm.func @cast_sycl_accessor_to_owner_less_base(
58+
// CHECK-SAME: [[SRC:%.*]]: !llvm.ptr<[[ACC1]]>) -> !llvm.ptr<[[OWNERLESSBASE:.*]]>
59+
// CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[ACC1]]> to !llvm.ptr<[[OWNERLESSBASE]]>
60+
// CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[OWNERLESSBASE]]>
61+
62+
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_accessor_1_i32_rw_gb>) -> memref<?x!sycl.owner_less_base>
63+
func.return %0: memref<?x!sycl.owner_less_base>
64+
}
65+
66+
// -----
67+
68+
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
69+
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
70+
!sycl_LocalAccessorBaseDevice_1_ = !sycl.LocalAccessorBaseDevice<[1], (!sycl_range_1_, !sycl_range_1_, !sycl_id_1_)>
71+
!sycl_local_accessor_base_1_i32_rw = !sycl.local_accessor_base<[1, i32, read_write], (!sycl_LocalAccessorBaseDevice_1_, memref<?xi32, 3>)>
72+
func.func @cast_sycl_local_accessor_base_to_accessor_common(%arg0: memref<?x!sycl_local_accessor_base_1_i32_rw>) -> memref<?x!sycl.accessor_common> {
73+
// CHECK-LABEL: llvm.func @cast_sycl_local_accessor_base_to_accessor_common(
74+
// CHECK-SAME: [[SRC:%.*]]: !llvm.ptr<[[LAB1:.*]]>) -> !llvm.ptr<[[COMMON]]>
75+
// CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[LAB1]]> to !llvm.ptr<[[COMMON]]>
76+
// CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[COMMON]]
77+
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_local_accessor_base_1_i32_rw>) -> memref<?x!sycl.accessor_common>
78+
func.return %0: memref<?x!sycl.accessor_common>
79+
}
80+
81+
!sycl_local_accessor_1_i32_rw = !sycl.local_accessor<[1, i32], (!sycl_local_accessor_base_1_i32_rw)>
82+
func.func @cast_sycl_local_accessor_to_local_accessor_base(%arg0: memref<?x!sycl_local_accessor_1_i32_rw>) -> memref<?x!sycl_local_accessor_base_1_i32_rw> {
83+
// CHECK-LABEL: llvm.func @cast_sycl_local_accessor_to_local_accessor_base(
84+
// CHECK-SAME: [[SRC:%.*]]: !llvm.ptr<[[LA1:.*]]>) -> !llvm.ptr<[[LAB1]]>
85+
// CHECK-NEXT: [[RES:%.*]] = llvm.bitcast [[SRC]] : !llvm.ptr<[[LA1]]> to !llvm.ptr<[[LAB1]]>
86+
// CHECK-NEXT: llvm.return [[RES]] : !llvm.ptr<[[LAB1]]
87+
88+
%0 = "sycl.cast"(%arg0) : (memref<?x!sycl_local_accessor_1_i32_rw>) -> memref<?x!sycl_local_accessor_base_1_i32_rw>
89+
func.return %0: memref<?x!sycl_local_accessor_base_1_i32_rw>
90+
}
91+
92+

0 commit comments

Comments
 (0)