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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
