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

[Justin Lebar via Phabricator via cfe-commits]

jlebar added a comment.

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.

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)?

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:

  __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.  :)


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

https://reviews.llvm.org/D56411