Skip to content

Commit 3a7c778

Browse files
jinge90jinge90
committed
Merge remote-tracking branch 'upstream/sycl' into simd_emulate_for_imf_libdevice
2 parents 595ce8c + 2e1c36f commit 3a7c778

25 files changed

+941
-343
lines changed

‎clang/lib/Basic/Targets/SPIR.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -230,7 +230,7 @@ class LLVM_LIBRARY_VISIBILITY SPIR64FPGATargetInfo : public SPIR64TargetInfo {
230230
public:
231231
SPIR64FPGATargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
232232
: SPIR64TargetInfo(Triple, Opts) {}
233-
virtual size_t getMaxBitIntWidth() const override { return 2048; }
233+
virtual size_t getMaxBitIntWidth() const override { return 4096; }
234234
};
235235

236236
// x86-32 SPIR Windows target

‎clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-host-intelfpga-bitint.cpp

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -3,23 +3,23 @@
33
// This test checks that we generate appropriate code for division
44
// operations of _BitInts of size greater than 128 bits, since it
55
// is allowed when -fintelfpga is enabled. The test uses a value of
6-
// 2048 for the bitsize as that is the maximum that is currently
6+
// 4096 for the bitsize as that is the maximum that is currently
77
// supported.
88

9-
// CHECK: define{{.*}} void @_Z3fooDB2048_S_(i2048* {{.*}} sret(i2048) align 8 %agg.result, i2048* {{.*}} byval(i2048) align 8 %[[ARG1:[0-9]+]], i2048* {{.*}} byval(i2048) align 8 %[[ARG2:[0-9]+]])
10-
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
11-
// CHECK: %[[VAR_A:a]].addr = alloca i2048, align 8
12-
// CHECK: %[[VAR_B:b]].addr = alloca i2048, align 8
13-
// CHECK: %[[VAR_A]] = load i2048, i2048* %[[ARG1]], align 8
14-
// CHECK: %[[VAR_B]] = load i2048, i2048* %[[ARG2]], align 8
15-
// CHECK: store i2048 %[[VAR_A]], i2048* %[[VAR_A]].addr, align 8
16-
// CHECK: store i2048 %[[VAR_B]], i2048* %[[VAR_B]].addr, align 8
17-
// CHECK: %[[TEMP1:[0-9]+]] = load i2048, i2048* %[[VAR_A]].addr, align 8
18-
// CHECK: %[[TEMP2:[0-9]+]] = load i2048, i2048* %[[VAR_B]].addr, align 8
19-
// CHECK: %div = sdiv i2048 %[[TEMP1]], %[[TEMP2]]
20-
// CHECK: store i2048 %div, i2048* %agg.result, align 8
21-
// CHECK: %[[RES:[0-9+]]] = load i2048, i2048* %agg.result, align 8
22-
// CHECK: store i2048 %[[RES]], i2048* %agg.result, align 8
9+
// CHECK: define{{.*}} void @_Z3fooDB4096_S_(i4096* {{.*}} sret(i4096) align 8 %agg.result, i4096* {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], i4096* {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]])
10+
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
11+
// CHECK: %[[VAR_A:a]].addr = alloca i4096, align 8
12+
// CHECK: %[[VAR_B:b]].addr = alloca i4096, align 8
13+
// CHECK: %[[VAR_A]] = load i4096, i4096* %[[ARG1]], align 8
14+
// CHECK: %[[VAR_B]] = load i4096, i4096* %[[ARG2]], align 8
15+
// CHECK: store i4096 %[[VAR_A]], i4096* %[[VAR_A]].addr, align 8
16+
// CHECK: store i4096 %[[VAR_B]], i4096* %[[VAR_B]].addr, align 8
17+
// CHECK: %[[TEMP1:[0-9]+]] = load i4096, i4096* %[[VAR_A]].addr, align 8
18+
// CHECK: %[[TEMP2:[0-9]+]] = load i4096, i4096* %[[VAR_B]].addr, align 8
19+
// CHECK: %div = sdiv i4096 %[[TEMP1]], %[[TEMP2]]
20+
// CHECK: store i4096 %div, i4096* %agg.result, align 8
21+
// CHECK: %[[RES:[0-9+]]] = load i4096, i4096* %agg.result, align 8
22+
// CHECK: store i4096 %[[RES]], i4096* %agg.result, align 8
2323
// CHECK: ret void
2424
return a / b;
2525
}

‎clang/test/CodeGenSYCL/no-opaque-ptrs-sycl-intelfpga-bitint.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -3,24 +3,24 @@
33
// This test checks that we generate appropriate code for division
44
// operations of _BitInts of size greater than 128 bits, since it
55
// is allowed when -fintelfpga is enabled. The test uses a value
6-
// of 2048 for the bitsize, the max that is currently supported.
6+
// of 4096 for the bitsize, the max that is currently supported.
77

88
#include "Inputs/sycl.hpp"
99

10-
// CHECK: define{{.*}} void @_Z3fooDB2048_S_(i2048 addrspace(4)* {{.*}} sret(i2048) align 8 %agg.result, i2048* {{.*}} byval(i2048) align 8 %[[ARG1:[0-9]+]], i2048* {{.*}} byval(i2048) align 8 %[[ARG2:[0-9]+]])
11-
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
12-
// CHECK: %[[VAR_A:a]] = load i2048, i2048* %[[ARG1]], align 8
13-
// CHECK: %[[VAR_B:b]] = load i2048, i2048* %[[ARG2]], align 8
14-
// CHECK: %[[RES:div]] = sdiv i2048 %[[VAR_A]], %[[VAR_B]]
15-
// CHECK: store i2048 %[[RES]], i2048 addrspace(4)* %agg.result, align 8
10+
// CHECK: define{{.*}} void @_Z3fooDB4096_S_(i4096 addrspace(4)* {{.*}} sret(i4096) align 8 %agg.result, i4096* {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], i4096* {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]])
11+
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
12+
// CHECK: %[[VAR_A:a]] = load i4096, i4096* %[[ARG1]], align 8
13+
// CHECK: %[[VAR_B:b]] = load i4096, i4096* %[[ARG2]], align 8
14+
// CHECK: %[[RES:div]] = sdiv i4096 %[[VAR_A]], %[[VAR_B]]
15+
// CHECK: store i4096 %[[RES]], i4096 addrspace(4)* %agg.result, align 8
1616
// CHECK: ret void
1717
return a / b;
1818
}
1919

2020
int main() {
2121
sycl::handler h;
2222
auto lambda = []() {
23-
_BitInt(2048) a, b = 3, c = 4;
23+
_BitInt(4096) a, b = 3, c = 4;
2424
a = foo(b, c);
2525
};
2626
h.single_task(lambda);

‎clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -3,22 +3,22 @@
33
// This test checks that we generate appropriate code for division
44
// operations of _BitInts of size greater than 128 bits, since it
55
// is allowed when -fintelfpga is enabled. The test uses a value of
6-
// 2048, the maximum bitsize that is currently supported.
6+
// 4096, the maximum bitsize that is currently supported.
77

8-
// CHECK: define{{.*}} void @_Z3fooDB2048_S_(ptr {{.*}} sret(i2048) align 8 %agg.result, ptr {{.*}} byval(i2048) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i2048) align 8 %[[ARG2:[0-9]+]])
9-
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
10-
// CHECK: %[[VAR_A:a]].addr = alloca i2048, align 8
11-
// CHECK: %[[VAR_B:b]].addr = alloca i2048, align 8
12-
// CHECK: %[[VAR_A]] = load i2048, ptr %[[ARG1]], align 8
13-
// CHECK: %[[VAR_B]] = load i2048, ptr %[[ARG2]], align 8
14-
// CHECK: store i2048 %[[VAR_A]], ptr %[[VAR_A]].addr, align 8
15-
// CHECK: store i2048 %[[VAR_B]], ptr %[[VAR_B]].addr, align 8
16-
// CHECK: %[[TEMP1:[0-9]+]] = load i2048, ptr %[[VAR_A]].addr, align 8
17-
// CHECK: %[[TEMP2:[0-9]+]] = load i2048, ptr %[[VAR_B]].addr, align 8
18-
// CHECK: %div = sdiv i2048 %[[TEMP1]], %[[TEMP2]]
19-
// CHECK: store i2048 %div, ptr %agg.result, align 8
20-
// CHECK: %[[RES:[0-9+]]] = load i2048, ptr %agg.result, align 8
21-
// CHECK: store i2048 %[[RES]], ptr %agg.result, align 8
8+
// CHECK: define{{.*}} void @_Z3fooDB4096_S_(ptr {{.*}} sret(i4096) align 8 %agg.result, ptr {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]])
9+
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
10+
// CHECK: %[[VAR_A:a]].addr = alloca i4096, align 8
11+
// CHECK: %[[VAR_B:b]].addr = alloca i4096, align 8
12+
// CHECK: %[[VAR_A]] = load i4096, ptr %[[ARG1]], align 8
13+
// CHECK: %[[VAR_B]] = load i4096, ptr %[[ARG2]], align 8
14+
// CHECK: store i4096 %[[VAR_A]], ptr %[[VAR_A]].addr, align 8
15+
// CHECK: store i4096 %[[VAR_B]], ptr %[[VAR_B]].addr, align 8
16+
// CHECK: %[[TEMP1:[0-9]+]] = load i4096, ptr %[[VAR_A]].addr, align 8
17+
// CHECK: %[[TEMP2:[0-9]+]] = load i4096, ptr %[[VAR_B]].addr, align 8
18+
// CHECK: %div = sdiv i4096 %[[TEMP1]], %[[TEMP2]]
19+
// CHECK: store i4096 %div, ptr %agg.result, align 8
20+
// CHECK: %[[RES:[0-9+]]] = load i4096, ptr %agg.result, align 8
21+
// CHECK: store i4096 %[[RES]], ptr %agg.result, align 8
2222
// CHECK: ret void
2323
return a / b;
2424
}

‎clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -3,24 +3,24 @@
33
// This test checks that we generate appropriate code for division
44
// operations of _BitInts of size greater than 128 bits, since it
55
// is allowed when -fintelfpga is enabled. The test uses a value of
6-
// 2048 for bitint size, the maximum that is currently supported.
6+
// 4096 for bitint size, the maximum that is currently supported.
77

88
#include "Inputs/sycl.hpp"
99

10-
// CHECK: define{{.*}} void @_Z3fooDB2048_S_(ptr addrspace(4) {{.*}} sret(i2048) align 8 %agg.result, ptr {{.*}} byval(i2048) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i2048) align 8 %[[ARG2:[0-9]+]])
11-
signed _BitInt(2048) foo(signed _BitInt(2048) a, signed _BitInt(2048) b) {
12-
// CHECK: %[[VAR_A:a]] = load i2048, ptr %[[ARG1]], align 8
13-
// CHECK: %[[VAR_B:b]] = load i2048, ptr %[[ARG2]], align 8
14-
// CHECK: %[[RES:div]] = sdiv i2048 %[[VAR_A]], %[[VAR_B]]
15-
// CHECK: store i2048 %[[RES]], ptr addrspace(4) %agg.result, align 8
10+
// CHECK: define{{.*}} void @_Z3fooDB4096_S_(ptr addrspace(4) {{.*}} sret(i4096) align 8 %agg.result, ptr {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]])
11+
signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) {
12+
// CHECK: %[[VAR_A:a]] = load i4096, ptr %[[ARG1]], align 8
13+
// CHECK: %[[VAR_B:b]] = load i4096, ptr %[[ARG2]], align 8
14+
// CHECK: %[[RES:div]] = sdiv i4096 %[[VAR_A]], %[[VAR_B]]
15+
// CHECK: store i4096 %[[RES]], ptr addrspace(4) %agg.result, align 8
1616
// CHECK: ret void
1717
return a / b;
1818
}
1919

2020
int main() {
2121
sycl::handler h;
2222
auto lambda = []() {
23-
_BitInt(2048) a, b = 3, c = 4;
23+
_BitInt(4096) a, b = 3, c = 4;
2424
a = foo(b, c);
2525
};
2626
h.single_task(lambda);

‎clang/test/Driver/sycl-offload.c

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -88,8 +88,8 @@
8888

8989
/// Check that -aux-triple is passed with -fsycl -fintelfpga
9090
// RUN: %clang -### -fsycl -fintelfpga %s 2>&1 \
91-
// RUN: | FileCheck -DARCH=x86_64 -DARCH2=spir64_fpga -check-prefix=CHK-SYCL-FPGA-AUX-TRIPLE %s
92-
// CHK-SYCL-FPGA-AUX-TRIPLE: clang{{.*}} "-cc1" "-triple" "[[ARCH]]-{{.*}}"{{.*}} "-aux-triple" "[[ARCH2]]-{{.*}}"{{.*}} "-fsycl-is-host"
91+
// RUN: | FileCheck -DARCH=spir64_fpga -check-prefix=CHK-SYCL-FPGA-AUX-TRIPLE %s
92+
// CHK-SYCL-FPGA-AUX-TRIPLE: clang{{.*}} "-cc1" "-triple" "{{.*}}"{{.*}} "-aux-triple" "[[ARCH]]-{{.*}}"{{.*}} "-fsycl-is-host"
9393
/// ###########################################################################
9494

9595
/// Validate SYCL option values

‎clang/test/SemaSYCL/sycl-intelfpga.cpp

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,13 +5,14 @@
55

66
// Tests that we do not issue errors for _Bitints of size greater than 128
77
// when -fintelfpga is enabled. The backend is expected to be able to handle
8-
// this. When -fintelfpga is not passed, we continue to diagnose.
8+
// this, upto a maximum size of 4096. When -fintelfpga is not passed,
9+
// we continue to diagnose size greater than 128.
910

10-
// device-intelfpga-error@+4 3{{signed _BitInt of bit sizes greater than 2048 not supported}}
11-
// host-intelfpga-error@+3 3{{signed _BitInt of bit sizes greater than 2048 not supported}}
11+
// device-intelfpga-error@+4 3{{signed _BitInt of bit sizes greater than 4096 not supported}}
12+
// host-intelfpga-error@+3 3{{signed _BitInt of bit sizes greater than 4096 not supported}}
1213
// device-error@+2 3{{signed _BitInt of bit sizes greater than 128 not supported}}
1314
// host-error@+1 3{{signed _BitInt of bit sizes greater than 128 not supported}}
14-
signed _BitInt(2049) foo(signed _BitInt(2049) a, signed _BitInt(2049) b) {
15+
signed _BitInt(4097) foo(signed _BitInt(4097) a, signed _BitInt(4097) b) {
1516
return a / b;
1617
}
1718
// device-error@+4 3{{signed _BitInt of bit sizes greater than 128 not supported}}

‎llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//===----------- ESIMDUtils.hpp - ESIMD t-forms-related utility functions ===//
1+
//===------------ ESIMDUtils.h - ESIMD utility functions ------------------===//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.
@@ -35,5 +35,15 @@ void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF,
3535
// Tells whether given function is a ESIMD kernel.
3636
bool isESIMDKernel(const Function &F);
3737

38+
/// Reports and error with the message \p Msg concatenated with the optional
39+
/// \p OptMsg if \p Condition is false.
40+
inline void assert_and_diag(bool Condition, StringRef Msg,
41+
StringRef OptMsg = "") {
42+
if (!Condition) {
43+
auto T = Twine(Msg) + OptMsg;
44+
llvm::report_fatal_error(T, true /* crash diagnostics */);
45+
}
46+
}
47+
3848
} // namespace esimd
3949
} // namespace llvm

‎llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,13 @@
1+
//===------------ ESIMDUtils.cpp - ESIMD utility functions ----------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// Utility functions for processing ESIMD code.
9+
//===----------------------------------------------------------------------===//
10+
111
#include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
212

313
#include "llvm/ADT/SmallPtrSet.h"

‎llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

Lines changed: 34 additions & 42 deletions
Original file line numberDiff line numberDiff line change
@@ -671,10 +671,8 @@ static const ESIMDIntrinDesc &getIntrinDesc(StringRef SrcSpelling) {
671671
const auto &Table = getIntrinTable();
672672
auto It = Table.find(SrcSpelling.str());
673673

674-
if (It == Table.end()) {
675-
Twine Msg("unknown ESIMD intrinsic: " + SrcSpelling);
676-
llvm::report_fatal_error(Msg, false /*no crash diag*/);
677-
}
674+
llvm::esimd::assert_and_diag(It != Table.end(),
675+
"unknown ESIMD intrinsic: ", SrcSpelling);
678676
return It->second;
679677
}
680678

@@ -926,7 +924,7 @@ struct UpdateUint64MetaDataToMaxValue {
926924
: M(M), Key(Key), NewVal(NewVal) {
927925
// Pre-select nodes for update to do less work in the '()' operator.
928926
llvm::NamedMDNode *GenXKernelMD = M.getNamedMetadata(GENX_KERNEL_METADATA);
929-
assert(GenXKernelMD && "invalid genx.kernels metadata");
927+
llvm::esimd::assert_and_diag(GenXKernelMD, "invalid genx.kernels metadata");
930928
for (auto Node : GenXKernelMD->operands()) {
931929
if (Node->getNumOperands() <= (unsigned)Key) {
932930
continue;
@@ -969,9 +967,8 @@ struct UpdateUint64MetaDataToMaxValue {
969967
static void translateSLMInit(CallInst &CI) {
970968
auto F = CI.getFunction();
971969
auto *ArgV = CI.getArgOperand(0);
972-
if (!isa<ConstantInt>(ArgV))
973-
llvm::report_fatal_error(llvm::Twine(__FILE__ " ") +
974-
"integral constant is expected for slm size");
970+
llvm::esimd::assert_and_diag(isa<ConstantInt>(ArgV), __FILE__,
971+
" integral constant is expected for slm size");
975972

976973
uint64_t NewVal = cast<llvm::ConstantInt>(ArgV)->getZExtValue();
977974
assert(NewVal != 0 && "zero slm bytes being requested");
@@ -985,10 +982,9 @@ static void translateSLMInit(CallInst &CI) {
985982
static void translateNbarrierInit(CallInst &CI) {
986983
auto F = CI.getFunction();
987984
auto *ArgV = CI.getArgOperand(0);
988-
if (!isa<ConstantInt>(ArgV))
989-
llvm::report_fatal_error(
990-
llvm::Twine(__FILE__ " ") +
991-
"integral constant is expected for named barrier count");
985+
llvm::esimd::assert_and_diag(
986+
isa<ConstantInt>(ArgV), __FILE__,
987+
" integral constant is expected for named barrier count");
992988

993989
auto NewVal = cast<llvm::ConstantInt>(ArgV)->getZExtValue();
994990
assert(NewVal != 0 && "zero named barrier count being requested");
@@ -1000,20 +996,18 @@ static void translateNbarrierInit(CallInst &CI) {
1000996
static void translatePackMask(CallInst &CI) {
1001997
using Demangler = id::ManglingParser<SimpleAllocator>;
1002998
Function *F = CI.getCalledFunction();
1003-
assert(F && "function to translate is invalid");
999+
llvm::esimd::assert_and_diag(F, "function to translate is invalid");
10041000

10051001
StringRef MnglName = F->getName();
10061002
Demangler Parser(MnglName.begin(), MnglName.end());
10071003
id::Node *AST = Parser.parse();
10081004

1009-
if (!AST || !Parser.ForwardTemplateRefs.empty()) {
1010-
Twine Msg("failed to demangle ESIMD intrinsic: " + MnglName);
1011-
llvm::report_fatal_error(Msg, false /*no crash diag*/);
1012-
}
1013-
if (AST->getKind() != id::Node::KFunctionEncoding) {
1014-
Twine Msg("bad ESIMD intrinsic: " + MnglName);
1015-
llvm::report_fatal_error(Msg, false /*no crash diag*/);
1016-
}
1005+
llvm::esimd::assert_and_diag(
1006+
AST && Parser.ForwardTemplateRefs.empty(),
1007+
"failed to demangle ESIMD intrinsic: ", MnglName);
1008+
llvm::esimd::assert_and_diag(AST->getKind() == id::Node::KFunctionEncoding,
1009+
"bad ESIMD intrinsic: ", MnglName);
1010+
10171011
auto *FE = static_cast<id::FunctionEncoding *>(AST);
10181012
llvm::LLVMContext &Context = CI.getContext();
10191013
Type *TTy = nullptr;
@@ -1042,19 +1036,17 @@ static void translatePackMask(CallInst &CI) {
10421036
static void translateUnPackMask(CallInst &CI) {
10431037
using Demangler = id::ManglingParser<SimpleAllocator>;
10441038
Function *F = CI.getCalledFunction();
1045-
assert(F && "function to translate is invalid");
1039+
llvm::esimd::assert_and_diag(F, "function to translate is invalid");
10461040
StringRef MnglName = F->getName();
10471041
Demangler Parser(MnglName.begin(), MnglName.end());
10481042
id::Node *AST = Parser.parse();
10491043

1050-
if (!AST || !Parser.ForwardTemplateRefs.empty()) {
1051-
Twine Msg("failed to demangle ESIMD intrinsic: " + MnglName);
1052-
llvm::report_fatal_error(Msg, false /*no crash diag*/);
1053-
}
1054-
if (AST->getKind() != id::Node::KFunctionEncoding) {
1055-
Twine Msg("bad ESIMD intrinsic: " + MnglName);
1056-
llvm::report_fatal_error(Msg, false /*no crash diag*/);
1057-
}
1044+
llvm::esimd::assert_and_diag(
1045+
AST && Parser.ForwardTemplateRefs.empty(),
1046+
"failed to demangle ESIMD intrinsic: ", MnglName);
1047+
llvm::esimd::assert_and_diag(AST->getKind() == id::Node::KFunctionEncoding,
1048+
"bad ESIMD intrinsic: ", MnglName);
1049+
10581050
auto *FE = static_cast<id::FunctionEncoding *>(AST);
10591051
llvm::LLVMContext &Context = CI.getContext();
10601052
Type *TTy = nullptr;
@@ -1194,8 +1186,9 @@ void translateFmuladd(CallInst *CI) {
11941186

11951187
// Translates an LLVM intrinsic to a form, digestable by the BE.
11961188
bool translateLLVMIntrinsic(CallInst *CI) {
1197-
Function *F = CI->getCalledFunction() ? CI->getCalledFunction() : nullptr;
1198-
assert(F && F->isIntrinsic());
1189+
Function *F = CI->getCalledFunction();
1190+
llvm::esimd::assert_and_diag(F && F->isIntrinsic(),
1191+
"malformed llvm intrinsic call");
11991192

12001193
switch (F->getIntrinsicID()) {
12011194
case Intrinsic::assume:
@@ -1277,7 +1270,8 @@ translateSpirvGlobalUses(LoadInst *LI, StringRef SpirvGlobalName,
12771270
NewInst = generateGenXCall(EEI, "group.count", true);
12781271
}
12791272

1280-
assert(NewInst && "Load from global SPIRV builtin was not translated");
1273+
llvm::esimd::assert_and_diag(
1274+
NewInst, "Load from global SPIRV builtin was not translated");
12811275
EEI->replaceAllUsesWith(NewInst);
12821276
InstsToErase.push_back(EEI);
12831277
}
@@ -1437,19 +1431,17 @@ static Function *createTestESIMDDeclaration(const ESIMDIntrinDesc &Desc,
14371431
static void translateESIMDIntrinsicCall(CallInst &CI) {
14381432
using Demangler = id::ManglingParser<SimpleAllocator>;
14391433
Function *F = CI.getCalledFunction();
1440-
assert(F && "function to translate is invalid");
1434+
llvm::esimd::assert_and_diag(F, "function to translate is invalid");
14411435
StringRef MnglName = F->getName();
14421436
Demangler Parser(MnglName.begin(), MnglName.end());
14431437
id::Node *AST = Parser.parse();
14441438

1445-
if (!AST || !Parser.ForwardTemplateRefs.empty()) {
1446-
Twine Msg("failed to demangle ESIMD intrinsic: " + MnglName);
1447-
llvm::report_fatal_error(Msg, false /*no crash diag*/);
1448-
}
1449-
if (AST->getKind() != id::Node::KFunctionEncoding) {
1450-
Twine Msg("bad ESIMD intrinsic: " + MnglName);
1451-
llvm::report_fatal_error(Msg, false /*no crash diag*/);
1452-
}
1439+
llvm::esimd::assert_and_diag(
1440+
AST && Parser.ForwardTemplateRefs.empty(),
1441+
"failed to demangle ESIMD intrinsic: ", MnglName);
1442+
llvm::esimd::assert_and_diag(AST->getKind() == id::Node::KFunctionEncoding,
1443+
"bad ESIMD intrinsic: ", MnglName);
1444+
14531445
auto *FE = static_cast<id::FunctionEncoding *>(AST);
14541446
id::StringView BaseNameV = FE->getName()->getBaseName();
14551447

‎sycl/include/CL/sycl.hpp

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -60,9 +60,7 @@
6060
#if SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO
6161
#include <sycl/ext/oneapi/backend/level_zero.hpp>
6262
#endif
63-
#include <sycl/ext/oneapi/bf16_storage_builtins.hpp>
6463
#include <sycl/ext/oneapi/device_global/properties.hpp>
65-
#include <sycl/ext/oneapi/experimental/builtins.hpp>
6664
#include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
6765
#include <sycl/ext/oneapi/filter_selector.hpp>
6866
#include <sycl/ext/oneapi/group_algorithm.hpp>

0 commit comments

Comments
 (0)