diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -11948,8 +11948,8 @@ /// if (diagIfOpenMPDeviceCode(Loc, diag::err_vla_unsupported)) /// return ExprError(); /// // Otherwise, continue parsing as normal. - SemaDiagnosticBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, - unsigned DiagID); + SemaDiagnosticBuilder + diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID, FunctionDecl *FD); /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current /// context is "used as host code". @@ -11965,17 +11965,19 @@ /// return ExprError(); /// // Otherwise, continue parsing as normal. SemaDiagnosticBuilder diagIfOpenMPHostCode(SourceLocation Loc, - unsigned DiagID); + unsigned DiagID, FunctionDecl *FD); - SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID); + SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID, + FunctionDecl *FD = nullptr); SemaDiagnosticBuilder targetDiag(SourceLocation Loc, - const PartialDiagnostic &PD) { - return targetDiag(Loc, PD.getDiagID()) << PD; + const PartialDiagnostic &PD, + FunctionDecl *FD = nullptr) { + return targetDiag(Loc, PD.getDiagID(), FD) << PD; } /// Check if the expression is allowed to be used in expressions for the /// offloading devices. - void checkDeviceDecl(const ValueDecl *D, SourceLocation Loc); + void checkDeviceDecl(ValueDecl *D, SourceLocation Loc); enum CUDAFunctionTarget { CFT_Device, diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -14,6 +14,7 @@ #include "UsedDeclVisitor.h" #include "clang/AST/ASTContext.h" #include "clang/AST/ASTDiagnostic.h" +#include "clang/AST/Decl.h" #include "clang/AST/DeclCXX.h" #include "clang/AST/DeclFriend.h" #include "clang/AST/DeclObjC.h" @@ -1733,11 +1734,12 @@ } } -Sema::SemaDiagnosticBuilder Sema::targetDiag(SourceLocation Loc, - unsigned DiagID) { +Sema::SemaDiagnosticBuilder +Sema::targetDiag(SourceLocation Loc, unsigned DiagID, FunctionDecl *FD) { + FD = FD ? FD : getCurFunctionDecl(); if (LangOpts.OpenMP) - return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID) - : diagIfOpenMPHostCode(Loc, DiagID); + return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID, FD) + : diagIfOpenMPHostCode(Loc, DiagID, FD); if (getLangOpts().CUDA) return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID) : CUDADiagIfHostCode(Loc, DiagID); @@ -1746,7 +1748,7 @@ return SYCLDiagIfDeviceCode(Loc, DiagID); return SemaDiagnosticBuilder(SemaDiagnosticBuilder::K_Immediate, Loc, DiagID, - getCurFunctionDecl(), *this); + FD, *this); } Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, unsigned DiagID, @@ -1765,15 +1767,14 @@ DiagID, getCurFunctionDecl(), *this); } - SemaDiagnosticBuilder DB = - getLangOpts().CUDAIsDevice - ? CUDADiagIfDeviceCode(Loc, DiagID) - : CUDADiagIfHostCode(Loc, DiagID); + SemaDiagnosticBuilder DB = getLangOpts().CUDAIsDevice + ? CUDADiagIfDeviceCode(Loc, DiagID) + : CUDADiagIfHostCode(Loc, DiagID); SetIsLastErrorImmediate(DB.isImmediate()); return DB; } -void Sema::checkDeviceDecl(const ValueDecl *D, SourceLocation Loc) { +void Sema::checkDeviceDecl(ValueDecl *D, SourceLocation Loc) { if (isUnevaluatedContext()) return; @@ -1791,13 +1792,17 @@ return; } + // Try to associate errors with the lexical context, if that is a function, or + // the value declaration otherwise. + FunctionDecl *FD = + isa(C) ? cast(C) : dyn_cast(D); auto CheckType = [&](QualType Ty) { if (Ty->isDependentType()) return; if (Ty->isExtIntType()) { if (!Context.getTargetInfo().hasExtIntType()) { - targetDiag(Loc, diag::err_device_unsupported_type) + targetDiag(Loc, diag::err_device_unsupported_type, FD) << D << false /*show bit size*/ << 0 /*bitsize*/ << Ty << Context.getTargetInfo().getTriple().str(); } @@ -1810,11 +1815,12 @@ !Context.getTargetInfo().hasFloat128Type()) || (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 && !Context.getTargetInfo().hasInt128Type())) { - targetDiag(Loc, diag::err_device_unsupported_type) + if (targetDiag(Loc, diag::err_device_unsupported_type, FD) << D << true /*show bit size*/ << static_cast(Context.getTypeSize(Ty)) << Ty - << Context.getTargetInfo().getTriple().str(); - targetDiag(D->getLocation(), diag::note_defined_here) << D; + << Context.getTargetInfo().getTriple().str()) + D->setInvalidDecl(); + targetDiag(D->getLocation(), diag::note_defined_here, FD) << D; } }; @@ -1826,6 +1832,8 @@ CheckType(ParamTy); CheckType(FPTy->getReturnType()); } + if (const auto *FPTy = dyn_cast(Ty)) + CheckType(FPTy->getReturnType()); } /// Looks through the macro-expansion chain for the given diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -9418,6 +9418,9 @@ } } + if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) + checkDeviceDecl(NewFD, D.getBeginLoc()); + if (!getLangOpts().CPlusPlus) { // Perform semantic checking on the function declaration. if (!NewFD->isInvalidDecl() && NewFD->isMain()) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -373,7 +373,7 @@ } if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) { - if (const auto *VD = dyn_cast(D)) + if (auto *VD = dyn_cast(D)) checkDeviceDecl(VD, Loc); if (!Context.getTargetInfo().isTLSSupported()) diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -1898,11 +1898,11 @@ } // anonymous namespace Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc, - unsigned DiagID) { + unsigned DiagID, + FunctionDecl *FD) { assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice && "Expected OpenMP device compilation."); - FunctionDecl *FD = getCurFunctionDecl(); SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop; if (FD) { FunctionEmissionStatus FES = getEmissionStatus(FD); @@ -1925,14 +1925,15 @@ } } - return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); + return SemaDiagnosticBuilder(Kind, Loc, DiagID, FD, *this); } Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc, - unsigned DiagID) { + unsigned DiagID, + FunctionDecl *FD) { assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice && "Expected OpenMP host compilation."); - FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl()); + FunctionEmissionStatus FES = getEmissionStatus(FD); SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop; switch (FES) { case FunctionEmissionStatus::Emitted: @@ -1948,7 +1949,7 @@ break; } - return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); + return SemaDiagnosticBuilder(Kind, Loc, DiagID, FD, *this); } static OpenMPDefaultmapClauseKind diff --git a/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp b/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp --- a/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp +++ b/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp @@ -39,10 +39,12 @@ }; #ifndef _ARCH_PPC -// expected-note@+1 {{'boo' defined here}} +// expected-error@+2 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} +// expected-note@+1 2{{'boo' defined here}} void boo(__float128 A) { return; } #else -// expected-note@+1 {{'boo' defined here}} +// expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} +// expected-note@+1 2{{'boo' defined here}} void boo(long double A) { return; } #endif #pragma omp declare target @@ -51,10 +53,11 @@ 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}} +// expected-error@+5 {{'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}} +// expected-error@+3 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} #endif +// expected-note@+1 {{called by 'foo'}} boo(0); return; } @@ -98,28 +101,49 @@ a = &b; } -// TODO: We should diagnose the return type and argument type here. +// expected-note@+2 {{'ld_return1a' defined here}} +// expected-error@+1 {{'ld_return1a' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} long double ld_return1a() { return 0; } +// expected-note@+2 {{'ld_arg1a' defined here}} +// expected-error@+1 {{'ld_arg1a' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void ld_arg1a(long double ld) {} // TODO: We should diagnose the return type and argument type here. typedef long double ld_ty; +// expected-note@+2 {{'ld_return1b' defined here}} +// expected-error@+1 {{'ld_return1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but device 'nvptx64-unknown-unknown' does not support it}} ld_ty ld_return1b() { return 0; } +// expected-note@+2 {{'ld_arg1b' defined here}} +// expected-error@+1 {{'ld_arg1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but device 'nvptx64-unknown-unknown' does not support it}} void ld_arg1b(ld_ty ld) {} +// TODO: These errors should not be emitted. +// expected-note@+2 {{'ld_return1c' defined here}} +// expected-error@+1 {{'ld_return1c' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} static long double ld_return1c() { return 0; } +// expected-note@+2 {{'ld_arg1c' defined here}} +// expected-error@+1 {{'ld_arg1c' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} static void ld_arg1c(long double ld) {} +// TODO: These errors should not be emitted. +// expected-note@+2 {{'ld_return1d' defined here}} +// expected-error@+1 {{'ld_return1d' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} inline long double ld_return1d() { return 0; } +// expected-note@+2 {{'ld_arg1d' defined here}} +// expected-error@+1 {{'ld_arg1d' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} inline void ld_arg1d(long double ld) {} +// expected-error@+2 {{'ld_return1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} // expected-note@+1 {{'ld_return1e' defined here}} static long double ld_return1e() { return 0; } +// expected-error@+2 {{'ld_arg1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} // expected-note@+1 {{'ld_arg1e' defined here}} static void ld_arg1e(long double ld) {} +// expected-error@+2 {{'ld_return1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} // expected-note@+1 {{'ld_return1f' defined here}} inline long double ld_return1f() { return 0; } +// expected-error@+2 {{'ld_arg1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} // expected-note@+1 {{'ld_arg1f' defined here}} inline void ld_arg1f(long double ld) {} @@ -152,46 +176,47 @@ } void external() { -// expected-error@+1 {{'ld_return1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void *p1 = reinterpret_cast(&ld_return1e); -// expected-error@+1 {{'ld_arg1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void *p2 = reinterpret_cast(&ld_arg1e); -// expected-error@+1 {{'ld_return1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void *p3 = reinterpret_cast(&ld_return1f); -// expected-error@+1 {{'ld_arg1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} void *p4 = reinterpret_cast(&ld_arg1f); void *p5 = reinterpret_cast(&ld_use3); void *p6 = reinterpret_cast(&ld_use4); } #ifndef _ARCH_PPC -// TODO: We should diagnose the return type and argument type here. +// expected-note@+2 {{'ld_return2a' defined here}} +// expected-error@+1 {{'ld_return2a' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} __float128 ld_return2a() { return 0; } +// expected-note@+2 {{'ld_arg2a' defined here}} +// expected-error@+1 {{'ld_arg2a' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}} void ld_arg2a(__float128 ld) {} -// TODO: We should diagnose the return type and argument type here. typedef __float128 fp128_ty; +// expected-note@+2 {{'ld_return2b' defined here}} +// expected-error@+1 {{'ld_return2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but device 'nvptx64-unknown-unknown' does not support it}} fp128_ty ld_return2b() { return 0; } +// expected-note@+2 {{'ld_arg2b' defined here}} +// expected-error@+1 {{'ld_arg2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but device 'nvptx64-unknown-unknown' does not support it}} void ld_arg2b(fp128_ty ld) {} #endif #pragma omp end declare target // TODO: There should not be an error here, dead_inline is never emitted. -// expected-note@+1 3{{'f' defined here}} +// expected-note@+1 {{'f' defined here}} inline long double dead_inline(long double f) { #pragma omp target map(f) -// TODO: We should not emit the same error message 3 times, here and elsewhere in this file. - // expected-error@+1 3{{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} + // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} f = 1; return f; } // TODO: There should not be an error here, dead_static is never emitted. -// expected-note@+1 3{{'f' defined here}} +// expected-note@+1 {{'f' defined here}} static long double dead_static(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}} + // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} f = 1; return f; } @@ -204,18 +229,18 @@ } #ifndef _ARCH_PPC -// expected-note@+1 3{{'f' defined here}} +// expected-note@+1 {{'f' defined here}} __float128 foo2(__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}} + // expected-error@+1 {{'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}} +// expected-note@+1 {{'f' defined here}} long double foo3(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}} + // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} f = 1; return f; }