LGTM.

LGTM.

In D58514#1435431, @wuhao5 wrote:

In D58514#1435296, @rjmccall wrote:

In D58514#1435228, @wuhao5 wrote:

Okay, so really just a block self-reference. We could really just add a feature for that that would avoid both the complexity and the expense of the self-capture dance.

Is there a plan to cover this case? or is it a legitimate use case that Clang should handle?

You are currently relying on something that ARC doesn't guarantee, so the client code should be fixed to explicitly copy the block. I think we would be happy to consider a proposal in the long run to allow blocks to self-reference more easily, which will effectively bypass the problem.

I am not sure if I follow here - is it not that the weak pointer holds a block that's in the stack but is supposed to be in the heap?

In D58514#1435228, @wuhao5 wrote:

Okay, so really just a block self-reference. We could really just add a feature for that that would avoid both the complexity and the expense of the self-capture dance.

Is there a plan to cover this case? or is it a legitimate use case that Clang should handle?

LGTM, thanks!

This patch LGTM.

Sorry, I've gotten trapped under a backlog and haven't had time to think much about this. I don't want to block progress here.

In D58514#1429758, @wuhao5 wrote:

Can I ask why you want a weak reference to a block in the first place? It seems basically useless — blocks can certainly appear in reference cycles, but I don't know why you'd ever try to break that cycle with the block instead of somewhere else.

The simplified version:

auto b = ^{

if (check) {
  dispatch_after(queue, 1, b);
} else {
 // done.
}

};
dispatch_after(queue, 1, b);

In D58514#1429652, @erik.pilkington wrote:

In D58514#1429606, @rjmccall wrote:

There is no way in the existing ABI for copying a block to cause other references to the block to become references to the heap block. We do do that for __block variables, but not for block objects themselves.

Do you think thats worth doing? We could add a forwarding pointer to the end of a block literal on the stack (after all the captures) and flip an unused bit to track it, preserving ABI. Seems like now that we're delaying _Block_copyies this might be a bigger issue.

In D58514#1429662, @wuhao5 wrote:

The specified user model of blocks is that they stay on the stack until they get copied, and there can be multiple such copies. ARC just automates that process. So the address of a block is not consistent before you've forced a copy.

I tend to agree that a better user model would have been for blocks to be allocated in one place, without any of this copying business, and for the compiler to make an intelligent decision about stack vs. heap based on how the block is used. That's the model we've used for closures in Swift. But that would've required the compiler to have a better ability to propagate information about how the block was used, which Clang isn't really set up for, and it would've required noescape annotations to be introduced and used reliably throughout the SDK, which seemed like a big request at the time. So it's not how it works.

There is no way in the existing ABI for copying a block to cause other references to the block to become references to the heap block. We do do that for __block variables, but not for block objects themselves.

I see - this makes sense. Right that I'd expect the compiler would know more about where the block is being used and make the variable consistent. my other worry is, although not realistically, that there can be other projects to use this weak/strong pointer trick to do a recursive block invocation. it becomes to me a bit counter-intuitive that I will need to know more about how block and where it should be copied, which currently we don't have to worry about it at all.

Right now we force an explicit copy before using it, but still like to request that this would be handled by Clang at some later point :)

In D58514#1429514, @wuhao5 wrote:

Hey guys, this is Hao, I am working with Chrome team to sort this issue out.

In D58514#1428606, @rjmccall wrote:

I remember this coming up 7-8 years ago. I intentionally decided against doing a copy/release when assigning to __weak because if the block wasn't already guaranteed to be copied then it was probably better to crash than to silently assign a value that's about to be deallocated. Note that copying the block we assign into wb doesn't change anything about the block stored in b.

I don't know why Chromium is building a weak reference to a block in the first place, but assuming they have a good reason to be doing it, they should fix their code to force a copy before forming a weak reference.

The culprit code is here

it's true that it can be copied before assigning to the weak var, and I understand the reasoning behind. however, my question is: just from the code itself, each variable has the proper scope and assignment, if the block copy happen automatically, just like what we should expect ARC would do, should it not mutate itself to something else. to be more precise, should the block assigned to the weak var be the same after the block is copied? (and in the code, the block should be moved to the heap after calling -addObject: a few line below.)

so in the end of day, as a user, should we expect the compiler would move the block from stack to heap in time and the variable we hold is consistent?

I remember this coming up 7-8 years ago. I intentionally decided against doing a copy/release when assigning to __weak because if the block wasn't already guaranteed to be copied then it was probably better to crash than to silently assign a value that's about to be deallocated. Note that copying the block we assign into wb doesn't change anything about the block stored in b.

Using internal linkage instead of private makes sense to me. Even if we wanted to use private linkage, it never made sense to be doing this with \01l instead of just setting the linkage properly. Maybe this was implemented before LLVM directly supported private linkage, but that hasn't been a constraint in 10+ years.

Looks great, thanks!

Thanks, LGTM.

Thanks, LGTM.

This test file will probably grow over time, so please add a CHECK-LABEL line to the test case to make sure you're checking the body of foo(). Also, you'll probably need to allow the name %ref.tmp to be stripped for this test to pass in release builds; please go ahead and test a more complete pattern, including the alloca and the store of 1 to it. Otherwise LGTM.

In the abstract, it would be nice to warn about casts that always have UB because the address spaces don't overlap; however, I'm worried about how people might be using address spaces in strange ways today, knowing that they *do* overlap, only in ways that the compiler isn't smart enough for. I think we should just be permissive in non-OpenCL mode.

This seems eminently reasonable to warn about. I think people might legitimately complain about extending an existing warning to a new set of contexts without putting it under a new warning flag, though. How annoying would it be to clone the warnings and put them in a subgroup? You could also change the diagnostics to say "converting the result of an assignment to bool" instead of "using the result of an assignment as a condition".

Hmm. Yeah, Embedded C allows these casts, so contra my previous comment, I think we can't make them ill-formed outside of OpenCL mode.

I think the right model is that a builtin operator candidate exists for all possible address-space qualifiers on the pointer/reference pointee, and it should be straightforward to only actually add candidates for the qualifier used by the corresponding operand.

Well, it won't have a guarantee that it won't see unused annotations, but alright, I'm fine with this.

In D56411#1410598, @yaxunl wrote:

In D56411#1410153, @rjmccall wrote:

In D56411#1406212, @yaxunl wrote:

I would like to fix the validation issue only and leave the overload resolution issue for future.

As I understand it, the "validation issue" is just that you'd like a diagnostic to be emitted when resolving the template argument in order to force SFINAE to pick a different template. I think that's actually just the overload-resolution issue.

Currently there are two host-ness related issues about function type template arguments:

1. when there are two or more candidates for the template argument, clang goes through host-ness based overloading resolution, which does not work properly
2. when there is only one candidate for the template argument, clang does not go through overloading resolution, therefore the first issue does not show up. However, clang still checks host-ness of template argument. As discussed before, clang should not check host-ness in non-evaluation or constant-evaluation context. Instead, clang should check host-ness in template instantiation.
   
   I refer the first issue as host-ness overloading resolution issue and the second issue as host-ness validation issue. They are related but separate.
   
   The first issue only happens when host-ness based overloading resolution is used. For applications which can be compiled with nvcc, this cannot happen, therefore it is less common and less urgent.
   
   The second issue can happen to applications which can be compiled with nvcc, therefore is more imminent.
   
   Fixing the second issue is relatively straightforward. It does not need to introduce new AST types for host-ness. Also it is orthogonal to fixing the first issue.

Thanks, LGTM.

A couple tweaks, but otherwise LGTM.

In D56411#1406212, @yaxunl wrote:

I would like to fix the validation issue only and leave the overload resolution issue for future.

Okay, one last minor request, then LGTM.

One comment, but otherwise LGTM.

I see, alright.

The correctness condition here is solely that we cannot allow the block literal to go out of scope before the variable that it is assigned to. (Local block literals with captures have lifetimes like C compound literals: until the end of the enclosing block, rather than the end of the full-expression.) So doing this blindly for assignments is problematic because you'd need to reason about relative lifetimes, but doing it for initialization should always be fine. I don't see any reason why address-taken-ness would matter as long as that condition holds.

I agree. There should just never be a copy expression if the capture is of reference type, and there should therefore be no need to special-case that in IRGen.

No, this LGTM.

Okay, but it's not great design to have a kind of overloading that can't be resolved to an exact intended declaration even by an explicit cast. That's why I think making *optional* host/device typing is a good idea. And I strongly want to caution you against doing language design by just incrementally hacking at the compiler to progressively make more test-cases work, which is what it feels like you're doing.

It is totally unreasonable, at the time you are resolving a template argument, to investigate how the corresponding template parameter is used within the template and use that to shape how the template argument is resolved. That is simply not how the C++ template model works. Given that CODA doesn't distinguish between host and device functions in the type system, if you are going to have a rule here, it has to be based on, at most, (1) the current semantic context (which may not even be a function), (2) the template being specialized, and (3) the declarations in the template-argument set.

LGTM, thanks!

LGTM.

But what we've just been talking about is not a validity rule, it's an overload-resolution rule. It's not *invalid* to use a device function as a template argument to a host function template (or to a class template, which of course is neither host nor device). All you need to do is to resolve otherwise-intractable overload ambiguities by matching with the host-ness of the current context, which there's probably already code to do for when an overload set is used as e.g. a function argument.

In D56411#1398328, @rjmccall wrote:

In D56411#1398291, @tra wrote:

That said, does CUDA have a general rule resolving __host__ vs. __device__ overloads based on context? And does it allow overloading based solely on __host__ vs. __device__?

NVCC does not. Clang does. See https://goo.gl/EXnymm for the details.

AFAICT, NVIDIA is starting to consider adopting Clang's approach:
http://lists.llvm.org/pipermail/cfe-dev/2018-November/060070.html (original message from Bryce apparently didn't make it to the cfe-dev archive)

Okay. Probably the template-argument rule ought to be the same as the address-of-function rule, which I assume means that there's a final pass that resolves ambiguities in favor of functions that can be used from the current context, to the extent that that's meaningful. It's hard to tell because that document does not appear to include a formal specification.

In D56411#1398291, @tra wrote:

That said, does CUDA have a general rule resolving __host__ vs. __device__ overloads based on context? And does it allow overloading based solely on __host__ vs. __device__?

NVCC does not. Clang does. See https://goo.gl/EXnymm for the details.

AFAICT, NVIDIA is starting to consider adopting Clang's approach:
http://lists.llvm.org/pipermail/cfe-dev/2018-November/060070.html (original message from Bryce apparently didn't make it to the cfe-dev archive)

In D58218#1398124, @jfb wrote:

In D58218#1398096, @rjmccall wrote:

Well, you can always make a variable with a more directly-applicable name than capturedByInit and update it as appropriate, like locIsByrefHeader.

Sounds good. I made it const too, to avoid inadvertent modification.

In D56411#1398097, @yaxunl wrote:

In D56411#1365878, @yaxunl wrote:

In D56411#1365745, @rjmccall wrote:

In D56411#1365727, @yaxunl wrote:

In D56411#1360010, @rjmccall wrote:

I think the diagnostic should come during instantiation when you find an evaluated use of a host function within a device function.

It seems the body of function template is checked only during parsing of the definition of the template itself. When a function
template is instantiated, the body of the instantiated function is not checked again.

No, that's not correct. However, it's checked somewhat differently, and it's possible that the existing diagnostic is not set up to fire along all common paths. Try moving the diagnostic to MarkFunctionReferenced, and note that OdrUse will be false in all the unevaluated contexts.

You are right. After I disable current diagnostic, I saw PerformPendingInstantiations at the end of parsing the TU, where the AST of the instantiated function is iterated and MarkFunctionReferenced is called. I will try to fix my patch as suggested. Thanks.

I got one concern. If we want to do overload resolution of function type template argument based on host or device, we need to do that before template instantiation, right?

e.g. we have two functions having the same name f and type, but one is __host__ and the other is __device__, and we pass it as a template argument to a template function g. We want to choose __device__ f if g itself is __device__ and __host__ f if g itself is __host__. If we want to do this we have to do the check before template instantiation, right?

Well, you can always make a variable with a more directly-applicable name than capturedByInit and update it as appropriate, like locIsByrefHeader.

Okay. Is it acceptable for the annotation to simply disappear in this case? I don't really understand the purposes of annotations well enough to judge.

Basically LGTM, especially if we need an emergency fix, but please consider addressing my comment before committing since I'd expect it to be straightforward to solve.

Sorry, this is deep in the loop analysis code, and I don't remember anything about this stuff, so I'm not sure I'm going to be a useful reviewer.

I really don't understand this well enough to review.

Good catch! Is there a reasonable way to just make FileCheck itself enforce this, or does that cause too many false positives with tests that are using e.g. x86_64 as the entire prefix?

Ah, thanks, seems like a reasonable fix.

LGTM, thanks!

Hmm. Richard, I've mostly let you stay out of this, but do you have any thoughts about the right manage to parse attributes here? We want to allow address_space attributes to be written in the method-qualifiers list, but when we do that, it's hard to avoid parsing arbitrary attributes and then potentially needing to move them to the declarator; also, doing that might mean propagating awkward information in to deal with of delayed attribute parsing.

SGTM, thanks!

LGTM, but why did I just get like a million different notifications of this from different services.

LGTM as well.

Also, why is this HIP-specific? I thought the toolchain was largely shared with CUDA and there were just a few runtime differences.

Two minor tweaks, but otherwise LGTM.

In D57768#1386862, @Anastasia wrote:

• SYCL seem to require adding tight dependencies from the standard libraries into the compiler because many language features are hidden behind library classes. This is not common for Clang. We had a discussion about this issue during the implementation of OpenCL C++ and it was decided not to go this route for upstream Clang. Can you explain your current approach to implement this? I think @rjmccall or @rsmith might need to be involved in this.

Moving parsed attributes between lists isn't unreasonable if that's what you have to do; we already do that when processing the ObjC ARC qualifiers. The ambiguity with function attributes is pretty much inherent.

Richard felt that this was an odd special case, and I was happy to use the old language-designer's dodge of banning something today so that we can decide to allow it tomorrow. This isn't SFINAE-able.

Okay.

Asking for a minor tweak, but feel free to commit with that.

I agree: if a test doesn't seem to require the GNU extensions, let's not use -std=gnu++XX. Otherwise this LGTM.