Skip to content

Commit 889506d

Browse files
victor-edsvictor-eds
authored
[SYCL-MLIR] Change sycl.call assembly format (#8273)
Make format more similar to the func.call format: `$FunctionName ( $Args ) attr-dict : functional-type($Args, results)` Signed-off-by: Victor Perez <[email protected]>
1 parent 0b07661 commit 889506d

File tree

9 files changed

+71
-41
lines changed

9 files changed

+71
-41
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
@@ -724,7 +724,7 @@ def SYCLCallOp : SYCL_Op<"call", [CallOpInterface]> {
724724
}];
725725

726726
let assemblyFormat = [{
727-
`(` $Args `)` attr-dict `:` functional-type($Args, results)
727+
$FunctionName `(` $Args `)` attr-dict `:` functional-type($Args, results)
728728
}];
729729
}
730730

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@
1212
func.func private @foo() -> (i32)
1313

1414
func.func @test() -> (i32) {
15-
%0 = sycl.call() {FunctionName = @foo, MangledFunctionName = @foo, TypeName = @accessor} : () -> i32
15+
%0 = sycl.call @foo() {MangledFunctionName = @foo, TypeName = @accessor} : () -> i32
1616
return %0 : i32
1717
}
1818

@@ -31,7 +31,7 @@ func.func private @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6target
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_) {
3333
// CHECK: llvm.call @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE({{.*}}) : ([[ARG_TYPES]]) -> ()
34-
sycl.call(%arg0, %arg1, %arg2, %arg3, %arg4) {FunctionName = @__init, 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_) -> ()
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_) -> ()
3535
return
3636
}
3737

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

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,36 @@ func.func @test_cast_accessor(%arg: memref<1x!sycl_accessor_1_i32_w_gb>) -> memr
3030
return %0 : memref<1x!sycl.accessor_common>
3131
}
3232

33+
// CHECK-LABEL: test_void_call
34+
func.func @test_void_call() {
35+
sycl.call @foo() {MangledFunctionName = @foov, TypeName = @A} : () -> ()
36+
return
37+
}
38+
39+
// CHECK-LABEL: test_void_call_with_arg
40+
func.func @test_void_call_with_arg(%arg_0: i32) {
41+
sycl.call @foo(%arg_0) {MangledFunctionName = @fooi, TypeName = @A} : (i32) -> ()
42+
return
43+
}
44+
45+
// CHECK-LABEL: test_i32_call
46+
func.func @test_i32_call() {
47+
%0 = sycl.call @bar() {MangledFunctionName = @barv, TypeName = @A} : () -> (i32)
48+
return
49+
}
50+
51+
// CHECK-LABEL: test_i32_call_with_arg
52+
func.func @test_i32_call_with_arg(%arg_0: i32) {
53+
%0 = sycl.call @bar(%arg_0) {MangledFunctionName = @bari, TypeName = @A} : (i32) -> (i32)
54+
return
55+
}
56+
57+
// CHECK-LABEL: test_i32_call_with_two_args
58+
func.func @test_i32_call_with_two_args(%arg_0: i32, %arg_1: i64) {
59+
%0 = sycl.call @bar(%arg_0, %arg_1) {MangledFunctionName = @baril, TypeName = @A} : (i32, i64) -> (i32)
60+
return
61+
}
62+
3363
// CHECK-LABEL: test_num_work_items
3464
func.func @test_num_work_items() -> !sycl_range_1_ {
3565
%0 = sycl.num_work_items() : () -> !sycl_range_1_

‎mlir-sycl/test/Transforms/inliner.mlir

Lines changed: 24 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -9,11 +9,11 @@
99

1010
// ALWAYS-INLINE-LABEL: func.func @caller() -> i32 {
1111
// ALWAYS-INLINE-NEXT: %c1_i32 = arith.constant 1 : i32
12-
// ALWAYS-INLINE-NEXT: %0 = sycl.call() {FunctionName = @inline_hint_callee_, MangledFunctionName = @inline_hint_callee, TypeName = @A} : () -> i32
13-
// ALWAYS-INLINE-NEXT: %1 = sycl.call() {FunctionName = @private_callee_, MangledFunctionName = @private_callee, TypeName = @A} : () -> i32
14-
// ALWAYS-INLINE-NEXT: %2 = sycl.call() {FunctionName = @small_callee_, MangledFunctionName = @small_callee, TypeName = @A} : () -> i32
15-
// ALWAYS-INLINE-NEXT: %3 = sycl.call() {FunctionName = @callee_, MangledFunctionName = @callee, TypeName = @A} : () -> i32
16-
// ALWAYS-INLINE-NEXT: %4 = sycl.call() {FunctionName = @gpu_func_callee_, MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
12+
// ALWAYS-INLINE-NEXT: %0 = sycl.call @inline_hint_callee_() {MangledFunctionName = @inline_hint_callee, TypeName = @A} : () -> i32
13+
// ALWAYS-INLINE-NEXT: %1 = sycl.call @private_callee_() {MangledFunctionName = @private_callee, TypeName = @A} : () -> i32
14+
// ALWAYS-INLINE-NEXT: %2 = sycl.call @small_callee_() {MangledFunctionName = @small_callee, TypeName = @A} : () -> i32
15+
// ALWAYS-INLINE-NEXT: %3 = sycl.call @callee_() {MangledFunctionName = @callee, TypeName = @A} : () -> i32
16+
// ALWAYS-INLINE-NEXT: %4 = sycl.call @gpu_func_callee_() {MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
1717
// ALWAYS-INLINE-NEXT: %5 = arith.addi %c1_i32, %0 : i32
1818
// ALWAYS-INLINE-NEXT: %6 = arith.addi %1, %2 : i32
1919
// ALWAYS-INLINE-NEXT: %7 = arith.addi %3, %4 : i32
@@ -40,8 +40,8 @@
4040
// INLINE-DAG: %c2_i32 = arith.constant 2 : i32
4141
// INLINE-DAG: %c3_i32 = arith.constant 3 : i32
4242
// INLINE-DAG: %c4_i32 = arith.constant 4 : i32
43-
// INLINE-NEXT: %0 = sycl.call() {FunctionName = @callee_, MangledFunctionName = @callee, TypeName = @A} : () -> i32
44-
// INLINE-NEXT: %1 = sycl.call() {FunctionName = @gpu_func_callee_, MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
43+
// INLINE-NEXT: %0 = sycl.call @callee_() {MangledFunctionName = @callee, TypeName = @A} : () -> i32
44+
// INLINE-NEXT: %1 = sycl.call @gpu_func_callee_() {MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
4545
// INLINE-NEXT: %2 = arith.addi %c1_i32, %c2_i32 : i32
4646
// INLINE-NEXT: %3 = arith.addi %c3_i32, %c4_i32 : i32
4747
// INLINE-NEXT: %4 = arith.addi %0, %1 : i32
@@ -60,12 +60,12 @@
6060
gpu.module @module {
6161

6262
func.func @caller() -> i32 {
63-
%res1 = sycl.call() {FunctionName = @"always_inline_callee_", MangledFunctionName = @always_inline_callee, TypeName = @A} : () -> i32
64-
%res2 = sycl.call() {FunctionName = @"inline_hint_callee_", MangledFunctionName = @inline_hint_callee, TypeName = @A} : () -> i32
65-
%res3 = sycl.call() {FunctionName = @"private_callee_", MangledFunctionName = @private_callee, TypeName = @A} : () -> i32
66-
%res4 = sycl.call() {FunctionName = @"small_callee_", MangledFunctionName = @small_callee, TypeName = @A} : () -> i32
67-
%res5 = sycl.call() {FunctionName = @"callee_", MangledFunctionName = @callee, TypeName = @A} : () -> i32
68-
%res6 = sycl.call() {FunctionName = @"gpu_func_callee_", MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
63+
%res1 = sycl.call @always_inline_callee_() {MangledFunctionName = @always_inline_callee, TypeName = @A} : () -> i32
64+
%res2 = sycl.call @inline_hint_callee_() {MangledFunctionName = @inline_hint_callee, TypeName = @A} : () -> i32
65+
%res3 = sycl.call @private_callee_() {MangledFunctionName = @private_callee, TypeName = @A} : () -> i32
66+
%res4 = sycl.call @small_callee_() {MangledFunctionName = @small_callee, TypeName = @A} : () -> i32
67+
%res5 = sycl.call @callee_() {MangledFunctionName = @callee, TypeName = @A} : () -> i32
68+
%res6 = sycl.call @gpu_func_callee_() {MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
6969
%add1 = arith.addi %res1, %res2 : i32
7070
%add2 = arith.addi %res3, %res4 : i32
7171
%add3 = arith.addi %res5, %res6 : i32
@@ -117,8 +117,8 @@ gpu.func @gpu_func_callee() -> i32 attributes {passthrough = ["alwaysinline"]} {
117117
// COM: Ensure a func.func can be inlined in a gpu.func caller iff the callee is 'alwaysinline'.
118118
// ALWAYS-INLINE-LABEL: gpu.func @caller() -> i32 {
119119
// ALWAYS-INLINE-NEXT: %c1_i32 = arith.constant 1 : i32
120-
// ALWAYS-INLINE-NEXT: %0 = sycl.call() {FunctionName = @callee_, MangledFunctionName = @callee, TypeName = @A} : () -> i32
121-
// ALWAYS-INLINE-NEXT: %1 = sycl.call() {FunctionName = @gpu_callee_, MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
120+
// ALWAYS-INLINE-NEXT: %0 = sycl.call @callee_() {MangledFunctionName = @callee, TypeName = @A} : () -> i32
121+
// ALWAYS-INLINE-NEXT: %1 = sycl.call @gpu_callee_() {MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
122122
// ALWAYS-INLINE-NEXT: %2 = arith.addi %c1_i32, %0 : i32
123123
// ALWAYS-INLINE-NEXT: %3 = arith.addi %2, %1 : i32
124124
// ALWAYS-INLINE-NEXT: gpu.return %3 : i32
@@ -127,8 +127,8 @@ gpu.func @gpu_func_callee() -> i32 attributes {passthrough = ["alwaysinline"]} {
127127
// COM: Ensure a func.func can be inlined in a gpu.func caller.
128128
// INLINE-LABEL: gpu.func @caller() -> i32 {
129129
// INLINE-NEXT: %c1_i32 = arith.constant 1 : i32
130-
// INLINE-NEXT: %0 = sycl.call() {FunctionName = @callee_, MangledFunctionName = @callee, TypeName = @A} : () -> i32
131-
// INLINE-NEXT: %1 = sycl.call() {FunctionName = @gpu_callee_, MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
130+
// INLINE-NEXT: %0 = sycl.call @callee_() {MangledFunctionName = @callee, TypeName = @A} : () -> i32
131+
// INLINE-NEXT: %1 = sycl.call @gpu_callee_() {MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
132132
// INLINE-NEXT: %2 = arith.addi %c1_i32, %0 : i32
133133
// INLINE-NEXT: %3 = arith.addi %2, %1 : i32
134134
// INLINE-NEXT: gpu.return %3 : i32
@@ -137,9 +137,9 @@ gpu.func @gpu_func_callee() -> i32 attributes {passthrough = ["alwaysinline"]} {
137137
gpu.module @module {
138138

139139
gpu.func @caller() -> i32 {
140-
%res1 = sycl.call() {FunctionName = @"inlinable_callee_", MangledFunctionName = @inlinable_callee, TypeName = @A} : () -> i32
141-
%res2 = sycl.call() {FunctionName = @"callee_", MangledFunctionName = @callee, TypeName = @A} : () -> i32
142-
%res3 = sycl.call() {FunctionName = @"gpu_callee_", MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
140+
%res1 = sycl.call @inlinable_callee_() {MangledFunctionName = @inlinable_callee, TypeName = @A} : () -> i32
141+
%res2 = sycl.call @callee_() {MangledFunctionName = @callee, TypeName = @A} : () -> i32
142+
%res3 = sycl.call @gpu_callee_() {MangledFunctionName = @gpu_func_callee, TypeName = @A} : () -> i32
143143
%res4 = arith.addi %res1, %res2 : i32
144144
%res5 = arith.addi %res4, %res3 : i32
145145
gpu.return %res5 : i32
@@ -176,7 +176,7 @@ gpu.func @gpu_func_callee() -> i32 attributes {passthrough = ["alwaysinline"]} {
176176
// INLINE-DAG: %c2_i32 = arith.constant 2 : i32
177177
// INLINE-DAG: %c1_i32_0 = arith.constant 1 : i32
178178
// INLINE-DAG: %c2_i32_1 = arith.constant 2 : i32
179-
// INLINE-DAG: %0 = sycl.call() {FunctionName = @callee_, MangledFunctionName = @callee, TypeName = @A} : () -> i32
179+
// INLINE-DAG: %0 = sycl.call @callee_() {MangledFunctionName = @callee, TypeName = @A} : () -> i32
180180
// INLINE-NEXT: %1 = arith.addi %c2_i32_1, %0 : i32
181181
// INLINE-NEXT: %2 = arith.addi %c1_i32_0, %1 : i32
182182
// INLINE-NEXT: %3 = arith.addi %c1_i32, %c2_i32 : i32
@@ -189,22 +189,22 @@ gpu.func @gpu_func_callee() -> i32 attributes {passthrough = ["alwaysinline"]} {
189189

190190
func.func private @inline_hint_callee() -> i32 attributes {passthrough = ["inlinehint"]} {
191191
%c_i32 = arith.constant 1 : i32
192-
%res1 = sycl.call() {FunctionName = @"private_callee_", MangledFunctionName = @private_callee, TypeName = @A} : () -> i32
192+
%res1 = sycl.call @private_callee_() {MangledFunctionName = @private_callee, TypeName = @A} : () -> i32
193193
%res2 = arith.addi %c_i32, %res1 : i32
194194
return %res2 : i32
195195
}
196196

197197
func.func private @private_callee() -> i32 {
198198
%c_i32 = arith.constant 2 : i32
199-
%res1 = sycl.call() {FunctionName = @"callee_", MangledFunctionName = @callee, TypeName = @A} : () -> i32
199+
%res1 = sycl.call @callee_() {MangledFunctionName = @callee, TypeName = @A} : () -> i32
200200
%res2 = arith.addi %c_i32, %res1 : i32
201201
return %res2 : i32
202202
}
203203

204204
func.func @callee() -> i32 {
205205
%c1 = arith.constant 1 : i32
206206
%c2 = arith.constant 2 : i32
207-
%res1 = sycl.call() {FunctionName = @"inline_hint_callee_", MangledFunctionName = @inline_hint_callee, TypeName = @A} : () -> i32
207+
%res1 = sycl.call @inline_hint_callee_() {MangledFunctionName = @inline_hint_callee, TypeName = @A} : () -> i32
208208
%add1 = arith.addi %c1, %c2 : i32
209209
%add2 = arith.addi %res1, %add1 : i32
210210
return %add2 : i32

‎mlir-sycl/test/Transforms/sycl-method-to-sycl-call.mlir

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@
2020
// CHECK-NEXT: %[[VAL_6:.*]] = "polygeist.pointer2memref"(%[[VAL_5]]) : (!llvm.ptr<!sycl_accessor_2_i32_rw_gb, 4>) -> memref<?x!sycl_accessor_2_i32_rw_gb, 4>
2121
// CHECK-NEXT: memref.store %[[VAL_1]], %[[ID_ALLOCA]]{{\[}}%[[VAL_2]]] : memref<1x!sycl_id_2_>
2222
// CHECK-NEXT: %[[ID_CAST:.*]] = memref.cast %[[ID_ALLOCA]] : memref<1x!sycl_id_2_> to memref<?x!sycl_id_2_>
23-
// CHECK-NEXT: %[[VAL_7:.*]] = sycl.call(%[[VAL_6]], %[[ID_CAST]]) {FunctionName = @"operator[]", MangledFunctionName = @_ZNK4sycl3_V18accessorIiLi2ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEixILi2EvEERiNS0_2idILi2EEE, TypeName = @accessor} : (memref<?x!sycl_accessor_2_i32_rw_gb, 4>, memref<?x!sycl_id_2_>) -> memref<?xi32, 4>
23+
// CHECK-NEXT: %[[VAL_7:.*]] = sycl.call @"operator[]"(%[[VAL_6]], %[[ID_CAST]]) {MangledFunctionName = @_ZNK4sycl3_V18accessorIiLi2ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEixILi2EvEERiNS0_2idILi2EEE, TypeName = @accessor} : (memref<?x!sycl_accessor_2_i32_rw_gb, 4>, memref<?x!sycl_id_2_>) -> memref<?xi32, 4>
2424
// CHECK-NEXT: return %[[VAL_7]] : memref<?xi32, 4>
2525
// CHECK-NEXT: }
2626

@@ -39,7 +39,7 @@ func.func @accessor_subscript_operator(%arg0: !sycl_accessor_2_i32_rw_gb, %arg1:
3939
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.addrspacecast %[[VAL_4]] : !llvm.ptr<!sycl_range_2_> to !llvm.ptr<!sycl_range_2_, 4>
4040
// CHECK-NEXT: %[[VAL_6:.*]] = "polygeist.pointer2memref"(%[[VAL_5]]) : (!llvm.ptr<!sycl_range_2_, 4>) -> memref<?x!sycl_range_2_, 4>
4141
// CHECK-NEXT: %[[VAL_7:.*]] = sycl.cast %[[VAL_6]] : memref<?x!sycl_range_2_, 4> to memref<?x!sycl_array_2_, 4>
42-
// CHECK-NEXT: %[[VAL_8:.*]] = sycl.call(%[[VAL_7]], %[[VAL_1]]) {FunctionName = @get, MangledFunctionName = @_ZNK4sycl3_V16detail5arrayILi2EE3getEi, TypeName = @array} : (memref<?x!sycl_array_2_, 4>, i32) -> i64
42+
// CHECK-NEXT: %[[VAL_8:.*]] = sycl.call @get(%[[VAL_7]], %[[VAL_1]]) {MangledFunctionName = @_ZNK4sycl3_V16detail5arrayILi2EE3getEi, TypeName = @array} : (memref<?x!sycl_array_2_, 4>, i32) -> i64
4343
// CHECK-NEXT: return %[[VAL_8]] : i64
4444
// CHECK-NEXT: }
4545

@@ -56,7 +56,7 @@ func.func @range_get(%arg0: !sycl_range_2_, %arg1: i32) -> i64 {
5656
// CHECK-NEXT: %[[VAL_3:.*]] = "polygeist.memref2pointer"(%[[VAL_2]]) : (memref<1x!sycl_range_2_>) -> !llvm.ptr<!sycl_range_2_>
5757
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.addrspacecast %[[VAL_3]] : !llvm.ptr<!sycl_range_2_> to !llvm.ptr<!sycl_range_2_, 4>
5858
// CHECK-NEXT: %[[VAL_5:.*]] = "polygeist.pointer2memref"(%[[VAL_4]]) : (!llvm.ptr<!sycl_range_2_, 4>) -> memref<?x!sycl_range_2_, 4>
59-
// CHECK-NEXT: %[[VAL_6:.*]] = sycl.call(%[[VAL_5]]) {FunctionName = @size, MangledFunctionName = @_ZNK4sycl3_V15rangeILi2EE4sizeEv, TypeName = @range} : (memref<?x!sycl_range_2_, 4>) -> i64
59+
// CHECK-NEXT: %[[VAL_6:.*]] = sycl.call @size(%[[VAL_5]]) {MangledFunctionName = @_ZNK4sycl3_V15rangeILi2EE4sizeEv, TypeName = @range} : (memref<?x!sycl_range_2_, 4>) -> i64
6060
// CHECK-NEXT: return %[[VAL_6]] : i64
6161
// CHECK-NEXT: }
6262

@@ -73,7 +73,7 @@ func.func @range_size(%arg0: !sycl_range_2_) -> i64 {
7373
// CHECK-NEXT: %[[VAL_3:.*]] = "polygeist.memref2pointer"(%[[VAL_2]]) : (memref<1x!sycl_item_1_>) -> !llvm.ptr<!sycl_item_1_>
7474
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.addrspacecast %[[VAL_3]] : !llvm.ptr<!sycl_item_1_> to !llvm.ptr<!sycl_item_1_, 4>
7575
// CHECK-NEXT: %[[VAL_5:.*]] = "polygeist.pointer2memref"(%[[VAL_4]]) : (!llvm.ptr<!sycl_item_1_, 4>) -> memref<?x!sycl_item_1_, 4>
76-
// CHECK-NEXT: %[[VAL_6:.*]] = sycl.call(%[[VAL_5]]) {FunctionName = @get_id, MangledFunctionName = @_ZNK4sycl3_V14itemILi1ELb1EE6get_idEv, TypeName = @item} : (memref<?x!sycl_item_1_, 4>) -> !sycl_id_1_
76+
// CHECK-NEXT: %[[VAL_6:.*]] = sycl.call @get_id(%[[VAL_5]]) {MangledFunctionName = @_ZNK4sycl3_V14itemILi1ELb1EE6get_idEv, TypeName = @item} : (memref<?x!sycl_item_1_, 4>) -> !sycl_id_1_
7777
// CHECK-NEXT: return %[[VAL_6]] : !sycl_id_1_
7878
// CHECK-NEXT: }
7979

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -655,7 +655,7 @@ SYCL_EXTERNAL void group_get_local_linear_range(sycl::group<1> group) {
655655
// CHECK-MLIR-NEXT: %0 = "polygeist.memref2pointer"(%arg0) : (memref<?x![[ITEM2]]>) -> !llvm.ptr<![[ITEM2]]>
656656
// CHECK-MLIR-NEXT: %1 = llvm.addrspacecast %0 : !llvm.ptr<![[ITEM2]]> to !llvm.ptr<![[ITEM2]], 4>
657657
// CHECK-MLIR-NEXT: %2 = "polygeist.pointer2memref"(%1) : (!llvm.ptr<![[ITEM2]], 4>) -> memref<?x![[ITEM2]], 4>
658-
// CHECK-MLIR-NEXT: %3 = sycl.call(%2, %2) {FunctionName = @"operator==", MangledFunctionName = @_ZNK4sycl3_V14itemILi2ELb1EEeqERKS2_, TypeName = @item} : (memref<?x![[ITEM2]], 4>, memref<?x![[ITEM2]], 4>) -> i1
658+
// CHECK-MLIR-NEXT: %3 = sycl.call @"operator=="(%2, %2) {MangledFunctionName = @_ZNK4sycl3_V14itemILi2ELb1EEeqERKS2_, TypeName = @item} : (memref<?x![[ITEM2]], 4>, memref<?x![[ITEM2]], 4>) -> i1
659659
// CHECK-MLIR-NEXT: return
660660
// CHECK-MLIR-NEXT: }
661661

@@ -678,7 +678,7 @@ SYCL_EXTERNAL void method_2(sycl::item<2, true> item) {
678678
// CHECK-MLIR-NEXT: %3 = "polygeist.memref2pointer"(%arg1) : (memref<?x!sycl_id_2_>) -> !llvm.ptr<!sycl_id_2_>
679679
// CHECK-MLIR-NEXT: %4 = llvm.addrspacecast %3 : !llvm.ptr<!sycl_id_2_> to !llvm.ptr<!sycl_id_2_, 4>
680680
// CHECK-MLIR-NEXT: %5 = "polygeist.pointer2memref"(%4) : (!llvm.ptr<!sycl_id_2_, 4>) -> memref<?x!sycl_id_2_, 4>
681-
// CHECK-MLIR-NEXT: %6 = sycl.call(%2, %5) {FunctionName = @"operator==", MangledFunctionName = @_ZNK4sycl3_V12idILi2EEeqERKS2_, TypeName = @id} : (memref<?x!sycl_id_2_, 4>, memref<?x!sycl_id_2_, 4>) -> i1
681+
// CHECK-MLIR-NEXT: %6 = sycl.call @"operator=="(%2, %5) {MangledFunctionName = @_ZNK4sycl3_V12idILi2EEeqERKS2_, TypeName = @id} : (memref<?x!sycl_id_2_, 4>, memref<?x!sycl_id_2_, 4>) -> i1
682682
// CHECK-MLIR-NEXT: return
683683
// CHECK-MLIR-NEXT: }
684684

@@ -702,7 +702,7 @@ SYCL_EXTERNAL void op_1(sycl::id<2> a, sycl::id<2> b) {
702702
// CHECK-MLIR-NEXT: %1 = "sycl.id.get"(%0, %c0_i32) {ArgumentTypes = [memref<?x!sycl_array_2_, 4>, i32], FunctionName = @get, MangledFunctionName = @_ZNK4sycl3_V16detail5arrayILi2EE3getEi, TypeName = @array} : (!sycl_id_2_, i32) -> i64
703703
// CHECK-MLIR-NEXT: %2 = "sycl.id.get"(%0, %c1_i32) {ArgumentTypes = [memref<?x!sycl_array_2_, 4>, i32], FunctionName = @get, MangledFunctionName = @_ZNK4sycl3_V16detail5arrayILi2EE3getEi, TypeName = @array} : (!sycl_id_2_, i32) -> i64
704704
// CHECK-MLIR-NEXT: %3 = arith.addi %1, %2 : i64
705-
// CHECK-MLIR-NEXT: %4 = sycl.call(%3) {FunctionName = @abs, MangledFunctionName = @_ZN4sycl3_V13absImEENSt9enable_ifIXsr6detail14is_ugenintegerIT_EE5valueES3_E4typeES3_} : (i64) -> i64
705+
// CHECK-MLIR-NEXT: %4 = sycl.call @abs(%3) {MangledFunctionName = @_ZN4sycl3_V13absImEENSt9enable_ifIXsr6detail14is_ugenintegerIT_EE5valueES3_E4typeES3_} : (i64) -> i64
706706
// CHECK-MLIR-NEXT: return
707707
// CHECK-MLIR-NEXT: }
708708

0 commit comments

Comments
 (0)