Skip to content

Commit 5764fbb

Browse files
whitneywhtsangwhitneywhtsang
committed
[SYCL-MLIR] Add support of initializing sycl::detail::Boolean
Signed-off-by: Tsang, Whitney <whitney.tsang@intel.com>
1 parent 30e9d48 commit 5764fbb

File tree

3 files changed

+60
-5
lines changed

3 files changed

+60
-5
lines changed

‎polygeist/tools/cgeist/Lib/CGExpr.cc‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -455,8 +455,11 @@ ValueCategory MLIRScanner::VisitCXXStdInitializerListExpr(
455455

456456
ArrayPtr = CommonArrayToPointer(ArrayPtr);
457457

458-
Res = Builder.create<LLVM::InsertValueOp>(Loc, Res,
459-
ArrayPtr.getValue(Builder), 0);
458+
Value ArrayPtrVal = ArrayPtr.getValue(Builder);
459+
if (auto ST = dyn_cast<LLVM::LLVMStructType>(Res.getType()))
460+
ArrayPtrVal = castToMemSpaceOfType(ArrayPtrVal, ST.getBody()[0]);
461+
462+
Res = Builder.create<LLVM::InsertValueOp>(Loc, Res, ArrayPtrVal, 0);
460463
Field++;
461464
auto ITy =
462465
Glob.getTypes().getMLIRType(Field->getType()).cast<mlir::IntegerType>();

‎polygeist/tools/cgeist/Lib/ValueCategory.cc‎

Lines changed: 9 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -44,9 +44,15 @@ ValueCategory::ValueCategory(mlir::Value val, bool isReference)
4444

4545
ValueCategory::ValueCategory(mlir::Value Val, mlir::Value Index)
4646
: val{Val}, isReference{true}, Index{Index} {
47-
assert(val.getType().isa<MemRefType>() &&
48-
val.getType().cast<MemRefType>().getElementType().isa<VectorType>() &&
49-
"Expecting memref of vector");
47+
assert(
48+
((val.getType().isa<MemRefType>() &&
49+
val.getType().cast<MemRefType>().getElementType().isa<VectorType>()) ||
50+
(val.getType().isa<LLVM::LLVMPointerType>() &&
51+
val.getType()
52+
.cast<LLVM::LLVMPointerType>()
53+
.getElementType()
54+
.isa<VectorType>())) &&
55+
"Expecting memref/pointer of vector");
5056
assert(Index.getType().isa<IntegerType>() && "Expecting integer index");
5157
}
5258

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

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
// RUN: clang++ -fsycl -fsycl-device-only -O0 -w -emit-mlir %s -o - | FileCheck %s
2+
3+
#include <sycl/sycl.hpp>
4+
5+
// CHECK-LABEL: func.func @_Z7vecinitv() -> !llvm.struct<(vector<4xi8>)>
6+
// CHECK-DAG: %[[VAL_0:.*]] = arith.constant 4 : i64
7+
// CHECK-DAG: %[[VAL_1:.*]] = arith.constant 1 : i8
8+
// CHECK-DAG: %[[VAL_2:.*]] = arith.constant 0 : i8
9+
// CHECK-DAG: %[[VAL_3:.*]] = arith.constant 1 : i64
10+
// CHECK-NEXT: %[[VAL_4:.*]] = llvm.alloca %[[VAL_3]] x !llvm.struct<(vector<4xi8>)> : (i64) -> !llvm.ptr<struct<(vector<4xi8>)>>
11+
// CHECK-NEXT: %[[VAL_5:.*]] = llvm.alloca %[[VAL_3]] x !llvm.struct<(vector<4xi8>)> : (i64) -> !llvm.ptr<struct<(vector<4xi8>)>>
12+
// CHECK-NEXT: %[[VAL_6:.*]] = llvm.alloca %[[VAL_3]] x !llvm.struct<(ptr<i8, 4>, i64)> : (i64) -> !llvm.ptr<struct<(ptr<i8, 4>, i64)>>
13+
// CHECK-NEXT: %[[VAL_7:.*]] = llvm.alloca %[[VAL_3]] x !llvm.array<4 x i8> : (i64) -> !llvm.ptr<array<4 x i8>>
14+
// CHECK-NEXT: %[[VAL_8:.*]] = llvm.alloca %[[VAL_3]] x !llvm.array<4 x i8> : (i64) -> !llvm.ptr<array<4 x i8>>
15+
// CHECK-NEXT: %[[VAL_9:.*]] = llvm.alloca %[[VAL_3]] x !llvm.struct<(vector<4xi8>)> : (i64) -> !llvm.ptr<struct<(vector<4xi8>)>>
16+
// CHECK-NEXT: %[[VAL_10:.*]] = llvm.getelementptr %[[VAL_8]][0, 0] : (!llvm.ptr<array<4 x i8>>) -> !llvm.ptr<i8>
17+
// CHECK-NEXT: llvm.store %[[VAL_2]], %[[VAL_10]] : !llvm.ptr<i8>
18+
// CHECK-NEXT: %[[VAL_11:.*]] = llvm.getelementptr %[[VAL_8]][0, 1] : (!llvm.ptr<array<4 x i8>>) -> !llvm.ptr<i8>
19+
// CHECK-NEXT: llvm.store %[[VAL_1]], %[[VAL_11]] : !llvm.ptr<i8>
20+
// CHECK-NEXT: %[[VAL_12:.*]] = llvm.getelementptr %[[VAL_8]][0, 2] : (!llvm.ptr<array<4 x i8>>) -> !llvm.ptr<i8>
21+
// CHECK-NEXT: llvm.store %[[VAL_2]], %[[VAL_12]] : !llvm.ptr<i8>
22+
// CHECK-NEXT: %[[VAL_13:.*]] = llvm.getelementptr %[[VAL_8]][0, 3] : (!llvm.ptr<array<4 x i8>>) -> !llvm.ptr<i8>
23+
// CHECK-NEXT: llvm.store %[[VAL_2]], %[[VAL_13]] : !llvm.ptr<i8>
24+
// CHECK-NEXT: %[[VAL_14:.*]] = llvm.load %[[VAL_8]] : !llvm.ptr<array<4 x i8>>
25+
// CHECK-NEXT: llvm.store %[[VAL_14]], %[[VAL_7]] : !llvm.ptr<array<4 x i8>>
26+
// CHECK-NEXT: %[[VAL_15:.*]] = llvm.mlir.undef : !llvm.struct<(ptr<i8, 4>, i64)>
27+
// CHECK-NEXT: %[[VAL_16:.*]] = llvm.getelementptr %[[VAL_7]][0, 0] : (!llvm.ptr<array<4 x i8>>) -> !llvm.ptr<i8>
28+
// CHECK-NEXT: %[[VAL_17:.*]] = llvm.addrspacecast %[[VAL_16]] : !llvm.ptr<i8> to !llvm.ptr<i8, 4>
29+
// CHECK-NEXT: %[[VAL_18:.*]] = llvm.insertvalue %[[VAL_17]], %[[VAL_15]][0] : !llvm.struct<(ptr<i8, 4>, i64)>
30+
// CHECK-NEXT: %[[VAL_19:.*]] = llvm.insertvalue %[[VAL_0]], %[[VAL_18]][1] : !llvm.struct<(ptr<i8, 4>, i64)>
31+
// CHECK-NEXT: %[[VAL_20:.*]] = llvm.addrspacecast %[[VAL_9]] : !llvm.ptr<struct<(vector<4xi8>)>> to !llvm.ptr<struct<(vector<4xi8>)>, 4>
32+
// CHECK-NEXT: llvm.store %[[VAL_19]], %[[VAL_6]] : !llvm.ptr<struct<(ptr<i8, 4>, i64)>>
33+
// CHECK-NEXT: sycl.call(%[[VAL_20]], %[[VAL_6]]) {FunctionName = @Boolean, MangledFunctionName = @_ZN4sycl3_V16detail7BooleanILi4EEC1ESt16initializer_listIaE, TypeName = @Boolean} : (!llvm.ptr<struct<(vector<4xi8>)>, 4>, !llvm.ptr<struct<(ptr<i8, 4>, i64)>>) -> ()
34+
// CHECK-NEXT: %[[VAL_21:.*]] = llvm.load %[[VAL_9]] : !llvm.ptr<struct<(vector<4xi8>)>>
35+
// CHECK-NEXT: llvm.store %[[VAL_21]], %[[VAL_5]] : !llvm.ptr<struct<(vector<4xi8>)>>
36+
// CHECK-NEXT: %[[VAL_22:.*]] = llvm.addrspacecast %[[VAL_4]] : !llvm.ptr<struct<(vector<4xi8>)>> to !llvm.ptr<struct<(vector<4xi8>)>, 4>
37+
// CHECK-NEXT: %[[VAL_23:.*]] = llvm.addrspacecast %[[VAL_5]] : !llvm.ptr<struct<(vector<4xi8>)>> to !llvm.ptr<struct<(vector<4xi8>)>, 4>
38+
// CHECK-NEXT: sycl.call(%[[VAL_22]], %[[VAL_23]]) {FunctionName = @Boolean, MangledFunctionName = @_ZN4sycl3_V16detail7BooleanILi4EEC1ERKS3_, TypeName = @Boolean} : (!llvm.ptr<struct<(vector<4xi8>)>, 4>, !llvm.ptr<struct<(vector<4xi8>)>, 4>) -> ()
39+
// CHECK-NEXT: %[[VAL_24:.*]] = llvm.load %[[VAL_4]] : !llvm.ptr<struct<(vector<4xi8>)>>
40+
// CHECK-NEXT: return %[[VAL_24]] : !llvm.struct<(vector<4xi8>)>
41+
// CHECK-NEXT: }
42+
43+
SYCL_EXTERNAL sycl::detail::Boolean<4> vecinit() {
44+
sycl::detail::Boolean<4> t{false, true, false, false};
45+
return t;
46+
}

0 commit comments

Comments
 (0)