Analysis that determines if a parallel region can reach another parallel region in any target region of the TU.
A new global var is emitted with the name of the kernel + "_nested_parallelism", which is either 0 or 1 depending on the result.
Details
Details
Diff Detail
Diff Detail
- Repository
- rG LLVM Github Monorepo
Event Timeline
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 4378 | Any reason you want to keep this variable? Why not add the condition directly into the ID? | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 4378 | Or probably better make NesterParallelism equal you the condition. | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 602 | Is it good enough to just have two states? Is it possible that we can't tell if a parallel region is nested or not, or do we just conservatively take it as true? | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 602 | What others states we might possibly have? The only way we can't tell if a parallel region is nested or not is when we have functions with parallel regions in another TU, for this case we conservatively take it as false, for all the other cases this approach should work. | |
| 4378 | Not really, we can remove the variable and add it to the IF statement | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 602 | If a function containing a parallel region can be reached from different paths, and the parallel region may or may not be nested, that would be the case right? What's more, I suppose we could do more aggressive optimization if we do not have nested parallelism, which means the conservative way would be to take it as a true. Did I misunderstand anything here? | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 602 | We start optimistic, so assuming no nested parallelism is right. | |
| 689 | Also known as | |
| 3353 | Use an int8 type, int1 is always tricky. | |
| 4379 | Don't use magic constants (5, above), that's why we have these variables (WrapperFunctionArgNo). Also, start at the wrapper, so ParallelRegion. The DepClassTy should be optional. | |
| llvm/test/Transforms/OpenMP/nested_parallelism.ll | ||
| 5 | Put the C code as a comment in here so it's easier to read the IR. | |
| 343 | A lot of this code is really not needed. Maybe start with a simpler example or simplify the IR. | |
| 418 | Remove all attributes. | |
| 441 | Remove the TBAA metadata and everything we don't need. !0 - !3 and !6 and !7 are needed. Adjust the offload_info and nvvm annotations, and module.flags accordingly. | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 602 | Shilei, are you referring to something like this: void foo () {
#pragma omp parallel
{}
}
void bar() {
#target
foo();
}
void zoo() {
#target parallel
foo();
}I believe the optimization is not from the perspective of the inner parallel region, but the kernel itself. So this would be consider nested. Rafa, Correct me if I am wrong. | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 602 | Okay then that makes sense. | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 4373 | NestedParallelism = !FnAA.ReachedKnownParallelRegions.empty() || !FnAA.ReachedUnknownParallelRegions.empty() | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 4373 | Rafa. I believe this comment by johannes has not been solved. The the if statement before the assignment. It can also be added to this line. | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 4373 | I think what he meant was to check if the AA was a valid state before reading NestedParallelism |= (!FnAA.ReachedKnownParallelRegions.empty() || !FnAA.ReachedUnknownParallelRegions.empty()) && FnAA.getState().isValidState() Johannes, please correct me if I’m wrong. | |
| llvm/lib/Transforms/IPO/OpenMPOpt.cpp | ||
|---|---|---|
| 4373 | You got the default wrong. If the state of FnAA is invalid we need to assume the worst. Right now you ignore invalid AAs, which is bad. So if Invalid, or if the other condition, then we assume nested parallelism. | |
Is it good enough to just have two states? Is it possible that we can't tell if a parallel region is nested or not, or do we just conservatively take it as true?