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

[Richard Smith - zygoloid via Phabricator via cfe-commits]

rsmith 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:
> 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.


Repository:
  rC Clang

https://reviews.llvm.org/D47757