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
Diff Detail
Event Timeline
Comment Actions
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 | ||
|---|---|---|
| 2 | Why do you want to run this with stats? I guess later you just use filecheck and verify the IR, right? | |
| 145 | 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. | |
Comment Actions
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 | ||
|---|---|---|
| 2 | 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. | |
| 145 | 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. | |
Comment Actions
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 | ||
|---|---|---|
| 2 | 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. | |
| 145 | The test case is useful, to check we do not accidentally hoist above something we are not allowed to. (=negative test) | |
Comment Actions
- 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.
Comment Actions
Added __tgt_target_data_begin to the commented original code, just to make the test cases easier to read.
| llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll | ||
|---|---|---|
| 2 | 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. | |
Comment Actions
- Fix heavyComputation1 test case.
- Passing all test cases through optimizer in order to remove some unnecessary code.
Comment Actions
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.
Why do you want to run this with stats? I guess later you just use filecheck and verify the IR, right?