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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked 10 inline comments as done.
yaxunl added inline comments.


================
Comment at: clang/lib/Sema/SemaCUDA.cpp:750
 
+bool Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
+                                  const sema::Capture &Capture) {
----------------
tra wrote:
> What does the return value mean? We don't seem to check it anyways.  If we don't care about the result, perhaps the function should be void.
> If we do, then it would be good to document its purpose and returned values and, probably, rename it to better indicate what is it it's supposed to check.
it should return void. fixed.


================
Comment at: clang/lib/Sema/SemaLambda.cpp:1783
+
+      if (LangOpts.CUDA && LangOpts.CUDAIsDevice)
+        CUDACheckLambdaCapture(CallOperator, From);
----------------
tra wrote:
> I would expect Sema-level diags to be produced during both host and device compilation. Some of the diags for HD may need to be postponed until after things have been CodeGen'ed, but the checks should happen nevertheless.
> 
> 
A host device function could be emitted in both host and device compilation. The way the deferred diags works is that they are only emitted when the function is sure to be emitted in a specific compilation. In host compilation, when a host device function is sure to be emitted, it is emitted as host function, therefore diags for host compilation and only diags for host compilation should be emitted. The same with device compilation.

This makes sense since we do not know if a host device function will be emitted in device compilation when we are doing host compilation, since to do that we have to go through the whole device compilation whereas currently device compilation and host compilation are separate process.

That said, when we emit diags for captures by reference, we should only emit them when the lambdas are emitted as device functions. When they are emitted as host functions in host compilation, these captures are valid and should not be diagnosed.


================
Comment at: clang/test/CodeGenCUDA/lambda.cu:53
+// DEV:  call void @_ZZ12test_resolvevENKUlvE_clEv
+// DEV-LABE: define internal void @_ZZ12test_resolvevENKUlvE_clEv
+// DEV:  call i32 @_Z10overloadedIiET_v
----------------
ashi1 wrote:
> There is a typo here, DEV-LABEL
fixed. thanks


================
Comment at: clang/test/SemaCUDA/lambda.cu:27
+
+    kernel<<<1,1>>>([&](){ hd(b); });
+    // dev-error at -1 {{capture host side class data member by this pointer in device or host device lambda function}}
----------------
pfultz2 wrote:
> Will this still produce diagnostics when the lambda is explicitly `__device__`? Maybe you could add a test case for that.
> 
> ```
> kernel<<<1,1>>>([&]() __device__ { hd(b); });
> ```
yes. added a test


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

https://reviews.llvm.org/D78655