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

Yaxun Liu via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Wed Jan 9 07:59:40 PST 2019


yaxunl added a comment.

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

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


Sorry I am not quite get it. bar() is a `__host__` function with definition, so clang does codegen for it. clang also does codegen for foo() since it has `__host__ __device__` attribute.


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

https://reviews.llvm.org/D56411





More information about the cfe-commits mailing list