yaxunl marked 4 inline comments as done.
yaxunl added inline comments.
================
Comment at: clang/lib/Sema/SemaCUDA.cpp:722
+ if (getLangOpts().HIP) {
+ Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ Method->addAttr(CUDAHostAttr::CreateImplicit(Context));
----------------
tra wrote:
> What about `__global__` lambdas? We probably don't want to add HD attributes
> on them here.
lambda is not allowed to be kernel. I will add a lit test for that.
================
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;
----------------
tra wrote:
> 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?
It works.
We need to think about this in device compilation. In device compilation,
global variable is a device variable, the lambda is a device host function,
therefore the lambda is accessing the real a, not the shadow.
In the host compilation, the lambda is not really called, therefore it is not
emitted.
I will update the lit test with these checks.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D78655/new/
https://reviews.llvm.org/D78655
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
