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