Skip to content

Commit 87e10ab

Browse files
whitneywhtsangetiotto
whitneywhtsang
authored and
etiotto
committed
Move sycl to llvm type converter lit test to mlir-sycl (#49)
`populateFuncToLLVMConversionPatterns` needs to be added to `ConvertSYCLToLLVMPass`, in order for the SYCL types in `func.func` to be converted. `mlir-sycl/test/Conversion/SYCLToLLVM/sycl-types-to-llvm.mlir` converts from SYCL MLIR Dialect to LLVM MLIR Dialect. `polygeist/tools/cgeist/Test/Verification/sycl/SYCLToLLVM/types.mlir` used to convert from SYCL MLIR Dialect to LLVM IR. `builtin.unrealized_conversion_cast` is no longer needed in `mlir-sycl/test/Conversion/SYCLToLLVM/func-ops-to-llvm.mlir` as SYCL types in `func.func` are now converted. Signed-off-by: Tsang, Whitney <[email protected]>
1 parent 1d867c2 commit 87e10ab

File tree

4 files changed

+52
-82
lines changed

4 files changed

+52
-82
lines changed

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,7 @@
1212

1313
#include "mlir/Conversion/SYCLToLLVM/SYCLToLLVMPass.h"
1414
#include "../PassDetail.h"
15+
#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h"
1516
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
1617
#include "mlir/Conversion/SYCLToLLVM/SYCLToLLVM.h"
1718
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
@@ -35,6 +36,7 @@ void ConvertSYCLToLLVMPass::runOnOperation() {
3536
RewritePatternSet patterns(context);
3637

3738
sycl::populateSYCLToLLVMConversionPatterns(converter, patterns);
39+
populateFuncToLLVMConversionPatterns(converter, patterns);
3840

3941
ConversionTarget target(*context);
4042
target.addIllegalDialect<sycl::SYCLDialect>();

‎mlir-sycl/test/Conversion/SYCLToLLVM/func-ops-to-llvm.mlir

Lines changed: 13 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,12 @@
1-
// RUN: sycl-mlir-opt -convert-sycl-to-llvm %s -split-input-file | FileCheck %s
1+
// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm -verify-diagnostics %s | FileCheck %s
22

33
//===-------------------------------------------------------------------------------------------------===//
44
// Constructors for sycl::id<n>::id()
55
//===-------------------------------------------------------------------------------------------------===//
66

77
// CHECK: llvm.func @_ZN2cl4sycl2idILi1EEC2Ev([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.1",.*]])
88
func.func @id1Ctor(%arg0: memref<?x!sycl.id<1>>) {
9-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_1_> to [[THIS_PTR_TYPE]]
10-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi1EEC2Ev(%0) : ([[THIS_PTR_TYPE]]) -> ()
9+
// CHECK: llvm.call @_ZN2cl4sycl2idILi1EEC2Ev({{.*}}) : ([[THIS_PTR_TYPE]]) -> ()
1110
sycl.constructor(%arg0) {Type = @id} : (memref<?x!sycl.id<1>>) -> ()
1211
return
1312
}
@@ -16,8 +15,7 @@ func.func @id1Ctor(%arg0: memref<?x!sycl.id<1>>) {
1615

1716
// CHECK: llvm.func @_ZN2cl4sycl2idILi2EEC2Ev([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.2",.*]])
1817
func.func @id2Ctor(%arg0: memref<?x!sycl.id<2>>) {
19-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_2_> to [[THIS_PTR_TYPE]]
20-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi2EEC2Ev(%0) : ([[THIS_PTR_TYPE]]) -> ()
18+
// CHECK: llvm.call @_ZN2cl4sycl2idILi2EEC2Ev({{.*}}) : ([[THIS_PTR_TYPE]]) -> ()
2119
sycl.constructor(%arg0) {Type = @id} : (memref<?x!sycl.id<2>>) -> ()
2220
return
2321
}
@@ -26,8 +24,7 @@ func.func @id2Ctor(%arg0: memref<?x!sycl.id<2>>) {
2624

2725
// CHECK: llvm.func @_ZN2cl4sycl2idILi3EEC2Ev([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.3",.*]])
2826
func.func @id3Ctor(%arg0: memref<?x!sycl.id<3>>) {
29-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_3_> to [[THIS_PTR_TYPE]]
30-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi3EEC2Ev(%0) : ([[THIS_PTR_TYPE]]) -> ()
27+
// CHECK: llvm.call @_ZN2cl4sycl2idILi3EEC2Ev({{.*}}) : ([[THIS_PTR_TYPE]]) -> ()
3128
sycl.constructor(%arg0) {Type = @id} : (memref<?x!sycl.id<3>>) -> ()
3229
return
3330
}
@@ -40,8 +37,7 @@ func.func @id3Ctor(%arg0: memref<?x!sycl.id<3>>) {
4037

4138
// CHECK: llvm.func @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.1",.*]], i64)
4239
func.func @id1CtorSizeT(%arg0: memref<?x!sycl.id<1>>, %arg1: i64) {
43-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_1_> to [[THIS_PTR_TYPE]]
44-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%0, %arg1) : ([[THIS_PTR_TYPE]], i64) -> ()
40+
// CHECK: llvm.call @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE({{.*}}, %arg5) : ([[THIS_PTR_TYPE]], i64) -> ()
4541
sycl.constructor(%arg0, %arg1) {Type = @id} : (memref<?x!sycl.id<1>>, i64) -> ()
4642
return
4743
}
@@ -51,8 +47,7 @@ func.func @id1CtorSizeT(%arg0: memref<?x!sycl.id<1>>, %arg1: i64) {
5147

5248
// CHECK: llvm.func @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeE([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.2",.*]], i64)
5349
func.func @id2CtorSizeT(%arg0: memref<?x!sycl.id<2>>, %arg1: i64) {
54-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_2_> to [[THIS_PTR_TYPE]]
55-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeE(%0, %arg1) : ([[THIS_PTR_TYPE]], i64) -> ()
50+
// CHECK: llvm.call @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeE({{.*}}, %arg5) : ([[THIS_PTR_TYPE]], i64) -> ()
5651
sycl.constructor(%arg0, %arg1) {Type = @id} : (memref<?x!sycl.id<2>>, i64) -> ()
5752
return
5853
}
@@ -61,8 +56,7 @@ func.func @id2CtorSizeT(%arg0: memref<?x!sycl.id<2>>, %arg1: i64) {
6156

6257
// CHECK: llvm.func @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeE([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.3",.*]], i64)
6358
func.func @id3CtorSizeT(%arg0: memref<?x!sycl.id<3>>, %arg1: i64) {
64-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_3_> to [[THIS_PTR_TYPE]]
65-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeE(%0, %arg1) : ([[THIS_PTR_TYPE]], i64) -> ()
59+
// CHECK: llvm.call @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeE({{.*}}, %arg5) : ([[THIS_PTR_TYPE]], i64) -> ()
6660
sycl.constructor(%arg0, %arg1) {Type = @id} : (memref<?x!sycl.id<3>>, i64) -> ()
6761
return
6862
}
@@ -75,8 +69,7 @@ func.func @id3CtorSizeT(%arg0: memref<?x!sycl.id<3>>, %arg1: i64) {
7569

7670
// CHECK: llvm.func @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeEm([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.1",.*]], i64, i64)
7771
func.func @id1CtorRange(%arg0: memref<?x!sycl.id<1>>, %arg1: i64, %arg2: i64) {
78-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_1_> to [[THIS_PTR_TYPE]]
79-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeEm(%0, %arg1, %arg2) : ([[THIS_PTR_TYPE]], i64, i64) -> ()
72+
// CHECK: llvm.call @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeEm({{.*}}, %arg5, %arg6) : ([[THIS_PTR_TYPE]], i64, i64) -> ()
8073
sycl.constructor(%arg0, %arg1, %arg2) {Type = @id} : (memref<?x!sycl.id<1>>, i64, i64) -> ()
8174
return
8275
}
@@ -85,8 +78,7 @@ func.func @id1CtorRange(%arg0: memref<?x!sycl.id<1>>, %arg1: i64, %arg2: i64) {
8578

8679
// CHECK: llvm.func @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeEm([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.2",.*]], i64, i64)
8780
func.func @id2CtorRange(%arg0: memref<?x!sycl.id<2>>, %arg1: i64, %arg2: i64) {
88-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_2_> to [[THIS_PTR_TYPE]]
89-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeEm(%0, %arg1, %arg2) : ([[THIS_PTR_TYPE]], i64, i64) -> ()
81+
// CHECK: llvm.call @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeEm({{.*}}, %arg5, %arg6) : ([[THIS_PTR_TYPE]], i64, i64) -> ()
9082
sycl.constructor(%arg0, %arg1, %arg2) {Type = @id} : (memref<?x!sycl.id<2>>, i64, i64) -> ()
9183
return
9284
}
@@ -95,8 +87,7 @@ func.func @id2CtorRange(%arg0: memref<?x!sycl.id<2>>, %arg1: i64, %arg2: i64) {
9587

9688
// CHECK: llvm.func @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeEm([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.3",.*]], i64, i64)
9789
func.func @id3CtorRange(%arg0: memref<?x!sycl.id<3>>, %arg1: i64, %arg2: i64) {
98-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_3_> to [[THIS_PTR_TYPE]]
99-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeEm(%0, %arg1, %arg2) : ([[THIS_PTR_TYPE]], i64, i64) -> ()
90+
// CHECK: llvm.call @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeEm({{.*}}, %arg5, %arg6) : ([[THIS_PTR_TYPE]], i64, i64) -> ()
10091
sycl.constructor(%arg0, %arg1, %arg2) {Type = @id} : (memref<?x!sycl.id<3>>, i64, i64) -> ()
10192
return
10293
}
@@ -109,8 +100,7 @@ func.func @id3CtorRange(%arg0: memref<?x!sycl.id<3>>, %arg1: i64, %arg2: i64) {
109100

110101
// CHECK: llvm.func @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeEmm([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.1",.*]], i64, i64, i64)
111102
func.func @id1CtorItem(%arg0: memref<?x!sycl.id<1>>, %arg1: i64, %arg2: i64, %arg3: i64) {
112-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_1_> to [[THIS_PTR_TYPE]]
113-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeEmm(%0, %arg1, %arg2, %arg3) : ([[THIS_PTR_TYPE]], i64, i64, i64) -> ()
103+
// CHECK: llvm.call @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeEmm({{.*}}, %arg5, %arg6, %arg7) : ([[THIS_PTR_TYPE]], i64, i64, i64) -> ()
114104
sycl.constructor(%arg0, %arg1, %arg2, %arg3) {Type = @id} : (memref<?x!sycl.id<1>>, i64, i64, i64) -> ()
115105
return
116106
}
@@ -119,8 +109,7 @@ func.func @id1CtorItem(%arg0: memref<?x!sycl.id<1>>, %arg1: i64, %arg2: i64, %ar
119109

120110
// CHECK: llvm.func @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeEmm([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.2",.*]], i64, i64, i64)
121111
func.func @id2CtorItem(%arg0: memref<?x!sycl.id<2>>, %arg1: i64, %arg2: i64, %arg3: i64) {
122-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_2_> to [[THIS_PTR_TYPE]]
123-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeEmm(%0, %arg1, %arg2, %arg3) : ([[THIS_PTR_TYPE]], i64, i64, i64) -> ()
112+
// CHECK: llvm.call @_ZN2cl4sycl2idILi2EEC2ILi2EEENSt9enable_ifIXeqT_Li2EEmE4typeEmm({{.*}}, %arg5, %arg6, %arg7) : ([[THIS_PTR_TYPE]], i64, i64, i64) -> ()
124113
sycl.constructor(%arg0, %arg1, %arg2, %arg3) {Type = @id} : (memref<?x!sycl.id<2>>, i64, i64, i64) -> ()
125114
return
126115
}
@@ -129,8 +118,7 @@ func.func @id2CtorItem(%arg0: memref<?x!sycl.id<2>>, %arg1: i64, %arg2: i64, %ar
129118

130119
// CHECK: llvm.func @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeEmm([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::id.3",.*]], i64, i64, i64)
131120
func.func @id3CtorItem(%arg0: memref<?x!sycl.id<3>>, %arg1: i64, %arg2: i64, %arg3: i64) {
132-
// CHECK: %0 = builtin.unrealized_conversion_cast %arg0 : memref<?x!sycl_id_3_> to [[THIS_PTR_TYPE]]
133-
// CHECK-NEXT: llvm.call @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeEmm(%0, %arg1, %arg2, %arg3) : ([[THIS_PTR_TYPE]], i64, i64, i64) -> ()
121+
// CHECK: llvm.call @_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeEmm({{.*}}, %arg5, %arg6, %arg7) : ([[THIS_PTR_TYPE]], i64, i64, i64) -> ()
134122
sycl.constructor(%arg0, %arg1, %arg2, %arg3) {Type = @id} : (memref<?x!sycl.id<3>>, i64, i64, i64) -> ()
135123
return
136124
}

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

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,37 @@
1+
// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm -verify-diagnostics %s | FileCheck %s
2+
3+
// CHECK: llvm.func @test_array.1(%arg0: !llvm.[[ARRAY_1:struct<"class.cl::sycl::detail::array.*", \(array<1 x i64>\)>]])
4+
// CHECK: llvm.func @test_array.2(%arg0: !llvm.[[ARRAY_2:struct<"class.cl::sycl::detail::array.*", \(array<2 x i64>\)>]])
5+
// CHECK: llvm.func @test_id(%arg0: !llvm.[[ID_1:struct<"class.cl::sycl::id.*", \(]][[ARRAY_1]][[SUFFIX:\)>]], %arg1: !llvm.[[ID_1]][[ARRAY_1]][[SUFFIX]])
6+
// CHECK: llvm.func @test_range.1(%arg0: !llvm.[[RANGE_1:struct<"class.cl::sycl::range.*", \(]][[ARRAY_1]][[SUFFIX]])
7+
// CHECK: llvm.func @test_range.2(%arg0: !llvm.[[RANGE_2:struct<"class.cl::sycl::range.*", \(]][[ARRAY_2]][[SUFFIX]])
8+
// CHECK: llvm.func @test_accessorImplDevice(%arg0: !llvm.[[ACCESSORIMPLDEVICE_1:struct<"class.cl::sycl::detail::AccessorImplDevice.*", \(]][[ID_1]][[ARRAY_1]][[SUFFIX]], [[RANGE_1]][[ARRAY_1]][[SUFFIX]], [[RANGE_1]][[ARRAY_1]][[SUFFIX]][[SUFFIX]])
9+
// CHECK: llvm.func @test_accessor.1(%arg0: !llvm.[[ACCESSOR_1:struct<"class.cl::sycl::accessor.*", \(]][[ACCESSORIMPLDEVICE_1]][[ID_1]][[ARRAY_1]][[SUFFIX]], [[RANGE_1]][[ARRAY_1]][[SUFFIX]], [[RANGE_1]][[ARRAY_1]][[SUFFIX]][[SUFFIX]], struct<(ptr<i32, 1>)>)>)
10+
// CHECK: llvm.func @test_accessor.2(%arg0: !llvm.[[ACCESSOR_2:struct<"class.cl::sycl::accessor.*", \(]][[ACCESSORIMPLDEVICE_2:struct<"class.cl::sycl::detail::AccessorImplDevice.*", \(]][[ID_2:struct<"class.cl::sycl::id.*", \(]][[ARRAY_2]][[SUFFIX]], [[RANGE_2]][[ARRAY_2]][[SUFFIX]], [[RANGE_2]][[ARRAY_2]][[SUFFIX]][[SUFFIX]], struct<(ptr<i64, 1>)>)>)
11+
12+
module {
13+
func.func @test_array.1(%arg0: !sycl.array<[1], (memref<1xi64>)>) {
14+
return
15+
}
16+
func.func @test_array.2(%arg0: !sycl.array<[2], (memref<2xi64>)>) {
17+
return
18+
}
19+
func.func @test_id(%arg0: !sycl.id<1>, %arg1: !sycl.id<1>) {
20+
return
21+
}
22+
func.func @test_range.1(%arg0: !sycl.range<1>) {
23+
return
24+
}
25+
func.func @test_range.2(%arg0: !sycl.range<2>) {
26+
return
27+
}
28+
func.func @test_accessorImplDevice(%arg0: !sycl.accessor_impl_device<[1], (!sycl.id<1>, !sycl.range<1>, !sycl.range<1>)>) {
29+
return
30+
}
31+
func.func @test_accessor.1(%arg0: !sycl.accessor<[1, i32, write, global_buffer], (!sycl.accessor_impl_device<[1], (!sycl.id<1>, !sycl.range<1>, !sycl.range<1>)>)>) {
32+
return
33+
}
34+
func.func @test_accessor.2(%arg0: !sycl.accessor<[2, i64, write, global_buffer], (!sycl.accessor_impl_device<[2], (!sycl.id<2>, !sycl.range<2>, !sycl.range<2>)>)>) {
35+
return
36+
}
37+
}

‎polygeist/tools/cgeist/Test/Verification/sycl/SYCLToLLVM/types.mlir

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

0 commit comments

Comments
 (0)