yaxunl marked an inline comment as done.
yaxunl added a comment.

In D78655#2003681 <https://reviews.llvm.org/D78655#2003681>, @hliao wrote:

> I though the goal of adding HD/D attributes for lambda is to make the static 
> check easier as lambda used in device code or device lambda is sensitive to 
> captures. Invalid capture may render error accidentally without static check, 
> says we capture host variable reference in a device lambda. That makes the 
> final code invalid. Allowing regular lambda to be used in global or device 
> function is considering harmful.


Inferring a lambda function by default as `__host__ __device__` does not mean 
skipping the check for harmful captures.

If we add such checks, it does not matter whether the `__host__ __device__` 
attribute is explicit or implicit, they go through the same check.

How to infer the device/host-ness of a lambda function is a usability issue. It 
is orthogonal to the issue of missing diagnostics about captures.

Forcing users to explicitly mark a lambda function as `__device__ __host__` 
itself does not help diagnose the harmful captures if such diags do not exist.

Let's think about a lambda function which captures references to host 
variables. If it is only used in host code, as in ordinary C++ host code. 
Marking it `host device` implicitly does not change anything, since it is not 
emitted in device code. If it is used in device code, it will likely cause mem 
fault at run time since currently we do not diagnose it. Does it help if we 
force users to mark it `__device__ __host__`? I don't think so. Users will just 
reluctantly add `__device__ __host__` to it and they still end up as mem fault. 
If we add the diagnostic about the harmful captures, it does not matter whether 
the `__device__ __host__` attribute is explicit or implicit, users get the same 
diagnostic about harmful captures. So the effect is the same. However, 
usability is improved if users do not need to add `__device__ __host__` by 
themselves.

Does inferring lambda function as `__device__ __host__` by default making the 
diagnostic about harmful captures more difficult? No. It should be the same for 
lambdas with explicit `__device__ __host__`. This needs to be a deferred 
diagnostic like those we already did, which are only emitted if the function is 
really emitted. It does not matter whether the device/host attrs are implicit 
or explicit.


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

Reply via email to