Skip to content

Commit 704451a

Browse files
victor-edsvictor-eds
authored
[SYCL-MLIR] Change sycl.constructor assembly format (#8274)
Define new assembly format closer to a func.call: `$TypeName ( $Args ) attr-dict : ( type($Args) )` Signed-off-by: Victor Perez <[email protected]>
1 parent 8ccbd90 commit 704451a

File tree

10 files changed

+96
-81
lines changed

10 files changed

+96
-81
lines changed

‎mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -644,7 +644,7 @@ def SYCLConstructorOp : SYCL_Op<"constructor", [CallOpInterface]> {
644644
}];
645645

646646
let assemblyFormat = [{
647-
`(` $Args `)` attr-dict `:` functional-type($Args, results)
647+
$TypeName `(` $Args `)` attr-dict `:` `(` type($Args) `)`
648648
}];
649649
}
650650

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -67,6 +67,9 @@ bool SYCLCastOp::areCastCompatible(TypeRange Inputs, TypeRange Outputs) {
6767
}
6868

6969
LogicalResult SYCLConstructorOp::verify() {
70+
if (getNumOperands() == 0)
71+
return emitOpError("A sycl::constructor must be passed the object to build "
72+
"as the first argument");
7073
auto MT = getOperand(0).getType().dyn_cast<MemRefType>();
7174
if (MT && isSYCLType(MT.getElementType()))
7275
return success();

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

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

‎mlir-sycl/test/Dialect/IR/SYCL/constructor.mlir

Lines changed: 0 additions & 35 deletions
This file was deleted.

‎mlir-sycl/test/Dialect/IR/SYCL/invalid.mlir

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,31 @@ func.func @test_cast_bad_shape(%arg: memref<1x!sycl_id_1_>) -> memref<2x!sycl_ar
2323

2424
// -----
2525

26+
func.func @test_not_arg_constructor() {
27+
// expected-error @+1 {{'sycl.constructor' op A sycl::constructor must be passed the object to build as the first argument}}
28+
"sycl.constructor"() {TypeName = @range, MangledFunctionName = @rangev} : () -> ()
29+
}
30+
31+
// -----
32+
33+
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
34+
35+
func.func @test_non_memref_arg_constructor(%range: !sycl_range_1_) {
36+
// expected-error @+1 {{'sycl.constructor' op The first argument of a sycl::constructor op has to be a MemRef to a SYCL type}}
37+
sycl.constructor @range(%range) {MangledFunctionName = @rangev} : (!sycl_range_1_)
38+
}
39+
40+
// -----
41+
42+
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
43+
44+
func.func @test_non_sycl_arg_constructor(%i: memref<1xi32>) {
45+
// expected-error @+1 {{'sycl.constructor' op The first argument of a sycl::constructor op has to be a MemRef to a SYCL type}}
46+
sycl.constructor @range(%i) {MangledFunctionName = @rangev} : (memref<1xi32>)
47+
}
48+
49+
// -----
50+
2651
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
2752

2853
func.func @test_num_work_items(%i: i32) -> !sycl_range_1_ {

‎mlir-sycl/test/Dialect/IR/SYCL/ops.mlir

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -60,6 +60,28 @@ func.func @test_i32_call_with_two_args(%arg_0: i32, %arg_1: i64) {
6060
return
6161
}
6262

63+
// CHECK-LABEL: test_constructor
64+
func.func @test_constructor(%id: memref<1x!sycl_id_1_>) {
65+
sycl.constructor @id(%id) {MangledFunctionName = @idv} : (memref<1x!sycl_id_1_>)
66+
return
67+
}
68+
69+
// CHECK-LABEL: test_constructor_with_arg
70+
func.func @test_constructor_with_arg(%range: memref<1x!sycl_range_1_>,
71+
%arg_0: i32) {
72+
sycl.constructor @range(%range, %arg_0) {MangledFunctionName = @rangei} : (memref<1x!sycl_range_1_>, i32)
73+
return
74+
}
75+
76+
// CHECK-LABEL: test_constructor_with_args
77+
func.func @test_constructor_with_args(%range: memref<1x!sycl_range_1_>,
78+
%arg_0: i32,
79+
%arg_1: i32,
80+
%arg_2: i32) {
81+
sycl.constructor @range(%range, %arg_0, %arg_1, %arg_2) {MangledFunctionName = @rangeiii} : (memref<1x!sycl_range_1_>, i32, i32, i32)
82+
return
83+
}
84+
6385
// CHECK-LABEL: test_num_work_items
6486
func.func @test_num_work_items() -> !sycl_range_1_ {
6587
%0 = sycl.num_work_items() : () -> !sycl_range_1_

‎polygeist/tools/cgeist/Test/Verification/sycl/constructors.cpp

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -109,7 +109,7 @@ extern "C" SYCL_EXTERNAL void cons_0(sycl::id<1> i, sycl::range<1> r) {
109109
// CHECK-NEXT: %3 = "polygeist.memref2pointer"(%alloca) : (memref<1x!sycl_id_2_>) -> !llvm.ptr<!sycl_id_2_>
110110
// CHECK-NEXT: %4 = llvm.addrspacecast %3 : !llvm.ptr<!sycl_id_2_> to !llvm.ptr<!sycl_id_2_, 4>
111111
// CHECK-NEXT: %5 = "polygeist.pointer2memref"(%4) : (!llvm.ptr<!sycl_id_2_, 4>) -> memref<?x!sycl_id_2_, 4>
112-
// CHECK-NEXT: sycl.constructor(%5) {MangledFunctionName = @_ZN4sycl3_V12idILi2EEC1Ev, TypeName = @id} : (memref<?x!sycl_id_2_, 4>) -> ()
112+
// CHECK-NEXT: sycl.constructor @id(%5) {MangledFunctionName = @_ZN4sycl3_V12idILi2EEC1Ev} : (memref<?x!sycl_id_2_, 4>)
113113

114114
// CHECK-LLVM-LABEL: define spir_func void @cons_1()
115115
// CHECK-LLVM-SAME: #[[FUNCATTRS]]
@@ -129,7 +129,7 @@ extern "C" SYCL_EXTERNAL void cons_1() {
129129
// CHECK-NEXT: %0 = "polygeist.memref2pointer"(%alloca) : (memref<1x!sycl_id_2_>) -> !llvm.ptr<!sycl_id_2_>
130130
// CHECK-NEXT: %1 = llvm.addrspacecast %0 : !llvm.ptr<!sycl_id_2_> to !llvm.ptr<!sycl_id_2_, 4>
131131
// CHECK-NEXT: %2 = "polygeist.pointer2memref"(%1) : (!llvm.ptr<!sycl_id_2_, 4>) -> memref<?x!sycl_id_2_, 4>
132-
// CHECK-NEXT: sycl.constructor(%2, %arg0, %arg1) {MangledFunctionName = @_ZN4sycl3_V12idILi2EEC1ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeEm, TypeName = @id} : (memref<?x!sycl_id_2_, 4>, i64, i64) -> ()
132+
// CHECK-NEXT: sycl.constructor @id(%2, %arg0, %arg1) {MangledFunctionName = @_ZN4sycl3_V12idILi2EEC1ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeEm} : (memref<?x!sycl_id_2_, 4>, i64, i64)
133133

134134
// CHECK-LLVM-LABEL: define spir_func void @cons_2(i64 noundef %0, i64 noundef %1)
135135
// CHECK-LLVM-SAME: #[[FUNCATTRS]]
@@ -151,7 +151,7 @@ extern "C" SYCL_EXTERNAL void cons_2(size_t a, size_t b) {
151151
// CHECK-NEXT: %3 = "polygeist.memref2pointer"(%arg0) : (memref<?x![[ITEM]]>) -> !llvm.ptr<![[ITEM]]>
152152
// CHECK-NEXT: %4 = llvm.addrspacecast %3 : !llvm.ptr<![[ITEM]]> to !llvm.ptr<![[ITEM]], 4>
153153
// CHECK-NEXT: %5 = "polygeist.pointer2memref"(%4) : (!llvm.ptr<![[ITEM]], 4>) -> memref<?x![[ITEM]], 4>
154-
// CHECK-NEXT: sycl.constructor(%2, %5) {MangledFunctionName = @_ZN4sycl3_V12idILi2EEC1ILi2ELb1EEERNSt9enable_ifIXeqT_Li2EEKNS0_4itemILi2EXT0_EEEE4typeE, TypeName = @id} : (memref<?x!sycl_id_2_, 4>, memref<?x![[ITEM]], 4>) -> ()
154+
// CHECK-NEXT: sycl.constructor @id(%2, %5) {MangledFunctionName = @_ZN4sycl3_V12idILi2EEC1ILi2ELb1EEERNSt9enable_ifIXeqT_Li2EEKNS0_4itemILi2EXT0_EEEE4typeE} : (memref<?x!sycl_id_2_, 4>, memref<?x![[ITEM]], 4>)
155155

156156
// CHECK-LLVM: define spir_func void @cons_3([[ITEM_TYPE:%"class.sycl::_V1::item.2.true"]]* noundef byval(%"class.sycl::_V1::item.2.true") align 8 [[ARG0:%.*]]) #[[FUNCATTRS]]
157157
// CHECK-LLVM-DAG: [[ID:%.*]] = alloca [[ID_TYPE:%"class.sycl::_V1::id.2"]]
@@ -172,7 +172,7 @@ extern "C" SYCL_EXTERNAL void cons_3(sycl::item<2, true> val) {
172172
// CHECK-NEXT: %3 = "polygeist.memref2pointer"(%arg0) : (memref<?x!sycl_id_2_>) -> !llvm.ptr<!sycl_id_2_>
173173
// CHECK-NEXT: %4 = llvm.addrspacecast %3 : !llvm.ptr<!sycl_id_2_> to !llvm.ptr<!sycl_id_2_, 4>
174174
// CHECK-NEXT: %5 = "polygeist.pointer2memref"(%4) : (!llvm.ptr<!sycl_id_2_, 4>) -> memref<?x!sycl_id_2_, 4>
175-
// CHECK-NEXT: sycl.constructor(%2, %5) {MangledFunctionName = @_ZN4sycl3_V12idILi2EEC1ERKS2_, TypeName = @id} : (memref<?x!sycl_id_2_, 4>, memref<?x!sycl_id_2_, 4>) -> ()
175+
// CHECK-NEXT: sycl.constructor @id(%2, %5) {MangledFunctionName = @_ZN4sycl3_V12idILi2EEC1ERKS2_} : (memref<?x!sycl_id_2_, 4>, memref<?x!sycl_id_2_, 4>)
176176

177177
// CHECK-LLVM: define spir_func void @cons_4([[ID_TYPE:%"class.sycl::_V1::id.2"]]* noundef byval(%"class.sycl::_V1::id.2") align 8 [[ARG0:%.*]]) #[[FUNCATTRS]]
178178
// CHECK-LLVM-DAG: [[ID:%.*]] = alloca [[ID_TYPE]]
@@ -187,7 +187,7 @@ extern "C" SYCL_EXTERNAL void cons_4(sycl::id<2> val) {
187187
// CHECK-LABEL: func.func @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1025ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(
188188
// CHECK-SAME: {{.*}}) attributes {[[SPIR_FUNCCC]], [[LINKONCE]], {{.*}}}
189189
// CHECK: [[I:%.*]] = "polygeist.subindex"(%arg0, %c0) : (memref<?x!sycl_accessor_1_i32_w_gb, 4>, index) -> memref<?x!sycl_accessor_impl_device_1_, 4>
190-
// CHECK: sycl.constructor([[I]], {{%.*}}, {{%.*}}, {{%.*}}) {MangledFunctionName = @_ZN4sycl3_V16detail18AccessorImplDeviceILi1EEC1ENS0_2idILi1EEENS0_5rangeILi1EEES7_, TypeName = @AccessorImplDevice} : (memref<?x!sycl_accessor_impl_device_1_, 4>, memref<?x!sycl_id_1_>, memref<?x!sycl_range_1_>, memref<?x!sycl_range_1_>) -> ()
190+
// CHECK: sycl.constructor @AccessorImplDevice([[I]], {{%.*}}, {{%.*}}, {{%.*}}) {MangledFunctionName = @_ZN4sycl3_V16detail18AccessorImplDeviceILi1EEC1ENS0_2idILi1EEENS0_5rangeILi1EEES7_} : (memref<?x!sycl_accessor_impl_device_1_, 4>, memref<?x!sycl_id_1_>, memref<?x!sycl_range_1_>, memref<?x!sycl_range_1_>)
191191

192192
// CHECK-LLVM-LABEL: define spir_func void @cons_5()
193193
// CHECK-LLVM-SAME: #[[FUNCATTRS]]
@@ -201,7 +201,7 @@ extern "C" SYCL_EXTERNAL void cons_5() {
201201

202202
// CHECK-LABEL: func.func @cons_6(
203203
// CHECK-SAME: %{{.*}}: i32
204-
// CHECK: sycl.constructor({{.*}}, {{.*}}) {MangledFunctionName = @[[VEC_SPLAT_CTR:.*]], TypeName = @vec} : (memref<?x!sycl_vec_i32_8_, 4>, memref<?xi32, 4>) -> ()
204+
// CHECK: sycl.constructor @vec({{.*}}, {{.*}}) {MangledFunctionName = @[[VEC_SPLAT_CTR:.*]]} : (memref<?x!sycl_vec_i32_8_, 4>, memref<?xi32, 4>)
205205
// CHECK: func.func @[[VEC_SPLAT_CTR]](%{{.*}}: memref<?x!sycl_vec_i32_8_, 4> {{{.*}}}, %{{.*}}: memref<?xi32, 4> {{{.*}}}) attributes {[[SPIR_FUNCCC]], [[LINKONCE]], {{.*}}}
206206
// CHECK: vector.splat %{{.*}} : vector<8xi32>
207207

@@ -218,7 +218,7 @@ extern "C" SYCL_EXTERNAL void cons_6(int Arg) {
218218

219219
// CHECK-LABEL: func.func @cons_7(
220220
// CHECK-SAME: %[[ARG0:.*]]: f32 {{{.*}}}, %[[ARG1:.*]]: f32 {{{.*}}}, %[[ARG2:.*]]: f32 {{{.*}}}, %[[ARG3:.*]]: f32 {{{.*}}})
221-
// CHECK: sycl.constructor(%{{.*}}, %[[ARG0]], %[[ARG1]], %[[ARG2]], %[[ARG3]]) {MangledFunctionName = @[[VEC_INITLIST_CTR:.*]], TypeName = @vec} : (memref<?x!sycl_vec_f32_4_, 4>, f32, f32, f32, f32) -> ()
221+
// CHECK: sycl.constructor @vec(%{{.*}}, %[[ARG0]], %[[ARG1]], %[[ARG2]], %[[ARG3]]) {MangledFunctionName = @[[VEC_INITLIST_CTR:.*]]} : (memref<?x!sycl_vec_f32_4_, 4>, f32, f32, f32, f32)
222222
// CHECK: func.func @[[VEC_INITLIST_CTR]](%{{.*}}: memref<?x!sycl_vec_f32_4_, 4> {{{.*}}}, %{{.*}}: f32 {{{.*}}}, %{{.*}}: f32 {{{.*}}}, %{{.*}}: f32 {{{.*}}}, %{{.*}}: f32 {{{.*}}}) attributes {[[SPIR_FUNCCC]], [[LINKONCE]], {{.*}}}
223223

224224
// CHECK-LLVM-LABEL: define spir_func void @cons_7(
@@ -231,7 +231,7 @@ extern "C" SYCL_EXTERNAL void cons_7(float A, float B, float C, float D) {
231231

232232
// CHECK-LABEL: func.func @cons_8(
233233
// CHECK-SAME: %[[ARG0:.*]]: memref<?x!sycl_vec_f64_16_, 4> {{{.*}}})
234-
// CHECK: sycl.constructor(%{{.*}}, %[[ARG0]]) {MangledFunctionName = @[[VEC_COPY_CTR:.*]], TypeName = @vec} : (memref<?x!sycl_vec_f64_16_, 4>, memref<?x!sycl_vec_f64_16_, 4>) -> ()
234+
// CHECK: sycl.constructor @vec(%{{.*}}, %[[ARG0]]) {MangledFunctionName = @[[VEC_COPY_CTR:.*]]} : (memref<?x!sycl_vec_f64_16_, 4>, memref<?x!sycl_vec_f64_16_, 4>)
235235
// CHECK: func.func @[[VEC_COPY_CTR]](%{{.*}}: memref<?x!sycl_vec_f64_16_, 4> {{{.*}}}, %{{.*}}: memref<?x!sycl_vec_f64_16_, 4> {{{.*}}}) attributes {[[SPIR_FUNCCC]], [[LINKONCE]], {{.*}}}
236236

237237
// CHECK-LLVM-LABEL: define spir_func void @cons_8(
@@ -244,7 +244,7 @@ extern "C" SYCL_EXTERNAL void cons_8(const sycl::vec<sycl::cl_double, 16> &Other
244244

245245
// CHECK-LABEL: func.func @cons_9(
246246
// CHECK-SAME: %[[ARG0:.*]]: vector<3xi8>
247-
// CHECK: sycl.constructor(%{{.*}}, %[[ARG0]]) {MangledFunctionName = @[[VEC_NATIVE_CTR:.*]], TypeName = @vec} : (memref<?x!sycl_vec_i8_3_, 4>, vector<3xi8>) -> ()
247+
// CHECK: sycl.constructor @vec(%{{.*}}, %[[ARG0]]) {MangledFunctionName = @[[VEC_NATIVE_CTR:.*]]} : (memref<?x!sycl_vec_i8_3_, 4>, vector<3xi8>)
248248
// CHECK: func.func @[[VEC_NATIVE_CTR]](%{{.*}}: memref<?x!sycl_vec_i8_3_, 4> {{{.*}}}, %{{.*}}: vector<3xi8> {{{.*}}}) attributes {[[SPIR_FUNCCC]], [[LINKONCE]], {{.*}}}
249249

250250
// CHECK-LLVM-LABEL: define spir_func void @cons_9(
@@ -257,7 +257,7 @@ extern "C" SYCL_EXTERNAL void cons_9(const sycl::vec<sycl::cl_char, 3>::vector_t
257257

258258
// CHECK-LABEL: func.func @cons_10(
259259
// CHECK-SAME: %[[ARG0:.*]]: memref<?x!sycl_vec_i64_8_, 4> {{{.*}}}, %[[ARG1:.*]]: memref<?x!sycl_vec_i64_4_, 4> {{{.*}}}, %[[ARG2:.*]]: memref<?x!sycl_vec_i64_2_, 4> {{{.*}}}, %{{.*}}: i64 {{{.*}}}, %{{.*}}: i64 {{{.*}}}) attributes {[[SPIR_FUNCCC]], [[LINKEXTERNAL]], {{.*}}}
260-
// CHECK: sycl.constructor(%3, %[[ARG0]], %[[ARG1]], %[[ARG2]], %6, %9) {MangledFunctionName = @[[VEC_INITLIST_VEC_CTR:.*]], TypeName = @vec} : (memref<?x!sycl_vec_i64_16_, 4>, memref<?x!sycl_vec_i64_8_, 4>, memref<?x!sycl_vec_i64_4_, 4>, memref<?x!sycl_vec_i64_2_, 4>, memref<?xi64, 4>, memref<?xi64, 4>) -> ()
260+
// CHECK: sycl.constructor @vec(%3, %[[ARG0]], %[[ARG1]], %[[ARG2]], %6, %9) {MangledFunctionName = @[[VEC_INITLIST_VEC_CTR:.*]]} : (memref<?x!sycl_vec_i64_16_, 4>, memref<?x!sycl_vec_i64_8_, 4>, memref<?x!sycl_vec_i64_4_, 4>, memref<?x!sycl_vec_i64_2_, 4>, memref<?xi64, 4>, memref<?xi64, 4>)
261261
// CHECK: func.func @[[VEC_INITLIST_VEC_CTR]](%{{.*}}: memref<?x!sycl_vec_i64_16_, 4> {{{.*}}}, %{{.*}}: memref<?x!sycl_vec_i64_8_, 4> {{{.*}}}, %{{.*}}: memref<?x!sycl_vec_i64_4_, 4> {{{.*}}}, %{{.*}}: memref<?x!sycl_vec_i64_2_, 4> {{{.*}}}, %{{.*}}: memref<?xi64, 4> {{{.*}}}, %{{.*}}: memref<?xi64, 4> {{{.*}}}) attributes {[[SPIR_FUNCCC]], [[LINKONCE]], {{.*}}}
262262

263263
// CHECK-LLVM-LABEL: define spir_func void @cons_10(

‎polygeist/tools/cgeist/Test/Verification/sycl/functions.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -155,7 +155,7 @@ SYCL_EXTERNAL void range_size(sycl::range<2> r) {
155155
// CHECK-MLIR-NEXT: %[[VAL_4:.*]] = "polygeist.memref2pointer"(%[[VAL_2]]) : (memref<1x!sycl_range_2_>) -> !llvm.ptr<!sycl_range_2_>
156156
// CHECK-MLIR-NEXT: %[[VAL_5:.*]] = llvm.addrspacecast %[[VAL_4]] : !llvm.ptr<!sycl_range_2_> to !llvm.ptr<!sycl_range_2_, 4>
157157
// CHECK-MLIR-NEXT: %[[VAL_6:.*]] = "polygeist.pointer2memref"(%[[VAL_5]]) : (!llvm.ptr<!sycl_range_2_, 4>) -> memref<?x!sycl_range_2_, 4>
158-
// CHECK-MLIR-NEXT: sycl.constructor(%[[VAL_6]], %[[VAL_3]]) {MangledFunctionName = @_ZN4sycl3_V15rangeILi2EEC1ERKS2_, TypeName = @range} : (memref<?x!sycl_range_2_, 4>, memref<?x!sycl_range_2_, 4>) -> ()
158+
// CHECK-MLIR-NEXT: sycl.constructor @range(%[[VAL_6]], %[[VAL_3]]) {MangledFunctionName = @_ZN4sycl3_V15rangeILi2EEC1ERKS2_} : (memref<?x!sycl_range_2_, 4>, memref<?x!sycl_range_2_, 4>)
159159
// CHECK-MLIR-NEXT: %[[VAL_7:.*]] = affine.load %[[VAL_2]][0] : memref<1x!sycl_range_2_>
160160
// CHECK-MLIR-NEXT: return %[[VAL_7]] : !sycl_range_2_
161161
// CHECK-MLIR-NEXT: }

‎polygeist/tools/cgeist/Test/Verification/sycl/issue7835.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@
1010
// CHECK-MLIR-SAME: %arg1: !llvm.ptr<!llvm.struct<(memref<?xi32, 1>)>> {llvm.align = 8 : i64, llvm.byval = !llvm.struct<(memref<?xi32, 1>)>, llvm.noundef})
1111
// CHECK-MLIR-SAME: kernel attributes {llvm.cconv = #llvm.cconv<spir_kernelcc>, llvm.linkage = #llvm.linkage<weak_odr>
1212

13-
// CHECK-MLIR: sycl.constructor(%4, %7) {MangledFunctionName = @_ZN4sycl3_V15rangeILi1EEC1ERKS2_, TypeName = @range} : (memref<?x!sycl_range_1_, 4>, memref<?x!sycl_range_1_, 4>) -> ()
13+
// CHECK-MLIR: sycl.constructor @range(%4, %7) {MangledFunctionName = @_ZN4sycl3_V15rangeILi1EEC1ERKS2_} : (memref<?x!sycl_range_1_, 4>, memref<?x!sycl_range_1_, 4>)
1414
// CHECK-MLIR: %8 = affine.load %alloca_0[0] : memref<1x!sycl_range_1_>
1515
// CHECK-MLIR: affine.store %8, %1[0] : memref<?x!sycl_range_1_>
1616
// CHECK-MLIR: %9 = "polygeist.subindex"(%cast_2, %c1) : (memref<?x!llvm.struct<(!sycl_range_1_, !llvm.struct<(memref<?xi32, 4>)>)>>, index) -> memref<?x!llvm.struct<(memref<?xi32, 4>)>>

‎polygeist/tools/cgeist/Test/Verification/sycl/sycl_extern_func.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@
1212
// CHECK-MLIR: call @cons_5() : () -> ()
1313

1414
// CHECK-MLIR-LABEL: func.func @cons_5() attributes {llvm.cconv = #llvm.cconv<spir_funccc>, llvm.linkage = #llvm.linkage<external>
15-
// CHECK-MLIR: sycl.constructor(%{{.*}})
15+
// CHECK-MLIR: sycl.constructor{{.*}}(%{{.*}})
1616
// CHECK-MLIR-NEXT: return
1717

1818
// CHECK-LLVM-LABEL: define spir_func void @cons_5()

0 commit comments

Comments
 (0)