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

[Yaxun Liu via Phabricator via cfe-commits]

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