[CUDA][HIP] Defer overloading resolution diagnostics for host device functions

[yaxunl]

Summary

In CUDA/HIP a function may become implicit host device function by
pragma or constexpr. A host device function is checked in both
host and device compilation. However it may be emitted only
on host or device side, therefore the diagnostics should be
deferred until it is known to be emitted.

Currently clang is only able to defer certain diagnostics. This causes
false alarms and limits the usefulness of host device functions.

This patch lets clang defer all overloading resolution diagnostics for host device functions.

An option -fgpu-defer-diag is added to control this behavior. By default
it is off.

It is NFC for other languages.

[yaxunl]

update for depending patch change

[tra]

It's an interesting idea and it may be needed to handle wider range of implicitly-HD functions.
However it's likely to have interesting consequences, not all of them desirable.
It may be worth considering hiding the new behavior behind a flag, make it optional at first, give it some soak time and only then make it the default.
Summoning @rsmith for the big picture input.

One side effect of this change is that we'll probably stop producing diagnostics for the code that is known to be wrong, but happens to be unused.

E.g.

__host__ __device__ static void hd() {
   no_such_type_t t;
}

We obviously can never compile this function on either side of the compilation and clang currently does diagnose it on both sides.
However, if all diagnostics in HD functions are postponed until codegen, this error will go unreported.

What if I have a syntax error in a HD function? I'd still want it to be diagnosed.
It's going to be interesting to see how it affects parsing subsequent code.
E.g https://godbolt.org/z/1vfPcc -- clang complains about the error in innocent due to an error in hd. If all diags issues in hd are postponed and the user only sees error: function definition is not allowed here, it will be rather confusing, because the innocent itself is perfectly fine.

Also, this patch may need more eyes on it. I'm sure there are more corner cases we may need to think about.
E.g. how would it affect SFINAE? Will it allow some things to succeed where they would've failed previously? If so, is it always the right thing to do?

We'll certainly want to have more tests that cover the errors that we still want to be produced.

[jdoerfert]

In D84364#2168036, @tra wrote:

E.g.

__host__ __device__ static void hd() {
   no_such_type_t t;
}

We obviously can never compile this function on either side of the compilation and clang currently does diagnose it on both sides.
However, if all diagnostics in HD functions are postponed until codegen, this error will go unreported.

I am not certain but I could imagine a scenario in which problematic code would trigger visible changes to the program behavior w/o an error ever being emitted. That seems really undesirable.

[yaxunl]

Added option -fgpu-defer-diag to control this new behavior. By default it is off.
Added test for syntax error.

[yaxunl]

In D84364#2168036, @tra wrote:

One side effect of this change is that we'll probably stop producing diagnostics for the code that is known to be wrong, but happens to be unused.

E.g.

__host__ __device__ static void hd() {
   no_such_type_t t;
}

We obviously can never compile this function on either side of the compilation and clang currently does diagnose it on both sides.
However, if all diagnostics in HD functions are postponed until codegen, this error will go unreported.

What if I have a syntax error in a HD function? I'd still want it to be diagnosed.
It's going to be interesting to see how it affects parsing subsequent code.
E.g https://godbolt.org/z/1vfPcc -- clang complains about the error in innocent due to an error in hd. If all diags issues in hd are postponed and the user only sees error: function definition is not allowed here, it will be rather confusing, because the innocent itself is perfectly fine.

A diagnostic will be deferred only if clang successfully parsed a function declaration and the current context is a function declaration.

In the above situation, clang fails to parse the function and the current context is not a function, therefore the error is emitted immediately, like before.

I have added a lit test for such situation.

Also, this patch may need more eyes on it. I'm sure there are more corner cases we may need to think about.
E.g. how would it affect SFINAE? Will it allow some things to succeed where they would've failed previously? If so, is it always the right thing to do?

We only defer diagnostics happen inside a host device function. My guess is that SFINAE will work as before since it happens not in a function scope, therefore it is not deferred. I will add some tests.

[yaxunl]

I'd like to clarify a little bit about deferring an error.

First it is localized to a function. If an error causes a function not parsed completely, it will be emitted immediately.

So if an error is deferred, it means clang at least parses that function containing it.

Second, if an entity caused an error which is deferred, it does not mean clang thinks gee, let's defer it, and thinks now we do not have error here. It is more like OK here is an error and we know it. We just save that message somewhere and do not print it out for now. clang knows that error happened and propagates the error status the same way as non-deferred diagnostics inside that function. In a sense, the deferred diagnostic and immediate diagnostic does not differ much.

[yaxunl]

Another thing we can do is to limit the diagnostic messages that can be deferred. For example, we only defer semantic diagnostics, or overloading resolution related diagnostics. According to previous experience, what bothers us most are the diagnostics triggered by the differences in function overloading on device and host sides. If a diagnostic tends to happen on both sides, e.g. syntax error, there is less value to defer it.

[jdoerfert]

FWIW, OpenMP does also defer some diagnostics. It feels like a mess we can't avoid. That means, I think there is merit in generalizing this. I haven't reviewed this in any detail though.

[yaxunl]

added lit test for SFINAE

[yaxunl]

update the lit test for SFINAE. make sure substitution failure does not incur error msg if there is valid substitution.

Since template is not allowed in local class, there is no need for test SFINAE inside host device function.

[tra]

I'm going to try the patch on our CUDA code and see how it fares. Stay tuned.

[tra]

In D84364#2170244, @tra wrote:

I'm going to try the patch on our CUDA code and see how it fares. Stay tuned.

The patch works well enough to build TensorFlow which is a good sign that existing working code should be OK -- we're making HD checks less strict, so already working code should not be affected.

The remaining concern is whether we're going to unintentionally allow something we should not have. Let me see if I can come up with some examples, in addition to the hdf_not_called()one I've commented on in the test file.

clang/test/SemaCUDA/deferred-all.cu

13–19

I think we do need diagnostics in this scenario, regardless of whether the function is called or not.
We would have emitted it for a regular inline function and I believe it should be the case here, too.

45

Nit: I'd rename type -> isA which would make its use in sfinae() more obvious.

[yaxunl]

In D84364#2170769, @tra wrote:

In D84364#2170244, @tra wrote:

I'm going to try the patch on our CUDA code and see how it fares. Stay tuned.

The patch works well enough to build TensorFlow which is a good sign that existing working code should be OK -- we're making HD checks less strict, so already working code should not be affected.

The remaining concern is whether we're going to unintentionally allow something we should not have. Let me see if I can come up with some examples, in addition to the hdf_not_called()one I've commented on in the test file.

I added a Deferrable bit to the diagnostics which can be specified in td files. This can be added to individual diagnostic defs or added to a bunch of diagnostic defs all together.

This field is used to control whether a diagnostic message can be deferred.

For now I enabled this bit for the overloading resolution diagnostics since these have been seen as false alarms in host device functions. We could let more diagnostics be deferrable if we see it is necessary.

clang/test/SemaCUDA/deferred-all.cu

13–19

Limited deferred diagnostics to a set of specified diagnostics in .td files.

45

done

[yaxunl]

Added Deferrable bit to diagnostics to control whether a diagnostic can be deferred.

[yaxunl]

I just saw bugzilla bug https://bugs.llvm.org/show_bug.cgi?id=46922

my patch https://reviews.llvm.org/D77954 is supposed to fix this issue. However since implicit host device functions often cause overloading resolution diagnostics on the device side which are not deferred, my patch caused regressions and was reverted several times. Currently it was still reverted.

I think to fix this issue we need to make overloading resolution diagnostics deferrable.

[tra]

In D84364#2176091, @yaxunl wrote:

I added a Deferrable bit to the diagnostics which can be specified in td files. This can be added to individual diagnostic defs or added to a bunch of diagnostic defs all together.

This field is used to control whether a diagnostic message can be deferred.

This may be a case of "too much, but not enough". It will be unnecessary for most of the diagnostics we have. Overload resolution is likely to be the primary beneficiary, inline asm and exceptions may be two other classes, but I can't think of anything else ATM.
At the same time it may not be enough, because we also need to take into account where and when particular diagnostic is emitted. I.e. the same diagnostic may need to be postponed when we emit it from CUDA code, yet we may want to *not* postpone it if it's in the code which has nothing to do with CUDA. E.g. C++ code which has oveloading-related error in an inline function which would not be codegen'ed. I would expect such error to be reported as it would be if the same function was compiled in plain C++ mode.

For now I enabled this bit for the overloading resolution diagnostics since these have been seen as false alarms in host device functions. We could let more diagnostics be deferrable if we see it is necessary.

I just saw bugzilla bug https://bugs.llvm.org/show_bug.cgi?id=46922
my patch https://reviews.llvm.org/D77954 is supposed to fix this issue. However since implicit host device functions often cause overloading resolution diagnostics on the device side which are not deferred, my patch caused regressions and was reverted several times. Currently it was still reverted.

I think to fix this issue we need to make overloading resolution diagnostics deferrable.

I'm missing something here. How would deferred diagnostics help with the issue in the bug 46922? The HD function there is codegen'ed on both sides, so the only difference postponing would make is that we'd emit the diagnostic a bit later.
Do you mean that postponed diags patch is not the fix, but rather a prerequisite for the overload resolution changes patch?

[yaxunl]

Defer overload resolution diags only if there are wrong-sided candidates.

[yaxunl]

In D84364#2201336, @tra wrote:

In D84364#2176091, @yaxunl wrote:

I added a Deferrable bit to the diagnostics which can be specified in td files. This can be added to individual diagnostic defs or added to a bunch of diagnostic defs all together.

This field is used to control whether a diagnostic message can be deferred.

This may be a case of "too much, but not enough". It will be unnecessary for most of the diagnostics we have. Overload resolution is likely to be the primary beneficiary, inline asm and exceptions may be two other classes, but I can't think of anything else ATM.
At the same time it may not be enough, because we also need to take into account where and when particular diagnostic is emitted. I.e. the same diagnostic may need to be postponed when we emit it from CUDA code, yet we may want to *not* postpone it if it's in the code which has nothing to do with CUDA. E.g. C++ code which has oveloading-related error in an inline function which would not be codegen'ed. I would expect such error to be reported as it would be if the same function was compiled in plain C++ mode.

I added a heuristics based rule for deferring overloading resolution related diagnostics. Basically clang will check if there are wrong-sided candidates when an overloading resolution error happens. This seems to be able to capture most of the host-ness related diagnostics which we want to defer whereas without affecting the diagnostics that are unrelated to host-ness.

For now I enabled this bit for the overloading resolution diagnostics since these have been seen as false alarms in host device functions. We could let more diagnostics be deferrable if we see it is necessary.

I just saw bugzilla bug https://bugs.llvm.org/show_bug.cgi?id=46922
my patch https://reviews.llvm.org/D77954 is supposed to fix this issue. However since implicit host device functions often cause overloading resolution diagnostics on the device side which are not deferred, my patch caused regressions and was reverted several times. Currently it was still reverted.

I think to fix this issue we need to make overloading resolution diagnostics deferrable.

I'm missing something here. How would deferred diagnostics help with the issue in the bug 46922? The HD function there is codegen'ed on both sides, so the only difference postponing would make is that we'd emit the diagnostic a bit later.
Do you mean that postponed diags patch is not the fix, but rather a prerequisite for the overload resolution changes patch?

Right. We need to favor same sided candidates for host device functions to fix that issue, the obstacle to implement the correct host/device based overloading resolution is that implicit host device function causes diagnostics in device compilation, which should be deferred but currently are not. With this patch we may be able to fix the overloading resolution issue properly.

[tra]

LGTM.

Nice!

To sum it up -- the patch introduces -fgpu-defer-diag flag which allows deferring overload resolution diagnostics, if overload set included candidates from both sides.

We may be deferring cases when we don't have to (e.g. df()->()callee2() should've errored out right away, even during host compilation, as there's no way it could ever be valid), but this approach is a good starting point. It affects only the interesting subset of diags, is not enabled by default, and can be refined further, if necessary.

clang/include/clang/Basic/LangOptions.def

244

The description does not seem to reflect reality. IIUIC only overload-related diagnostic messages will be deferred.

[yaxunl]

In D84364#2255572, @tra wrote:

LGTM.

Nice!

To sum it up -- the patch introduces -fgpu-defer-diag flag which allows deferring overload resolution diagnostics, if overload set included candidates from both sides.

We may be deferring cases when we don't have to (e.g. df()->()callee2() should've errored out right away, even during host compilation, as there's no way it could ever be valid), but this approach is a good starting point. It affects only the interesting subset of diags, is not enabled by default, and can be refined further, if necessary.

Thanks Artem. Would you please also review https://reviews.llvm.org/D84362 since this patch depends on that. Thanks.

clang/include/clang/Basic/LangOptions.def

244

will fix when committing

[hctim]

Looks like this patch broke the MSan buildbots, PTAL (repro instructions https://github.com/google/sanitizers/wiki/SanitizerBotReproduceBuild):

http://lab.llvm.org:8011/builders/sanitizer-x86_64-linux-fast/builds/46239/steps/check-clang%20msan/logs/stdio

FAIL: Clang :: SemaCUDA/deferred-oeverload.cu (11308 of 26387)
******************** TEST 'Clang :: SemaCUDA/deferred-oeverload.cu' FAILED ********************
Script:
--
: 'RUN: at line 1';   /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang -cc1 -internal-isystem /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/lib/clang/12.0.0/include -nostdsysteminc -fcuda-is-device -fsyntax-only -verify=dev,com /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/test/SemaCUDA/deferred-oeverload.cu    -std=c++11 -fgpu-defer-diag
: 'RUN: at line 3';   /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang -cc1 -internal-isystem /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/lib/clang/12.0.0/include -nostdsysteminc -fsyntax-only -verify=host,com /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/test/SemaCUDA/deferred-oeverload.cu    -std=c++11 -fgpu-defer-diag
--
Exit Code: 77

Command Output (stderr):
--
==41680==WARNING: MemorySanitizer: use-of-uninitialized-value
    #0 0xddfbf73 in clang::OverloadCandidateSet::CompleteCandidates(clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11484:9
    #1 0xde4e7ca in clang::OverloadCandidateSet::NoteCandidates(std::__1::pair<clang::SourceLocation, clang::PartialDiagnostic>, clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, llvm::StringRef, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11529:9
    #2 0xde6316f in FinishOverloadedCallExpr(clang::Sema&, clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, clang::OverloadCandidateSet*, clang::OverloadCandidate**, clang::OverloadingResult, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:12959:19
    #3 0xde6296b in clang::Sema::BuildOverloadedCallExpr(clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:13032:10
    #4 0xd4e126b in clang::Sema::BuildCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaExpr.cpp:6378:16
    #5 0xd53e073 in clang::Sema::ActOnCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaExpr.cpp:6275:7
    #6 0xcaa957f in clang::Parser::ParsePostfixExpressionSuffix(clang::ActionResult<clang::Expr*, true>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:2066:23
    #7 0xcaaf03e in clang::Parser::ParseCastExpression(clang::Parser::CastParseKind, bool, bool&, clang::Parser::TypeCastState, bool, bool*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:1811:9
    #8 0xcaa5f3e in clang::Parser::ParseCastExpression(clang::Parser::CastParseKind, bool, clang::Parser::TypeCastState, bool, bool*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:681:20
    #9 0xcaa15f3 in clang::Parser::ParseAssignmentExpression(clang::Parser::TypeCastState) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:173:20
    #10 0xcaa13ad in clang::Parser::ParseExpression(clang::Parser::TypeCastState) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:124:18
    #11 0xcbc8288 in clang::Parser::ParseExprStatement(clang::Parser::ParsedStmtContext) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:446:19
    #12 0xcbc2183 in clang::Parser::ParseStatementOrDeclarationAfterAttributes(llvm::SmallVector<clang::Stmt*, 32u>&, clang::Parser::ParsedStmtContext, clang::SourceLocation*, clang::Parser::ParsedAttributesWithRange&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:234:12
    #13 0xcbc119d in clang::Parser::ParseStatementOrDeclaration(llvm::SmallVector<clang::Stmt*, 32u>&, clang::Parser::ParsedStmtContext, clang::SourceLocation*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:106:20
    #14 0xcbdd830 in clang::Parser::ParseCompoundStatementBody(bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:1098:11
    #15 0xcbe0b67 in clang::Parser::ParseFunctionStatementBody(clang::Decl*, clang::Parser::ParseScope&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:2259:21
    #16 0xc9c26d4 in clang::Parser::ParseFunctionDefinition(clang::ParsingDeclarator&, clang::Parser::ParsedTemplateInfo const&, clang::Parser::LateParsedAttrList*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1375:10
    #17 0xca17b5b in clang::Parser::ParseDeclGroup(clang::ParsingDeclSpec&, clang::DeclaratorContext, clang::SourceLocation*, clang::Parser::ForRangeInit*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseDecl.cpp:1924:27
    #18 0xc9bf6ee in clang::Parser::ParseDeclOrFunctionDefInternal(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec&, clang::AccessSpecifier) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1135:10
    #19 0xc9bde93 in clang::Parser::ParseDeclarationOrFunctionDefinition(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec*, clang::AccessSpecifier) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1151:12
    #20 0xc9bb56b in clang::Parser::ParseExternalDeclaration(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:971:12
    #21 0xc9b5d18 in clang::Parser::ParseTopLevelDecl(clang::OpaquePtr<clang::DeclGroupRef>&, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:716:12
    #22 0xc9a49cf in clang::ParseAST(clang::Sema&, bool, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseAST.cpp:158:20
    #23 0x8de7de0 in clang::FrontendAction::Execute() /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Frontend/FrontendAction.cpp:950:8
    #24 0x8ce176d in clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Frontend/CompilerInstance.cpp:984:33
    #25 0x905ccb9 in clang::ExecuteCompilerInvocation(clang::CompilerInstance*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp:278:25
    #26 0xb9bc9d in cc1_main(llvm::ArrayRef<char const*>, char const*, void*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/cc1_main.cpp:240:15
    #27 0xb93aa6 in ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/driver.cpp:330:12
    #28 0xb928ad in main /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/driver.cpp:407:12
    #29 0x7f3d1d0dc09a in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x2409a)
    #30 0xb15209 in _start (/b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang-12+0xb15209)

SUMMARY: MemorySanitizer: use-of-uninitialized-value /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11484:9 in clang::OverloadCandidateSet::CompleteCandidates(clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>)
Exiting

--

********************
Testing:  0.. 10.. 20.. 30.. 40.. 50.. 60.. 70.. 80.. 90.. 
********************
Failed Tests (1):
  Clang :: SemaCUDA/deferred-oeverload.cu

[yaxunl]

In D84364#2279684, @hctim wrote:

Looks like this patch broke the MSan buildbots, PTAL (repro instructions https://github.com/google/sanitizers/wiki/SanitizerBotReproduceBuild):

http://lab.llvm.org:8011/builders/sanitizer-x86_64-linux-fast/builds/46239/steps/check-clang%20msan/logs/stdio

FAIL: Clang :: SemaCUDA/deferred-oeverload.cu (11308 of 26387)
******************** TEST 'Clang :: SemaCUDA/deferred-oeverload.cu' FAILED ********************
Script:
--
: 'RUN: at line 1';   /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang -cc1 -internal-isystem /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/lib/clang/12.0.0/include -nostdsysteminc -fcuda-is-device -fsyntax-only -verify=dev,com /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/test/SemaCUDA/deferred-oeverload.cu    -std=c++11 -fgpu-defer-diag
: 'RUN: at line 3';   /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang -cc1 -internal-isystem /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/lib/clang/12.0.0/include -nostdsysteminc -fsyntax-only -verify=host,com /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/test/SemaCUDA/deferred-oeverload.cu    -std=c++11 -fgpu-defer-diag
--
Exit Code: 77

Command Output (stderr):
--
==41680==WARNING: MemorySanitizer: use-of-uninitialized-value
    #0 0xddfbf73 in clang::OverloadCandidateSet::CompleteCandidates(clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11484:9
    #1 0xde4e7ca in clang::OverloadCandidateSet::NoteCandidates(std::__1::pair<clang::SourceLocation, clang::PartialDiagnostic>, clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, llvm::StringRef, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11529:9
    #2 0xde6316f in FinishOverloadedCallExpr(clang::Sema&, clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, clang::OverloadCandidateSet*, clang::OverloadCandidate**, clang::OverloadingResult, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:12959:19
    #3 0xde6296b in clang::Sema::BuildOverloadedCallExpr(clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:13032:10
    #4 0xd4e126b in clang::Sema::BuildCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaExpr.cpp:6378:16
    #5 0xd53e073 in clang::Sema::ActOnCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaExpr.cpp:6275:7
    #6 0xcaa957f in clang::Parser::ParsePostfixExpressionSuffix(clang::ActionResult<clang::Expr*, true>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:2066:23
    #7 0xcaaf03e in clang::Parser::ParseCastExpression(clang::Parser::CastParseKind, bool, bool&, clang::Parser::TypeCastState, bool, bool*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:1811:9
    #8 0xcaa5f3e in clang::Parser::ParseCastExpression(clang::Parser::CastParseKind, bool, clang::Parser::TypeCastState, bool, bool*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:681:20
    #9 0xcaa15f3 in clang::Parser::ParseAssignmentExpression(clang::Parser::TypeCastState) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:173:20
    #10 0xcaa13ad in clang::Parser::ParseExpression(clang::Parser::TypeCastState) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:124:18
    #11 0xcbc8288 in clang::Parser::ParseExprStatement(clang::Parser::ParsedStmtContext) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:446:19
    #12 0xcbc2183 in clang::Parser::ParseStatementOrDeclarationAfterAttributes(llvm::SmallVector<clang::Stmt*, 32u>&, clang::Parser::ParsedStmtContext, clang::SourceLocation*, clang::Parser::ParsedAttributesWithRange&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:234:12
    #13 0xcbc119d in clang::Parser::ParseStatementOrDeclaration(llvm::SmallVector<clang::Stmt*, 32u>&, clang::Parser::ParsedStmtContext, clang::SourceLocation*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:106:20
    #14 0xcbdd830 in clang::Parser::ParseCompoundStatementBody(bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:1098:11
    #15 0xcbe0b67 in clang::Parser::ParseFunctionStatementBody(clang::Decl*, clang::Parser::ParseScope&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:2259:21
    #16 0xc9c26d4 in clang::Parser::ParseFunctionDefinition(clang::ParsingDeclarator&, clang::Parser::ParsedTemplateInfo const&, clang::Parser::LateParsedAttrList*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1375:10
    #17 0xca17b5b in clang::Parser::ParseDeclGroup(clang::ParsingDeclSpec&, clang::DeclaratorContext, clang::SourceLocation*, clang::Parser::ForRangeInit*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseDecl.cpp:1924:27
    #18 0xc9bf6ee in clang::Parser::ParseDeclOrFunctionDefInternal(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec&, clang::AccessSpecifier) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1135:10
    #19 0xc9bde93 in clang::Parser::ParseDeclarationOrFunctionDefinition(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec*, clang::AccessSpecifier) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1151:12
    #20 0xc9bb56b in clang::Parser::ParseExternalDeclaration(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:971:12
    #21 0xc9b5d18 in clang::Parser::ParseTopLevelDecl(clang::OpaquePtr<clang::DeclGroupRef>&, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:716:12
    #22 0xc9a49cf in clang::ParseAST(clang::Sema&, bool, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseAST.cpp:158:20
    #23 0x8de7de0 in clang::FrontendAction::Execute() /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Frontend/FrontendAction.cpp:950:8
    #24 0x8ce176d in clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Frontend/CompilerInstance.cpp:984:33
    #25 0x905ccb9 in clang::ExecuteCompilerInvocation(clang::CompilerInstance*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp:278:25
    #26 0xb9bc9d in cc1_main(llvm::ArrayRef<char const*>, char const*, void*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/cc1_main.cpp:240:15
    #27 0xb93aa6 in ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/driver.cpp:330:12
    #28 0xb928ad in main /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/driver.cpp:407:12
    #29 0x7f3d1d0dc09a in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x2409a)
    #30 0xb15209 in _start (/b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang-12+0xb15209)

SUMMARY: MemorySanitizer: use-of-uninitialized-value /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11484:9 in clang::OverloadCandidateSet::CompleteCandidates(clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>)
Exiting

--

********************
Testing:  0.. 10.. 20.. 30.. 40.. 50.. 60.. 70.. 80.. 90.. 
********************
Failed Tests (1):
  Clang :: SemaCUDA/deferred-oeverload.cu

Thanks. Reverted it. I will fix it.

[yaxunl]

In D84364#2279684, @hctim wrote:

Looks like this patch broke the MSan buildbots, PTAL (repro instructions https://github.com/google/sanitizers/wiki/SanitizerBotReproduceBuild):

http://lab.llvm.org:8011/builders/sanitizer-x86_64-linux-fast/builds/46239/steps/check-clang%20msan/logs/stdio

FAIL: Clang :: SemaCUDA/deferred-oeverload.cu (11308 of 26387)
******************** TEST 'Clang :: SemaCUDA/deferred-oeverload.cu' FAILED ********************
Script:
--
: 'RUN: at line 1';   /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang -cc1 -internal-isystem /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/lib/clang/12.0.0/include -nostdsysteminc -fcuda-is-device -fsyntax-only -verify=dev,com /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/test/SemaCUDA/deferred-oeverload.cu    -std=c++11 -fgpu-defer-diag
: 'RUN: at line 3';   /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang -cc1 -internal-isystem /b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/lib/clang/12.0.0/include -nostdsysteminc -fsyntax-only -verify=host,com /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/test/SemaCUDA/deferred-oeverload.cu    -std=c++11 -fgpu-defer-diag
--
Exit Code: 77

Command Output (stderr):
--
==41680==WARNING: MemorySanitizer: use-of-uninitialized-value
    #0 0xddfbf73 in clang::OverloadCandidateSet::CompleteCandidates(clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11484:9
    #1 0xde4e7ca in clang::OverloadCandidateSet::NoteCandidates(std::__1::pair<clang::SourceLocation, clang::PartialDiagnostic>, clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, llvm::StringRef, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11529:9
    #2 0xde6316f in FinishOverloadedCallExpr(clang::Sema&, clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, clang::OverloadCandidateSet*, clang::OverloadCandidate**, clang::OverloadingResult, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:12959:19
    #3 0xde6296b in clang::Sema::BuildOverloadedCallExpr(clang::Scope*, clang::Expr*, clang::UnresolvedLookupExpr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:13032:10
    #4 0xd4e126b in clang::Sema::BuildCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaExpr.cpp:6378:16
    #5 0xd53e073 in clang::Sema::ActOnCallExpr(clang::Scope*, clang::Expr*, clang::SourceLocation, llvm::MutableArrayRef<clang::Expr*>, clang::SourceLocation, clang::Expr*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaExpr.cpp:6275:7
    #6 0xcaa957f in clang::Parser::ParsePostfixExpressionSuffix(clang::ActionResult<clang::Expr*, true>) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:2066:23
    #7 0xcaaf03e in clang::Parser::ParseCastExpression(clang::Parser::CastParseKind, bool, bool&, clang::Parser::TypeCastState, bool, bool*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:1811:9
    #8 0xcaa5f3e in clang::Parser::ParseCastExpression(clang::Parser::CastParseKind, bool, clang::Parser::TypeCastState, bool, bool*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:681:20
    #9 0xcaa15f3 in clang::Parser::ParseAssignmentExpression(clang::Parser::TypeCastState) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:173:20
    #10 0xcaa13ad in clang::Parser::ParseExpression(clang::Parser::TypeCastState) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseExpr.cpp:124:18
    #11 0xcbc8288 in clang::Parser::ParseExprStatement(clang::Parser::ParsedStmtContext) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:446:19
    #12 0xcbc2183 in clang::Parser::ParseStatementOrDeclarationAfterAttributes(llvm::SmallVector<clang::Stmt*, 32u>&, clang::Parser::ParsedStmtContext, clang::SourceLocation*, clang::Parser::ParsedAttributesWithRange&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:234:12
    #13 0xcbc119d in clang::Parser::ParseStatementOrDeclaration(llvm::SmallVector<clang::Stmt*, 32u>&, clang::Parser::ParsedStmtContext, clang::SourceLocation*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:106:20
    #14 0xcbdd830 in clang::Parser::ParseCompoundStatementBody(bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:1098:11
    #15 0xcbe0b67 in clang::Parser::ParseFunctionStatementBody(clang::Decl*, clang::Parser::ParseScope&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseStmt.cpp:2259:21
    #16 0xc9c26d4 in clang::Parser::ParseFunctionDefinition(clang::ParsingDeclarator&, clang::Parser::ParsedTemplateInfo const&, clang::Parser::LateParsedAttrList*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1375:10
    #17 0xca17b5b in clang::Parser::ParseDeclGroup(clang::ParsingDeclSpec&, clang::DeclaratorContext, clang::SourceLocation*, clang::Parser::ForRangeInit*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseDecl.cpp:1924:27
    #18 0xc9bf6ee in clang::Parser::ParseDeclOrFunctionDefInternal(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec&, clang::AccessSpecifier) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1135:10
    #19 0xc9bde93 in clang::Parser::ParseDeclarationOrFunctionDefinition(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec*, clang::AccessSpecifier) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:1151:12
    #20 0xc9bb56b in clang::Parser::ParseExternalDeclaration(clang::Parser::ParsedAttributesWithRange&, clang::ParsingDeclSpec*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:971:12
    #21 0xc9b5d18 in clang::Parser::ParseTopLevelDecl(clang::OpaquePtr<clang::DeclGroupRef>&, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/Parser.cpp:716:12
    #22 0xc9a49cf in clang::ParseAST(clang::Sema&, bool, bool) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Parse/ParseAST.cpp:158:20
    #23 0x8de7de0 in clang::FrontendAction::Execute() /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Frontend/FrontendAction.cpp:950:8
    #24 0x8ce176d in clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Frontend/CompilerInstance.cpp:984:33
    #25 0x905ccb9 in clang::ExecuteCompilerInvocation(clang::CompilerInstance*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/FrontendTool/ExecuteCompilerInvocation.cpp:278:25
    #26 0xb9bc9d in cc1_main(llvm::ArrayRef<char const*>, char const*, void*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/cc1_main.cpp:240:15
    #27 0xb93aa6 in ExecuteCC1Tool(llvm::SmallVectorImpl<char const*>&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/driver.cpp:330:12
    #28 0xb928ad in main /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/tools/driver/driver.cpp:407:12
    #29 0x7f3d1d0dc09a in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x2409a)
    #30 0xb15209 in _start (/b/sanitizer-x86_64-linux-fast/build/llvm_build_msan/bin/clang-12+0xb15209)

SUMMARY: MemorySanitizer: use-of-uninitialized-value /b/sanitizer-x86_64-linux-fast/build/llvm-project/clang/lib/Sema/SemaOverload.cpp:11484:9 in clang::OverloadCandidateSet::CompleteCandidates(clang::Sema&, clang::OverloadCandidateDisplayKind, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation, llvm::function_ref<bool (clang::OverloadCandidate&)>)
Exiting

--

********************
Testing:  0.. 10.. 20.. 30.. 40.. 50.. 60.. 70.. 80.. 90.. 
********************
Failed Tests (1):
  Clang :: SemaCUDA/deferred-oeverload.cu

I am trying to reproduce this issue by building clang with -fsanitize=memory. However I got memory sanitizer failures with llvm-tblgen:

[422/2415] Building Options.inc...
FAILED: cd /home/yaxunl/git/llvm/build-msan && /home/yaxunl/git/llvm/build-msan/bin/llvm-tblgen -gen-opt-parser-defs -I /home/yaxunl/git/llvm/llvm/tools/clang/include/clang/Driver -I/home/yaxunl/git/llvm/llvm/tools/clang/include -I/home/yaxunl/git/llvm/build-msan/tools/clang/include -I/home/yaxunl/git/llvm/build-msan/include -I/home/yaxunl/git/llvm/llvm/include /home/yaxunl/git/llvm/llvm/tools/clang/include/clang/Driver/Options.td --write-if-changed -o tools/clang/include/clang/Driver/Options.inc -d tools/clang/include/clang/Driver/Options.inc.d

30130==WARNING: MemorySanitizer: use-of-uninitialized-value

#0 0xbdfe6a in find /home/yaxunl/git/llvm/llvm/include/llvm/ADT/StringRef.h:320:11
#1 0xbdfe6a in llvm::StringRef::split(llvm::SmallVectorImpl<llvm::StringRef>&, char, int, bool) const /home/yaxunl/git/llvm/llvm/lib/Support/StringRef.cpp:344:20
#2 0xbee57c in llvm::Triple::Triple(llvm::Twine const&) /home/yaxunl/git/llvm/llvm/lib/Support/Triple.cpp:747:19
#3 0xc29d25 in llvm::sys::getProcessTriple() /home/yaxunl/git/llvm/llvm/lib/Support/Host.cpp:1597:10
#4 0xb8704e in ParseCommandLineOptions /home/yaxunl/git/llvm/llvm/lib/Support/CommandLine.cpp:1337:17
#5 0xb8704e in llvm::cl::ParseCommandLineOptions(int, char const* const*, llvm::StringRef, llvm::raw_ostream*, char const*, bool) /home/yaxunl/git/llvm/llvm/lib/Support/CommandLine.cpp:1312:24
#6 0xb27a53 in main /home/yaxunl/git/llvm/llvm/utils/TableGen/TableGen.cpp:280:3
#7 0x7f4569cc883f in __libc_start_main /build/glibc-e6zv40/glibc-2.23/csu/../csu/libc-start.c:291
#8 0x42a7a8 in _start (/mnt/sdc1/home/yaxunl/git/llvm/build-msan/bin/llvm-tblgen+0x42a7a8)

Is this some known issue? Thanks.

[hctim]

Not a known issue - no, but MSan doesn't play nice with uninistrumented libraries (including things like libcxx) - and so it can be tricky to ensure your build is properly sanitized, which is why I'd recommend the build script :).

[yaxunl]

In D84364#2282458, @hctim wrote:

Not a known issue - no, but MSan doesn't play nice with uninistrumented libraries (including things like libcxx) - and so it can be tricky to ensure your build is properly sanitized, which is why I'd recommend the build script :).

Thanks for the suggestion. I reproduced the issue with the build script and fixed it.

[shafik]

clang/lib/Sema/SemaDecl.cpp

14543

Do you have tests that cover this case?

This section of code came up while I was investigating this bug report: https://github.com/llvm/llvm-project/issues/48974

I tried replacing hasUncompilableErrorOccurred() with hasAnyUnrecoverableErrorsInThisFunction() and it fixes the reported bug and it does not cause any tests to fail with check-clang but it looks like hasUncompilableErrorOccurred() is doing a lot more checks and so I expected some tests to fail when I swapped it out.

If we are missing test coverage we should add some.

clang/include/clang/Basic/LangOptions.def

 
 LANGOPT(CUDAIsDevice      , 1, 0, "compiling for CUDA device")
 LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "allowing variadic functions in CUDA device code")
 LANGOPT(CUDAHostDeviceConstexpr, 1, 1, "treating unattributed constexpr functions as __host__ __device__")
 LANGOPT(CUDADeviceApproxTranscendentals, 1, 0, "using approximate transcendental functions")
 LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code")
 LANGOPT(GPUAllowDeviceInit, 1, 0, "allowing device side global init functions for HIP")
 LANGOPT(GPUMaxThreadsPerBlock, 32, 1024, "default max threads per block for kernel launch bounds for HIP")
+LANGOPT(GPUDeferDiag, 1, 0, "defer all semantic diagnostic messages in host device functions for CUDA/HIP")
 
 LANGOPT(SYCL              , 1, 0, "SYCL")
 LANGOPT(SYCLIsDevice      , 1, 0, "Generate code for SYCL device")
 LANGOPT(SYCLVersion       , 32, 0, "Version of the SYCL standard used")
 
 LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")
 
 LANGOPT(SizedDeallocation , 1, 0, "sized deallocation")

clang/include/clang/Driver/Options.td

 def hip_version_EQ : Joined<["--"], "hip-version=">,
   HelpText<"HIP version in the format of major.minor.patch">;
 def fhip_dump_offload_linker_script : Flag<["-"], "fhip-dump-offload-linker-script">,
   Group<f_Group>, Flags<[NoArgumentUnused, HelpHidden]>;
 defm hip_new_launch_api : OptInFFlag<"hip-new-launch-api",
   "Use", "Don't use", " new kernel launching API for HIP">;
 defm gpu_allow_device_init : OptInFFlag<"gpu-allow-device-init",
   "Allow", "Don't allow", " device side init function in HIP">;
+defm gpu_defer_diag : OptInFFlag<"gpu-defer-diag",
+  "Defer", "Don't defer", " all semantic diagnostics in host device functions"
+  " for CUDA/HIP">;
 def gpu_max_threads_per_block_EQ : Joined<["--"], "gpu-max-threads-per-block=">,
   Flags<[CC1Option]>,
   HelpText<"Default max threads per block for kernel launch bounds for HIP">;
 def libomptarget_nvptx_path_EQ : Joined<["--"], "libomptarget-nvptx-path=">, Group<i_Group>,
   HelpText<"Path to libomptarget-nvptx libraries">;
 def dD : Flag<["-"], "dD">, Group<d_Group>, Flags<[CC1Option]>,
   HelpText<"Print macro definitions in -E mode in addition to normal output">;
 def dI : Flag<["-"], "dI">, Group<d_Group>, Flags<[CC1Option]>,

clang/include/clang/Sema/Sema.h

public:
   /// in template instantiation) to avoid stack overflow.
   void runWithSufficientStackSpace(SourceLocation Loc,
                                    llvm::function_ref<void()> Fn);
 
   /// Helper class that creates diagnostics with optional
   /// template instantiation stacks.
   ///
   /// This class provides a wrapper around the basic DiagnosticBuilder
-  /// class that emits diagnostics. SemaDiagnosticBuilder is
+  /// class that emits diagnostics. ImmediateDiagBuilder is
   /// responsible for emitting the diagnostic (as DiagnosticBuilder
   /// does) and, if the diagnostic comes from inside a template
   /// instantiation, printing the template instantiation stack as
   /// well.
-  class SemaDiagnosticBuilder : public DiagnosticBuilder {
+  class ImmediateDiagBuilder : public DiagnosticBuilder {
     Sema &SemaRef;
     unsigned DiagID;
 
   public:
-    SemaDiagnosticBuilder(DiagnosticBuilder &DB, Sema &SemaRef, unsigned DiagID)
-      : DiagnosticBuilder(DB), SemaRef(SemaRef), DiagID(DiagID) { }
+    const static bool IsDiagBuilder = false;
+    ImmediateDiagBuilder(DiagnosticBuilder &DB, Sema &SemaRef, unsigned DiagID)
+        : DiagnosticBuilder(DB), SemaRef(SemaRef), DiagID(DiagID) {}
+    ImmediateDiagBuilder(DiagnosticBuilder &&DB, Sema &SemaRef, unsigned DiagID)
+        : DiagnosticBuilder(DB), SemaRef(SemaRef), DiagID(DiagID) {}
 
     // This is a cunning lie. DiagnosticBuilder actually performs move
     // construction in its copy constructor (but due to varied uses, it's not
     // possible to conveniently express this as actual move construction). So
     // the default copy ctor here is fine, because the base class disables the
-    // source anyway, so the user-defined ~SemaDiagnosticBuilder is a safe no-op
+    // source anyway, so the user-defined ~ImmediateDiagBuilder is a safe no-op
     // in that case anwyay.
-    SemaDiagnosticBuilder(const SemaDiagnosticBuilder&) = default;
+    ImmediateDiagBuilder(const ImmediateDiagBuilder &) = default;
 
-    ~SemaDiagnosticBuilder() {
+    ~ImmediateDiagBuilder() {
       // If we aren't active, there is nothing to do.
       if (!isActive()) return;
 
       // Otherwise, we need to emit the diagnostic. First flush the underlying
       // DiagnosticBuilder data, and clear the diagnostic builder itself so it
       // won't emit the diagnostic in its own destructor.
       //
       // This seems wasteful, in that as written the DiagnosticBuilder dtor will
       // do its own needless checks to see if the diagnostic needs to be
       // emitted. However, because we take care to ensure that the builder
       // objects never escape, a sufficiently smart compiler will be able to
       // eliminate that code.
       FlushCounts();
       Clear();
 
       // Dispatch to Sema to emit the diagnostic.
       SemaRef.EmitCurrentDiagnostic(DiagID);
     }
 
     /// Teach operator<< to produce an object of the correct type.
-    template<typename T>
-    friend const SemaDiagnosticBuilder &operator<<(
-        const SemaDiagnosticBuilder &Diag, const T &Value) {
+    template <typename T>
+    friend const ImmediateDiagBuilder &
+    operator<<(const ImmediateDiagBuilder &Diag, const T &Value) {
       const DiagnosticBuilder &BaseDiag = Diag;
       BaseDiag << Value;
       return Diag;
     }
-    const static bool IsDiagBuilder = false;
   };
 
-  /// Emit a diagnostic.
-  SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID) {
-    DiagnosticBuilder DB = Diags.Report(Loc, DiagID);
-    return SemaDiagnosticBuilder(DB, *this, DiagID);
+  /// A generic diagnostic builder for errors which may or may not be deferred.
+  ///
+  /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch)
+  /// which are not allowed to appear inside __device__ functions and are
+  /// allowed to appear in __host__ __device__ functions only if the host+device
+  /// function is never codegen'ed.
+  ///
+  /// To handle this, we use the notion of "deferred diagnostics", where we
+  /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed.
+  ///
+  /// This class lets you emit either a regular diagnostic, a deferred
+  /// diagnostic, or no diagnostic at all, according to an argument you pass to
+  /// its constructor, thus simplifying the process of creating these "maybe
+  /// deferred" diagnostics.
+  class SemaDiagnosticBuilder {
+  public:
+    enum Kind {
+      /// Emit no diagnostics.
+      K_Nop,
+      /// Emit the diagnostic immediately (i.e., behave like Sema::Diag()).
+      K_Immediate,
+      /// Emit the diagnostic immediately, and, if it's a warning or error, also
+      /// emit a call stack showing how this function can be reached by an a
+      /// priori known-emitted function.
+      K_ImmediateWithCallStack,
+      /// Create a deferred diagnostic, which is emitted only if the function
+      /// it's attached to is codegen'ed.  Also emit a call stack as with
+      /// K_ImmediateWithCallStack.
+      K_Deferred
+    };
+
+    SemaDiagnosticBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
+                          FunctionDecl *Fn, Sema &S);
+    SemaDiagnosticBuilder(SemaDiagnosticBuilder &&D);
+    SemaDiagnosticBuilder(const SemaDiagnosticBuilder &) = default;
+    ~SemaDiagnosticBuilder();
+
+    bool isImmediate() const { return ImmediateDiag.hasValue(); }
+
+    /// Convertible to bool: True if we immediately emitted an error, false if
+    /// we didn't emit an error or we created a deferred error.
+    ///
+    /// Example usage:
+    ///
+    ///   if (SemaDiagnosticBuilder(...) << foo << bar)
+    ///     return ExprError();
+    ///
+    /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
+    /// want to use these instead of creating a SemaDiagnosticBuilder yourself.
+    operator bool() const { return isImmediate(); }
+
+    template <typename T>
+    friend const SemaDiagnosticBuilder &
+    operator<<(const SemaDiagnosticBuilder &Diag, const T &Value) {
+      if (Diag.ImmediateDiag.hasValue())
+        *Diag.ImmediateDiag << Value;
+      else if (Diag.PartialDiagId.hasValue())
+        Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second
+            << Value;
+      return Diag;
     }
 
+    friend const SemaDiagnosticBuilder &
+    operator<<(const SemaDiagnosticBuilder &Diag, const PartialDiagnostic &PD) {
+      if (Diag.ImmediateDiag.hasValue())
+        PD.Emit(*Diag.ImmediateDiag);
+      else if (Diag.PartialDiagId.hasValue())
+        Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second = PD;
+      return Diag;
+    }
+
+    void AddFixItHint(const FixItHint &Hint) const {
+      if (ImmediateDiag.hasValue())
+        ImmediateDiag->AddFixItHint(Hint);
+      else if (PartialDiagId.hasValue())
+        S.DeviceDeferredDiags[Fn][*PartialDiagId].second.AddFixItHint(Hint);
+    }
+
+    friend ExprResult ExprError(const SemaDiagnosticBuilder &) {
+      return ExprError();
+    }
+    friend StmtResult StmtError(const SemaDiagnosticBuilder &) {
+      return StmtError();
+    }
+    operator ExprResult() const { return ExprError(); }
+    operator StmtResult() const { return StmtError(); }
+    operator TypeResult() const { return TypeError(); }
+    operator DeclResult() const { return DeclResult(true); }
+    operator MemInitResult() const { return MemInitResult(true); }
+
+  private:
+    Sema &S;
+    SourceLocation Loc;
+    unsigned DiagID;
+    FunctionDecl *Fn;
+    bool ShowCallStack;
+
+    // Invariant: At most one of these Optionals has a value.
+    // FIXME: Switch these to a Variant once that exists.
+    llvm::Optional<ImmediateDiagBuilder> ImmediateDiag;
+    llvm::Optional<unsigned> PartialDiagId;
+  };
+  using DiagBuilderT = SemaDiagnosticBuilder;
+
+  /// Is the last error level diagnostic immediate. This is used to determined
+  /// whether the next info diagnostic should be immediate.
+  bool IsLastErrorImmediate = true;
+
+  /// Emit a diagnostic.
+  SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID);
+
   /// Emit a partial diagnostic.
-  SemaDiagnosticBuilder Diag(SourceLocation Loc, const PartialDiagnostic& PD);
+  SemaDiagnosticBuilder Diag(SourceLocation Loc, const PartialDiagnostic &PD);
 
   /// Build a partial diagnostic.
   PartialDiagnostic PDiag(unsigned DiagID = 0); // in SemaInternal.h
 
+  /// Whether uncompilable error has occurred. This includes error happens
+  /// in deferred diagnostics.
+  bool hasUncompilableErrorOccurred() const;
+
   bool findMacroSpelling(SourceLocation &loc, StringRef name);
 
   /// Get a string to suggest for zero-initialization of a type.
   std::string
   getFixItZeroInitializerForType(QualType T, SourceLocation Loc) const;
   std::string getFixItZeroLiteralForType(QualType T, SourceLocation Loc) const;
 
   /// Calls \c Lexer::getLocForEndOfToken()
public:
   /// known-emitted callers (plus the location of the call).
   ///
   /// Functions that we can tell a priori must be emitted aren't added to this
   /// map.
   llvm::DenseMap</* Callee = */ CanonicalDeclPtr<FunctionDecl>,
                  /* Caller = */ FunctionDeclAndLoc>
       DeviceKnownEmittedFns;
 
-  /// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be
-  /// deferred.
-  ///
-  /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch)
-  /// which are not allowed to appear inside __device__ functions and are
-  /// allowed to appear in __host__ __device__ functions only if the host+device
-  /// function is never codegen'ed.
-  ///
-  /// To handle this, we use the notion of "deferred diagnostics", where we
-  /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed.
-  ///
-  /// This class lets you emit either a regular diagnostic, a deferred
-  /// diagnostic, or no diagnostic at all, according to an argument you pass to
-  /// its constructor, thus simplifying the process of creating these "maybe
-  /// deferred" diagnostics.
-  class DeviceDiagBuilder {
-  public:
-    enum Kind {
-      /// Emit no diagnostics.
-      K_Nop,
-      /// Emit the diagnostic immediately (i.e., behave like Sema::Diag()).
-      K_Immediate,
-      /// Emit the diagnostic immediately, and, if it's a warning or error, also
-      /// emit a call stack showing how this function can be reached by an a
-      /// priori known-emitted function.
-      K_ImmediateWithCallStack,
-      /// Create a deferred diagnostic, which is emitted only if the function
-      /// it's attached to is codegen'ed.  Also emit a call stack as with
-      /// K_ImmediateWithCallStack.
-      K_Deferred
-    };
-
-    DeviceDiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID,
-                      FunctionDecl *Fn, Sema &S);
-    DeviceDiagBuilder(DeviceDiagBuilder &&D);
-    DeviceDiagBuilder(const DeviceDiagBuilder &) = default;
-    ~DeviceDiagBuilder();
-
-    /// Convertible to bool: True if we immediately emitted an error, false if
-    /// we didn't emit an error or we created a deferred error.
-    ///
-    /// Example usage:
-    ///
-    ///   if (DeviceDiagBuilder(...) << foo << bar)
-    ///     return ExprError();
-    ///
-    /// But see CUDADiagIfDeviceCode() and CUDADiagIfHostCode() -- you probably
-    /// want to use these instead of creating a DeviceDiagBuilder yourself.
-    operator bool() const { return ImmediateDiag.hasValue(); }
-
-    template <typename T>
-    friend const DeviceDiagBuilder &operator<<(const DeviceDiagBuilder &Diag,
-                                               const T &Value) {
-      if (Diag.ImmediateDiag.hasValue())
-        *Diag.ImmediateDiag << Value;
-      else if (Diag.PartialDiagId.hasValue())
-        Diag.S.DeviceDeferredDiags[Diag.Fn][*Diag.PartialDiagId].second
-            << Value;
-      return Diag;
-    }
-
-  private:
-    Sema &S;
-    SourceLocation Loc;
-    unsigned DiagID;
-    FunctionDecl *Fn;
-    bool ShowCallStack;
-
-    // Invariant: At most one of these Optionals has a value.
-    // FIXME: Switch these to a Variant once that exists.
-    llvm::Optional<SemaDiagnosticBuilder> ImmediateDiag;
-    llvm::Optional<unsigned> PartialDiagId;
-  };
-
-  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
-  /// is "used as device code".
+  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
+  /// context is "used as device code".
   ///
-  /// - If CurContext is a __host__ function, does not emit any diagnostics.
+  /// - If CurContext is a __host__ function, does not emit any diagnostics
+  ///   unless \p EmitOnBothSides is true.
   /// - If CurContext is a __device__ or __global__ function, emits the
   ///   diagnostics immediately.
   /// - If CurContext is a __host__ __device__ function and we are compiling for
   ///   the device, creates a diagnostic which is emitted if and when we realize
   ///   that the function will be codegen'ed.
   ///
   /// Example usage:
   ///
   ///  // Variable-length arrays are not allowed in CUDA device code.
   ///  if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget())
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
-  DeviceDiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
+  SemaDiagnosticBuilder CUDADiagIfDeviceCode(SourceLocation Loc,
+                                             unsigned DiagID,
+                                             bool EmitOnBothSides = false);
 
-  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context
-  /// is "used as host code".
+  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
+  /// context is "used as host code".
   ///
   /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched.
-  DeviceDiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID);
+  SemaDiagnosticBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID,
+                                           bool EmitOnBothSides = false);
 
-  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
   /// context is "used as device code".
   ///
   /// - If CurContext is a `declare target` function or it is known that the
   /// function is emitted for the device, emits the diagnostics immediately.
   /// - If CurContext is a non-`declare target` function and we are compiling
   ///   for the device, creates a diagnostic which is emitted if and when we
   ///   realize that the function will be codegen'ed.
   ///
   /// Example usage:
   ///
   ///  // Variable-length arrays are not allowed in NVPTX device code.
   ///  if (diagIfOpenMPDeviceCode(Loc, diag::err_vla_unsupported))
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
-  DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID);
+  SemaDiagnosticBuilder diagIfOpenMPDeviceCode(SourceLocation Loc,
+                                               unsigned DiagID);
 
-  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+  /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
   /// context is "used as host code".
   ///
   /// - If CurContext is a `declare target` function or it is known that the
   /// function is emitted for the host, emits the diagnostics immediately.
   /// - If CurContext is a non-host function, just ignore it.
   ///
   /// Example usage:
   ///
   ///  // Variable-length arrays are not allowed in NVPTX device code.
   ///  if (diagIfOpenMPHostode(Loc, diag::err_vla_unsupported))
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
-  DeviceDiagBuilder diagIfOpenMPHostCode(SourceLocation Loc, unsigned DiagID);
+  SemaDiagnosticBuilder diagIfOpenMPHostCode(SourceLocation Loc,
+                                             unsigned DiagID);
 
-  DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
+  SemaDiagnosticBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
+  SemaDiagnosticBuilder targetDiag(SourceLocation Loc,
+                                   const PartialDiagnostic &PD) {
+    return targetDiag(Loc, PD.getDiagID()) << PD;
+  }
 
   /// 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,
public:
   /// Describes the reason a calling convention specification was ignored, used
   /// for diagnostics.
   enum class CallingConventionIgnoredReason {
     ForThisTarget = 0,
     VariadicFunction,
     ConstructorDestructor,
     BuiltinFunction
   };
-  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+  /// Creates a SemaDiagnosticBuilder 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);
+  SemaDiagnosticBuilder 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

clang/lib/Driver/ToolChains/Cuda.cpp

if (DeviceOffloadingKind == Action::OFK_Cuda) {
 
     if (DriverArgs.hasFlag(options::OPT_fcuda_approx_transcendentals,
                            options::OPT_fno_cuda_approx_transcendentals, false))
       CC1Args.push_back("-fcuda-approx-transcendentals");
 
     if (DriverArgs.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
                            false))
       CC1Args.push_back("-fgpu-rdc");
+
+    if (DriverArgs.hasFlag(options::OPT_fgpu_defer_diag,
+                           options::OPT_fno_gpu_defer_diag, false))
+      CC1Args.push_back("-fgpu-defer-diag");
   }
 
   if (DriverArgs.hasArg(options::OPT_nogpulib))
     return;
 
   std::string LibDeviceFile = CudaInstallation.getLibDeviceFile(GpuArch);
 
   if (LibDeviceFile.empty()) {

clang/lib/Driver/ToolChains/HIP.cpp

std::string ArgStr =
         std::string("--gpu-max-threads-per-block=") + MaxThreadsPerBlock.str();
     CC1Args.push_back(DriverArgs.MakeArgStringRef(ArgStr));
   }
 
   if (DriverArgs.hasFlag(options::OPT_fgpu_allow_device_init,
                          options::OPT_fno_gpu_allow_device_init, false))
     CC1Args.push_back("-fgpu-allow-device-init");
 
+  if (DriverArgs.hasFlag(options::OPT_fgpu_defer_diag,
+                         options::OPT_fno_gpu_defer_diag, false))
+    CC1Args.push_back("-fgpu-defer-diag");
+
   CC1Args.push_back("-fcuda-allow-variadic-functions");
 
   // Default to "hidden" visibility, as object level linking will not be
   // supported for the foreseeable future.
   if (!DriverArgs.hasArg(options::OPT_fvisibility_EQ,
                          options::OPT_fvisibility_ms_compat)) {
     CC1Args.append({"-fvisibility", "hidden"});
     CC1Args.push_back("-fapply-global-visibility-to-externs");

clang/lib/Frontend/CompilerInvocation.cpp

if (Args.hasArg(OPT_fcuda_is_device))
     Opts.CUDAIsDevice = 1;
 
   if (Args.hasArg(OPT_fcuda_allow_variadic_functions))
     Opts.CUDAAllowVariadicFunctions = 1;
 
   if (Args.hasArg(OPT_fno_cuda_host_device_constexpr))
     Opts.CUDAHostDeviceConstexpr = 0;
 
+  if (Args.hasArg(OPT_fgpu_defer_diag))
+    Opts.GPUDeferDiag = 1;
+
   if (Opts.CUDAIsDevice && Args.hasArg(OPT_fcuda_approx_transcendentals))
     Opts.CUDADeviceApproxTranscendentals = 1;
 
   Opts.GPURelocatableDeviceCode = Args.hasArg(OPT_fgpu_rdc);
   if (Args.hasArg(OPT_fgpu_allow_device_init)) {
     if (Opts.HIP)
       Opts.GPUAllowDeviceInit = 1;
     else

clang/lib/Sema/AnalysisBasedWarnings.cpp

if (Diags.getIgnoreAllWarnings() ||
       (Diags.getSuppressSystemWarnings() &&
        S.SourceMgr.isInSystemHeader(D->getLocation())))
     return;
 
   // For code in dependent contexts, we'll do this at instantiation time.
   if (cast<DeclContext>(D)->isDependentContext())
     return;
 
-  if (Diags.hasUncompilableErrorOccurred()) {
+  if (S.hasUncompilableErrorOccurred()) {
     // Flush out any possibly unreachable diagnostics.
     flushDiagnostics(S, fscope);
     return;
   }
 
   const Stmt *Body = D->getBody();
   assert(Body);
 

clang/lib/Sema/Sema.cpp

void Sema::EmitCurrentDiagnostic(unsigned DiagID) {
   // If this is not a note, and we're in a template instantiation
   // that is different from the last template instantiation where
   // we emitted an error, print a template instantiation
   // backtrace.
   if (!DiagnosticIDs::isBuiltinNote(DiagID))
     PrintContextStack();
 }
 
-Sema::SemaDiagnosticBuilder
-Sema::Diag(SourceLocation Loc, const PartialDiagnostic& PD) {
-  SemaDiagnosticBuilder Builder(Diag(Loc, PD.getDiagID()));
-  PD.Emit(Builder);
+Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc,
+                                       const PartialDiagnostic &PD) {
+  return Diag(Loc, PD.getDiagID()) << PD;
+}
 
-  return Builder;
+bool Sema::hasUncompilableErrorOccurred() const {
+  if (getDiagnostics().hasUncompilableErrorOccurred())
+    return true;
+  auto *FD = dyn_cast<FunctionDecl>(CurContext);
+  if (!FD)
+    return false;
+  auto Loc = DeviceDeferredDiags.find(FD);
+  if (Loc == DeviceDeferredDiags.end())
+    return false;
+  for (auto PDAt : Loc->second) {
+    if (DiagnosticIDs::isDefaultMappingAsError(PDAt.second.getDiagID()))
+      return true;
+  }
+  return false;
 }
 
 // Print notes showing how we can reach FD starting from an a priori
 // known-callable function.
 static void emitCallStackNotes(Sema &S, FunctionDecl *FD) {
   auto FnIt = S.DeviceKnownEmittedFns.find(FD);
   while (FnIt != S.DeviceKnownEmittedFns.end()) {
     // Respect error limit.

 // (usually by way of CUDADiagIfDeviceCode, CUDADiagIfHostCode, or
 // CheckCUDACall), we first check if the current function is known-emitted.  If
 // so, we immediately output the diagnostic.
 //
 // Otherwise, we "defer" the diagnostic.  It sits in Sema::DeviceDeferredDiags
 // until we discover that the function is known-emitted, at which point we take
 // it out of this map and emit the diagnostic.
 
-Sema::DeviceDiagBuilder::DeviceDiagBuilder(Kind K, SourceLocation Loc,
-                                           unsigned DiagID, FunctionDecl *Fn,
-                                           Sema &S)
+Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(Kind K, SourceLocation Loc,
+                                                   unsigned DiagID,
+                                                   FunctionDecl *Fn, Sema &S)
     : S(S), Loc(Loc), DiagID(DiagID), Fn(Fn),
       ShowCallStack(K == K_ImmediateWithCallStack || K == K_Deferred) {
   switch (K) {
   case K_Nop:
     break;
   case K_Immediate:
   case K_ImmediateWithCallStack:
-    ImmediateDiag.emplace(S.Diag(Loc, DiagID));
+    ImmediateDiag.emplace(
+        ImmediateDiagBuilder(S.Diags.Report(Loc, DiagID), S, DiagID));
     break;
   case K_Deferred:
     assert(Fn && "Must have a function to attach the deferred diag to.");
     auto &Diags = S.DeviceDeferredDiags[Fn];
     PartialDiagId.emplace(Diags.size());
     Diags.emplace_back(Loc, S.PDiag(DiagID));
     break;
   }
 }
 
-Sema::DeviceDiagBuilder::DeviceDiagBuilder(DeviceDiagBuilder &&D)
+Sema::SemaDiagnosticBuilder::SemaDiagnosticBuilder(SemaDiagnosticBuilder &&D)
     : S(D.S), Loc(D.Loc), DiagID(D.DiagID), Fn(D.Fn),
       ShowCallStack(D.ShowCallStack), ImmediateDiag(D.ImmediateDiag),
       PartialDiagId(D.PartialDiagId) {
   // Clean the previous diagnostics.
   D.ShowCallStack = false;
   D.ImmediateDiag.reset();
   D.PartialDiagId.reset();
 }
 
-Sema::DeviceDiagBuilder::~DeviceDiagBuilder() {
+Sema::SemaDiagnosticBuilder::~SemaDiagnosticBuilder() {
   if (ImmediateDiag) {
     // Emit our diagnostic and, if it was a warning or error, output a callstack
     // if Fn isn't a priori known-emitted.
     bool IsWarningOrError = S.getDiagnostics().getDiagnosticLevel(
                                 DiagID, Loc) >= DiagnosticsEngine::Warning;
     ImmediateDiag.reset(); // Emit the immediate diag.
     if (IsWarningOrError && ShowCallStack)
       emitCallStackNotes(S, Fn);
   } else {
     assert((!PartialDiagId || ShowCallStack) &&
            "Must always show call stack for deferred diags.");
   }
 }
 
-Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
+Sema::SemaDiagnosticBuilder Sema::targetDiag(SourceLocation Loc,
+                                             unsigned DiagID) {
   if (LangOpts.OpenMP)
     return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID)
                                    : diagIfOpenMPHostCode(Loc, DiagID);
   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,
+  return SemaDiagnosticBuilder(SemaDiagnosticBuilder::K_Immediate, Loc, DiagID,
                                getCurFunctionDecl(), *this);
 }
 
+Sema::SemaDiagnosticBuilder Sema::Diag(SourceLocation Loc, unsigned DiagID) {
+  if (!getLangOpts().CUDA || !LangOpts.GPUDeferDiag)
+    return SemaDiagnosticBuilder(SemaDiagnosticBuilder::K_Immediate, Loc,
+                                 DiagID, getCurFunctionDecl(), *this);
+
+  SemaDiagnosticBuilder DB =
+      getLangOpts().CUDAIsDevice
+          ? CUDADiagIfDeviceCode(Loc, DiagID, /*EmitOnBothSides=*/true)
+          : CUDADiagIfHostCode(Loc, DiagID, /*EmitOnBothSides=*/true);
+
+  if (Diags.getDiagnosticIDs()->isDefaultMappingAsError(DiagID))
+    IsLastErrorImmediate = DB.isImmediate();
+  return DB;
+}
+
 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.

clang/lib/Sema/SemaAttr.cpp

void Sema::DiagnoseUnterminatedPragmaPack() {
   if (PackStack.Stack.empty())
     return;
   bool IsInnermost = true;
   for (const auto &StackSlot : llvm::reverse(PackStack.Stack)) {
     Diag(StackSlot.PragmaPushLocation, diag::warn_pragma_pack_no_pop_eof);
     // The user might have already reset the alignment, so suggest replacing
     // the reset with a pop.
     if (IsInnermost && PackStack.CurrentValue == PackStack.DefaultValue) {
-      DiagnosticBuilder DB = Diag(PackStack.CurrentPragmaLocation,
-                                  diag::note_pragma_pack_pop_instead_reset);
+      auto DB = Diag(PackStack.CurrentPragmaLocation,
       SourceLocation FixItLoc = Lexer::findLocationAfterToken(
           PackStack.CurrentPragmaLocation, tok::l_paren, SourceMgr, LangOpts,
           /*SkipTrailing=*/false);
       if (FixItLoc.isValid())
         DB << FixItHint::CreateInsertion(FixItLoc, "pop");
     }
     IsInnermost = false;
   }

clang/lib/Sema/SemaCUDA.cpp

 
 void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
   if (getLangOpts().CUDAIsDevice && VD->isConstexpr() &&
       (VD->isFileVarDecl() || VD->isStaticDataMember())) {
     VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
   }
 }
 
-Sema::DeviceDiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
-                                                   unsigned DiagID) {
+Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
+                                                       unsigned DiagID,
+                                                       bool EmitOnBothSides) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
-  DeviceDiagBuilder::Kind DiagKind = [this] {
+  SemaDiagnosticBuilder::Kind DiagKind = [&] {
+    if (!isa<FunctionDecl>(CurContext))
+      return SemaDiagnosticBuilder::K_Immediate;
     switch (CurrentCUDATarget()) {
     case CFT_Global:
     case CFT_Device:
-      return DeviceDiagBuilder::K_Immediate;
+      return SemaDiagnosticBuilder::K_Immediate;
     case CFT_HostDevice:
       // An HD function counts as host code if we're compiling for host, and
       // device code if we're compiling for device.  Defer any errors in device
       // mode until the function is known-emitted.
-      if (getLangOpts().CUDAIsDevice) {
-        return (getEmissionStatus(cast<FunctionDecl>(CurContext)) ==
-                FunctionEmissionStatus::Emitted)
-                   ? DeviceDiagBuilder::K_ImmediateWithCallStack
-                   : DeviceDiagBuilder::K_Deferred;
-      }
-      return DeviceDiagBuilder::K_Nop;
-
+      if (!getLangOpts().CUDAIsDevice)
+        return SemaDiagnosticBuilder::K_Nop;
+      if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
+        return SemaDiagnosticBuilder::K_Immediate;
+                 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
+                 : SemaDiagnosticBuilder::K_Deferred;
     default:
-      return DeviceDiagBuilder::K_Nop;
+      return EmitOnBothSides ? SemaDiagnosticBuilder::K_Immediate
+                             : SemaDiagnosticBuilder::K_Nop;
     }
   }();
-  return DeviceDiagBuilder(DiagKind, Loc, DiagID,
+  return SemaDiagnosticBuilder(DiagKind, Loc, DiagID,
                                dyn_cast<FunctionDecl>(CurContext), *this);
 }
 
-Sema::DeviceDiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
-                                                 unsigned DiagID) {
+Sema::SemaDiagnosticBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
+                                                     unsigned DiagID,
+                                                     bool EmitOnBothSides) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
-  DeviceDiagBuilder::Kind DiagKind = [this] {
+  SemaDiagnosticBuilder::Kind DiagKind = [&] {
+    if (!isa<FunctionDecl>(CurContext))
+      return SemaDiagnosticBuilder::K_Immediate;
     switch (CurrentCUDATarget()) {
     case CFT_Host:
-      return DeviceDiagBuilder::K_Immediate;
+      return SemaDiagnosticBuilder::K_Immediate;
     case CFT_HostDevice:
       // An HD function counts as host code if we're compiling for host, and
       // device code if we're compiling for device.  Defer any errors in device
       // mode until the function is known-emitted.
       if (getLangOpts().CUDAIsDevice)
-        return DeviceDiagBuilder::K_Nop;
-
+        return SemaDiagnosticBuilder::K_Nop;
+      if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
+        return SemaDiagnosticBuilder::K_Immediate;
       return (getEmissionStatus(cast<FunctionDecl>(CurContext)) ==
               FunctionEmissionStatus::Emitted)
-                 ? DeviceDiagBuilder::K_ImmediateWithCallStack
-                 : DeviceDiagBuilder::K_Deferred;
+                 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
+                 : SemaDiagnosticBuilder::K_Deferred;
     default:
-      return DeviceDiagBuilder::K_Nop;
+      return EmitOnBothSides ? SemaDiagnosticBuilder::K_Immediate
+                             : SemaDiagnosticBuilder::K_Nop;
     }
   }();
-  return DeviceDiagBuilder(DiagKind, Loc, DiagID,
+  return SemaDiagnosticBuilder(DiagKind, Loc, DiagID,
                                dyn_cast<FunctionDecl>(CurContext), *this);
 }
 
 bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   assert(Callee && "Callee may not be null.");
 
   auto &ExprEvalCtx = ExprEvalContexts.back();
   if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated())
     return true;
 
   // FIXME: Is bailing out early correct here?  Should we instead assume that
   // the caller is a global initializer?
   FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
   if (!Caller)
     return true;
 
   // If the caller is known-emitted, mark the callee as known-emitted.
   // Otherwise, mark the call in our call graph so we can traverse it later.
   bool CallerKnownEmitted =
       getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted;
-  DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee,
+  SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
                                           CallerKnownEmitted] {
     switch (IdentifyCUDAPreference(Caller, Callee)) {
     case CFP_Never:
     case CFP_WrongSide:
       assert(Caller && "Never/wrongSide calls require a non-null caller");
       // If we know the caller will be emitted, we know this wrong-side call
       // will be emitted, so it's an immediate error.  Otherwise, defer the
       // error until we know the caller is emitted.
-      return CallerKnownEmitted ? DeviceDiagBuilder::K_ImmediateWithCallStack
-                                : DeviceDiagBuilder::K_Deferred;
+      return CallerKnownEmitted
+                 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
+                 : SemaDiagnosticBuilder::K_Deferred;
     default:
-      return DeviceDiagBuilder::K_Nop;
+      return SemaDiagnosticBuilder::K_Nop;
     }
   }();
 
-  if (DiagKind == DeviceDiagBuilder::K_Nop)
+  if (DiagKind == SemaDiagnosticBuilder::K_Nop)
     return true;
 
   // Avoid emitting this error twice for the same location.  Using a hashtable
   // like this is unfortunate, but because we must continue parsing as normal
   // after encountering a deferred error, it's otherwise very tricky for us to
   // ensure that we only emit this deferred error once.
   if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second)
     return true;
 
-  DeviceDiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
+  SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
       << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
   if (!Callee->getBuiltinID())
-    DeviceDiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl,
-                      Caller, *this)
+    SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
+                          diag::note_previous_decl, Caller, *this)
         << Callee;
-  return DiagKind != DeviceDiagBuilder::K_Immediate &&
-         DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack;
+  return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
+         DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
 }
 
 // Check the wrong-sided reference capture of lambda for CUDA/HIP.
 // A lambda function may capture a stack variable by reference when it is
 // defined and uses the capture by reference when the lambda is called. When
 // the capture and use happen on different sides, the capture is invalid and
 // should be diagnosed.
 void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
   // only if the lambda structure is populated by a host function then passed
   // to and called in a device function or kernel.
   bool CalleeIsDevice = Callee->hasAttr<CUDADeviceAttr>();
   bool CallerIsHost =
       !Caller->hasAttr<CUDAGlobalAttr>() && !Caller->hasAttr<CUDADeviceAttr>();
   bool ShouldCheck = CalleeIsDevice && CallerIsHost;
   if (!ShouldCheck || !Capture.isReferenceCapture())
     return;
-  auto DiagKind = DeviceDiagBuilder::K_Deferred;
+  auto DiagKind = SemaDiagnosticBuilder::K_Deferred;
   if (Capture.isVariableCapture()) {
-    DeviceDiagBuilder(DiagKind, Capture.getLocation(),
+    SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
                           diag::err_capture_bad_target, Callee, *this)
         << Capture.getVariable();
   } else if (Capture.isThisCapture()) {
-    DeviceDiagBuilder(DiagKind, Capture.getLocation(),
+    SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
                           diag::err_capture_bad_target_this_ptr, Callee, *this);
   }
   return;
 }
 
 void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
     return;

clang/lib/Sema/SemaDecl.cpp

 
       MarkBaseAndMemberDestructorsReferenced(Destructor->getLocation(),
                                              Destructor->getParent());
     }
 
     // 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 (getDiagnostics().hasUncompilableErrorOccurred() ||
+    if (hasUncompilableErrorOccurred() ||
         getDiagnostics().getSuppressAllDiagnostics()) {
       DiscardCleanupsInEvaluationContext();
     }
-    if (!getDiagnostics().hasUncompilableErrorOccurred() &&
+    if (!hasUncompilableErrorOccurred() &&
         !isa<FunctionTemplateDecl>(dcl)) {
       // Since the body is valid, issue any analysis-based warnings that are
       // enabled.
       ActivePolicy = &WP;
     }
 
     if (!IsInstantiation && FD && FD->isConstexpr() && !FD->isInvalidDecl() &&
         !CheckConstexprFunctionDefinition(FD, CheckConstexprKind::Diagnose))

 
   if (!IsInstantiation)
     PopDeclContext();
 
   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 (getDiagnostics().hasUncompilableErrorOccurred()) {
+  if (hasUncompilableErrorOccurred()) {
     DiscardCleanupsInEvaluationContext();
   }
 
   if (LangOpts.OpenMP || LangOpts.CUDA || LangOpts.SYCLIsDevice) {
     auto ES = getEmissionStatus(FD);
     if (ES == Sema::FunctionEmissionStatus::Emitted ||
         ES == Sema::FunctionEmissionStatus::Unknown)
       DeclsToCheckForDeferredDiags.push_back(FD);

       checkThisInStaticMemberFunctionAttributes(Method);
 }
 
 /// ImplicitlyDefineFunction - An undeclared identifier was used in a function
 /// call, forming a call to an implicitly defined function (per C99 6.5.1p2).
 NamedDecl *Sema::ImplicitlyDefineFunction(SourceLocation Loc,
                                           IdentifierInfo &II, Scope *S) {
   // Find the scope in which the identifier is injected and the corresponding
   // DeclContext.
   // FIXME: C89 does not say what happens if there is no enclosing block scope.
   // In that case, we inject the declaration into the translation unit scope
   // instead.
   Scope *BlockScope = S;
   while (!BlockScope->isCompoundStmtScope() && BlockScope->getParent())
     BlockScope = BlockScope->getParent();
 
   Scope *ContextScope = BlockScope;

clang/lib/Sema/SemaExprObjC.cpp

static void applyCocoaAPICheck(Sema &S, const ObjCMessageExpr *Msg,
                                               const NSAPI &, edit::Commit &)) {
   SourceLocation MsgLoc = Msg->getExprLoc();
   if (S.Diags.isIgnored(DiagID, MsgLoc))
     return;
 
   SourceManager &SM = S.SourceMgr;
   edit::Commit ECommit(SM, S.LangOpts);
   if (refactor(Msg,*S.NSAPIObj, ECommit)) {
-    DiagnosticBuilder Builder = S.Diag(MsgLoc, DiagID)
-                        << Msg->getSelector() << Msg->getSourceRange();
+    auto Builder = S.Diag(MsgLoc, DiagID)
     // FIXME: Don't emit diagnostic at all if fixits are non-commitable.
     if (!ECommit.isCommitable())
       return;
     for (edit::Commit::edit_iterator
            I = ECommit.edit_begin(), E = ECommit.edit_end(); I != E; ++I) {
       const edit::Commit::Edit &Edit = *I;
       switch (Edit.Kind) {
       case edit::Commit::Act_Insert:
if (Method && Method->isDirectMethod()) {
     }
 
     // Under ARC, self can't be assigned, and doing a direct call to `self`
     // when it's a Class is hence safe.  For other cases, we can't trust `self`
     // is what we think it is, so we reject it.
     if (ReceiverType->isObjCClassType() && !isImplicit &&
         !(Receiver->isObjCSelfExpr() && getLangOpts().ObjCAutoRefCount)) {
       {
-        DiagnosticBuilder Builder =
-            Diag(Receiver->getExprLoc(),
+        auto Builder = Diag(Receiver->getExprLoc(),
                             diag::err_messaging_class_with_direct_method);
         if (Receiver->isObjCSelfExpr()) {
           Builder.AddFixItHint(FixItHint::CreateReplacement(
               RecRange, Method->getClassInterface()->getName()));
         }
       }
       Diag(Method->getLocation(), diag::note_direct_method_declared_at)
           << Method->getDeclName();
     }
 
     if (SuperLoc.isValid()) {
       {
-        DiagnosticBuilder Builder =
+        auto Builder =
             Diag(SuperLoc, diag::err_messaging_super_with_direct_method);
         if (ReceiverType->isObjCClassType()) {
           Builder.AddFixItHint(FixItHint::CreateReplacement(
               SuperLoc, Method->getClassInterface()->getName()));
         } else {
           Builder.AddFixItHint(FixItHint::CreateReplacement(SuperLoc, "self"));
         }
       }

 bool Sema::isKnownName(StringRef name) {
   if (name.empty())
     return false;
   LookupResult R(*this, &Context.Idents.get(name), SourceLocation(),
                  Sema::LookupOrdinaryName);
   return LookupName(R, TUScope, false);
 }
 
-static void addFixitForObjCARCConversion(Sema &S,
-                                         DiagnosticBuilder &DiagB,
-                                         Sema::CheckedConversionKind CCK,
-                                         SourceLocation afterLParen,
-                                         QualType castType,
-                                         Expr *castExpr,
-                                         Expr *realCast,
-                                         const char *bridgeKeyword,
-                                         const char *CFBridgeName) {
+template <typename DiagBuilderT>
+static void addFixitForObjCARCConversion(
+    Sema &S, DiagBuilderT &DiagB, Sema::CheckedConversionKind CCK,
+    SourceLocation afterLParen, QualType castType, Expr *castExpr,
+    Expr *realCast, const char *bridgeKeyword, const char *CFBridgeName) {
   // We handle C-style and implicit casts here.
   switch (CCK) {
   case Sema::CCK_ImplicitConversion:
   case Sema::CCK_ForBuiltinOverloadedOp:
   case Sema::CCK_CStyleCast:
   case Sema::CCK_OtherCast:
     break;
   case Sema::CCK_FunctionalCast:
S.Diag(loc, diag::err_arc_cast_requires_bridge)
       << castRange
       << castExpr->getSourceRange();
     bool br = S.isKnownName("CFBridgingRelease");
     ACCResult CreateRule =
       ARCCastChecker(S.Context, exprACTC, castACTC, true).Visit(castExpr);
     assert(CreateRule != ACC_bottom && "This cast should already be accepted.");
     if (CreateRule != ACC_plusOne)
     {
-      DiagnosticBuilder DiagB =
-        (CCK != Sema::CCK_OtherCast) ? S.Diag(noteLoc, diag::note_arc_bridge)
-                              : S.Diag(noteLoc, diag::note_arc_cstyle_bridge);
+      auto DiagB = (CCK != Sema::CCK_OtherCast)
+                       ? S.Diag(noteLoc, diag::note_arc_bridge)
 
       addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen,
                                    castType, castExpr, realCast, "__bridge ",
                                    nullptr);
     }
     if (CreateRule != ACC_plusZero)
     {
-      DiagnosticBuilder DiagB =
-        (CCK == Sema::CCK_OtherCast && !br) ?
-          S.Diag(noteLoc, diag::note_arc_cstyle_bridge_transfer) << castExprType :
-          S.Diag(br ? castExpr->getExprLoc() : noteLoc,
+      auto DiagB = (CCK == Sema::CCK_OtherCast && !br)
+                       ? S.Diag(noteLoc, diag::note_arc_cstyle_bridge_transfer)
+                             << castExprType
+                       : S.Diag(br ? castExpr->getExprLoc() : noteLoc,
                                 diag::note_arc_bridge_transfer)
                              << castExprType << br;
 
       addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen,
                                    castType, castExpr, realCast, "__bridge_transfer ",
                                    br ? "CFBridgingRelease" : nullptr);
     }
 
     return;
   }
S.Diag(loc, diag::err_arc_cast_requires_bridge)
       << castType
       << castRange
       << castExpr->getSourceRange();
     ACCResult CreateRule =
       ARCCastChecker(S.Context, exprACTC, castACTC, true).Visit(castExpr);
     assert(CreateRule != ACC_bottom && "This cast should already be accepted.");
     if (CreateRule != ACC_plusOne)
     {
-      DiagnosticBuilder DiagB =
-      (CCK != Sema::CCK_OtherCast) ? S.Diag(noteLoc, diag::note_arc_bridge)
-                               : S.Diag(noteLoc, diag::note_arc_cstyle_bridge);
+      auto DiagB = (CCK != Sema::CCK_OtherCast)
+                       ? S.Diag(noteLoc, diag::note_arc_bridge)
       addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen,
                                    castType, castExpr, realCast, "__bridge ",
                                    nullptr);
     }
     if (CreateRule != ACC_plusZero)
     {
-      DiagnosticBuilder DiagB =
-        (CCK == Sema::CCK_OtherCast && !br) ?
-          S.Diag(noteLoc, diag::note_arc_cstyle_bridge_retained) << castType :
-          S.Diag(br ? castExpr->getExprLoc() : noteLoc,
+      auto DiagB = (CCK == Sema::CCK_OtherCast && !br)
+                       ? S.Diag(noteLoc, diag::note_arc_cstyle_bridge_retained)
+                             << castType
+                       : S.Diag(br ? castExpr->getExprLoc() : noteLoc,
                                 diag::note_arc_bridge_retained)
                              << castType << br;
 
       addFixitForObjCARCConversion(S, DiagB, CCK, afterLParen,
                                    castType, castExpr, realCast, "__bridge_retained ",
                                    br ? "CFBridgingRetain" : nullptr);
     }
 
     return;
   }

clang/lib/Sema/SemaOpenMP.cpp

 /// Status of the function emission on the host/device.
 enum class FunctionEmissionStatus {
   Emitted,
   Discarded,
   Unknown,
 };
 } // anonymous namespace
 
-Sema::DeviceDiagBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
+Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
                                                          unsigned DiagID) {
   assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
          "Expected OpenMP device compilation.");
 
   FunctionDecl *FD = getCurFunctionDecl();
-  DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
+  SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop;
   if (FD) {
     FunctionEmissionStatus FES = getEmissionStatus(FD);
     switch (FES) {
     case FunctionEmissionStatus::Emitted:
-      Kind = DeviceDiagBuilder::K_Immediate;
+      Kind = SemaDiagnosticBuilder::K_Immediate;
       break;
     case FunctionEmissionStatus::Unknown:
       Kind = isOpenMPDeviceDelayedContext(*this)
-                 ? DeviceDiagBuilder::K_Deferred
-                 : DeviceDiagBuilder::K_Immediate;
+                 ? SemaDiagnosticBuilder::K_Deferred
+                 : SemaDiagnosticBuilder::K_Immediate;
       break;
     case FunctionEmissionStatus::TemplateDiscarded:
     case FunctionEmissionStatus::OMPDiscarded:
-      Kind = DeviceDiagBuilder::K_Nop;
+      Kind = SemaDiagnosticBuilder::K_Nop;
       break;
     case FunctionEmissionStatus::CUDADiscarded:
       llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation");
       break;
     }
   }
 
-  return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
+  return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
 }
 
-Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
+Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
                                                        unsigned DiagID) {
   assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
          "Expected OpenMP host compilation.");
   FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl());
-  DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
+  SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop;
   switch (FES) {
   case FunctionEmissionStatus::Emitted:
-    Kind = DeviceDiagBuilder::K_Immediate;
+    Kind = SemaDiagnosticBuilder::K_Immediate;
     break;
   case FunctionEmissionStatus::Unknown:
-    Kind = DeviceDiagBuilder::K_Deferred;
+    Kind = SemaDiagnosticBuilder::K_Deferred;
     break;
   case FunctionEmissionStatus::TemplateDiscarded:
   case FunctionEmissionStatus::OMPDiscarded:
   case FunctionEmissionStatus::CUDADiscarded:
-    Kind = DeviceDiagBuilder::K_Nop;
+    Kind = SemaDiagnosticBuilder::K_Nop;
     break;
   }
 
-  return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
+  return SemaDiagnosticBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
 }
 
 static OpenMPDefaultmapClauseKind
 getVariableCategoryFromDecl(const LangOptions &LO, const ValueDecl *VD) {
   if (LO.OpenMP <= 45) {
     if (VD->getType().getNonReferenceType()->isScalarType())
       return OMPC_DEFAULTMAP_scalar;
     return OMPC_DEFAULTMAP_aggregate;

clang/lib/Sema/SemaSYCL.cpp

 #include "clang/Sema/SemaDiagnostic.h"
 
 using namespace clang;
 
 // -----------------------------------------------------------------------------
 // SYCL device specific diagnostics implementation
 // -----------------------------------------------------------------------------
 
-Sema::DeviceDiagBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc,
+Sema::SemaDiagnosticBuilder 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] {
+  SemaDiagnosticBuilder::Kind DiagKind = [this, FD] {
     if (!FD)
-      return DeviceDiagBuilder::K_Nop;
+      return SemaDiagnosticBuilder::K_Nop;
     if (getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted)
-      return DeviceDiagBuilder::K_ImmediateWithCallStack;
-    return DeviceDiagBuilder::K_Deferred;
+      return SemaDiagnosticBuilder::K_ImmediateWithCallStack;
+    return SemaDiagnosticBuilder::K_Deferred;
   }();
-  return DeviceDiagBuilder(DiagKind, Loc, DiagID, FD, *this);
+  return SemaDiagnosticBuilder(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;
+  SemaDiagnosticBuilder::Kind DiagKind = SemaDiagnosticBuilder::K_Nop;
 
-  return DiagKind != DeviceDiagBuilder::K_Immediate &&
-         DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack;
+  return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
+         DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
 }

clang/lib/Sema/SemaStmt.cpp

if (!CaseListIsErroneous && !CaseListIsIncomplete && !HasConstantCond &&
         }
       }
 
       if (TheDefaultStmt && UnhandledNames.empty() && ED->isClosedNonFlag())
         Diag(TheDefaultStmt->getDefaultLoc(), diag::warn_unreachable_default);
 
       // Produce a nice diagnostic if multiple values aren't handled.
       if (!UnhandledNames.empty()) {
-        DiagnosticBuilder DB = Diag(CondExpr->getExprLoc(),
-                                    TheDefaultStmt ? diag::warn_def_missing_case
+        auto DB = Diag(CondExpr->getExprLoc(), TheDefaultStmt
+                                                   ? diag::warn_def_missing_case
                                                    : diag::warn_missing_case)
-                               << (int)UnhandledNames.size();
 
         for (size_t I = 0, E = std::min(UnhandledNames.size(), (size_t)3);
              I != E; ++I)
           DB << UnhandledNames[I];
       }
 
       if (!hasCasesNotInSwitch)
         SS->setAllEnumCasesCovered();

clang/lib/Sema/SemaStmtAsm.cpp

for (unsigned i = NumOutputs, e = NumOutputs + NumInputs; i != e; i++) {
     if (!Ty->isVoidType() || !Info.allowsMemory())
       if (RequireCompleteType(InputExpr->getBeginLoc(), Exprs[i]->getType(),
                               diag::err_dereference_incomplete_type))
         return StmtError();
 
     unsigned Size = Context.getTypeSize(Ty);
     if (!Context.getTargetInfo().validateInputSize(FeatureMap,
                                                    Literal->getString(), Size))
-      return StmtResult(
-          targetDiag(InputExpr->getBeginLoc(), diag::err_asm_invalid_input_size)
-          << Info.getConstraintStr());
+      return targetDiag(InputExpr->getBeginLoc(),
+                        diag::err_asm_invalid_input_size)
+             << Info.getConstraintStr();
   }
 
   // Check that the clobbers are valid.
   for (unsigned i = 0; i != NumClobbers; i++) {
     StringLiteral *Literal = Clobbers[i];
     assert(Literal->isAscii());
 
     StringRef Clobber = Literal->getString();

clang/lib/Sema/SemaTemplateInstantiate.cpp

Sema::InstantiatingTemplate::InstantiatingTemplate(
     SourceLocation PointOfInstantiation, SourceRange InstantiationRange,
     Decl *Entity, NamedDecl *Template, ArrayRef<TemplateArgument> TemplateArgs,
     sema::TemplateDeductionInfo *DeductionInfo)
     : SemaRef(SemaRef) {
   // Don't allow further instantiation if a fatal error and an uncompilable
   // error have occurred. Any diagnostics we might have raised will not be
   // visible, and we do not need to construct a correct AST.
   if (SemaRef.Diags.hasFatalErrorOccurred() &&
-      SemaRef.Diags.hasUncompilableErrorOccurred()) {
+      SemaRef.hasUncompilableErrorOccurred()) {
     Invalid = true;
     return;
   }
   Invalid = CheckInstantiationDepth(PointOfInstantiation, InstantiationRange);
   if (!Invalid) {
     CodeSynthesisContext Inst;
     Inst.Kind = Kind;
     Inst.PointOfInstantiation = PointOfInstantiation;

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

if (auto Name = D->getDeclName()) {
       Result = findInstantiationOf(Context, D,
                                    ParentDC->decls_begin(),
                                    ParentDC->decls_end());
     }
 
     if (!Result) {
       if (isa<UsingShadowDecl>(D)) {
         // UsingShadowDecls can instantiate to nothing because of using hiding.
-      } else if (Diags.hasUncompilableErrorOccurred()) {
+      } else if (hasUncompilableErrorOccurred()) {
         // We've already complained about some ill-formed code, so most likely
         // this declaration failed to instantiate. There's no point in
         // complaining further, since this is normal in invalid code.
         // FIXME: Use more fine-grained 'invalid' tracking for this.
       } else if (IsBeingInstantiated) {
         // The class in which this member exists is currently being
         // instantiated, and we haven't gotten around to instantiating this
         // member yet. This can happen when the code uses forward declarations

clang/lib/Sema/SemaTemplateVariadic.cpp

for (unsigned I = 0, N = Unexpanded.size(); I != N; ++I) {
 
     if (Name && NamesKnown.insert(Name).second)
       Names.push_back(Name);
 
     if (Unexpanded[I].second.isValid())
       Locations.push_back(Unexpanded[I].second);
   }
 
-  DiagnosticBuilder DB = Diag(Loc, diag::err_unexpanded_parameter_pack)
-                         << (int)UPPC << (int)Names.size();
+  auto DB = Diag(Loc, diag::err_unexpanded_parameter_pack)
   for (size_t I = 0, E = std::min(Names.size(), (size_t)2); I != E; ++I)
     DB << Names[I];
 
   for (unsigned I = 0, N = Locations.size(); I != N; ++I)
     DB << SourceRange(Locations[I]);
   return true;
 }
 

clang/lib/Sema/SemaType.cpp

if (fileInfo.getFileCharacteristic() != SrcMgr::C_User &&
     return FileID();
   }
 
   return file;
 }
 
 /// Creates a fix-it to insert a C-style nullability keyword at \p pointerLoc,
 /// taking into account whitespace before and after.
-static void fixItNullability(Sema &S, DiagnosticBuilder &Diag,
+template <typename DiagBuilderT>
+static void fixItNullability(Sema &S, DiagBuilderT &Diag,
                              SourceLocation PointerLoc,
                              NullabilityKind Nullability) {
   assert(PointerLoc.isValid());
   if (PointerLoc.isMacroID())
     return;
 
   SourceLocation FixItLoc = S.getLocForEndOfToken(PointerLoc);
   if (!FixItLoc.isValid() || FixItLoc == PointerLoc)

clang/test/SemaCUDA/deferred-all.cu

+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify=dev,com %s \
+// RUN:   -fgpu-defer-diag
+// RUN: %clang_cc1 -fsyntax-only -verify=host,com %s \
+// RUN:   -fgpu-defer-diag
+
+#include "Inputs/cuda.h"
+
+__device__ void callee(int);
+__host__ void callee(float); // host-note 2{{candidate function}}
+__host__ void callee(double); // host-note 2{{candidate function}}
+
+// Check no diagnostics for this function since it is never
+// called.
+
+inline __host__ __device__ void hdf_not_called() {
+  callee(1);
+  bad_line
+}
+
+// When emitted on device, there is syntax error.
+// When emitted on host, there is ambiguity and syntax error.
+  
+inline __host__ __device__ void hdf_called() {
+  callee(1); // host-error {{call to 'callee' is ambiguous}}
+  bad_line // com-error {{use of undeclared identifier 'bad_line'}}
+}
+
+// This is similar to the above but is always emitted on
+// both sides.
+
+__host__ __device__ void hdf_always_emitted() {
+  callee(1); // host-error {{call to 'callee' is ambiguous}}
+  bad_line // com-error {{use of undeclared identifier 'bad_line'}}
+}
+
+void hf() {
+ hdf_called(); // host-note {{called by 'hf'}}
+}
+ 
+__device__ void df() {
+ hdf_called(); // dev-note {{called by 'df'}}
+}
+
+// If a syntax error causes a function not declared, it cannot
+// be deferred.
+
+inline __host__ __device__ void bad_func() { // com-note {{to match this '{'}}
+// com-error {{expected '}'}}