[openmp] Emit deferred diag only when device compilation presents

[weiwang]

Summary

There is no need to check for deferred diag when device compilation or target is
not given. This results in considerable build time improvement in some cases.

[weiwang]

Our internal codebase never uses the target directive. Once the deferred diags is bypassed, we observed 18% e2e build time improvement.

[lebedev.ri]

In D109175#2980333, @weiwang wrote:

Our internal codebase never uses the target directive. Once the deferred diags is bypassed, we observed 18% e2e build time improvement.

Is that with -fopenmp or without?
That seems, kinda a lot more than i would have expected,
perhaps there are some other ways to reduce the overhead other than this approach?

[weiwang]

In D109175#2980446, @lebedev.ri wrote:

In D109175#2980333, @weiwang wrote:

Our internal codebase never uses the target directive. Once the deferred diags is bypassed, we observed 18% e2e build time improvement.

Is that with -fopenmp or without?
That seems, kinda a lot more than i would have expected,
perhaps there are some other ways to reduce the overhead other than this approach?

This is with -fopenmp and no other omp related flags. I'd prefer a more generic way of fixing this, but right now this seems to be most direct way.

[wenlei]

In D109175#2980446, @lebedev.ri wrote:

In D109175#2980333, @weiwang wrote:

Our internal codebase never uses the target directive. Once the deferred diags is bypassed, we observed 18% e2e build time improvement.

Is that with -fopenmp or without?
That seems, kinda a lot more than i would have expected,
perhaps there are some other ways to reduce the overhead other than this approach?

Yeah, though the slow down in build time is blocking us from moving to newer llvm for many projects. Currently I think it makes sense to give user the option to decide whether they want faster build or better diagnostics. Of course if the slow down is fully addressed in the future, the switch could be removed too.

[jdoerfert]

Why do we need this flag, is the absence of -fopenmp-targets not sufficient?

[weiwang]

In D109175#2980744, @jdoerfert wrote:

Why do we need this flag, is the absence of -fopenmp-targets not sufficient?

Just double checked, this is the full omp related options currently in use:

"-fopenmp"
"-fopenmp-version=31"
"-fopenmp-version=31"
"-fopenmp-cuda-parallel-target-regions"

We saw a huge number of DECLS_TO_CHECK_FOR_DEFERRED_DIAGS records. I don't know if this has anything to do with omp version being 31, since prior 5.0, everything is available on host.

[jdoerfert]

In D109175#2980806, @weiwang wrote:

In D109175#2980744, @jdoerfert wrote:

Why do we need this flag, is the absence of -fopenmp-targets not sufficient?

Just double checked, this is the full omp related options currently in use:

"-fopenmp"
"-fopenmp-version=31"
"-fopenmp-version=31"
"-fopenmp-cuda-parallel-target-regions"

We saw a huge number of DECLS_TO_CHECK_FOR_DEFERRED_DIAGS records. I don't know if this has anything to do with omp version being 31, since prior 5.0, everything is available on host.

I don't think we are selective right now. As I was saying, disable deferred parsing if fopenmp-targets is missing, no need for this option.

[weiwang]

In D109175#2980900, @jdoerfert wrote:

In D109175#2980806, @weiwang wrote:

In D109175#2980744, @jdoerfert wrote:

Why do we need this flag, is the absence of -fopenmp-targets not sufficient?

Just double checked, this is the full omp related options currently in use:

"-fopenmp"
"-fopenmp-version=31"
"-fopenmp-version=31"
"-fopenmp-cuda-parallel-target-regions"

We saw a huge number of DECLS_TO_CHECK_FOR_DEFERRED_DIAGS records. I don't know if this has anything to do with omp version being 31, since prior 5.0, everything is available on host.

I don't think we are selective right now. As I was saying, disable deferred parsing if fopenmp-targets is missing, no need for this option.

Sure I can certainly make the change. To make sure I understand you correctly: if -fopenmp-targets (or maybe fopenmp-is-device too) is not given from cmdline, we can just skip the deferred diags like this option does?

[jdoerfert]

In D109175#2980905, @weiwang wrote:

In D109175#2980900, @jdoerfert wrote:

In D109175#2980806, @weiwang wrote:

In D109175#2980744, @jdoerfert wrote:

Why do we need this flag, is the absence of -fopenmp-targets not sufficient?

Just double checked, this is the full omp related options currently in use:

"-fopenmp"
"-fopenmp-version=31"
"-fopenmp-version=31"
"-fopenmp-cuda-parallel-target-regions"

We saw a huge number of DECLS_TO_CHECK_FOR_DEFERRED_DIAGS records. I don't know if this has anything to do with omp version being 31, since prior 5.0, everything is available on host.

I don't think we are selective right now. As I was saying, disable deferred parsing if fopenmp-targets is missing, no need for this option.

Sure I can certainly make the change. To make sure I understand you correctly: if -fopenmp-targets (or maybe fopenmp-is-device too) is not given from cmdline, we can just skip the deferred diags like this option does?

I thought so, @ABataev wdyt?

[ABataev]

In D109175#2981054, @jdoerfert wrote:

In D109175#2980905, @weiwang wrote:

In D109175#2980900, @jdoerfert wrote:

In D109175#2980806, @weiwang wrote:

In D109175#2980744, @jdoerfert wrote:

Why do we need this flag, is the absence of -fopenmp-targets not sufficient?

Just double checked, this is the full omp related options currently in use:

"-fopenmp"
"-fopenmp-version=31"
"-fopenmp-version=31"
"-fopenmp-cuda-parallel-target-regions"

We saw a huge number of DECLS_TO_CHECK_FOR_DEFERRED_DIAGS records. I don't know if this has anything to do with omp version being 31, since prior 5.0, everything is available on host.

I don't think we are selective right now. As I was saying, disable deferred parsing if fopenmp-targets is missing, no need for this option.

Sure I can certainly make the change. To make sure I understand you correctly: if -fopenmp-targets (or maybe fopenmp-is-device too) is not given from cmdline, we can just skip the deferred diags like this option does?

I thought so, @ABataev wdyt?

Yes, deferred diags are used only for target-dependent compilation, so should be enough to check if IsTargetSpecified is false.

[weiwang]

update as discussed.

Changed a few affected tests to explicitly add target triple and added a case where deferred diags are skipped.

[yaxunl]

I agree with Johannes and Alexey that deferred diags are only needed when LangOpts.OMPTargetTriples.empty(). However, I am not sure whether it is only needed in device compilation.

For other offloading languages like CUDA/HIP it is needed in both device and host compilation.

[jdoerfert]

In D109175#2986782, @yaxunl wrote:

I agree with Johannes and Alexey that deferred diags are only needed when LangOpts.OMPTargetTriples.empty(). However, I am not sure whether it is only needed in device compilation.

For other offloading languages like CUDA/HIP it is needed in both device and host compilation.

Technically, we might even want to delay in host only mode for OpenMP, but that is something we can revisit (e.g., by dynamically setting a flag based on the directives we've seen).
@yaxunl Should we for now check if there is any associated offload job?

[weiwang]

In D109175#2987048, @jdoerfert wrote:

In D109175#2986782, @yaxunl wrote:

I agree with Johannes and Alexey that deferred diags are only needed when LangOpts.OMPTargetTriples.empty(). However, I am not sure whether it is only needed in device compilation.

For other offloading languages like CUDA/HIP it is needed in both device and host compilation.

Technically, we might even want to delay in host only mode for OpenMP, but that is something we can revisit (e.g., by dynamically setting a flag based on the directives we've seen).
@yaxunl Should we for now check if there is any associated offload job?

Shall we go ahead and get this change in and think about more longer term solution later?

[yaxunl]

In D109175#2989823, @weiwang wrote:

In D109175#2987048, @jdoerfert wrote:

In D109175#2986782, @yaxunl wrote:

I agree with Johannes and Alexey that deferred diags are only needed when LangOpts.OMPTargetTriples.empty(). However, I am not sure whether it is only needed in device compilation.

For other offloading languages like CUDA/HIP it is needed in both device and host compilation.

Technically, we might even want to delay in host only mode for OpenMP, but that is something we can revisit (e.g., by dynamically setting a flag based on the directives we've seen).
@yaxunl Should we for now check if there is any associated offload job?

Shall we go ahead and get this change in and think about more longer term solution later?

LGTM. This patch should be sufficient to limit deferred diags to OpenMP with offloading. Device compilation is covered by OpenMPIsDevice and host compilation is covered by !LangOpts.OMPTargetTriples.empty(). I will leave the decision to Johannes.

[weiwang]

In D109175#2989997, @yaxunl wrote:

In D109175#2989823, @weiwang wrote:

In D109175#2987048, @jdoerfert wrote:

In D109175#2986782, @yaxunl wrote:

I agree with Johannes and Alexey that deferred diags are only needed when LangOpts.OMPTargetTriples.empty(). However, I am not sure whether it is only needed in device compilation.

For other offloading languages like CUDA/HIP it is needed in both device and host compilation.

Technically, we might even want to delay in host only mode for OpenMP, but that is something we can revisit (e.g., by dynamically setting a flag based on the directives we've seen).
@yaxunl Should we for now check if there is any associated offload job?

Shall we go ahead and get this change in and think about more longer term solution later?

LGTM. This patch should be sufficient to limit deferred diags to OpenMP with offloading. Device compilation is covered by OpenMPIsDevice and host compilation is covered by !LangOpts.OMPTargetTriples.empty(). I will leave the decision to Johannes.

Thanks. @jdoerfert. Could you approve this?

[jdoerfert]

LGTM

clang/lib/Sema/SemaDecl.cpp

void Sema::AddInitializerToDecl(Decl *RealDecl, Expr *Init, bool DirectInit) {
   if (CXXDirectInit) {
     assert(DirectInit && "Call-style initializer must be direct init.");
     VDecl->setInitStyle(VarDecl::CallInit);
   } else if (DirectInit) {
     // This must be list-initialization. No other way is direct-initialization.
     VDecl->setInitStyle(VarDecl::ListInit);
   }
 
-  if (LangOpts.OpenMP && VDecl->isFileVarDecl())
+  if (LangOpts.OpenMP &&
+      (LangOpts.OpenMPIsDevice || !LangOpts.OMPTargetTriples.empty()) &&
+      VDecl->isFileVarDecl())
     DeclsToCheckForDeferredDiags.insert(VDecl);
   CheckCompleteVariableDeclaration(VDecl);
 }
 
 /// ActOnInitializerError - Given that there was an error parsing an
 /// initializer for the given declaration, try to return to some form
 /// of sanity.
 void Sema::ActOnInitializerError(Decl *D) {
Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
   PopFunctionScopeInfo(ActivePolicy, dcl);
   // If any errors have occurred, clear out any temporaries that may have
   // been leftover. This ensures that these temporaries won't be picked up for
   // deletion in some later function.
   if (hasUncompilableErrorOccurred()) {
     DiscardCleanupsInEvaluationContext();
   }
 
-  if (FD && (LangOpts.OpenMP || LangOpts.CUDA || LangOpts.SYCLIsDevice)) {
+  if (FD && ((LangOpts.OpenMP && (LangOpts.OpenMPIsDevice ||
+                                  !LangOpts.OMPTargetTriples.empty())) ||
+             LangOpts.CUDA || LangOpts.SYCLIsDevice)) {
     auto ES = getEmissionStatus(FD);
     if (ES == Sema::FunctionEmissionStatus::Emitted ||
         ES == Sema::FunctionEmissionStatus::Unknown)
       DeclsToCheckForDeferredDiags.insert(FD);
   }
 
   return dcl;
 }

clang/test/OpenMP/declare_target_messages.cpp

 // RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp -fopenmp-targets=x86_64-apple-macos10.7.0 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
 // RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,dev5 -fopenmp -fopenmp-is-device -fopenmp-targets=x86_64-apple-macos10.7.0 -aux-triple x86_64-apple-macos10.7.0 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
 
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fnoopenmp-use-tls -ferror-limit 100 -o - %s
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-is-device -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-targets=x86_64-apple-macos10.7.0 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-is-device -fopenmp-targets=x86_64-apple-macos10.7.0 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
 // RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp-version=45 -fopenmp-simd -fnoopenmp-use-tls -ferror-limit 100 -o - %s
 
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5 -fopenmp -fnoopenmp-use-tls -ferror-limit 100 -o - %s
 #pragma omp end declare target // expected-error {{unexpected OpenMP directive '#pragma omp end declare target'}}
 
 int a, b, z; // omp5-error {{variable captured in declare target region must appear in a to clause}}
 __thread int t; // expected-note {{defined as threadprivate or thread local}}
 
 #pragma omp declare target . // expected-error {{expected '(' after 'declare target'}}
 
 #pragma omp declare target

clang/test/SemaCUDA/openmp-target.cu

 // RUN: %clang_cc1 -triple x86_64 -verify=expected,dev \
 // RUN:            -verify-ignore-unexpected=note \
-// RUN:            -fopenmp -fopenmp-version=50 -o - %s
+// RUN:            -fopenmp -fopenmp-version=50 -fopenmp-targets=amdgcn-amd-amdhsa -o - %s
 // RUN: %clang_cc1 -triple x86_64 -verify -verify-ignore-unexpected=note\
-// RUN:            -fopenmp -fopenmp-version=50 -o - -x c++ %s
+// RUN:            -fopenmp -fopenmp-version=50 -fopenmp-targets=amdgcn-amd-amdhsa -o - -x c++ %s
 // RUN: %clang_cc1 -triple x86_64 -verify=dev -verify-ignore-unexpected=note\
 // RUN:            -fcuda-is-device -o - %s
 
 #if __CUDA__
 #include "Inputs/cuda.h"
 __device__ void cu_devf();
 #endif