Use correct flag naming.

Rebase.

Gentle ping.

Use correct capitalisation in helper.

Address review feedback around explicit use of std::.

In D155856#4644379, @yaxunl wrote:

Are these passes called when compiling C++ to bitcode? If so, then device functions not called by kernels in the same module will be removed, right? Then we cannot support calling device functions in a different TU. Should these passes be moved to the llvm codegen pipelines so that they will only be called post-linking?

Fix some wording to further clarify this is an HIP + AMDGPU exclusive extension.

Correctly reflect AMDGPU-only nature of the extension.

Add back clarification about RDC.

Fix typo, add release notes reflecting the HIP-only nature.

Define the interpose macro iff in interpose mode.

Gentle ping.

Remove noise.

Update per review comments.

Use --hipstdpar instead of -hipstdpar in the test.

Updating this to reflect the outcome of the RFC, which is that we're adding it as a HIP extension exclusively. Added support for handling thread_local, which is unsupported.

Updating to reflect the outcome of the RFC, which is that this will be added as a HIP extension exclusively.

Updating this to reflect the outcome of the RFC, which is that it will be added as a HIP extension exclusively.

Updating this to reflect the outcome of the RFC, which is that this will be added as a HIP extension exclusively.

Updating this to reflect the outcome of the RFC, which is that this will be added as a HIP extension exclusively.

Updating this to reflect the outcome of the RFC, which is that this shall be a HIP only extension. As such, documentation lives within the HIP Support master document; the other patches in the series will be updated accordingly.

Add handling of thread_local variables that are used by accelerator reachable functions: since this is currently unsupported, we error out with as tidy a diagnostic as possible.

Remove unneeded cast, the dynamic case already emitted a generic pointer to typeinfo

Update test.

LGTM, thanks.

In D157452#4586554, @rjmccall wrote:

The path of least resistance here is that IRGen should just insert conversions from the global AS to the default AS as part of evaluating typeid. I haven't looked at it closely, but that seems to be what this patch is doing.

However, std::type_info is an interesting special case in that we actually know statically that all objects of that type will be allocated in the global AS, so there's really no reason to pass around pointers in the default AS; std::type_info * should just default to being in the global AS. It'd be a non-trivial feature in support of a somewhat uncommonly-used C++ feature, and I can't tell how best to spend your time, *but*... if you're so inclined, I think it would make a somewhat compelling feature to be able to declare a default AS for a class type, which your target could then adopt in the headers for std::type_info.

Rework the patch as the proposed approach was unsound; keep typeid generic.

Add support for handling incompatible ASM.

Switch to __ASM prefix.

Fix typo.

Remove noise / unintended file. Add support for dealing with unsupported ASM.

Fix typo.

Add support for handling certain cases of unambiguously accelerator unsupported ASM i.e. cases where constraints are clearly mismatched. When that happens, we instead emit an ASM__stdpar_unsupported stub which takes as its single argument the constexpr string value of the ASM block. Later, in the AcceleratorCodeSelection pass, if such a stub is reachable from an accelerator callable, we error out and print the offending ASM alongside the location.

In D157452#4573076, @yaxunl wrote:

It is a little concerning how far the global address will spread further. Compared to handling user-defined global variables, we keep the global address to its definition in the IR and any use of it will use the generic pointer addrcasted from its definition. This simplifies things a lot since the AST is not aware of the global address. Should we reconsider the handling of type id and type info here with a similar approach to how the user-defined global variables are handled? Or we are confident that the effect of global address can be confined.

In D155850#4570336, @efriedma wrote:

LGTM (but please don't merge until we reach consensus on the overall feature)

Extend handling of unsupported builtins to include dealing with the target attribute.

Remove confusing guidance around mGPU.

Update diff, handling of globals was wrong in that it'd mess up extern __shared__ cases.

Remove noise, correct style.

Rebase.

Ping?

In D156713#4549890, @foad wrote:

In D156713#4547539, @b-sumner wrote:

In D156713#4547505, @arsenm wrote:

In D156713#4547504, @b-sumner wrote:

In D156713#4547434, @foad wrote:

Do we want a separate GWS feature? Or a combined "hasGDSAndGWS" feature?

We do need a feature that include GWS. GDS was dropped before GWS in GFX9, but compute never used GDS. So GDSandGWS would probably work.

I didn't realize GDS was dropped in gfx9, should split the two then

Sorry for mangling this. GDS was dropped in gfx90a. GWS was not dropped in gfx9.

What about gfx909 and gfx90c?

In the current patch I have assumed gfx908 and gfx909 both have GDS, and gfx90a and gfx90c (and gfx94x) lack it.

Use pointee's address space for the check.

Address review comments.

In D156539#4547814, @rjmccall wrote:

In D156539#4546834, @AlexVlx wrote:

In D156539#4542836, @rjmccall wrote:

We should probably write this code to work properly in case we add a target that makes __builtin_alloca return a pointer in the private address space. Could you recover the target AS from the type of the expression instead of assuming LangAS::Default?

I believe it should work as written (it is possible that I misunderstand the objection, case in which I do apologise). The issue is precisely that for some targets (amdgcn being one) alloca returns a pointer to private, which doesn't compose with the default AS being unspecified / nonexistent, for some languages (e.g. HIP, C/C++ etc.). For the OCL case @arsenm mentions, I believe that we make a choice based on LangOpts.OpenCLGenericAddressSpace, and the code above would only be valid from the OCL language perspective iff that is true. I've put together a short Godbolt that captures these (and the bug the patch should fix, please see the bitcasts between ASes in the non-OCL case): https://gcc.godbolt.org/z/sYGK76zqv.

An address space conversion is required in IR if there is a difference between the address space of the pointer type formally returned by the call to __builtin_alloca and the address space produced by the alloca operation in IR. If Sema has set the type of __builtin_alloca to formally return something in the stack address space, no conversion is required. What I'm saying that I'd like you to not directly refer to LangAS::Default in this code, because you're assuming that __builtin_alloca is always returning a pointer that's formally in LangAS::Default. Please recover the target address space from the type of the expression.

Additionally, in IRGen we allow the target to hook address-space conversions; please call performAddrSpaceCast instead of directly emitting an addrspacecast instruction.

Incorporated review feedback. Updated test.

In D156539#4542836, @rjmccall wrote:

We should probably write this code to work properly in case we add a target that makes __builtin_alloca return a pointer in the private address space. Could you recover the target AS from the type of the expression instead of assuming LangAS::Default?

Exploit the fact that ROCm/AMDGPU does LTCG at the moment and for the foreseeable future by moving the accelerator code selection pass to later.

Add support for dealing with builtins that are unavailable on the target.

This adds more ecumenical handling of unsupported builtins, as per the review discussion (a suffixed equivalent stub is emitted instead); it's paired with an associated change in accelerator code selection pass, where the actual check for these stubs occurs. I've also adjusted where the latter pass gets added to the opt pipeline, for the AMDGCN target; for the latter it's better, for the moment, to run it later because we essentially do LTCG, and therefore can unambiguously determine reachability by operating on the full module.

Ping @rjmccall or @efriedma for a non AMD perspective, if possible.

In D155870#4535425, @yaxunl wrote:

__dynamic_cast is part of standard C++ library. If we ever implement it for GPU, chances are we will use libc++abi with the same signature as other targets, i.e., the 2nd and 3rd arguments are generic pointers.

I feel it is safer to do an address space cast when calling the function instead of changing its signature.

Gentle ping.

Rebase.

@yaxunl interesting point - are you worried about cases where due to missing inlining / const prop an indirect call site that can be replaced with a direct one would remain indirect? I think the problem in that case would actually be different, in that possibly reachable functions would not be identified as such and would be erroneously removed. I'm not sure there's any case where we'd fail to remove a meant to be unreachable function. We can definitely go with the __clang_unsupported approach, but I think I'd prefer these to be compile time errors rather than remarks + runtime printf, not in the least because printf adds some overhead. A way to ensure we don't "miss a spot" might be to check after removal for any remaining unsupported builtins, instead of doing it during reachability computation (this is coupled with the special naming from the prior post).

Removed some accidental noise.

Rebase.

Ping.

Rebase.

Thank you @yaxunl. @rjmccall / @efriedma any input on this? I'd like to try landing it next week to unblock some additional work. Thanks!

In D153092#4452070, @yaxunl wrote:

This could be a good chance to switch VT to constant address space instead of global address space. AFAIK if a target has global addr space they usually also has constant addr space since they usually support OpenCL or CUDA/HIP. Is there any reason we cannot introduce a CGM.ConstantGlobalsInt8PtrTy and use it for VT instead?

In D153092#4448514, @yaxunl wrote:

In D153092#4447445, @AlexVlx wrote:

Fixed issue found via internal testing (thanks @yaxunl).

Can we add a test to cover the regression found via internal testing? Thanks.