[PATCH] D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl added a comment.

In D56411#1349364 <https://reviews.llvm.org/D56411#1349364>, @jlebar wrote:

> Without reading the patch in detail (sorry) but looking mainly at the testcase: It looks like we're not checking how overloading and `__host__ __device__` functions play into this.  Maybe there are some additional edge-cases to explore/check.


will add test for `__host__` `__device__`.

> Just some examples:
> 
> Will we DTRT and parse `bar` call as calling the `device` overload of `bar` in
> 
>   __host__ void bar() {}
>   __device__ int bar() { return 0; }
>   __host__ __device__ void foo() { int x = bar(); }
>   template <void (*devF)()> __global__ void kernel() { devF();}
>   
>   kernel<foo>();
> 
> 
> ?  Also will we know that we don't have to codegen `foo` for host (so `foo` is actually able to do things that only device functions can)?

we DTRT for this case. Here `__host__` bar needs to return int since foo() expects that. will add a test for that.

> Another one: How should the following template be instantiated?
> 
>   __host__ constexpr int n() { return 0; }
>   __device__ constexpr int n() { return 1; }
>   template <int> __global__ void kernel() {}
>   
>   kernel<n()>
> 
> 
> Presumably the call to `n` should be the host one?  That seems correct to me, but then it's pretty odd that a function pointer template argument would point to a *device* function.  Maybe that's the right thing, but I bet I can come up with something weird, like:

I think n() should be resolved in the containing function context. n itself is not template argument. the result of n() is.

> 
> 
>   __host__ void bar() {}
>   __device__ int bar() { return 0; }
>   __device__ auto baz() -> decltype(foo<n()>()) {} // which n() does it call?  Presumably host, but:
>   __device__ auto baz() -> decltype(bar()) {}  // does baz return void or int?  Presumably...the device one, int?
> 
> 
> Now mix in templates and sizeof and...yeah.  Rife for opportunities.  :)

I think this example is different from the issue which this patch tries to address. In the case of function type template parameter, it is less controversial about host/device resolution. The function argument is supposed to be called by the function template, therefore its host/device attribute should be consistent with the function template. Whereas in the above decltype example, such reqirement does not exist. Therefore I tend to suggest we keep things as they are, i.e., bar is host/device resolved in its containing function context.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D56411/new/

https://reviews.llvm.org/D56411