[llvm-branch-commits] [clang] 0d14528 - [OpenMP] Delay more diagnostics of potentially non-emitted code
Tom Stellard via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Feb 15 11:58:41 PST 2021
Author: Johannes Doerfert
Date: 2021-02-15T11:57:51-08:00
New Revision: 0d14528f8082ba12e34f41cbf931a7e7425b694a
URL: https://github.com/llvm/llvm-project/commit/0d14528f8082ba12e34f41cbf931a7e7425b694a
DIFF: https://github.com/llvm/llvm-project/commit/0d14528f8082ba12e34f41cbf931a7e7425b694a.diff
LOG: [OpenMP] Delay more diagnostics of potentially non-emitted code
Even code in target and declare target regions might not be emitted.
With this patch we delay more diagnostics and use laziness and linkage
to determine if a function is emitted (for the device). Note that we
still eagerly emit diagnostics for target regions, unfortunately, see
the TODO for the reason.
This hopefully fixes PR48933.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D95928
(cherry picked from commit 1dd66e6111a8247c6c7931143251c0cf1442b905)
Added:
Modified:
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/nvptx_allocate_messages.cpp
clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
Removed:
################################################################################
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 6457c6d024cf..1f7ab49ccdd7 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -18332,42 +18332,51 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
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<OMPDeclareTargetDeclAttr::DevTypeTy> 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<OMPDeclareTargetDeclAttr::DevTypeTy> 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<OMPDeclareTargetDeclAttr::DevTypeTy> 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,
@@ -18381,17 +18390,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
(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
index 596aa4b7b5c1..4063c185388d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -1884,8 +1884,7 @@ void Sema::popOpenMPFunctionRegion(const FunctionScopeInfo *OldFSI) {
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 @@ Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
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
index a4d78b6ab588..9a61da73eb39 100644
--- a/clang/test/OpenMP/nvptx_allocate_messages.cpp
+++ b/clang/test/OpenMP/nvptx_allocate_messages.cpp
@@ -81,8 +81,7 @@ int main () {
#endif // DEVICE && !REQUIRES
#pragma omp allocate(b)
#if defined(DEVICE) && !defined(REQUIRES)
-// expected-note at +3 {{in instantiation of function template specialization 'foo<int>' requested here}}
-// expected-note at +2 {{called by 'main'}}
+// expected-note at +2 2{{called by 'main'}}
#endif // DEVICE && !REQUIRES
return (foo<int>() + bar());
}
diff --git a/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp b/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
index c71615d2521f..87ea00a90822 100644
--- a/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
+++ b/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
@@ -52,6 +52,7 @@ int maini1() {
#pragma omp target map(tofrom \
: a, b)
{
+ // expected-note at +1 {{called by 'maini1'}}
S s(a);
static long aaa = 23;
a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); // 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
index 1b89a891887d..a319c78f73c5 100644
--- a/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
+++ b/clang/test/OpenMP/nvptx_unsupported_type_messages.cpp
@@ -81,18 +81,12 @@ void baz1() {
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 at +1 {{'b' defined here}}
long double *a, b = 0;
-// expected-error at +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 at +1 {{'b' defined here}}
long double *a, b = 0;
-// expected-error at +1 {{'b' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
a = &b;
}
template<bool>
@@ -108,7 +102,6 @@ long double ld_return1a() { return 0; }
// expected-error at +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 at +2 {{'ld_return1b' defined here}}
// expected-error at +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 @@ ld_ty ld_return1b() { return 0; }
// expected-error at +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 at +2 {{'ld_return1c' defined here}}
-// expected-error at +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 at +2 {{'ld_arg1c' defined here}}
-// expected-error at +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 at +2 {{'ld_return1d' defined here}}
-// expected-error at +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 at +2 {{'ld_arg1d' defined here}}
-// expected-error at +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 at +2 {{'ld_return1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
// expected-note at +1 {{'ld_return1e' defined here}}
static long double ld_return1e() { return 0; }
-// expected-error at +2 {{'ld_arg1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
// expected-note at +1 {{'ld_arg1e' defined here}}
static void ld_arg1e(long double ld) {}
-// expected-error at +2 {{'ld_return1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
// expected-note at +1 {{'ld_return1f' defined here}}
inline long double ld_return1f() { return 0; }
-// expected-error at +2 {{'ld_arg1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
// expected-note at +1 {{'ld_arg1f' defined here}}
inline void ld_arg1f(long double ld) {}
inline void ld_use1() {
-// expected-note at +1 {{'ld' defined here}}
long double ld = 0;
-// TODO: We should not diagnose this as the function is dead.
-// expected-error at +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 at +1 {{'ld' defined here}}
long double ld = 0;
-// TODO: We should not diagnose this as the function is dead.
-// expected-error at +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 @@ static void ld_use4() {
}
void external() {
+// expected-error at +1 {{'ld_return1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
void *p1 = reinterpret_cast<void*>(&ld_return1e);
+// expected-error at +1 {{'ld_arg1e' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
void *p2 = reinterpret_cast<void*>(&ld_arg1e);
+// expected-error at +1 {{'ld_return1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
void *p3 = reinterpret_cast<void*>(&ld_return1f);
+// expected-error at +1 {{'ld_arg1f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
void *p4 = reinterpret_cast<void*>(&ld_arg1f);
+// TODO: The error message "called by" is not great.
+// expected-note at +1 {{called by 'external'}}
void *p5 = reinterpret_cast<void*>(&ld_use3);
+// expected-note at +1 {{called by 'external'}}
void *p6 = reinterpret_cast<void*>(&ld_use4);
}
More information about the llvm-branch-commits
mailing list