Page MenuHomePhabricator

[OpenMPOpt] Test case 1 - Latency Hiding for Host to Device Memory Transfers
ClosedPublic

Authored by hamax97 on May 15 2020, 4:40 PM.

Details

Summary

First test case that has FIXMEs showing how the runtime function __tgt_target_data_begin should be split into issue and wait functions, and those functions should be moved forwards and backwards, respectively.

Diff Detail

Event Timeline

hamax97 created this revision.May 15 2020, 4:40 PM

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.

llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll
3

Why do you want to run this with stats? I guess later you just use filecheck and verify the IR, right?

146

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.

hamax97 marked 2 inline comments as done.May 16 2020, 10:32 AM

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?. Also, making the API thinner would help a lot I guess. Actually, I don't fully understand what the baseptr array is for, could you give me a simple explanation of its usage?.

llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll
3

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.

146

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.

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.

llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll
3

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.

146

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 ;)

hamax97 updated this revision to Diff 265802.May 22 2020, 2:43 PM
  • Fixed illegal movement: now there's one negative test case (issue cannot be moved) and one positive (with noalias) where the issue can be moved.
  • Added data transfer only test case.
hamax97 updated this revision to Diff 265929.May 24 2020, 7:39 AM

Added __tgt_target_data_begin to the commented original code, just to make the test cases easier to read.

jdoerfert added inline comments.May 26 2020, 5:31 PM
llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll
3

No error output and asserts needed. We should however already use filecheck to check for the outut we care about. You can use llvm/utils/update_test_checks.py on this file. Once we optimize this we can nicely see the effect in the diff.

hamax97 updated this revision to Diff 267762.Jun 1 2020, 4:58 PM

Added FileCheck constructs

jdoerfert accepted this revision.Jun 1 2020, 5:44 PM

LGTM. One nit: rerun the update test check with --function-signature

This revision is now accepted and ready to land.Jun 1 2020, 5:44 PM
hamax97 updated this revision to Diff 275563.Jul 5 2020, 12:43 PM
hamax97 added a reviewer: sstefan1.
  • Fix heavyComputation1 test case.
  • Passing all test cases through optimizer in order to remove some unnecessary code.

This can be closed, right?

This can be closed, right?

Yes, it's already merged.

sstefan1 closed this revision.Aug 12 2020, 1:50 PM

Closed with: https://github.com/llvm/llvm-project/commit/6f0d99d2b9b3b8ae96dd91c8988cc067b9c9afb9

BTW if you add: Differential Revision: https://reviews.llvm.org/D##### to the commit message it should automatically close and link the commit to your profile.