[PATCH] D47757: [Sema] Produce diagnostics when unavailable aligned allocation/deallocation functions are called

[Artem Belevich via Phabricator via cfe-commits]

tra added inline comments.


================
Comment at: test/SemaCUDA/call-host-fn-from-device.cu:88
 __host__ __device__ void class_specific_delete(T *t, U *u) {
-  delete t; // ok, call sized device delete even though host has preferable non-sized version
+  delete t; // expected-error {{reference to __host__ function 'operator delete' in __host__ __device__ function}}
   delete u; // ok, call non-sized HD delete rather than sized D delete
----------------
rsmith wrote:
> rsmith wrote:
> > tra wrote:
> > > The C++ magic is way above my paygrade, but as far as CUDA goes this is a regression, compared to what nvcc does. This code in NVCC produced a warning and clang should not error out at this point in time either as 
> > > it's not an error to call a host function from HD unless we use HD in a host function, and we would not know how it's used until later. I think the error should be postponed until codegen. 
> > > 
> > > 
> > > 
> > > 
> > We're in `-fcuda-is-device` mode, so IIUC it's correct to reject a call to a host function here (because `__host__ __device__` is treated as basically meaning `__device__` in that mode for the purpose of checking whether a call is valid), right?
> > 
> > However, the comment suggests that the intent was that this would instead call the device version. Did that actually previously happen (in which case this patch is somehow affecting overload resolution and should be fixed), or is the comment prior to this patch wrong and we were silently calling a host function from a device function (in which case this patch is fine, but we should add a FIXME here to select the device delete function if we think that's appropriate)?
> OK, I see from prior review comments (that phab is helpfully hiding from view) that this is just adding a diagnostic and the overload resolution behavior is unchanged. So I think this change is correct. @tra, can you confirm? My testing shows that
> 
> ```
> __host__ void f(); __host__ __device__ void g() { f(); }
> ```
> 
> is accepted by default but rejected in `-fcuda-is-device` mode, which is consistent with the behavior after this patch is applied.
"-fcuda-is-device" does not necessarily mean that the __host__ __device__ function will be used.

In the example above the error is indeed correct, as g() is considered to be externally visible and we will attempt to generate code for it, and we can't call f() on device.

However, if you make it static, there should be no error:
```
__host__ void f(); 
static __host__ __device__ void g() { f(); }
```

CUDA is somewhat weird when it comes to what's considered available and what is not.
If you want some of the gory details, see D12453 and https://goo.gl/EXnymm

@jlebar has details on how we handle the errors in cases like that:
https://github.com/llvm-mirror/clang/blob/master/lib/Sema/SemaCUDA.cpp#L590



Repository:
  rC Clang

https://reviews.llvm.org/D47757