[Clang][Flang][OpenMP] `if` clause for `teams` construct has no effect on accelerator
- Dominant language
- LLVM
- Stars
- 40.5k
- Forks
- 18.7k
- PR merge metrics
- PR metrics pending
Description
This is a follow-up to #205345, #207444
---
The OpenMP v5.2 and v6.0 specification mention the following addition to the OpenMP teams construct:
> When an `if` clause is present on a `teams` construct and the `if` clause expression evaluates to `false`, the number of formed teams is one. The use of a variable in an `if` clause expression of a `teams` construct causes an implicit reference to the variable in all enclosing constructs. The `if` clause expression is evaluated in the context outside of the `teams` construct.
Prior to #207444, this was not respected for Clang for host-side execution. Here, calls to `__kmp_push_num_teams` were missing.
However, accelerator execution is still affected by this, and does not respect the number of teams. In https://github.com/llvm/llvm-project/pull/207444#issuecomment-4912406134, I took a look at the generated IR code by Clang. While calls to `__kmp_push_num_threads` are inserted, the call is essentially a no-op. Instead, the number of leagues for the created team are set in the call to `__tgt_target_kernel`. This is not handled yet.
In Clang, [`CGOpenMPRuntime::getNumTeamsExprForTargetDirective`](https://github.com/llvm/llvm-project/blob/ef1b46f29293db7b16dee598fb1e8f2a62555c33/clang/lib/CodeGen/CGOpenMPRuntime.cpp#L6474) only checks for `OMPNumTeamsClause`, not for a potential `OMPIfClause`. I assume that a similar fix to #207444 should suffice here. In Flang, [`emitTargetCall`](https://github.com/llvm/llvm-project/blob/642b85388cca2114e467c58c02fe79b32f7b221f/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp#L9712) checks for an if condition, but only for the `target` construct itself. Handling for a potential `if` clause on the teams construct is not implemented.
This leads to the following incorrect number of teams printed:
```fortran
program main
use omp_lib
implicit none
integer :: teams
!$omp target teams if(teams: .false.) map(tofrom: teams)
teams = omp_get_num_teams()
!$omp end target teams
write (*,*) teams
end program main
```
```console
$ flang -v 2>&1 | head -n 1
flang version 23.0.0git (https://github.com/llvm/llvm-project.git c2f5040f47c8e7991b36cea3afc0b855956a5475
$ flang -fopenmp --offload-arch=gfx1101 test.f90 -fopenmp-version=52
warning: OpenMP support for version 52 in flang is still incomplete [-Wexperimental-option]
warning: OpenMP support for version 52 in flang is still incomplete [-Wexperimental-option]
warning: OpenMP support for version 52 in flang is still incomplete [-Wexperimental-option]
$ ./a.out
216
```
and a similar result for Clang:
```c
#include
#include
int main(void) {
int teams = 0;
int parallel = 0;
#pragma omp target teams distribute parallel for device(0) if(parallel: 0) if(teams: 0) map(tofrom: teams)
for(int i = 0; i < 1000; ++i)
{
teams = omp_get_num_teams();
parallel = omp_get_num_threads();
}
printf("Num Leagues %d\n", teams);
printf("Num Threads %d\n", parallel);
}
```
```console
$ clang -v 2>&1 | head -n 1
clang version 23.0.0git (https://github.com/llvm/llvm-project.git c2f5040f47c8e7991b36cea3afc0b855956a5475)
$ clang -fopenmp --offload-arch=gfx1101 test.c -O0 -fopenmp-version=52
$ ./a.out
Num Leagues 216
Num Threads 0
```
Here, the `parallel` construct is properly ignored, whereas the `teams` construct still creates 216 teams instead of 1.
Contributor guide
Research direction
Reproduce the issue with the provided Clang C and Flang Fortran examples, then inspect clang/lib/CodeGen/CGOpenMPRuntime.cpp at getNumTeamsExprForTargetDirective and llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp at emitTargetCall. Compare the existing handling of if clauses and verify that an if(false) teams construct produces one team for accelerator execution.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp, fortran
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Clearly specified
- Newbie friendliness
- 52/100