In D101630#2799425, @yaxunl wrote:

But how do we control emitting LLVM IR with or without bundle? -emit-llvm -emit-gpu-object or -emit-llvm -emit-gpu-bundle? -emit-* is usually for specifying a specific file type.

In D101630#2798975, @yaxunl wrote:

For sure we will need -fgpu-bundle-device-output to control bundling of intermediate files. Then adding -emit-gpu-object and -emit-gpu-bundle may be redundant and can cause confusion. What if users specify -c -fgpu-bundle-device-output -emit-gpu-object or -c -fno-gpu-bundle-device-output -emit-gpu-bundle? To me a single option -fgpu-bundle-device-output to control all device output seems cleaner.

LGTM for CUDA.

Still LGTM.

I'm done with testing. The patch does not seem to break anything obvious. Tensorflow builds and works.

In D101630#2792052, @yaxunl wrote:

I think for intermediate outputs e.g. preprocessor expansion, IR, and assembly, probably it makes sense not to bundle by default.

LGTM. I would like to test the patch on our code first. Please wait a bit before landing the patch. I should be able to have the results tomorrow.

In D101630#2787714, @yaxunl wrote:

How does nvcc --genco behave when there are multiple GPU arch's? Does it output a fat binary containing multiple ISA's? Also, does it support device-only compilation for intermediate outputs?

Context for the changes in this patch: https://reviews.llvm.org/D97340#2775477

Overall looks good, though I've got one more question.

In D101630#2777346, @yaxunl wrote:

In D101630#2748513, @tra wrote:

How about this:
If the user explicitly specified --cuda-host-only or --cuda-device-only, then by default only allow producing the natural output format, unless a bundled output is requested by an option. This should keep existing users working.
If the compilation is done without explicitly requested sub-compilation(s), then bundle the output by default. This should keep the GPU-unaware tools like ccache happy as they would always get the single output they expect.

WDYT?

--cuda-host-only always have one output, therefore there is no point of bundle its output. We only need to decide the proper behavior of --cuda-device-only.

In D102936#2774686, @jwakely wrote:

You can't use __GLIBCXX__ this way. It will be different for different snapshots from the gcc-11 branch. Some distros are already shipping gcc-11 snapshots with later dates.

I would just check RELEASE == 11. If __failed_assertion is present, you'll rename it. If it's not present, nothing gets renamed but it works anyway.

Check only _GLIBCXX_RELEASE

Fixed typo in push/pop macro name.

I've verified that Tensorflow still builds with this patch and that the patch does fix the regressions we've seen.
If you could land this patch soon, that would be appreciated.

In D102801#2771664, @yaxunl wrote:

In the updated patch I have a simpler solution which is easier to explain to the users. Basically we classify variables by how they are emitted: device side only, host side only, both sides as different entities (e.g. default constexpr var), and both sides as unified entity (e.g. managed var). For variables emitted on both sides as separate entities, we have limited knowledge and we limit what we can do for them. I think users should understand the compiler's limitation in such cases. And they can easily workaround that by making the variable explicitly device variable.

This patch does not appear to fix the second regression introduced by the D102237.

In D102237#2769475, @yaxunl wrote:

In D102237#2767538, @tra wrote:

Here's a slightly simpler reproducer: https://godbolt.org/z/rW6P9e37s

I have a fix for this: https://reviews.llvm.org/D102801

Tentative LGTM as we need it to fix the regression soon.

Here's a slightly simpler reproducer: https://godbolt.org/z/rW6P9e37s

Here's one example reproducer: https://godbolt.org/z/77M596W89
It's rather hairy, but should be usable for further debugging.

In D102237#2767131, @yaxunl wrote:

Sam, this patch has apparently triggered some unwanted side effects. I'm still reducing the failures to something that could be used for debugging, but the rough symptoms are:

I'll land this patch along with D100124

LGTM with a test nit.

In D102306#2760775, @msearles wrote:

In D102306#2758166, @JonChesterfield wrote:

Is there documentation for the mapping from product names to gfx numbers?

See the Processors section of https://llvm.org/docs/AMDGPUUsage.html

In effect this patch applies __host__ __device__ to a subset of the standard library headers and whatever headers *they* happen to include. While it may happen to work, I'm not at all confident that it does not create interesting issues.

In D100124#2757251, @steffenlarsen wrote:

@tra Thanks a ton for the review! This is my first LLVM patch so I only know as much as the Code Review documentation tells me. Is there a process for chasing up additional reviews?

LGTM in general.
Perhaps it would make sense to combine this patch with D102237 as both patches are changing the same code for the same reason, just for slightly different kinds of variables.

LGTM for CUDA. This matches the intent of deferred diags -- we only emit them if we get to generate the code for the sources that triggered them, so they should not show up for the false constexpr branches.

LGTM with few nits.

In D101630#2744861, @yaxunl wrote:

[snip] it is the convention for compiler to have one output.
The compilation is like a pipeline. If we break it into stages, users would expect to use the output from one stage as input for the next stage. This is possible only if there is one output.
Also, when users do not want the output to be bundled, it is usually for debugging or special purposes. Users need to know the naming convention of the multiple outputs. I think it is justifiable to enable this by an option.

In D101630#2730273, @yaxunl wrote:

How about an option -fhip-bundle-device-output. If it is on, device output is bundled no matter how many GPU arch there are. By default it is on.

What will happen with this patch in the following scenarios:

• --offload_arch=A -S -o out.s
• --offload_arch=A --offload-arch=B -S -o out.s

CUDA compilation currently errors out if -o is used when more than one output would be produced.
E.g.

% bin/clang++ -x cuda --offload-arch=sm_60 --offload-arch=sm_70 --cuda-path=$HOME/local/cuda-10.2  zz.cu -c -E 
#... preprocessed output from host and 2 GPU compilations is printed out

In D100060#2726888, @ashi1 wrote:

@tra have you seen this exception in your CUDA buildbots before?
https://lab.llvm.org/staging/#/builders/152/builds/3

LGTM.

Few nits. LGTM overall.

In D99997#2709907, @JonChesterfield wrote:

Tagging @jdoerfert as HIP CI may be of interest. I wonder if an openmp directory added here could be persuaded to run on nvptx or amdgpu hardware, based on what happens to be in the bot.

In D100124#2707731, @steffenlarsen wrote:

Do you know if any existing code already uses the __nvvm_* builtins for cp.async? In other words, does nvcc provide them already or is it something we're free to name as we wish? I do not see any relevant intrinsics mentioned in NVVM IR spec: https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html and I don't think NVCC's builtins are publicly documented anywhere.

I don't know of any yet. We will be using these in the relatively near future, but we can still change them no problem. However, the intrinsic and builtin naming for NVVM and NVPTX seems a bit inconsistent so it may be a long discussion (or maybe not.)

General style nit: the script has very inconsistent quoting for the variables. They are quoted in some places but not others.

Do you know if any existing code already uses the __nvvm_* builtins for cp.async? In other words, does nvcc provide them already or is it something we're free to name as we wish?
I do not see any relevant intrinsics mentioned in NVVM IR spec: https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html and I don't think NVCC's builtins are publicly documented anywhere.

The planned new option for offloading will be a more generic solution, however, I expect it will take time to develop and be adopted.

LGTM. I'll leave the patch approval to @gkistanova .

LGTM.

LGTM overall, modulo few test and naming nits.

In D100404#2701902, @t-tye wrote:

In D100404#2701630, @palves wrote:

Sure, it will work for me, though FWIW I think it's the worse option of the 3 potential solutions, and I'll be puzzled if we end with that UI.
Can anyone else speak up and state your opinions, please?

I agree with @palves that this seems the least consistent of the options. Is there an objection to to the more consistent and more intuitive alternative?

LGTM overall.

Small test nit. LGTM otherwise.

Overall the patch looks good. We may still need to tweak intrinsic properties later, but this is a good starting point.