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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits