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