[PATCH] D17103: [CUDA] Don't crash when trying to printf a non-scalar object.

[Justin Lebar via cfe-commits]

jlebar added a comment.

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


http://reviews.llvm.org/D17103