Landed ir r365913

LGTM. Thank you!

In D63756#1557658, @cdevadas wrote:

Hi Sam,
The compiler generates metadata for the first 48 bytes. I compiled a sample code and verified it. The backend does nothing for the extra bytes now.
I will soon submit the backend patch to generate the new metadata.

Overall looks good. Thank you for making the change.

SGTM in principle. Folding the stubs would be bad as their addresses are implicitly used to identify the kernels to launch.

LGTM, but I'm not familiar with the use of the stuff under Testing/. I've added @zturner who did some work there.

LGTM. This is a cleaner way to provide stub name tweaks.

In D63335#1544324, @hliao wrote:

I think debugger does have sufficient information to deal with this and that would be the right place to deal with the issue.

em, I did push the later as well, :(. OK, I will simplify the patch to change any functionality but move the calculation of device name into a common interface. So that, vendor could adjust that internally with minimal change. OK?

In D63335#1544315, @hliao wrote:

Sorry, I still don't think I understand the reasons for this change. The stub and the kernel do have a different name now. I don't quite get it why the debugger can differentiate the names when they differ by prefix, but can't when they differ by suffix. It sounds like an attempt to work around a problem somewhere else.

Could you talk to the folks requesting the change and get more details on what exactly we need to do here and, more importantly, why.

But, after unmangling, debugger still could match both as they are almost identical excep the final variants, like clone. The debugger will set all locations matching that specified kernel name.

In D63335#1544026, @hliao wrote:

Is it OK for us to mangle __device_stub __ as the nested name into the original one, says, we prepend _ZN15__device_stub__E, so that we have _ZN15__device_stub__E10kernelfuncIiEvv

and

$ c++filt _ZN15__device_stub__E10kernelfuncIiEvv
__device_stub__(kernelfunc<int>, void, void)

In D63335#1544019, @hliao wrote:

In D63335#1543854, @tra wrote:

In D63335#1543845, @hliao wrote:

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

I understand that you want a different name for the stub. My question is why the ".stub" suffix was not sufficient and how does having a prefix instead helps? Making the name un-demangleable is undesirable, IMO. There should be a good reason to justify it.

it's based on debugger people told me, with ".stub", the debugger still could find it match the original device kernel even though it could find both of them. But, they want to match the original one only and leave the stub one intentionally unmatched.

In D63335#1543845, @hliao wrote:

it's requested from debugger people. they don't want to the host-side stub could match the device-side kernel function name. the previous scheme cannot prevent that.

Is there particular reason you need to switch to this naming scheme?

Yay! Thank you. Lack of this has stopped me when I tried gn last time. I'll try again once this patch lands.

In D62393#1542098, @__simt__ wrote:

In D62393#1542042, @tra wrote:

C++ volatile will give you that. You also rely on atomicity. C++ volatile does not guarantee that, even if NVPTX does happen to. It's a mere coincidence. What if next PTX revision provides a better way to implement C++ volatile without providing atomicity? Then your code will all of a sudden break in new and exciting ways.

I think it's completely healthy and fair to be skeptical that we'll do the right thing by you, but FWIW the changes I want to make to CUDA volatile after we ship std::atomic<T> go in the other direction. And the legacy requirements are pretty entrenched here, too.

In D62393#1541976, @ABataev wrote:

In D62393#1541969, @tra wrote:

@reames , @tra , @Hahnfeld , @jlebar , @chandlerc, I see that this was discussed in D50391 (in terms of using PTX's volatile to provide atomic lowering), but it's not clear to me that using volatile with atomic semantics in LLVM is something we guarantee will work (even in CUDA mode). I imagine that we'd need to lower these in terms of relaxed atomics in Clang for this to really be guaranteed to work correctly.

Do I understand it correctly that the argument is about using C++ volatile with the assumption that those will map to ld.volatile/st.volatile, which happen to provide sufficiently strong guarantees to be equivalent of LLVM's atomic monotonic operations?

If that's the case, then I would agree that it's an implementation detail one should not be relying on. If atomicity is needed, we should figure out a way to express that in C++.

In practice, the standard C++ library support in cuda is somewhat limited. I don't think atomic<> works, so volatile might be a 'happens to work' short-term workaround. If that's the case, then there should a comment describing what's going on here and TODO to fix it when better support for atomics on GPU is available.

Another option would be to use atomic*() functions provided by CUDA, but those do have limited functionality on older GPUs.

Yet another alternative is to explicitly use ld.volatile/st.volatile from inline assembly. I don't think we have it plumbed as clang builtin.

Artem, thanks for your comment. But we need not only atomicity, bht also we need to command the ptxas compiler to not optimize accesses to parallelLevel.

@reames , @tra , @Hahnfeld , @jlebar , @chandlerc, I see that this was discussed in D50391 (in terms of using PTX's volatile to provide atomic lowering), but it's not clear to me that using volatile with atomic semantics in LLVM is something we guarantee will work (even in CUDA mode). I imagine that we'd need to lower these in terms of relaxed atomics in Clang for this to really be guaranteed to work correctly.

So, in short, what you're saying is that lambda type may leak into the mangled name of a __global__ function and ne need to ensure that the mangled name is identical for both host and device, hence the need for consistent naming of lambdas.

So, the only thing this patch appears to do is make everything with this attribute uninitialized on device side and give protected visibility.
If I understand it correctly, you're using the attribute in order to construct something that's sort of opposite of the currently used device vars with host-side shadows. Only now the real variable lives on the host side and it's device side that gets the 'shadow' copy. Do I understand it correctly?

Syntactically the patch looks OK to me, but I think the purpose and meaning of the builtin type should be documented in more details. Based on this patch alone it's not clear at all what it's supposed to be used for and how.

In D62603#1521979, @yaxunl wrote:

I think static __device__ globals would fall into the same category -- nominally they should not be visible outside of device-side object file, but in practice we do need to make them visible from the host side of the same TU.

Are you sure nvcc support accessing static __device__ variables in host code? That would be expensive to implement.

In D62603#1521792, @hliao wrote:

that should assume that variable is not declared with static. that's also the motivation of this patch.

NVCC also allows that: https://godbolt.org/z/t78RvM

BTW, that code posted looks quite weird to me, how the code could make sense by return a pointer of device variable? or a pointer of shadow host variable?

Note for the future -- it would be great if we could finish discussing the patch before landing it.
I would still like to see the host-side test.

In D62603#1521484, @hliao wrote:

thanks, but that static __device__ variable won't have shadow in host anymore.

I'd add a comment with a brief explanation for the const variant and a TODO() to remove it.

LGTM.

In D61765#1500457, @gtbercea wrote:

This won't affect CUDA in any way, all we have added is OpenMP specific.

In D61765#1500309, @gtbercea wrote:

As soon as libc++ the limits header included in

__clang_cuda_cmath.h:15
``` is not found:

__clang_cuda_cmath.h:15:10: fatal error: 'limits' file not found
#include <limits>

Not even CUDA works actually so I'm not sure what the best answer to this problem is.

In D61765#1499957, @jdoerfert wrote:

Two small changes and then it is fine with me. @tra ?

In D61458#1488550, @hliao wrote:

In D61458#1488523, @tra wrote:

Perhaps we should allow this in all unevaluated contexts?
I.e. int s = sizeof(foo(x)); should also work.

good point, do we have a dedicated context for sizeof? that make the checking easier.

Perhaps we should allow this in all unevaluated contexts?
I.e. int s = sizeof(foo(x)); should also work.

LGTM, but I've added @rsmith who is way more familiar with this code.

In D60907#1484756, @jdoerfert wrote:

I actually don't want to preinclude anything and my arguments are (mostly) for the OpenMP offloading code path not necessarily Cuda.
Maybe to clarify, what I want is:

1. Make sure the clang/Headers/math.h is found first if math.h is included.
2. Use a scheme similar to the one described https://reviews.llvm.org/D47849#1483653 in clang/Headers/math.h
3. Only add math.h function overloads in our math.h. <- This is debatable

+1 to Hal's comments.

A kernel functions in CUDA is actually two different functions. One is the real kernel we compile for the GPU, another is a host-side stub that launches the device-side kernel.

FYI, I have almost-ready set of patches to implement missing bits of sm_75 support, including this change:
https://reviews.llvm.org/D60279

LGTM. Thank you for cleaning this up.

In D60727#1467241, @echristo wrote:

Nick: Can you do some archaeology on the original patch and find out if there was supposed to be something supported here?

Art: Any thoughts?

@arsenm Matt, FYI, this patch seems to be a continuation of D59863 you've commented on.

It looks like you are solving two problems here.
a) you want to create multiple device passes for the same GPU, but with different options.
b) you may want to pass different compiler options to different device compilations.
The patch effectively hard-codes {gpu, options} tuple into --offloading-target-id variant.
Is that correct?

• Converted class to struct+function as Tim suggested.

In D60220#1456430, @emankov wrote:

Oooh, sorry, but I've just pushed the fix. But with the following words: "Add missing long double device functions' declarations. Provide only declarations to prevent any use of long double on the device side, because CUDA does not support long double on the device side."

One more thing -- perhaps the long double declarations should be put under #ifndef _MSC_VER in all the files to make the change unobservable on non-windows platforms.

Thank you for fixing this!

• Added PTX64 to the list of builtins' constraints.

• Fixed minor issues with parameters of the new builtins:
  • __imma*_st_c_i32 builtins have 'const int * src'
  • __bmma_m8n8k128_mma_xor_popc_b1 does not have 'satf' argument.

Cleaned up mma test generation.

• Enabled .satf for s4/u4.