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

[Justin Lebar via Phabricator via cfe-commits]

jlebar added a comment.

  __host__ void bar() {}
  __device__ int bar() { return 0; }
  __host__ __device__ void foo() { int x = bar(); }
  template <void (*devF)()> __global__ void kernel() { devF();}
  
  kernel<foo>();



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

`__host__ bar()` should not need to return int if `foo` is inline (or templated), because then we should never codegen `foo` for host.  I guess my question is, we should be sure that `kernel<foo>()` does not force an inline/templated `foo` to be codegen'ed for host.  (Sorry that wasn't more clear before.)

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

Yes, that's a fair way to think about it.  It just is a bit weird that in this context `&n` refers to one function but `n()` refers to another.  Maybe that's unavoidable.  :shrug:

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

Agreed.

> Therefore I tend to suggest we keep things as they are, i.e., bar is host/device resolved in its containing function context.

I'm not sure what is the containing function context in these examples, since all of the definitions don't have a containing function.

Currently `baz()` returns void, but it sort of seems to me like the decltype should morally be executed within a `__device__` context?

Anyway I know much of this is a distraction from your patch.  So long as we have `__host__ __device__` tests I'm happy here.


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

https://reviews.llvm.org/D56411