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

Artem Belevich via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Mon Jun 22 12:54:40 PDT 2020


tra added a comment.

In D78655#2107190 <https://reviews.llvm.org/D78655#2107190>, @yaxunl wrote:

> It seems we can only promote non-capturing lambdas, no matter whether it has enclosing function or not.


I'd be OK with promoting only non-capturing lambdas until we figure out a consistent way to deal with the capturing ones.
Or we can promote captured ones, too and rely on postponed diags to guard against producing wrong-side code. We may need to improve that a bit anyways. E.g. what should we do if I write something like this:

  __device__ int dv;
  int hv;
  
  __host__ __device__ int hd() {
    return [a = dv, b=hv](){ return a + b;}();
  }

https://godbolt.org/z/op_FE6 -- NVCC complains about not being able to access hv in a device function (which makes sense considering that it converts HD -> D after source splitting, but clang happily allows capturing both variables (but will likely fail during ptxas due to the fact that there will be no hv on device side).


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

https://reviews.llvm.org/D78655





More information about the cfe-commits mailing list