Skip to content

Commit 8775b9f

Browse files
victor-edsvictor-eds
authored
[SYCL-MLIR] Drop some non-standard types from the SYCL dialect (#8370)
Drop types that: 1. Are not defined in the SYCL spec; 2. Are not bases of other types. Dropped types are: - AssertHappened; - BFloat16; - GetOp; - GetScalarOp; - TupleCopyAssignableValueHolder; - TupleValueHolder. --------- Signed-off-by: Victor Perez <[email protected]>
1 parent 8097273 commit 8775b9f

File tree

10 files changed

+37
-311
lines changed

10 files changed

+37
-311
lines changed

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

Lines changed: 0 additions & 41 deletions
Original file line numberDiff line numberDiff line change
@@ -179,34 +179,13 @@ def SYCL_ArrayType : SYCL_Type<"Array", "array"> {
179179
let assemblyFormat = "`<` `[` $dimension `]` `,` `(` $body `)` `>`";
180180
}
181181

182-
def SYCL_AssertHappenedType : SYCL_Type<"AssertHappened", "assert_happened"> {
183-
let parameters = (ins ArrayRefParameter<"mlir::Type">:$body);
184-
let assemblyFormat = "`<` `(` $body `)` `>`";
185-
}
186-
187182
def SYCL_AtomicType : SYCL_Type<"Atomic", "atomic"> {
188183
let parameters = (ins "::mlir::Type":$dataType,
189184
"mlir::sycl::AccessAddrSpace":$addrSpace,
190185
ArrayRefParameter<"mlir::Type">:$body);
191186
let assemblyFormat = "`<` `[` $dataType `,` custom<AccessAddrSpace>($addrSpace) `]` `,` `(` $body `)` `>`";
192187
}
193188

194-
def SYCL_BFloat16Type : SYCL_Type<"BFloat16", "bfloat16"> {
195-
let parameters = (ins ArrayRefParameter<"mlir::Type">:$body);
196-
let assemblyFormat = "`<` `(` $body `)` `>`";
197-
}
198-
199-
def SYCL_GetOpType : SYCL_Type<"GetOp", "get_op"> {
200-
let parameters = (ins "::mlir::Type":$dataType);
201-
let assemblyFormat = "`<` $dataType `>`";
202-
}
203-
204-
def SYCL_GetScalarOpType : SYCL_Type<"GetScalarOp", "get_scalar_op"> {
205-
let parameters = (ins "::mlir::Type":$dataType,
206-
ArrayRefParameter<"mlir::Type">:$body);
207-
let assemblyFormat = "`<` `[` $dataType `]` `,` `(` $body `)` `>`";
208-
}
209-
210189
def SYCL_GroupType : SYCL_Type<"Group", "group"> {
211190
let parameters = (ins "unsigned":$dimension,
212191
ArrayRefParameter<"mlir::Type">:$body);
@@ -331,22 +310,6 @@ def SYCL_SwizzledVecType : SYCL_Type<"SwizzledVec", "swizzled_vec"> {
331310
ArrayRefParameter<"mlir::Type">:$body);
332311
let assemblyFormat = "`<` `[` $vecType `,` $indexes `]` `,` `(` $body `)` `>`";
333312
}
334-
335-
def SYCL_TupleCopyAssignableValueHolderType
336-
: SYCL_Type<"TupleCopyAssignableValueHolder", "tuple_copy_assignable_value_holder",
337-
[SYCLInheritanceTypeTrait<"TupleValueHolderType">]> {
338-
let parameters = (ins "::mlir::Type":$dataType,
339-
"bool":$isTriviallyCopyAssignable,
340-
ArrayRefParameter<"mlir::Type">:$body);
341-
let assemblyFormat = "`<` `[` $dataType `,` $isTriviallyCopyAssignable `]` `,` `(` $body `)` `>`";
342-
}
343-
344-
def SYCL_TupleValueHolderType : SYCL_Type<"TupleValueHolder", "tuple_value_holder"> {
345-
let parameters = (ins "::mlir::Type":$dataType,
346-
ArrayRefParameter<"mlir::Type">:$body);
347-
let assemblyFormat = "`<` `[` $dataType `]` `,` `(` $body `)` `>`";
348-
}
349-
350313
def SYCL_VecType : SYCL_Type<"Vec", "vec"> {
351314
let parameters = (ins "::mlir::Type":$dataType,
352315
"int":$numElements,
@@ -362,8 +325,6 @@ def AccessorMemRef : MemRefOf<[SYCL_AccessorType]>;
362325
def AccessorSubscriptMemRef : MemRefOf<[SYCL_AccessorSubscriptType]>;
363326
def ArrayMemRef : MemRefOf<[SYCL_ArrayType]>;
364327
def AtomicMemRef : MemRefOf<[SYCL_AtomicType]>;
365-
def GetOpMemRef : MemRefOf<[SYCL_GetOpType]>;
366-
def GetScalarOpMemRef : MemRefOf<[SYCL_GetScalarOpType]>;
367328
def GroupMemRef : MemRefOf<[SYCL_GroupType]>;
368329
def IDMemRef : MemRefOf<[SYCL_IDType]>;
369330
def ItemBaseMemRef : MemRefOf<[SYCL_ItemBaseType]>;
@@ -391,8 +352,6 @@ def SYCLMemref : AnyTypeOf<[
391352
AccessorSubscriptMemRef,
392353
ArrayMemRef,
393354
AtomicMemRef,
394-
GetOpMemRef,
395-
GetScalarOpMemRef,
396355
GroupMemRef,
397356
IDMemRef,
398357
ItemBaseMemRef,

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

Lines changed: 0 additions & 64 deletions
Original file line numberDiff line numberDiff line change
@@ -148,13 +148,6 @@ static Optional<Type> convertArrayType(sycl::ArrayType type,
148148
return structTy;
149149
}
150150

151-
/// Converts SYCL AssertHappened type to LLVM type.
152-
static Optional<Type> convertAssertHappenedType(sycl::AssertHappenedType type,
153-
LLVMTypeConverter &converter) {
154-
return convertBodyType("struct.sycl::_V1::detail::AssertHappened",
155-
type.getBody(), converter);
156-
}
157-
158151
/// Converts SYCL atomic type to LLVM type.
159152
static Optional<Type> convertAtomicType(sycl::AtomicType type,
160153
LLVMTypeConverter &converter) {
@@ -163,26 +156,6 @@ static Optional<Type> convertAtomicType(sycl::AtomicType type,
163156
return convertBodyType("class.sycl::_V1::atomic", type.getBody(), converter);
164157
}
165158

166-
/// Converts SYCL bfloat16 type to LLVM type.
167-
static Optional<Type> convertBFloat16Type(sycl::BFloat16Type type,
168-
LLVMTypeConverter &converter) {
169-
return convertBodyType("class.sycl::_V1::ext::oneapi::bfloat16",
170-
type.getBody(), converter);
171-
}
172-
173-
/// Converts SYCL GetOp type to LLVM type.
174-
static Optional<Type> convertGetOpType(sycl::GetOpType type,
175-
LLVMTypeConverter &converter) {
176-
return getI8Struct("class.sycl::_V1::detail::GetOp", converter);
177-
}
178-
179-
/// Converts SYCL GetScalarOp type to LLVM type.
180-
static Optional<Type> convertGetScalarOpType(sycl::GetScalarOpType type,
181-
LLVMTypeConverter &converter) {
182-
return convertBodyType("class.sycl::_V1::detail::GetScalarOp", type.getBody(),
183-
converter);
184-
}
185-
186159
/// Converts SYCL group type to LLVM type.
187160
static Optional<Type> convertGroupType(sycl::GroupType type,
188161
LLVMTypeConverter &converter) {
@@ -327,24 +300,6 @@ static Optional<Type> convertSwizzledVecType(sycl::SwizzledVecType type,
327300
return convertBodyType("class.sycl::_V1::detail::SwizzleOp", type.getBody(),
328301
converter);
329302
}
330-
331-
/// Converts SYCL TupleCopyAssignableValueHolder type to LLVM type.
332-
static Optional<Type> convertTupleCopyAssignableValueHolderType(
333-
sycl::TupleCopyAssignableValueHolderType type,
334-
LLVMTypeConverter &converter) {
335-
return convertBodyType(
336-
"struct.sycl::_V1::detail::TupleCopyAssignableValueHolder",
337-
type.getBody(), converter);
338-
}
339-
340-
/// Converts SYCL TupleValueHolder type to LLVM type.
341-
static Optional<Type>
342-
convertTupleValueHolderType(sycl::TupleValueHolderType type,
343-
LLVMTypeConverter &converter) {
344-
return convertBodyType("struct.sycl::_V1::detail::TupleValueHolder",
345-
type.getBody(), converter);
346-
}
347-
348303
/// Converts SYCL vec type to LLVM type.
349304
static Optional<Type> convertVecType(sycl::VecType type,
350305
LLVMTypeConverter &converter) {
@@ -550,21 +505,9 @@ void mlir::sycl::populateSYCLToLLVMTypeConversion(
550505
typeConverter.addConversion([&](sycl::ArrayType type) {
551506
return convertArrayType(type, typeConverter);
552507
});
553-
typeConverter.addConversion([&](sycl::AssertHappenedType type) {
554-
return convertAssertHappenedType(type, typeConverter);
555-
});
556508
typeConverter.addConversion([&](sycl::AtomicType type) {
557509
return convertAtomicType(type, typeConverter);
558510
});
559-
typeConverter.addConversion([&](sycl::BFloat16Type type) {
560-
return convertBFloat16Type(type, typeConverter);
561-
});
562-
typeConverter.addConversion([&](sycl::GetOpType type) {
563-
return convertGetOpType(type, typeConverter);
564-
});
565-
typeConverter.addConversion([&](sycl::GetScalarOpType type) {
566-
return convertGetScalarOpType(type, typeConverter);
567-
});
568511
typeConverter.addConversion([&](sycl::GroupType type) {
569512
return convertGroupType(type, typeConverter);
570513
});
@@ -621,13 +564,6 @@ void mlir::sycl::populateSYCLToLLVMTypeConversion(
621564
typeConverter.addConversion([&](sycl::SwizzledVecType type) {
622565
return convertSwizzledVecType(type, typeConverter);
623566
});
624-
typeConverter.addConversion(
625-
[&](sycl::TupleCopyAssignableValueHolderType type) {
626-
return convertTupleCopyAssignableValueHolderType(type, typeConverter);
627-
});
628-
typeConverter.addConversion([&](sycl::TupleValueHolderType type) {
629-
return convertTupleValueHolderType(type, typeConverter);
630-
});
631567
typeConverter.addConversion(
632568
[&](sycl::VecType type) { return convertVecType(type, typeConverter); });
633569
}

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

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -60,10 +60,6 @@ bool SYCLCastOp::areCastCompatible(TypeRange Inputs, TypeRange Outputs) {
6060
return InputElemType
6161
.hasTrait<SYCLInheritanceTypeTrait<OwnerLessBaseType>::Trait>();
6262
})
63-
.template Case<TupleValueHolderType>([&](auto) {
64-
return InputElemType
65-
.hasTrait<SYCLInheritanceTypeTrait<TupleValueHolderType>::Trait>();
66-
})
6763
.Default(false);
6864
}
6965

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

Lines changed: 6 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -135,8 +135,7 @@ SYCLOpAsmInterface::getAlias(mlir::Type Type, llvm::raw_ostream &OS) const {
135135
<< Ty.getCurrentDimension() << "_";
136136
return AliasResult::OverridableAlias;
137137
})
138-
.Case<mlir::sycl::AssertHappenedType, mlir::sycl::BFloat16Type,
139-
mlir::sycl::KernelHandlerType, mlir::sycl::StreamType>(
138+
.Case<mlir::sycl::KernelHandlerType, mlir::sycl::StreamType>(
140139
[&](auto Ty) {
141140
OS << "sycl_" << decltype(Ty)::getMnemonic() << "_";
142141
return AliasResult::FinalAlias;
@@ -170,30 +169,23 @@ SYCLOpAsmInterface::getAlias(mlir::Type Type, llvm::raw_ostream &OS) const {
170169
<< "_" << Ty.getType();
171170
return AliasResult::FinalAlias;
172171
})
172+
.Case<mlir::sycl::MaximumType, mlir::sycl::MinimumType>([&](auto Ty) {
173+
OS << "sycl_" << decltype(Ty)::getMnemonic() << "_" << Ty.getDataType()
174+
<< "_";
175+
return AliasResult::FinalAlias;
176+
})
173177
.Case<mlir::sycl::MultiPtrType>([&](auto Ty) {
174178
OS << "sycl_" << decltype(Ty)::getMnemonic() << "_" << Ty.getDataType()
175179
<< "_" << mlir::sycl::accessAddressSpaceAsString(Ty.getAddrSpace())
176180
<< "_";
177181
return AliasResult::OverridableAlias;
178182
})
179-
.Case<mlir::sycl::GetScalarOpType, mlir::sycl::MinimumType,
180-
mlir::sycl::MaximumType, mlir::sycl::TupleValueHolderType>(
181-
[&](auto Ty) {
182-
OS << "sycl_" << decltype(Ty)::getMnemonic() << "_"
183-
<< Ty.getDataType() << "_";
184-
return AliasResult::FinalAlias;
185-
})
186183
.Case<mlir::sycl::SwizzledVecType>([&](auto Ty) {
187184
const auto VecTy = Ty.getVecType();
188185
OS << "sycl_" << decltype(Ty)::getMnemonic() << "_"
189186
<< VecTy.getDataType() << "_" << VecTy.getNumElements() << "_";
190187
return AliasResult::OverridableAlias;
191188
})
192-
.Case<mlir::sycl::TupleCopyAssignableValueHolderType>([&](auto Ty) {
193-
OS << "sycl_" << decltype(Ty)::getMnemonic() << "_" << Ty.getDataType()
194-
<< "_";
195-
return AliasResult::OverridableAlias;
196-
})
197189
.Case<mlir::sycl::VecType>([&](auto Ty) {
198190
OS << "sycl_" << decltype(Ty)::getMnemonic() << "_" << Ty.getDataType()
199191
<< "_" << Ty.getNumElements() << "_";

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

Lines changed: 2 additions & 47 deletions
Original file line numberDiff line numberDiff line change
@@ -87,14 +87,6 @@ func.func @test_OwnerLessBase(%arg0: !sycl.owner_less_base) {
8787

8888
// -----
8989

90-
!sycl_assert_happened_ = !sycl.assert_happened<(i32, !llvm.array<257 x i8>, !llvm.array<257 x i8>, !llvm.array<129 x i8>, i32, i64, i64, i64, i64, i64, i64)>
91-
// CHECK: llvm.func @test_assert_happened(%arg0: !llvm.[[ASSERT_HAPPENED:struct<"struct.sycl::_V1::detail::AssertHappened", \(i32, array<257 x i8>, array<257 x i8>, array<129 x i8>, i32, i64, i64, i64, i64, i64, i64\)>]]) {
92-
func.func @test_assert_happened(%arg0: !sycl_assert_happened_) {
93-
return
94-
}
95-
96-
// -----
97-
9890
!sycl_atomic_i32_1_ = !sycl.atomic<[i32,1], (memref<?xi32, 1>)>
9991
!sycl_atomic_f32_3_ = !sycl.atomic<[f32,3], (memref<?xf32, 3>)>
10092
// CHECK: llvm.func @test_atomic(%arg0: !llvm.[[ATOMIC1:struct<"class.sycl::_V1::atomic", \(struct<\(ptr<f32, 3>, ptr<f32, 3>, i64, array<1 x i64>, array<1 x i64>\)>\)>]], %arg1: !llvm.[[ATOMIC1:struct<"class.sycl::_V1::atomic.1", \(struct<\(ptr<i32, 1>, ptr<i32, 1>, i64, array<1 x i64>, array<1 x i64>\)>\)>]]) {
@@ -104,29 +96,6 @@ func.func @test_atomic(%arg0: !sycl_atomic_f32_3_, %arg1: !sycl_atomic_i32_1_) {
10496

10597
// -----
10698

107-
!sycl_bfloat16_ = !sycl.bfloat16<(i16)>
108-
// CHECK: llvm.func @test_bfloat16(%arg0: !llvm.[[BFLOAT16:struct<"class.sycl::_V1::ext::oneapi::bfloat16", \(i16\)>]]) {
109-
func.func @test_bfloat16(%arg0: !sycl_bfloat16_) {
110-
return
111-
}
112-
113-
// -----
114-
115-
!sycl_get_scalar_op_i32_ = !sycl.get_scalar_op<[i32], (i32)>
116-
// CHECK: llvm.func @test_get_scalar_op(%arg0: !llvm.[[GET_SCALAR_OP:struct<"class.sycl::_V1::detail::GetScalarOp.*", \(i32\)>]])
117-
func.func @test_get_scalar_op(%arg0: !sycl_get_scalar_op_i32_) {
118-
return
119-
}
120-
121-
// -----
122-
123-
// CHECK: llvm.func @test_get_op(%arg0: !llvm.[[GET_OP:struct<"class.sycl::_V1::detail::GetOp", \(i8\)>]])
124-
func.func @test_get_op(%arg0: !sycl.get_op<i32>) {
125-
return
126-
}
127-
128-
// -----
129-
13099
!sycl_id_1_ = !sycl.id<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
131100
!sycl_range_1_ = !sycl.range<[1], (!sycl.array<[1], (memref<1xi64, 4>)>)>
132101
!sycl_item_base_1_ = !sycl.item_base<[1, true], (!sycl_range_1_, !sycl_id_1_, !sycl_id_1_)>
@@ -252,29 +221,15 @@ func.func @test_sub_group(%arg0: !sycl.sub_group) {
252221
func.func @test_stream(%arg0: !sycl_stream_) {
253222
return
254223
}
255-
256-
// -----
257-
258-
!sycl_tuple_value_holder_i32_ = !sycl.tuple_value_holder<[i32], (i32)>
259-
!sycl_tuple_copy_assignable_value_holder_i32_ = !sycl.tuple_copy_assignable_value_holder<[i32, true], (!sycl.tuple_value_holder<[i32], (i32)>)>
260-
// CHECK: llvm.func @test_tuple_value_holder(%arg0: !llvm.[[TUPLE_VALUE_HOLDER:struct<"struct.sycl::_V1::detail::TupleValueHolder", \(i32\)>]]) {
261-
func.func @test_tuple_value_holder(%arg0: !sycl_tuple_value_holder_i32_) {
262-
return
263-
}
264-
// CHECK: llvm.func @test_tuple_copy_assignable_value_holder(%arg0: !llvm.[[TUPLE_COPY_ASSIGNABLE_VALUE_HOLDER:struct<"struct.sycl::_V1::detail::TupleCopyAssignableValueHolder", \(]][[TUPLE_VALUE_HOLDER]][[SUFFIX]]) {
265-
func.func @test_tuple_copy_assignable_value_holder(%arg0: !sycl_tuple_copy_assignable_value_holder_i32_) {
266-
return
267-
}
268-
269224
// -----
270225

271226
!sycl_vec_f32_4_ = !sycl.vec<[f32, 4], (vector<4xf32>)>
272227
// CHECK: llvm.func @test_vec(%arg0: !llvm.[[VEC:struct<"class.sycl::_V1::vec", \(vector<4xf32>\)>]])
273228
func.func @test_vec(%arg0: !sycl_vec_f32_4_) {
274229
return
275230
}
276-
!sycl_swizzled_vec_f32_4_ = !sycl.swizzled_vec<[!sycl_vec_f32_4_, 0, 2], (memref<?x!sycl_vec_f32_4_, 4>, !sycl.get_op<i8>, !sycl.get_op<i8>)>
277-
// CHECK: llvm.func @test_swizzled_vec(%arg0: !llvm.[[SWIZZLED_VEC:struct<"class.sycl::_V1::detail::SwizzleOp"]], (struct<(ptr<[[VEC]], 4>, ptr<[[VEC]], 4>, i64, array<1 x i64>, array<1 x i64>)>, [[GET_OP]], [[GET_OP]][[SUFFIX]]) {
231+
!sycl_swizzled_vec_f32_4_ = !sycl.swizzled_vec<[!sycl_vec_f32_4_, 0, 2], (memref<?x!sycl_vec_f32_4_, 4>, !llvm.struct<(i8)>, !llvm.struct<(i8)>)>
232+
// CHECK: llvm.func @test_swizzled_vec(%arg0: !llvm.[[SWIZZLED_VEC:struct<"class.sycl::_V1::detail::SwizzleOp"]], (struct<(ptr<[[VEC]], 4>, ptr<[[VEC]], 4>, i64, array<1 x i64>, array<1 x i64>)>, [[GET_OP:struct<\(i8\)>]], [[GET_OP]][[SUFFIX]]) {
278233
func.func @test_swizzled_vec(%arg0: !sycl_swizzled_vec_f32_4_) {
279234
return
280235
}

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

Lines changed: 2 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -166,26 +166,6 @@ func.func @_Z7group_2N2cl4sycl5groupILi2EEE(%arg0: !sycl_group_2_) attributes {l
166166

167167
// -----
168168

169-
////////////////////////////////////////////////////////////////////////////////
170-
// GET_OP
171-
////////////////////////////////////////////////////////////////////////////////
172-
173-
// CHECK: func @get_op(%arg0: !sycl.get_op<i32>)
174-
func.func @get_op(%arg0: !sycl.get_op<i32>) attributes {llvm.linkage = #llvm.linkage<external>} {
175-
return
176-
}
177-
178-
////////////////////////////////////////////////////////////////////////////////
179-
// GET_SCALAR_OP
180-
////////////////////////////////////////////////////////////////////////////////
181-
182-
!sycl_get_scalar_op_i32_ = !sycl.get_scalar_op<[i32], (i32)>
183-
184-
// CHECK: func @get_scalar_op(%arg0: !sycl_get_scalar_op_i32_)
185-
func.func @get_scalar_op(%arg0: !sycl_get_scalar_op_i32_) attributes {llvm.linkage = #llvm.linkage<external>} {
186-
return
187-
}
188-
189169
////////////////////////////////////////////////////////////////////////////////
190170
// VEC
191171
////////////////////////////////////////////////////////////////////////////////
@@ -206,8 +186,8 @@ func.func @vec_1(%arg0: !sycl_vec_f32_4_) attributes {llvm.linkage = #llvm.linka
206186
// SWIZZLED_VEC
207187
////////////////////////////////////////////////////////////////////////////////
208188

209-
!sycl_swizzled_vec_i32_2_ = !sycl.swizzled_vec<[!sycl_vec_i32_2_, 0, 1], (memref<?x!sycl_vec_i32_2_, 4>, !sycl.get_op<i8>, !sycl.get_op<i8>)>
210-
!sycl_swizzled_vec_i32_2_1 = !sycl.swizzled_vec<[!sycl_vec_i32_2_, 0, 1], (memref<?x!sycl_vec_i32_2_, 4>, !sycl_swizzled_vec_i32_2_, !sycl_get_scalar_op_i32_)>
189+
!sycl_swizzled_vec_i32_2_ = !sycl.swizzled_vec<[!sycl_vec_i32_2_, 0, 1], (memref<?x!sycl_vec_i32_2_, 4>, !llvm.struct<(i8)>, !llvm.struct<(i8)>)>
190+
!sycl_swizzled_vec_i32_2_1 = !sycl.swizzled_vec<[!sycl_vec_i32_2_, 0, 1], (memref<?x!sycl_vec_i32_2_, 4>, !sycl_swizzled_vec_i32_2_, !llvm.struct<(i32)>)>
211191

212192
// CHECK: func @swizzled_vec_0(%arg0: !sycl_swizzled_vec_i32_2_)
213193
func.func @swizzled_vec_0(%arg0: !sycl_swizzled_vec_i32_2_) attributes {llvm.linkage = #llvm.linkage<external>} {

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

Lines changed: 5 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1371,21 +1371,16 @@ mlir::Type CodeGenTypes::getMLIRType(clang::QualType QT, bool *ImplicitRef,
13711371
const auto TypeName = RD->getName();
13721372
if (TypeName == "accessor" || TypeName == "accessor_common" ||
13731373
TypeName == "AccessorImplDevice" || TypeName == "AccessorSubscript" ||
1374-
TypeName == "array" || TypeName == "AssertHappened" ||
1375-
TypeName == "atomic" || TypeName == "bfloat16" ||
1376-
TypeName == "GetOp" || TypeName == "GetScalarOp" ||
1377-
TypeName == "group" || TypeName == "h_item" || TypeName == "id" ||
1378-
TypeName == "item" || TypeName == "ItemBase" ||
1379-
TypeName == "kernel_handler" ||
1374+
TypeName == "array" || TypeName == "atomic" || TypeName == "group" ||
1375+
TypeName == "h_item" || TypeName == "id" || TypeName == "item" ||
1376+
TypeName == "ItemBase" || TypeName == "kernel_handler" ||
13801377
TypeName == "LocalAccessorBaseDevice" ||
13811378
TypeName == "local_accessor_base" || TypeName == "local_accessor" ||
13821379
TypeName == "maximum" || TypeName == "minimum" ||
13831380
TypeName == "multi_ptr" || TypeName == "nd_item" ||
13841381
TypeName == "nd_range" || TypeName == "OwnerLessBase" ||
1385-
TypeName == "range" || TypeName == "stream" ||
1386-
TypeName == "sub_group" || TypeName == "SwizzleOp" ||
1387-
TypeName == "TupleCopyAssignableValueHolder" ||
1388-
TypeName == "TupleValueHolder" || TypeName == "vec")
1382+
TypeName == "range" || TypeName == "SwizzleOp" ||
1383+
TypeName == "stream" || TypeName == "sub_group" || TypeName == "vec")
13891384
return getSYCLType(RT, *this);
13901385

13911386
assert((AllowUndefinedSYCLTypes ||

0 commit comments

Comments
 (0)