Fznamznon updated this revision to Diff 266066.
Fznamznon added a comment.
Herald added a subscriber: sstefan1.
Re-implemented diagnostic itself, now only usages of declarations
with unsupported types are diagnosed.
Generalized approach between OpenMP and SYCL.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D74387/new/
https://reviews.llvm.org/D74387
Files:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Sema/Sema.h
clang/lib/Sema/CMakeLists.txt
clang/lib/Sema/Sema.cpp
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaDeclCXX.cpp
clang/lib/Sema/SemaExpr.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/SemaSYCL.cpp
clang/lib/Sema/SemaType.cpp
clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
clang/test/SemaSYCL/float128.cpp
Index: clang/test/SemaSYCL/float128.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/float128.cpp
@@ -0,0 +1,96 @@
+// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsycl -fsycl-is-device -fsyntax-only %s
+
+typedef __float128 BIGTY;
+
+template <class T>
+class Z {
+public:
+ // expected-note@+1 {{'field' defined here}}
+ T field;
+ // expected-note@+1 2{{'field1' defined here}}
+ __float128 field1;
+ using BIGTYPE = __float128;
+ // expected-note@+1 {{'bigfield' defined here}}
+ BIGTYPE bigfield;
+};
+
+void host_ok(void) {
+ __float128 A;
+ int B = sizeof(__float128);
+ Z<__float128> C;
+ C.field1 = A;
+}
+
+void usage() {
+ // expected-note@+1 3{{'A' defined here}}
+ __float128 A;
+ Z<__float128> C;
+ // expected-error@+2 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ C.field1 = A;
+ // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__float128') type support, but device 'spir64' does not support it}}
+ C.bigfield += 1.0;
+
+ // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ auto foo1 = [=]() {
+ __float128 AA;
+ // expected-note@+2 {{'BB' defined here}}
+ // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ auto BB = A;
+ // expected-error@+1 {{'BB' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ BB += 1;
+ };
+
+ // expected-note@+1 {{called by 'usage'}}
+ foo1();
+}
+
+template <typename t>
+void foo2(){};
+
+// expected-note@+3 {{'P' defined here}}
+// expected-error@+2 {{'P' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+// expected-note@+1 2{{'foo' defined here}}
+__float128 foo(__float128 P) { return P; }
+
+template <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+ // expected-note@+1 5{{called by 'kernel}}
+ kernelFunc();
+}
+
+int main() {
+ // expected-note@+1 {{'CapturedToDevice' defined here}}
+ __float128 CapturedToDevice = 1;
+ host_ok();
+ kernel<class variables>([=]() {
+ decltype(CapturedToDevice) D;
+ // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ auto C = CapturedToDevice;
+ Z<__float128> S;
+ // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ S.field1 += 1;
+ // expected-error@+1 {{'field' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ S.field = 1;
+ });
+
+ kernel<class functions>([=]() {
+ // expected-note@+1 2{{called by 'operator()'}}
+ usage();
+ // expected-note@+1 {{'BBBB' defined here}}
+ BIGTY BBBB;
+ // expected-note@+3 {{called by 'operator()'}}
+ // expected-error@+2 2{{'foo' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+ // expected-error@+1 {{'BBBB' requires 128 bit size 'BIGTY' (aka '__float128') type support, but device 'spir64' does not support it}}
+ auto A = foo(BBBB);
+ });
+
+ kernel<class ok>([=]() {
+ Z<__float128> S;
+ foo2<__float128>();
+ auto A = sizeof(CapturedToDevice);
+ });
+
+ return 0;
+}
Index: clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
===================================================================
--- clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
+++ clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
@@ -7,18 +7,23 @@
struct T {
char a;
#ifndef _ARCH_PPC
+ // expected-note@+1 {{'f' defined here}}
__float128 f;
#else
+ // expected-note@+1 {{'f' defined here}}
long double f;
#endif
char c;
T() : a(12), f(15) {}
#ifndef _ARCH_PPC
-// expected-error@+4 {{host requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+5 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
#else
-// expected-error@+2 {{host requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+3 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
#endif
- T &operator+(T &b) { f += b.a; return *this;}
+ T &operator+(T &b) {
+ f += b.a;
+ return *this;
+ }
};
struct T1 {
@@ -27,19 +32,36 @@
__int128 f1;
char c;
T1() : a(12), f(15) {}
- T1 &operator/(T1 &b) { f /= b.a; return *this;}
+ T1 &operator/(T1 &b) {
+ f /= b.a;
+ return *this;
+ }
};
+#ifndef _ARCH_PPC
+// expected-note@+1 {{'boo' defined here}}
+void boo(__float128 A) { return; }
+#else
+// expected-note@+1 {{'boo' defined here}}
+void boo(long double A) { return; }
+#endif
#pragma omp declare target
T a = T();
T f = a;
void foo(T a = T()) {
a = a + f; // expected-note {{called by 'foo'}}
+#ifndef _ARCH_PPC
+// expected-error@+4 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+#else
+// expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+#endif
+ boo(0);
return;
}
T bar() {
return T();
}
+
void baz() {
T t = bar();
}
@@ -56,3 +78,34 @@
T1 t = bar1();
}
#pragma omp end declare target
+
+#ifndef _ARCH_PPC
+// expected-note@+1 3{{'f' defined here}}
+__float128 foo1(__float128 f) {
+#pragma omp target map(f)
+ // expected-error@+1 3{{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+ f = 1;
+ return f;
+}
+#else
+// expected-note@+1 3{{'f' defined here}}
+long double foo1(long double f) {
+#pragma omp target map(f)
+ // expected-error@+1 3{{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+ f = 1;
+ return f;
+}
+#endif
+
+// Allow all sorts of stuff on host
+#ifndef _ARCH_PPC
+__float128 q, b;
+__float128 c = q + b;
+#else
+long double q, b;
+long double c = q + b;
+#endif
+
+void hostFoo() {
+ boo(c - b);
+}
Index: clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
===================================================================
--- clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
+++ clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp
@@ -71,11 +71,3 @@
}
#pragma omp end declare target
-BIGTYPE foo(BIGTYPE f) {
-#pragma omp target map(f)
- f = 1;
- return f;
-}
-
-// CHECK: define weak void @__omp_offloading_{{.+}}foo{{.+}}_l75([[BIGTYPE:.+]]*
-// CHECK: store [[BIGTYPE]] {{0xL00000000000000003FFF000000000000|0xM3FF00000000000000000000000000000}}, [[BIGTYPE]]* %
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ clang/lib/Sema/SemaType.cpp
@@ -1530,6 +1530,7 @@
break;
case DeclSpec::TST_float128:
if (!S.Context.getTargetInfo().hasFloat128Type() &&
+ !S.getLangOpts().SYCLIsDevice &&
!(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "__float128";
Index: clang/lib/Sema/SemaSYCL.cpp
===================================================================
--- /dev/null
+++ clang/lib/Sema/SemaSYCL.cpp
@@ -0,0 +1,49 @@
+//===- SemaSYCL.cpp - Semantic Analysis for SYCL constructs ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// This implements Semantic Analysis for SYCL constructs.
+//===----------------------------------------------------------------------===//
+
+#include "clang/Sema/Sema.h"
+#include "clang/Sema/SemaDiagnostic.h"
+
+using namespace clang;
+
+// -----------------------------------------------------------------------------
+// SYCL device specific diagnostics implementation
+// -----------------------------------------------------------------------------
+
+Sema::DeviceDiagBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc,
+ unsigned DiagID) {
+ assert(getLangOpts().SYCLIsDevice &&
+ "Should only be called during SYCL compilation");
+ FunctionDecl *FD = dyn_cast<FunctionDecl>(getCurLexicalContext());
+ DeviceDiagBuilder::Kind DiagKind = [this, FD] {
+ if (!FD)
+ return DeviceDiagBuilder::K_Nop;
+ if (getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted)
+ return DeviceDiagBuilder::K_ImmediateWithCallStack;
+ return DeviceDiagBuilder::K_Deferred;
+ }();
+ return DeviceDiagBuilder(DiagKind, Loc, DiagID, FD, *this);
+}
+
+bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) {
+ assert(getLangOpts().SYCLIsDevice &&
+ "Should only be called during SYCL compilation");
+ assert(Callee && "Callee may not be null.");
+
+ // Errors in unevaluated context don't need to be generated,
+ // so we can safely skip them.
+ if (isUnevaluatedContext() || isConstantEvaluated())
+ return true;
+
+ DeviceDiagBuilder::Kind DiagKind = DeviceDiagBuilder::K_Nop;
+
+ return DiagKind != DeviceDiagBuilder::K_Immediate &&
+ DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack;
+}
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -1832,23 +1832,28 @@
unsigned DiagID) {
assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
"Expected OpenMP device compilation.");
- FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl());
+
+ FunctionDecl *FD = getCurFunctionDecl();
DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
- switch (FES) {
- case FunctionEmissionStatus::Emitted:
- Kind = DeviceDiagBuilder::K_Immediate;
- break;
- case FunctionEmissionStatus::Unknown:
- Kind = isOpenMPDeviceDelayedContext(*this) ? DeviceDiagBuilder::K_Deferred
- : DeviceDiagBuilder::K_Immediate;
- break;
- case FunctionEmissionStatus::TemplateDiscarded:
- case FunctionEmissionStatus::OMPDiscarded:
- Kind = DeviceDiagBuilder::K_Nop;
- break;
- case FunctionEmissionStatus::CUDADiscarded:
- llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation");
- break;
+ if (FD) {
+ FunctionEmissionStatus FES = getEmissionStatus(FD);
+ switch (FES) {
+ case FunctionEmissionStatus::Emitted:
+ Kind = DeviceDiagBuilder::K_Immediate;
+ break;
+ case FunctionEmissionStatus::Unknown:
+ Kind = isOpenMPDeviceDelayedContext(*this)
+ ? DeviceDiagBuilder::K_Deferred
+ : DeviceDiagBuilder::K_Immediate;
+ break;
+ case FunctionEmissionStatus::TemplateDiscarded:
+ case FunctionEmissionStatus::OMPDiscarded:
+ Kind = DeviceDiagBuilder::K_Nop;
+ break;
+ case FunctionEmissionStatus::CUDADiscarded:
+ llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation");
+ break;
+ }
}
return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
@@ -1877,21 +1882,6 @@
return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
}
-void Sema::checkOpenMPDeviceExpr(const Expr *E) {
- assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
- "OpenMP device compilation mode is expected.");
- QualType Ty = E->getType();
- if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
- ((Ty->isFloat128Type() ||
- (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) &&
- !Context.getTargetInfo().hasFloat128Type()) ||
- (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
- !Context.getTargetInfo().hasInt128Type()))
- targetDiag(E->getExprLoc(), diag::err_omp_unsupported_type)
- << static_cast<unsigned>(Context.getTypeSize(Ty)) << Ty
- << Context.getTargetInfo().getTriple().str() << E->getSourceRange();
-}
-
static OpenMPDefaultmapClauseKind
getVariableCategoryFromDecl(const LangOptions &LO, const ValueDecl *VD) {
if (LO.OpenMP <= 45) {
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -293,6 +293,9 @@
if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD))
return true;
+
+ if (getLangOpts().SYCLIsDevice && !checkSYCLDeviceFunction(Loc, FD))
+ return true;
}
if (auto *MD = dyn_cast<CXXMethodDecl>(D)) {
@@ -352,6 +355,10 @@
diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
+ if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+ if (const auto *VD = dyn_cast<ValueDecl>(D))
+ checkDeviceDecl(VD, Loc);
+
if (isa<ParmVarDecl>(D) && isa<RequiresExprBodyDecl>(D->getDeclContext()) &&
!isUnevaluatedContext()) {
// C++ [expr.prim.req.nested] p3
@@ -13510,14 +13517,6 @@
}
}
- // Diagnose operations on the unsupported types for OpenMP device compilation.
- if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice) {
- if (Opc != BO_Assign && Opc != BO_Comma) {
- checkOpenMPDeviceExpr(LHSExpr);
- checkOpenMPDeviceExpr(RHSExpr);
- }
- }
-
switch (Opc) {
case BO_Assign:
ResultTy = CheckAssignmentOperands(LHS.get(), RHS, OpLoc, QualType());
@@ -14130,12 +14129,6 @@
<< Input.get()->getSourceRange());
}
}
- // Diagnose operations on the unsupported types for OpenMP device compilation.
- if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice) {
- if (UnaryOperator::isIncrementDecrementOp(Opc) ||
- UnaryOperator::isArithmeticOp(Opc))
- checkOpenMPDeviceExpr(InputExpr);
- }
switch (Opc) {
case UO_PreInc:
@@ -16394,6 +16387,9 @@
if (getLangOpts().CUDA)
CheckCUDACall(Loc, Func);
+ if (getLangOpts().SYCLIsDevice)
+ checkSYCLDeviceFunction(Loc, Func);
+
// If we need a definition, try to create one.
if (NeedDefinition && !Func->getBody()) {
runWithSufficientStackSpace(Loc, [&] {
Index: clang/lib/Sema/SemaDeclCXX.cpp
===================================================================
--- clang/lib/Sema/SemaDeclCXX.cpp
+++ clang/lib/Sema/SemaDeclCXX.cpp
@@ -14915,6 +14915,9 @@
MarkFunctionReferenced(ConstructLoc, Constructor);
if (getLangOpts().CUDA && !CheckCUDACall(ConstructLoc, Constructor))
return ExprError();
+ if (getLangOpts().SYCLIsDevice &&
+ !checkSYCLDeviceFunction(ConstructLoc, Constructor))
+ return ExprError();
return CheckForImmediateInvocation(
CXXConstructExpr::Create(
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -14423,7 +14423,7 @@
DiscardCleanupsInEvaluationContext();
}
- if (LangOpts.OpenMP || LangOpts.CUDA) {
+ if (LangOpts.OpenMP || LangOpts.CUDA || LangOpts.SYCLIsDevice) {
auto ES = getEmissionStatus(FD);
if (ES == Sema::FunctionEmissionStatus::Emitted ||
ES == Sema::FunctionEmissionStatus::Unknown)
@@ -18103,6 +18103,11 @@
Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
bool Final) {
+ // SYCL functions can be template, so we check if they have appropriate
+ // attribute prior to checking if it is a template.
+ if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelAttr>())
+ return FunctionEmissionStatus::Emitted;
+
// Templates are emitted when they're instantiated.
if (FD->isDependentContext())
return FunctionEmissionStatus::TemplateDiscarded;
Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -1698,10 +1698,56 @@
if (getLangOpts().CUDA)
return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID)
: CUDADiagIfHostCode(Loc, DiagID);
+
+ if (getLangOpts().SYCLIsDevice)
+ return SYCLDiagIfDeviceCode(Loc, DiagID);
+
return DeviceDiagBuilder(DeviceDiagBuilder::K_Immediate, Loc, DiagID,
getCurFunctionDecl(), *this);
}
+void Sema::checkDeviceDecl(const ValueDecl *D, SourceLocation Loc) {
+ if (isUnevaluatedContext())
+ return;
+
+ Decl *C = cast<Decl>(getCurLexicalContext());
+
+ // Memcpy operations for structs containing a member with unsupported type
+ // are ok, though.
+ if (const auto *MD = dyn_cast<CXXMethodDecl>(C)) {
+ if ((MD->isCopyAssignmentOperator() || MD->isMoveAssignmentOperator()) &&
+ MD->isTrivial())
+ return;
+
+ if (const auto *Ctor = dyn_cast<CXXConstructorDecl>(MD))
+ if (Ctor->isCopyOrMoveConstructor() && Ctor->isTrivial())
+ return;
+ }
+
+ QualType Ty = D->getType();
+ auto CheckType = [&](QualType Ty) {
+ if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
+ ((Ty->isFloat128Type() ||
+ (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) &&
+ !Context.getTargetInfo().hasFloat128Type()) ||
+ (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
+ !Context.getTargetInfo().hasInt128Type())) {
+ targetDiag(Loc, diag::err_device_unsupported_type)
+ << D << static_cast<unsigned>(Context.getTypeSize(Ty)) << Ty
+ << Context.getTargetInfo().getTriple().str();
+ targetDiag(D->getLocation(), diag::note_defined_here) << D;
+ }
+ };
+
+ CheckType(Ty);
+
+ if (const auto *FPTy = dyn_cast<FunctionProtoType>(Ty)) {
+ for (const auto &ParamTy : FPTy->param_types())
+ CheckType(ParamTy);
+ CheckType(FPTy->getReturnType());
+ }
+}
+
/// Looks through the macro-expansion chain for the given
/// location, looking for a macro expansion with the given name.
/// If one is found, returns true and sets the location to that
Index: clang/lib/Sema/CMakeLists.txt
===================================================================
--- clang/lib/Sema/CMakeLists.txt
+++ clang/lib/Sema/CMakeLists.txt
@@ -61,6 +61,7 @@
SemaStmt.cpp
SemaStmtAsm.cpp
SemaStmtAttr.cpp
+ SemaSYCL.cpp
SemaTemplate.cpp
SemaTemplateDeduction.cpp
SemaTemplateInstantiate.cpp
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -9867,10 +9867,6 @@
/// Pop OpenMP function region for non-capturing function.
void popOpenMPFunctionRegion(const sema::FunctionScopeInfo *OldFSI);
- /// Check if the expression is allowed to be used in expressions for the
- /// OpenMP devices.
- void checkOpenMPDeviceExpr(const Expr *E);
-
/// Checks if a type or a declaration is disabled due to the owning extension
/// being disabled, and emits diagnostic messages if it is disabled.
/// \param D type or declaration to be checked.
@@ -11650,6 +11646,10 @@
DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
+ /// Check if the expression is allowed to be used in expressions for the
+ /// offloading devices.
+ void checkDeviceDecl(const ValueDecl *D, SourceLocation Loc);
+
enum CUDAFunctionTarget {
CFT_Device,
CFT_Global,
@@ -12392,6 +12392,40 @@
ConstructorDestructor,
BuiltinFunction
};
+ /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+ /// context is "used as device code".
+ ///
+ /// - If CurLexicalContext is a kernel function or it is known that the
+ /// function will be emitted for the device, emits the diagnostics
+ /// immediately.
+ /// - If CurLexicalContext is a function and we are compiling
+ /// for the device, but we don't know that this function will be codegen'ed
+ /// for devive yet, creates a diagnostic which is emitted if and when we
+ /// realize that the function will be codegen'ed.
+ ///
+ /// Example usage:
+ ///
+ /// Diagnose __float128 type usage only from SYCL device code if the current
+ /// target doesn't support it
+ /// if (!S.Context.getTargetInfo().hasFloat128Type() &&
+ /// S.getLangOpts().SYCLIsDevice)
+ /// SYCLDiagIfDeviceCode(Loc, diag::err_type_unsupported) << "__float128";
+ DeviceDiagBuilder SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
+
+ /// Check whether we're allowed to call Callee from the current context.
+ ///
+ /// - If the call is never allowed in a semantically-correct program
+ /// emits an error and returns false.
+ ///
+ /// - If the call is allowed in semantically-correct programs, but only if
+ /// it's never codegen'ed, creates a deferred diagnostic to be emitted if
+ /// and when the caller is codegen'ed, and returns true.
+ ///
+ /// - Otherwise, returns true without emitting any diagnostics.
+ ///
+ /// Adds Callee to DeviceCallGraph if we don't know if its caller will be
+ /// codegen'ed yet.
+ bool checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee);
};
/// RAII object that enters a new expression evaluation context.
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10204,8 +10204,8 @@
"expected loop invariant expression or '<invariant1> * %0 + <invariant2>' kind of expression">;
def err_omp_wrong_dependency_iterator_type : Error<
"expected an integer or a pointer type of the outer loop counter '%0' for non-rectangular nests">;
-def err_omp_unsupported_type : Error <
- "host requires %0 bit size %1 type support, but device '%2' does not support it">;
+def err_device_unsupported_type : Error <
+ "%0 requires %1 bit size %2 type support, but device '%3' does not support it">;
def err_omp_lambda_capture_in_declare_target_not_to : Error<
"variable captured in declare target region must appear in a to clause">;
def err_omp_device_type_mismatch : Error<
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
