I spoke offline with tra@ about this; he was confused why I was coming down so hard here.

So if we can drop cuda 8, great. If some of the maintainers have a customer that requires cuda 8 to continue working with a recent llvm - which is sad but plausible - we're better off supporting that use case in tree than forcing an out of tree fork.

No, I don't try to add a workaround for the bug in LLVM since there are no problems in LLVM optimizations. The same code works correctly with Cuda9.2. But I would like to add it for cuda8 since we need to fully support it unless we drop the support for cuda8.

Hi, jumping in here. I haven't read the whole bug, apologies.

At [nvcc] from CUDA 10, that's not acceptable as we are declaring two functions only differ from the return type. It seems CUDA attributes do not contribute to the function signature. clang is quite different here.

Here's one for you:

Does XLA:GPU built with this change pass the exhaustive test you added for kRound after we revert the change that made kRound call into libdevice?

I uh... I also think this is an @rsmith question, I have no idea.

Looks great, one nit and one unfortunately substantial comment, sorry I didn't think about it earlier.

Thank you for the patch!

By default it is still treated as error, therefore no behavior change of clang.

I agree with Art. The fact that nvcc allows this is broken.

LGTM modulo jfb's outstanding comment.

That is much less scary! Thanks.

Actually, I wonder if we should have some protections built in here, because e.g. if filename contains "/../../" we could end up rm'ing an arbitrary directory on the machine.

Thanks!

Still LGTM.

I'm not totally comfortable with this code, but this change looks reasonable to me.

Strangely this didn't show up in my inbox until you pinged it. Not sure why. :-/

This seems reasonable to me but I'm not familiar enough with the x86 backend to feel comfortable signing off on it.

My main comment is, in most or all of these places where we s/isVolatile/isVolatile && isAtomic/, it's not going to be obvious to a reader whether isAtomic is really necessary, or just part of a conservative check.

LGTM, mostly nits.

Well this is heroic.

What's the status of this?

Without reading the patch in detail (sorry) but looking mainly at the testcase: It looks like we're not checking how overloading and __host__ __device__ functions play into this. Maybe there are some additional edge-cases to explore/check.

Thanks a lot for proofreading my totally untested code. :)

Address fedor.sergeev's comments.

Sorry for missing tra's ping earlier, I get a lot of HIP email traffic that's 99% inactionable by me, so I didn't notice my username in tra's earlier email.

+tra in the hopes that perhaps he's comfortable reviewing this (sorry that I'm not).

I defer to the others here.

Just to check, the notion is that it's OK if I report a sm version less than what I end up running on?

I'm not confident that the GUIDS remain stable, either. I have no data to tell one way or another, though.

You don't want to pass in the GPU-deadbeef form from nvidia-smi -L so we don't have this problem?

Looks right to me (other than the missing constexpr in C++14 onwards). Though this is subtle enough that I suspect the only way to know for sure is to try it.

In D47757#1142886, @ahatanak wrote:

I mean ToT clang (without my patch applied) seems to select the non-sized host version 'T::operator delete(void*)'.

@jlebar, is the change I made to call-host-fn-from-device.cu correct?

@rsmith friendly ping on this one.

In D48151#1133954, @rsmith wrote:

LGTM

@rsmith friendly ping on this one.

@rsmith friendly ping on this review.