This looks like one of the things we should *not* do as it affects correctness -- non-trivial constructor may be arbitrarily complex and the per-TU flag to enable this behavior is way too coarse, IMO.
On the other hand, I can believe that someone somewhere did write the code and relies to NVCC accepting it.

This is, or is very similar to, the problem that the host/device overloading addresses in CUDA.

Rebased on updated D59389

• Addressed Tim's comments.

• Addressed Tim's comments.

The intent is to avoid unintentional clashes with the preprocessor macros the user may have defined.
https://reviews.llvm.org/rL260647

In D59423#1431151, @emankov wrote:

Perhaps for consistency sake it would be better to replace __sptr -> __s and __cptr -> __c.

Well, it came from NVIDIA code, you know, I mean all those double underscores.

___ stands out as a sore thumb and raises unnecessary questions -- "why does it have three underscores, while __cptr is fine with two?".
Perhaps for consistency sake it would be better to replace __sptr -> __s and __cptr -> __c.
Given that we're just passing the args through, we don't really need to have ptr here.

Thank you for landing the workaround!

This patch fixes the bug in the ptxas tool with the processing of bytes

OK. I'll file my own and will add a comment that it may be a known issue.

Do you have NVIDIA's bug number for the crasher?

Something about DWARF debug info is making ptxas upset to the point of SIGSEV'ing. :-/

The change looks OK as far as regular CUDA is concerned.
That said, I'm not quite familiar with the use of bundling/unbundling actions and you should probably get someone who uses/depends on them to take a look. I think OpenMP uses them. Perhaps @ABataev would be the right person.

LGTM.

E.g.:

namespace {
__host__ __device__ a() {
  int prev;
  __asm__ __volatile__("mov %0, 0" : "=a" (prev)::);
  return prev;
}

__host__ __device__ b() {
  int prev;
  return prev;
}

} //namespace

Ideally we should always emit uninitialized diagnostics for b, but never for a in both host and device compilation modes.
I think we may want to propagate assignment from the inline asm statement -- we may not know the meaning of the constraint, but we do know which argument gets used/modified by the asm statement. Perhaps we can construct a fake GCCAsmStmt but bail out before we attempt to validate the asm string.

But it is going to be emitted for b() if b() is really used on the host or on the device.

Hi Artem, I think we can just delay emission of this warning to solve this problem.

There's a new quirk we've ran into after this patch landed. Consider this code:

int foo() {
  int prev;
  __asm__ __volatile__("whatever" : "=a" (prev)::);
  return prev;
}

Could you add a test to demonstrate the issue and verify that it is fixed?

In D58518#1406274, @t-tye wrote:

To clarify, I am saying that the stub does have a different name since it is conceptually part of the implementation of doing the call to the device function implementation, and is not in fact the the device function being called itself. However, when we generate code for a function that is present on both the host and device, both copies of the code are for the same source level function and so can have the same symbol name (which was a question that was asked)

In D58518#1406202, @t-tye wrote:

Yes this relates to supporting the debugger.

For the same function being present on both host and device, having the same name is correct as the debugger must set a breakpoint at both places. This is similar to needing to set a breakpoint at every place a function is inlined.

My guess is that this is needed because HIP debugger can see symbols from both host and device executables at the same time. Is that so?

Thank you.

In D56411#1400300, @rjmccall wrote:

Okay, but it's not great design to have a kind of overloading that can't be resolved to an exact intended declaration even by an explicit cast. That's why I think making *optional* host/device typing is a good idea. And I strongly want to caution you against doing language design by just incrementally hacking at the compiler to progressively make more test-cases work, which is what it feels like you're doing.

That said, does CUDA have a general rule resolving __host__ vs. __device__ overloads based on context? And does it allow overloading based solely on __host__ vs. __device__?

Thank you. LGTM.

LGTM. I've added jlebar@ as he's originally written the code.

In D57829#1387416, @yaxunl wrote:

In D57829#1387412, @tra wrote:

Could you elaborate on why you want to disable this metadata? I think the original idea of llvm.linker.options was that it should be ignored if the back-end does not support it.

If backend does not support it, it goes to TargetLoweringObjectFileELF::emitModuleMetadata and causes codegen to fail.

Could you elaborate on why you want to disable this metadata? I think the original idea of llvm.linker.options was that it should be ignored if the back-end does not support it.

Made a comment more readable.

Make the function object local.

Updated ASTMatchers unit test.

Addressed Justin's comments.

Addressed Justin's comments.

I would be ok with reusing that option, as long as it's documented that there is a difference in terms of how it can be used.

Updated the comment about SDKVersion use.

Would that be OK to use target_sdk_version to pass *CUDA* SDK version to the CC1 compilations?
I have upcoming changes that need to know the version to generate correct glue IR for CUDA. The driver currently figures out detected CUDA version in lib/Driver/ToolChains/Cuda.cpp and I could use -target-sdk-version to pass it on to CC1 instances.

In D57259#1372234, @stella.stamenova wrote:

I haven't had time to look into it in details but it looks like either this change or https://reviews.llvm.org/D57186 broke the windows bot.

http://lab.llvm.org:8011/builders/lldb-x64-windows-ninja/builds/1030
http://lab.llvm.org:8011/builders/lldb-x64-windows-ninja/builds/1030/steps/test/logs/stdio

Fixed a typo.

Nice. Thank you for adding the tests.
LGTM.

The tests you have only check the end result and do not directly verify the new functionality the patch adds.

You probably want to remove libcalls.patch from the patch diff.

In D55145#1331301, @denzp wrote:

In D55145#1324230, @tra wrote:

Here the behaviour of unsupported instructions is changed to match the behaviour of explicit intrinsics calls.

If I remember correctly, the crash was only happening for instructions lowered to libcall. Do we crash in any other case?
If libcall is the only source of failure, would it make more sense to bail out of this function early? We're going to fail in this case anyways.

You are right, normally it would make sense. But in this case, the end goal is to support these libcalls.
It would be easier to achieve when both instructions and intrinsics are behaving consistently. And I'm afraid, bailing them out, would be an extra work that will have to be dropped when the final libcall solution will come (means, very-very soon).

jlebar@ LGTM'ed via email.
Landed in rL348662