<div dir="ltr"><div dir="ltr"><div dir="ltr">> So, actually, I wonder if that's not the right answer. We generally allow different overloads to have different return types. What if, for example, the return type on the host is __float128 and on the device it's `MyLongFloatTy`?<div></div></div><div dir="ltr"><br></div><div>The problem is that conceptually compiling for host/device does not create a new set of overloads.</div><div><br></div><div>When we compile for (say) host, we build a full AST for all functions, including device functions, and that AST must pass sema checks.  This is significant for example because when compiling for device we need to know which kernel templates were instantiated on the host side, so we know which kernels to emit.</div><div><br></div><div>Here's a contrived example.</div><div><br></div><div>```</div><div><span style="color:rgb(80,0,80)"> __host__ int8 bar();</span><br style="color:rgb(80,0,80)"><span style="color:rgb(80,0,80)">__device__ int16 bar();</span><br style="color:rgb(80,0,80)"><span style="color:rgb(80,0,80)">__host__ __device__ auto foo() -> decltype(bar()) {}</span><br></div><div><span style="color:rgb(80,0,80)"><br></span></div><div><span style="color:rgb(80,0,80)">template <int N> </span><span style="color:rgb(80,0,80)">__global__ kernel();</span></div><div><span style="color:rgb(80,0,80)"><br></span></div><div><span style="color:rgb(80,0,80)">void launch_kernel() {</span></div><div><span style="color:rgb(80,0,80)">  kernel<sizeof(decltype(foo()))><<<...>>>();<br>}</span></div><div>```</div><div><br></div><div>This template instantiation had better be the same when compiling for host and device.</div><div><br></div><div>That's contrived, but consider this much simpler case:</div><div><br></div><div>```</div><div>void host_fn() {</div><div>  static_assert(sizeof(decltype(foo())) == sizeof(int8));</div><div>}</div><div>```</div><div><br></div><div>If we let foo return int16 in device mode, this static_assert will fail when compiling in *device* mode even though host_fn is never called on the device.  <a href="https://gcc.godbolt.org/z/gYq901">https://gcc.godbolt.org/z/gYq901</a></div><div><br></div><div>Why are we doing sema checks on the host code when compiling for device?  See contrived example above, we need quite a bit of info about the host code to infer those templates.</div></div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Thu, May 2, 2019 at 7:05 PM Hal Finkel via Phabricator <<a href="mailto:reviews@reviews.llvm.org">reviews@reviews.llvm.org</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">hfinkel added a comment.<br>
<br>
In D61458#1488970 <<a href="https://reviews.llvm.org/D61458#1488970" rel="noreferrer" target="_blank">https://reviews.llvm.org/D61458#1488970</a>>, @jlebar wrote:<br>
<br>
> Here's one for you:<br>
><br>
>   __host__ float bar();<br>
>   __device__ int bar();<br>
>   __host__ __device__ auto foo() -> decltype(bar()) {}<br>
><br>
><br>
> What is the return type of `foo`?  :)<br>
><br>
> I don't believe the right answer is, "float when compiling for host, int when compiling for device."<br>
<br>
<br>
So, actually, I wonder if that's not the right answer. We generally allow different overloads to have different return types. What if, for example, the return type on the host is __float128 and on the device it's `MyLongFloatTy`?<br>
<br>
> I'd be happy if we said this was an error, so long as it's well-defined what exactly we're disallowing.  But I bet @rsmith can come up with substantially more evil testcases than this.<br>
<br>
<br>
<br>
<br>
Repository:<br>
  rG LLVM Github Monorepo<br>
<br>
CHANGES SINCE LAST ACTION<br>
  <a href="https://reviews.llvm.org/D61458/new/" rel="noreferrer" target="_blank">https://reviews.llvm.org/D61458/new/</a><br>
<br>
<a href="https://reviews.llvm.org/D61458" rel="noreferrer" target="_blank">https://reviews.llvm.org/D61458</a><br>
<br>
<br>
<br>
</blockquote></div>