In D88250#2293346, @tra wrote:

In D88250#2293270, @jlebar wrote:

I know it comes in a separate change, but can we add a check to the test-suite?

Will do.

In D88250#2293270, @jlebar wrote:

I know it comes in a separate change, but can we add a check to the test-suite?

Fixed compatibility with pre-c++11

In D84362#2283078, @tra wrote:

It's possible. Unfortunately it's only triggered by our internal tool and it's hard to create a public reproducer for it. I'll debug and try to fix it on Monday.

In D84362#2283065, @yaxunl wrote:

I think this is probably a different issue. The issue reported in D84364 was introduced by change in D84364.

In D84362#2282992, @yaxunl wrote:

The fix is for the change in D84364. It has no effect on the change in this review. Are you sure the issue you saw is due to change in this review instead of change in D84364?

In D84362#2282890, @yaxunl wrote:

I have a fix for the issue reported in D84364. Would you like to try? Thanks.

In D87325#2280313, @ashi1 wrote:

@tra , I tried to invoke clang with -Xclang -emit-pch, but the -x hip path doesn't know about the precompile phase. It does pass -emit-pch to the cc1 command, but is overridden by the -emit-obj default flag to cc1 in compiler phase. Also, from there it will go through backend, assembler, linker, which we will need to disable if we want this method to work. Here is the experiment:

root@e6915ef660c7:~/llvm-project/build_rel# ./bin/clang++ -x hip --cuda-device-only --cuda-gpu-arch=gfx803 -Xclang -emit-pch a.hip -o a.hip.pch -ccc-print-bindings

If you dd -S and remove -ccc-print-bindingsthis command does produce a PCH.

In D87791#2279864, @yaxunl wrote:

It is requested by our debugger team, so it should work with amdgpu.

In D84362#2279688, @tra wrote:

Apparently this patch triggers compiler crashes on some of our code. I'll try to create a reproducer, but it would be great to revert the patch for now.

In D87791#2278417, @yaxunl wrote:

Therefore in either case there is no need to rename the intermediate .o files since they are temporary files which have unique names.

The .dwo files are not temporary files. They are supposed to be shipped with .o files for debugging info.

Apparently this patch triggers compiler crashes on some of our code. I'll try to create a reproducer, but it would be great to revert the patch for now.

Does this naming scheme the same as used for .o files? We may want to keep them in sync.

In D87325#2277717, @ashi1 wrote:

Having -emit-pch in the clang driver is useful because it doesn't require users to specify standard C++ include paths, clang include paths, and CUDA/HIP wrapper headers needed by CC1. That is error prone for the user.

This is a very nice cleanup. Thank you, Sam.

In D87325#2277316, @ashi1 wrote:

In D87325#2271676, @tra wrote:

Can you elaborate on the use case of PCH files for CUDA/HIP?

I believe one use-case for PCH is for common include headers such as hip_runtime.h which is being re-used in many application source files. To improve the performance, we can pre-compile the header and re-use it during online compilation.

In D84362#2275029, @yaxunl wrote:

In D84362#2274790, @yaxunl wrote:

There are use patterns expecting PartialDiagnosticInst << X << Y to continue to be a PartialDiagnostic&, e.g.

PartialDiagnosticAt PDAt(SourceLoc, PartialDiagnosticInst << X << Y);

However if we derive PartialDiagnostic and DiagnosticBuilder from a base class DiagnosticBuilderBase which implements the << operators, PartialDiagnosticInst << X << Y will become a DiagnosticBuilderBase&, then we can no longer write the above code.

That's one reason I use templates to implement << operators.

Do we want to sacrifice this convenience?

Can you elaborate on the use case of PCH files for CUDA/HIP?

So, the idea here is to do some sort of duck-typing and allow DiagBuilder to work with both DiagnosticBuilder and PartialDiagnostic.

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.

In D86376#2239391, @yaxunl wrote:

For example, in HIP program, there is a kernel void foo(int*). If a C++ program wants to launch it, the desirable way is

void foo(int*);
hipLaunchKernel(foo, grids, blocks, args, shmem, stream);

Due to the prefixed kernel stub name, currently the users have to use

void __device_stub_foo(int*);
hipLaunchKernel(__device_stub_foo, grids, blocks, args, shmem, stream);

In D86376#2236501, @yaxunl wrote:

My previous measurements did not warming up, which caused some one time overhead due to device initialization and loading of device binary. With warm up, the call of __hipPushCallConfigure/__hipPopCallConfigure takes about 19 us. Based on the trace from rocprofile, the time spent inside these functions can be ignored. Most of the time is spent making the calls. These functions stay in a shared library, which may be the reason why they take such long time. Making them always_inline may get rid of the overhead, however, that would require exposing internal data structures.

In D86376#2234719, @yaxunl wrote:

This patch appears to be somewhere in the gray area to me. My prior experience with CUDA suggests that it will make little to no difference. On the other hand, AMD GPUs may be different enough to prove me wrong. Without specific evidence, I still can't tell what's the case here.

Sorry, the overhead due to __hipPushConfigure/__hipPopConfigure is about 60 us. The typical kernel launching latency is about 500us, therefore the improvement is around 10%.

I'm OK with how the patch is implemented.
I'm still on the fence regarding whether it should be implemented.

How much does this inlining buy you in practice? I.e. what's a typical launch latency before/after the patch? For CUDA, config push/pop is negligible compared to the cost of actually launching the kernel on the GPU. It is measurable if the launch is asynchronous, but queueing kernels fast, does not help all that much in the long run -- you eventually have to run those kernels on the GPU, so in most cases you're just spend a bit more time idling while waiting for the queued kernels to finish. To be beneficial, you'll need a finely balanced CPU/GPU workload and that's rather hard to achieve. Not to the point where the minor savings here would be meaningful. I would assume the situation on AMD GPUs is not that different.

Separated blocks of RUN commands with an empty line to make navigation easier.

Couple of minor nits. LGTM otherwise.

Looks good in general. Mostly C++ style comments below.

The test is no longer sticking out: https://gist.github.com/Artem-B/d0b05c2e98a49158c02de23f7f4f0279

In D84068#2204713, @arsenm wrote:

If we ship them with clang, who/where/how builds them?
If they come from ROCm packages, how would those packages add stuff into *clang* install directory? Resource dir is a rather awkward location if contents may be expected to change routinely.

Symlinks. I've been building the device libraries as part of LLVM_EXTERNAL_PROJECTS, and think this should be the preferred way to build and package the libraries. This is how compiler-rt is packaged on linux distributions. The compiler-rt binaries are a separate package symlinked into the resource directory locations. I'm not sure what you mean exactly by change routinely, the libraries should be an implementation detail invisible to users, not something they should be directly relying on. Only clang actually knows how to use them correctly and every other user is buggy

What if I have multiple ROCm versions installed? Which one should provide the bitcode in the resource dir?

These should be treated as an integral part of clang, and not something to mix and match. Each rocm version should have its own copy of the device libraries. It only happens to work most of the time if you mismatch these, and this isn't a guaranteed property.

The fix is here https://reviews.llvm.org/D85686

We can restrict externalization to constant variables with explicit 'constant' attributes only, which should fix this issue.

Sam, just a FYI that the patch has a couple of unintended consequences.

In D85551#2203711, @jdoerfert wrote:

Wow, cool. I imagine it was hard to split this given the manual check lines. We really need to start using the upgrade scripts here.

I'm fine with this, @ABataev WDYT?

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.

In D85276#2199655, @yaxunl wrote:

Do we need to disable pgo and coverage mapping for device compilation? Or it is already disabled?

Also fixed the same bug in nearbyint().

LGTM for CUDA.

What is expected to happen to device statics in anonymous name spaces? It may be worth adding them to the tests.

LGTM, modulo couple of nits.

Also, do we need the header at all?
It would be much easier to just get clang itself to add normalized macros without trying to reconstruct them from the existing macros.

I think Sam's approach is reasonable.

I'm not sure it's particularly useful, to be honest. CUDA code still needs to be compatible with NVCC so it can't be used in portable code like TF or other currently used CUDA libraries.
It could be useful internally, though, so I'm fine with it for that purpose.

It's a good point. Perhaps this is one of the cases where we should *not* follow nvcc.
We can't have our cake (preserve static behavior) and eat it (treat it as non-static in case something on the host side may decide to use an API which uses symbol names). Something's got to give. While we could make it work in some cases, I don't think we can make it work consistently.
I think it would be reasonable to restrict APIs that access symbols by name to be applicable to visible symbols only.

Updated directory structure.

Phabricator seems to be confused by the renames. The change just moved terraform -> terraform/buildbot-mlir-nvidia
The directory structure looks like this now:

Really updated directory structure.

Updated directory structure.

In D84258#2172337, @kuhnel wrote:

I would not run multiple containers on one VM. As you said, k8s cannot share one GPUs across containers. I would rather create one "build slave" per VM (or a group of "build slaves" each in a separate VM) in buildbot and then have that VM(s) execute a set of "builders". We could have an m:n mapping of "build slaves" and "builders".

My mlir-nvidia builder is not very picky. It would probably run on any of your machines as long as it has an Nvidia card. Sorry about the non-inclusive wording here, but that's what buildbot calls them in the UI.

In D80858#2170781, @hliao wrote:

The problem is not whether we have solution to tell them but when we need to add that. Not all static device variables need to be visible to the host side. Externalizing them adds the overhead for the linker and may pose additional restrictions on aggressive optimizations. Do we have to support every ambiguous usage in the burden of the compiler change?

In D84364#2170244, @tra wrote:

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

In D80858#2170604, @hliao wrote:

In D80858#2170547, @tra wrote:

Would it work if we generate a globally unique visible aliases for the static vars and use the alias' name to register device-side entities without changing their visibility?

We still need to define how a static device variable should be visible on the host side.

Would it work if we generate a globally unique visible aliases for the static vars and use the alias' name to register device-side entities without changing their visibility?

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

In D83503#2168092, @gkistanova wrote:

Ok. Let’s see how this will work in the reality.

Updated status.py with the new builder names.

In D84256#2168058, @mehdi_amini wrote:

I don't have this context so I was just seeing this revision alone which setup google/mlir/docker/ and still has google/docker.
It makes sense if after your stack of revisions if there is no google/docker left.