[PATCH] D80051: [OpenMPOpt] Test case 1 - Latency Hiding for Host to Device Memory Transfers

[Johannes Doerfert via Phabricator via llvm-commits]

jdoerfert added a comment.

In D80051#2040176 <https://reviews.llvm.org/D80051#2040176>, @hamax97 wrote:

> In D80051#2039931 <https://reviews.llvm.org/D80051#2039931>, @jdoerfert wrote:
>
> > This is great! Thanks for starting. I left two comments but this looks pretty good to me.
> >  We will later need positive and negative test cases. We can make those smaller though,
> >  you just have the memory transfer request, no target region. We want to verify we move it
> >  properly, e.g., when it appears in loops, in conditionals, there are memory accesses
> >  before/after, ...
> >
> > Btw. if it helps, we can also create a "thinner" version of the API that takes less arguments
> >  and is applicable only in some often occurring cases, e.g., we can remove the baseptr array
> >  and use it only if ptr and baseptr have the same entries anyway.
>
>
> Do you think I should create all those test cases before starting to code?.


You do not have to but we need them eventually. To make it easier to write and read them, just include a call like `__tgt_target_data_begin` and pass on arguments, e.g.

  void (void *arg, ...) {
    ; some interesting code you want the issue to interact with
    __tgt_target_data_begin(..., arg, ...);
    ; some interesting code you want the wait to interact with
  }



> Also, making the API thinner would help a lot I guess.

A thinner API will mostly make the tests cases smaller but conceptually I don't think it will make much of a difference.

> Actually, I don't fully understand what the `baseptr` array is for, could you give me a simple explanation of its usage?.

So if you write `#pragma omp target data map(to:A[x:y])` it means you map from `&A[x]` to `&A[x+y]`. As there are special rules about partially mapped objects we need to store also `A`, the `baseptr` of the expression. `&A[x]` will be `ptr` and `y` will be `size`.



================
Comment at: llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll:2
+; RUN: opt -S -openmpopt -stats < %s 2>&1
+; REQUIRES: asserts
+
----------------
hamax97 wrote:
> jdoerfert wrote:
> > Why do you want to run this with stats? I guess later you just use filecheck and verify the IR, right?
> I thought it would be useful to perhaps count the number of split runtime calls, and how much they were moved forward and backwards. What do you think?. I have no problem not doing it.
We will eventually check the resulting IR which contains all the information. We basically compare against the expected splits and movements so any derivation needs to be flagged. That said, having a test case to verify our statistic tracking works and one to verify the remarks (=feedback to the user) work, is important. Though, both should be separate test cases with appropriate names.


================
Comment at: llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll:145
+;}
+define dso_local i32 @heavyComputation2(double* %a, i32 %size) {
+entry:
----------------
hamax97 wrote:
> jdoerfert wrote:
> > For the issue to be hoisted above the `rand` call the pointer needs to be noalias or we need better information about `rand`. As it is, `rand` might modify `a` which prevents moving.
> Oh yes, you're right. So, this test case is only useful if `*a` is `noalias`, or I could add some code between the `target data map()` and the `target teams` which would enable moving the `wait` function. This latter way would avoid the use of `noalias` which is not commonly used I guess. But maybe it's a good idea to have a test case where `noalias` is used, it would be a nice optimization.
The test case is useful, to check we do not accidentally hoist above something we are not allowed to. (=negative test)
You want tests with interesting code on both sides to verify movement in both directions, yes. We will need more tests for coverage later.
Noalias is not uncommon either. People use `__restrict/restrict` and the Attributor can derive `noalias` even without.
Generally speaking, we want all kinds of test cases to cover all kinds of situations. Usual approach, copy this function and add `noalias` to one of them ;)


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

https://reviews.llvm.org/D80051