hfinkel added a comment.

In http://reviews.llvm.org/D17103#349254, @jlebar wrote:

> In http://reviews.llvm.org/D17103#349245, @hfinkel wrote:
>
> > In http://reviews.llvm.org/D17103#349182, @jlebar wrote:
> >
> > > Yeah, I have no idea what's the right thing to do here.  We can always 
> > > pass a null pointer, that's easy.  David, Reid, do you know what is the 
> > > correct behavior?
> >
> >
> > I think we need to diagnose / reject this during semantic analysis (and 
> > then put a reasonable assert in the backend).
>
>
> Two things.
>
> a) That doesn't seem to be what we do in regular C++.  It will happily let 
> you pass a Struct in with only a warning.


Yes, but it also can be legally lowered and does not crash.

> b) At the moment, we don't have the capability to do a proper semantic 
> analysis of this.  The issue is, when doing sema checking of __host__ 
> __device__ functions, we don't know whether the function will end up being 
> codegen'ed for device.  And the semantics of cuda are that it's OK to do 
> things that are illegal in device mode from __host__ __device__ functions, so 
> long as you never codegen those functions for the device.

> 

> We have a plan to address (b) (basically, when doing sema checking, buffer 
> any errors we would emit if we were to codegen for device; then we can emit 
> all those errors right before codegen), but it's a much bigger thing.  Until 
> then, we need to do *something* other than crash here, even if we add 
> additional sema checking for plain __device__ fns.


Interesting dilemma. In the mean time, you can call CGM.ErrorUnsupported 
instead of removing arguments.


http://reviews.llvm.org/D17103



_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to