llvm / llvm/llvm-project

[Clang][Flang][OpenMP] `if` clause for `teams` construct has no effect on accelerator

Open
#208,672 1 comment 0 reactions 0 assignees View on GitHub
clang flang openmp
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.