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

[Michael Liao via Phabricator via cfe-commits]

hliao added a comment.

In D78655#2003837 <https://reviews.llvm.org/D78655#2003837>, @yaxunl wrote:

> 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.


Not only the capture is an issue, like a regular function, lambda could also access non-local variables/functions. Without marking proper HD attributes explicitly, it's difficult or expensive to statically check the use of them. If we have to guess that attributes based on how they are called, we may find cases where conflicting results may be derived depending on call sites. It would be quite confusing if developers struggle to solve one but trigger another one or more. From the other perspective, a lambda is just another function and should have consistent rule for its usage, resolution, and etc.


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

https://reviews.llvm.org/D78655