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 @@ -18333,42 +18333,51 @@ if (FD->isDependentContext()) return FunctionEmissionStatus::TemplateDiscarded; - FunctionEmissionStatus OMPES = FunctionEmissionStatus::Unknown; + // Check whether this function is an externally visible definition. + auto IsEmittedForExternalSymbol = [this, FD]() { + // We have to check the GVA linkage of the function's *definition* -- if we + // only have a declaration, we don't know whether or not the function will + // be emitted, because (say) the definition could include "inline". + FunctionDecl *Def = FD->getDefinition(); + + return Def && !isDiscardableGVALinkage( + getASTContext().GetGVALinkageForFunction(Def)); + }; + if (LangOpts.OpenMPIsDevice) { + // In OpenMP device mode we will not emit host only functions, or functions + // we don't need due to their linkage. Optional DevTy = OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl()); - if (DevTy.hasValue()) { + // DevTy may be changed later by + // #pragma omp declare target to(*) device_type(*). + // Therefore DevTyhaving no value does not imply host. The emission status + // will be checked again at the end of compilation unit with Final = true. + if (DevTy.hasValue()) if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host) - OMPES = FunctionEmissionStatus::OMPDiscarded; - else if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost || - *DevTy == OMPDeclareTargetDeclAttr::DT_Any) { - OMPES = FunctionEmissionStatus::Emitted; - } - } - } else if (LangOpts.OpenMP) { - // In OpenMP 4.5 all the functions are host functions. - if (LangOpts.OpenMP <= 45) { - OMPES = FunctionEmissionStatus::Emitted; - } else { - Optional DevTy = - OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl()); - // In OpenMP 5.0 or above, DevTy may be changed later by - // #pragma omp declare target to(*) device_type(*). Therefore DevTy - // having no value does not imply host. The emission status will be - // checked again at the end of compilation unit. - if (DevTy.hasValue()) { - if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) { - OMPES = FunctionEmissionStatus::OMPDiscarded; - } else if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host || - *DevTy == OMPDeclareTargetDeclAttr::DT_Any) - OMPES = FunctionEmissionStatus::Emitted; - } else if (Final) - OMPES = FunctionEmissionStatus::Emitted; - } - } - if (OMPES == FunctionEmissionStatus::OMPDiscarded || - (OMPES == FunctionEmissionStatus::Emitted && !LangOpts.CUDA)) - return OMPES; + return FunctionEmissionStatus::OMPDiscarded; + // If we have an explicit value for the device type, or we are in a target + // declare context, we need to emit all extern and used symbols. + if (isInOpenMPDeclareTargetContext() || DevTy.hasValue()) + if (IsEmittedForExternalSymbol()) + return FunctionEmissionStatus::Emitted; + // Device mode only emits what it must, if it wasn't tagged yet and needed, + // we'll omit it. + if (Final) + return FunctionEmissionStatus::OMPDiscarded; + } else if (LangOpts.OpenMP > 45) { + // In OpenMP host compilation prior to 5.0 everything was an emitted host + // function. In 5.0, no_host was introduced which might cause a function to + // be ommitted. + Optional DevTy = + OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl()); + if (DevTy.hasValue()) + if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) + return FunctionEmissionStatus::OMPDiscarded; + } + + if (Final && LangOpts.OpenMP && !LangOpts.CUDA) + return FunctionEmissionStatus::Emitted; if (LangOpts.CUDA) { // When compiling for device, host functions are never emitted. Similarly, @@ -18382,17 +18391,7 @@ (T == Sema::CFT_Device || T == Sema::CFT_Global)) return FunctionEmissionStatus::CUDADiscarded; - // Check whether this function is externally visible -- if so, it's - // known-emitted. - // - // We have to check the GVA linkage of the function's *definition* -- if we - // only have a declaration, we don't know whether or not the function will - // be emitted, because (say) the definition could include "inline". - FunctionDecl *Def = FD->getDefinition(); - - if (Def && - !isDiscardableGVALinkage(getASTContext().GetGVALinkageForFunction(Def)) - && (!LangOpts.OpenMP || OMPES == FunctionEmissionStatus::Emitted)) + if (IsEmittedForExternalSymbol()) return FunctionEmissionStatus::Emitted; } 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 @@ -1884,8 +1884,7 @@ static bool isOpenMPDeviceDelayedContext(Sema &S) { assert(S.LangOpts.OpenMP && S.LangOpts.OpenMPIsDevice && "Expected OpenMP device compilation."); - return !S.isInOpenMPTargetExecutionDirective() && - !S.isInOpenMPDeclareTargetContext(); + return !S.isInOpenMPTargetExecutionDirective(); } namespace { @@ -1911,6 +1910,13 @@ Kind = SemaDiagnosticBuilder::K_Immediate; break; case FunctionEmissionStatus::Unknown: + // TODO: We should always delay diagnostics here in case a target + // region is in a function we do not emit. However, as the + // current diagnostics are associated with the function containing + // the target region and we do not emit that one, we would miss out + // on diagnostics for the target region itself. We need to anchor + // the diagnostics with the new generated function *or* ensure we + // emit diagnostics associated with the surrounding function. Kind = isOpenMPDeviceDelayedContext(*this) ? SemaDiagnosticBuilder::K_Deferred : SemaDiagnosticBuilder::K_Immediate; diff --git a/clang/test/OpenMP/nvptx_allocate_messages.cpp b/clang/test/OpenMP/nvptx_allocate_messages.cpp --- a/clang/test/OpenMP/nvptx_allocate_messages.cpp +++ b/clang/test/OpenMP/nvptx_allocate_messages.cpp @@ -81,8 +81,7 @@ #endif // DEVICE && !REQUIRES #pragma omp allocate(b) #if defined(DEVICE) && !defined(REQUIRES) -// expected-note@+3 {{in instantiation of function template specialization 'foo' requested here}} -// expected-note@+2 {{called by 'main'}} +// expected-note@+2 2{{called by 'main'}} #endif // DEVICE && !REQUIRES return (foo() + bar()); } diff --git a/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp b/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp --- a/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp +++ b/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp @@ -52,6 +52,7 @@ #pragma omp target map(tofrom \ : a, b) { + // expected-note@+1 {{called by 'maini1'}} S s(a); static long aaa = 23; a = foo() + bar() + b + c + d + aa + aaa + FA(); // expected-note{{called by 'maini1'}} 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 @@ -81,18 +81,12 @@ T1 t = bar1(); } -// TODO: We should not emit an error for dead functions we do not emit. inline void dead_inline_declare_target() { -// expected-note@+1 {{'b' defined here}} long double *a, b = 0; -// expected-error@+1 {{'b' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} a = &b; } -// TODO: We should not emit an error for dead functions we do not emit. static void dead_static_declare_target() { -// expected-note@+1 {{'b' defined here}} long double *a, b = 0; -// expected-error@+1 {{'b' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} a = &b; } template @@ -108,7 +102,6 @@ // 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}} @@ -117,48 +110,28 @@ // 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) {} inline void ld_use1() { -// expected-note@+1 {{'ld' defined here}} long double ld = 0; -// TODO: We should not diagnose this as the function is dead. -// expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} ld += 1; } static void ld_use2() { -// expected-note@+1 {{'ld' defined here}} long double ld = 0; -// TODO: We should not diagnose this as the function is dead. -// expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}} ld += 1; } @@ -176,11 +149,18 @@ } 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); +// TODO: The error message "called by" is not great. +// expected-note@+1 {{called by 'external'}} void *p5 = reinterpret_cast(&ld_use3); +// expected-note@+1 {{called by 'external'}} void *p6 = reinterpret_cast(&ld_use4); }