llvm dev - Mar 2021 - [RFC] The `implements` attribute, or how to swap functions statically but late

The basic idea is to provide a way to replace a callee statically 
*inside* the LLVM-core
pipeline [0]. The initial use case was the missing translation of 
certain (math) intrinsics
by the NVPTX backend, however, we believe there are other use cases that 
would benefit from
a generic model, one is described below as well as the benefits over 
alternatives.

---

Use cases

1a) NVPTX + fast-math + math functions

The original use case was discussed here [1] and the problem is shown 
here [2].
In a nutshell, `llvm.sin` (and friends) are not supported by the backend 
and it
will yell at you. This is unfortunate but could arguably be resolved in 
a lot of
different ways (below). Note that there is not necessarily a 1:1 mapping 
for all
intrinsics and we might need to write a wrapper which would probably be in
clang/lib/Header with other CUDA wrapper code.

Considered solutions we did not pick:
a) Teach the backend explicitly about the mapping `llvm.sin` -> `__nv_sin`.
b) Teach clang not to emit the intrinsics in the first place.
c) Use an existing overload "hack", e.g., __asm__.

Let's discuss those first.
a) Works, is very limited to this use case, the mapping is not with the 
declarations
    and definitions which makes the maintenance less appealing.
b) Seems easy enough, we loose all the benefits of intrinsics though, 
more on that
    later.
c) Right now, IR isn't happy with defining intrinsics so there is little 
we could do
    better than b) here I think, short of allowing intrinsic definitions 
(which is an
    option).

1b) User can overwrite the implementation of any intrinsic, regardless 
if the backend
     supports it or not. They still get the intrinsic benefits during 
the compilation.
     This is not limited to NVPTX and math intrinsics but also opens the 
door for cross
     platform tooling, e.g., Tool XYZ will emit intrinsics our backed 
doesn't support
     but we can map them to implementations for our architecture.

1c) You can prototype/test different lowerings for functions and 
intrinsics easily by
     linking in a bitcode file (which can be compiled form C*).

2) Function specifications and implementations

I would like to embed a function specification with the implementation, e.g.
```
void vec_add(double *A, double *B, int N) {
   for (int i = 0; i < N; ++i)
     A[i] += B[i];
}
__attribute__((implements("vec_add")))
void vec_add_impl(double *A, double *B, int N) {
   #pragma omp parallel for
for (int i = 0; i < N; ++i)
     A[i] += B[i];
}
```

I think that would open up cool possibilities:
A) Replace only the first N specifications with implementations, allows 
to improve debugging.
B) Use the specification for the first part of the optimization pipeline 
so you
    can derive facts the implementation "hides". E.g., `llvm.sin`
calls
could be
    annotated as `readnone` while the `__nv_sin` call contains inline 
assembly and
    acts as a optimization barrier. Or, in the above example, it is 
obvious from the
    "specification" that `B` is readonly but the implementation could 
arbitrarily hide this,
    e.g., the implementation might be a runtime call.
C) Program verification, test synthesis, etc. would be possible on a new 
level in a defined way.
    By keeping implementation and specification close together and 
written in the same
    language we might get somewhere in this area.

---

Thoughts

I initially thought about the reverse attribute, so `vec_add` in the 
above example would have
`__attribute__((specifies(vec_add_impl))` for the same reasons (A-C). 
The problem is that I'm
not sure how to attach anything to an intrinsic. Some `__asm__` trickery 
might help but it's
unclear that this is better than the `implements` version. Maybe we want 
both for convenience
with the `implements` version being able to handle intrinsics.

~ Johannes


[0] https://reviews.llvm.org/D98516
[1]https://lists.llvm.org/pipermail/llvm-dev/2021-March/149117.html
[2] https://godbolt.org/z/PxsEWs


-- 
───────────────────
∽ Johannes (he/his)

Super in favor of this, both for ease of GPU codegen, potential BLAS
optimizations, as well as upstream users like Julia potentially being able
to take advantage of LLVM's optimizations for sin/etc.

Given the potential for multiple implementations to exist, is it worthwhile
adding some notion of a priority or perhaps a suggested target?

On Fri, Mar 19, 2021 at 8:10 PM Johannes Doerfert <
johannesdoerfert at gmail.com> wrote:
> The basic idea is to provide a way to replace a callee statically
> *inside* the LLVM-core
> pipeline [0]. The initial use case was the missing translation of
> certain (math) intrinsics
> by the NVPTX backend, however, we believe there are other use cases that
> would benefit from
> a generic model, one is described below as well as the benefits over
> alternatives.
>
> ---
>
> Use cases
>
> 1a) NVPTX + fast-math + math functions
>
> The original use case was discussed here [1] and the problem is shown
> here [2].
> In a nutshell, `llvm.sin` (and friends) are not supported by the backend
> and it
> will yell at you. This is unfortunate but could arguably be resolved in
> a lot of
> different ways (below). Note that there is not necessarily a 1:1 mapping
> for all
> intrinsics and we might need to write a wrapper which would probably be in
> clang/lib/Header with other CUDA wrapper code.
>
> Considered solutions we did not pick:
> a) Teach the backend explicitly about the mapping `llvm.sin` ->
`__nv_sin`.
> b) Teach clang not to emit the intrinsics in the first place.
> c) Use an existing overload "hack", e.g., __asm__.
>
> Let's discuss those first.
> a) Works, is very limited to this use case, the mapping is not with the
> declarations
>     and definitions which makes the maintenance less appealing.
> b) Seems easy enough, we loose all the benefits of intrinsics though,
> more on that
>     later.
> c) Right now, IR isn't happy with defining intrinsics so there is
little
> we could do
>     better than b) here I think, short of allowing intrinsic definitions
> (which is an
>     option).
>
> 1b) User can overwrite the implementation of any intrinsic, regardless
> if the backend
>      supports it or not. They still get the intrinsic benefits during
> the compilation.
>      This is not limited to NVPTX and math intrinsics but also opens the
> door for cross
>      platform tooling, e.g., Tool XYZ will emit intrinsics our backed
> doesn't support
>      but we can map them to implementations for our architecture.
>
> 1c) You can prototype/test different lowerings for functions and
> intrinsics easily by
>      linking in a bitcode file (which can be compiled form C*).
>
> 2) Function specifications and implementations
>
> I would like to embed a function specification with the implementation,
> e.g.
> ```
> void vec_add(double *A, double *B, int N) {
>    for (int i = 0; i < N; ++i)
>      A[i] += B[i];
> }
> __attribute__((implements("vec_add")))
> void vec_add_impl(double *A, double *B, int N) {
>    #pragma omp parallel for
> for (int i = 0; i < N; ++i)
>      A[i] += B[i];
> }
> ```
>
> I think that would open up cool possibilities:
> A) Replace only the first N specifications with implementations, allows
> to improve debugging.
> B) Use the specification for the first part of the optimization pipeline
> so you
>     can derive facts the implementation "hides". E.g., `llvm.sin`
calls
> could be
>     annotated as `readnone` while the `__nv_sin` call contains inline
> assembly and
>     acts as a optimization barrier. Or, in the above example, it is
> obvious from the
>     "specification" that `B` is readonly but the implementation
could
> arbitrarily hide this,
>     e.g., the implementation might be a runtime call.
> C) Program verification, test synthesis, etc. would be possible on a new
> level in a defined way.
>     By keeping implementation and specification close together and
> written in the same
>     language we might get somewhere in this area.
>
> ---
>
> Thoughts
>
> I initially thought about the reverse attribute, so `vec_add` in the
> above example would have
> `__attribute__((specifies(vec_add_impl))` for the same reasons (A-C).
> The problem is that I'm
> not sure how to attach anything to an intrinsic. Some `__asm__` trickery
> might help but it's
> unclear that this is better than the `implements` version. Maybe we want
> both for convenience
> with the `implements` version being able to handle intrinsics.
>
> ~ Johannes
>
>
> [0] https://reviews.llvm.org/D98516
> [1]https://lists.llvm.org/pipermail/llvm-dev/2021-March/149117.html
> [2] https://godbolt.org/z/PxsEWs
>
>
> --
> ───────────────────
> ∽ Johannes (he/his)
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20210319/c12b4802/attachment.html>