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