[PATCH] D78655: [HIP] Let lambda be host device by default

Artem Belevich via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Wed Apr 22 11:58:25 PDT 2020


tra added a reviewer: rsmith.
tra added a subscriber: rsmith.
tra added a comment.

Summoning @rsmith as I'm sure that there are interesting corner cases in lambda handling that we didn't consider.

Making lambdas implicitly HD will make it easier to write the code which can't be instantiated on one side of the compilation. That's probably observable via SFINAE, but I can't tell whether that matters.
By default I'd rather err on handling lambdas the same way as we do regular user-authored functions.



================
Comment at: clang/lib/Sema/SemaCUDA.cpp:722
+  if (getLangOpts().HIP) {
+    Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+    Method->addAttr(CUDAHostAttr::CreateImplicit(Context));
----------------
What about `__global__` lambdas? We probably don't want to add HD attributes on them here.


================
Comment at: clang/test/CodeGenCUDA/lambda.cu:16-26
+template<class F>
+__global__ void g(F f) { f(); }
+
+template<class F>
+void h(F f) { g<<<1,1>>>(f); }
+
+__device__ int a;
----------------
The test example may not be doing what it's seemingly supposed to be doing:
https://cppinsights.io/s/3a5c42ff

`h()` gets a temporary host-side object which keeps the reference to `a` and that reference will actually point to the host-side shadow of the actual device-side `a`. When you get to execute `g` it's `this` may not be very usable on device side and thus `f.operator()` will probably not work.

Alas, we currently have no diagnostics for that kind of error.

Change it to a non-capturing lambda, perhaps?


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

https://reviews.llvm.org/D78655





More information about the cfe-commits mailing list