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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
