[PATCH] D93189: Introduce the `!nocapture` metadata and "nocapture_use" operand bundle

[Johannes Doerfert via Phabricator via llvm-commits]

jdoerfert added a comment.

Thanks for the input. The details of the design are a bit in flux given the llvm-dev thread but I responded to the more conceptual parts below.

In D93189#2540135 <https://reviews.llvm.org/D93189#2540135>, @reames wrote:

> This isn't a comment on the proposed design so much as how the proposed design is described.  (I intend to respond to your llvm-dev thread... at some point.)
>
> I think you need to revise the specification wording into something more operational.  As currently described, you define metadata in terms of existing analyzes which aren't themselves well defined.  (e.g. What *precisely* does it mean for something to be captured?)  What operations are allowed and/or disallowed by the existence of the metadata/bundle?  For the ones which are disallowed, is that full UB?  Producing poison?  Something else?
>
> Entirely seriously, I'm not sure I could implement the current specification.  In particular, I don't believe I could take this specification and implement it in another compiler without *very* close examination of LLVM's current implementation details.  As one of the people who will end up maintaining this, having to back reference the implementation of a N-year old compiler to understand current semantics will rapidly become very problematic.

I don't know what you mean. Implement what? TBH, I'm also not sure what the description below does better, from a conceptual standpoint. Maybe it would help me to understand which parts of the proposed wording is problematic and why?

> Here's an attempt to take what you've written for the !nocapture metadata and translate it into something a bit more operation.  Your welcome to take this as either a starting point, or as simply an illustration of what I meant by operational.
>
> The existence of the !nocapture metadata on the instruction indicates that the location stored to is not captured.  That is, there exists a small set of frame local copies of points which can be used to access that location, and that if the optimizer can identify all of them, it can use that information to contribute to tracking all copies of the stored pointer.

(I saw later you have frame local on the defect list, this was my response initially) Generally, I don't think we have a concept of frame local or we should go there. It is also not true for the use case described here, the users might be in other "frames". I also think the wording swapped who is the "recipient" of the `nocapture` property.

Now with regards to content: Storing into a location doesn't affect the "capture property" of that location. Storing a pointer into a location affects the pointers "capture property".
So, to use your sentence as a base it should read: `The existence of the !nocapture metadata on the instruction indicates that the stored pointer is not captured.`

I explicitly don't want to require anyone to identify the access locations because they might not be exposed. Let's assume `bulk_memset` defined below is in a library and we know what it does.
We can now mark the stores of the pointers into the `ptrs` argument as `!nocapture` if we also add a `"nocapture_use"` to the call site of `bulk_memset`.

  /// lib_bulk.so
  ///{
  void bulk_memset(int nargs, char** ptrs, size_t* sizes, char c) {
    // For example we could do it in parallel:
    #pragma omp parallel for 
    for (int i = 0; i < nargs; ++i) {
      memset(ptrs[i], c, sizes[i]);
    }
  }
  ///}
  
  // User code, original:
  char foo(char *a, char *b) {
    *a = 1;
    *b = 2;
    char *ptrs[2];
    ptrs[0] = a;             // a is captured now
    ptrs[1] = b;             // b is captured now
    int sizes[2] = {1, 1};
    bulk_memset(2, ptrs, sizes, 0); 
    return *a + *b;          // needs to be 0!
  }
  
  // User code, optimized:
  void foo(char /* nocapture */ *a, char *  /* nocapture */ b) {
    *a = 1;
    *b = 2;
    char *ptrs[2];
    ptrs[0] /* !nocapture */ = a;  // a is *not* captured now
    ptrs[1] /* !nocapture */ = b;  // b is *not* captured now
    int sizes[2] = {1, 1};
    bulk_memset(2, ptrs, sizes, 0) ["nocapture_use"(a, b)]; 
    return *a + *b;          // needs to be 0!
  }

This is pretty much the use case we have as part of https://bugs.llvm.org/show_bug.cgi?id=48475 (among other places).
Basically, the API requires us to pass something in memory but we know all we do at the other end is to load it from
that memory and use it without "capturing" it. The pthread_create example I have in the commit message shows how this
can also be used if we see both sides of the API and can deduce the property.

> If the storage location has been captured, execution of a store w/the !nocapture metadata is immediate UB.  Storing a previously captured pointer into an uncaptured location is well defined.  That is, the metadata says nothing of the capture status for the stored value.

But that is what it should do, right? I'm confused why we would need to say something about the captured status of the location we store into at all.

> On it's own, a !nocapture store says nothing about whether the contents of storage are later loaded, and the pointer stored later propagated.  However, !nocapture stores and the nocapture_use operand bundle can be combined to provide this stronger fact.
>
> Defect List -- These are things I know are wrong with my attempt.
>
> "frame local" needs defined, and might not be quite what we want.
>
> "captured" is loosely defined, and needs revision.
>
> After my own attempt, one further suggestion on structuring the langref changes.  I think we need to introduce a definitional section defining capture, and spelling out implications of that, then having the inline description of the metadata be purely structural and link back.  Doing it all inline in the store section seems really forced.

A section about "capture" seems reasonable to me. We refer to that concept also in https://llvm.org/docs/LangRef.html#deoptimization-operand-bundles.

> (Mildly relevant - I have an old review which takes the approach I described towards specification on a vaguely related topic.  Might be useful as a reference.  https://reviews.llvm.org/D52192)

This is really interesting. thanks for the link. I'll probably circle back to that one sometime soon.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D93189