| Issue |
208672
|
| Summary |
[Clang][Flang][OpenMP] `if` clause for `teams` construct has no effect on accelerator
|
| Labels |
|
| Assignees |
|
| Reporter |
Thyre
|
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 <omp.h>
#include <stdio.h>
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.
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
