[PATCH] D78655: [CUDA][HIP] Let non-caputuring lambda be host device

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked 2 inline comments as done.
yaxunl added a comment.

In D78655#2019290 <https://reviews.llvm.org/D78655#2019290>, @rsmith wrote:

> There are two behaviors that seem to make sense:
>
> - Treat lambdas as implicitly HD (like constexpr functions) in all CUDA / HIP language modes. I don't think it makes sense for lambdas to become implicitly HD in C++17 simply because they become implicitly `constexpr`, nor for their HDness to depend on whether their parameter types happen to be literal types, etc. So in C++17, where lambdas are constexpr whenever they can be, the logical behavior would seem to be that lambdas are implicitly HD. And then for consistency with that, I'd expect them to be implicitly HD across all language modes.
> - Implicitly give lambdas the same HD-ness as the enclosing function (if there is one).
>
>   I would think the best choice may be to do both of these things: if there is an enclosing function, inherit its host or device attributes. And if not, then treat the lambda as implicitly HD. A slight variation on that, that might be better: lambdas with no lambda-capture are implicitly HD; lambdas with any lambda-capture (which must therefore have an enclosing function) inherit the enclosing function's HDness.
>
>   (Note that if we go this way, it makes no difference if there are reference captures, because they're always references on the same "side".)


Sorry for the delay.  I updated the patch as you suggested:

- lambdas without enclosing function are implicitly HD
- lambdas with no lambda-capture are implicitly HD
- lambdas with any lambda-capture (which must therefore have an enclosing function) inherit the enclosing function's HDness.





================
Comment at: clang/test/SemaCUDA/lambda.cu:25-35
+  kernel<<<1,1>>>([&](){ hd(); });
+  // expected-note at -1 {{in instantiation of function template specialization 'kernel<(lambda at}}
+  // expected-note at -2 {{candidate function not viable: call to __host__ function from __global__ function}}
+
+  kernel<<<1,1>>>([=, &b](){ hd(); });
+  // expected-note at -1 {{in instantiation of function template specialization 'kernel<(lambda at}}
+  // expected-note at -2 {{candidate function not viable: call to __host__ function from __global__ function}}
----------------
tra wrote:
> We may need a better diagnostic for this. Here we've correctly rejected captured lambdas, but the diagnostic is a generic 'can't use that'.
> If would be helpful to let user know that we can't use that because of the capturing lambdas.
added more information about lambda function to the diagnostic message


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

https://reviews.llvm.org/D78655