Thyre wrote:
I had another look at this fix, especially checking how the fix interacts with
the `target` construct.
For pure host-side regions, e.g.:
```c
#pragma omp target teams if(false)
#pragma omp teams if(false)
#pragma omp teams if(teams: false)
```
this works as expected and only one team is created.
This is also reflected in the generated IR:
```llvm
; Function Attrs: noinline norecurse nounwind optnone uwtable
define internal void @__omp_offloading_803_31c3b3d_main_l6(ptr noundef nonnull
align 4 dereferenceable(4) %teams, ptr noalias noundef %dyn_ptr) #1 {
entry:
%teams.addr = alloca ptr, align 8
%dyn_ptr.addr = alloca ptr, align 8
%0 = call i32 @__kmpc_global_thread_num(ptr @1)
store ptr %teams, ptr %teams.addr, align 8
store ptr %dyn_ptr, ptr %dyn_ptr.addr, align 8
%1 = load ptr, ptr %teams.addr, align 8, !nonnull !0, !align !12
call void @__kmpc_push_num_teams(ptr @1, i32 %0, i32 1, i32 0)
call void (ptr, i32, ptr, ...) @__kmpc_fork_teams(ptr @1, i32 1, ptr
@__omp_offloading_803_31c3b3d_main_l6.omp_outlined, ptr %1)
ret void
}
```
For offloaded code however, the call to
[`__kmpc_fork_teams`](https://github.com/llvm/llvm-project/blob/2b080317511d3285c7376e25ee9668edcc41c06e/openmp/device/src/Parallelism.cpp#L308)
is a no-op.
Setting the number of teams is done differently. A user code:
```c
#pragma omp target teams if(teams: false)
```
generates a call to:
```llvm
%21 = call i32 @__tgt_target_kernel(ptr @1, i64 0, i32 0, i32 0, ptr
@.__omp_offloading_803_31c3b3d_main_l6.region_id, ptr %kernel_args)
```
where the number of teams is passed as the third argument.
This is not handled in this PR, hence a code like this:
```c
#include <omp.h>
#include <stdio.h>
int main(void) {
int teams = 0;
#pragma omp target teams device(0) if(teams: 0) map(tofrom: teams)
{
teams = omp_get_num_teams();
}
printf("Num teams %d\n", teams);
}
```
will yield: `Num teams 216` on my workstation.
I haven't spent much time looking into this yet, but I'd propose to handle this
in a follow-up PR. I'll change the PR title and description to make clear that
this only addresses host execution.
https://github.com/llvm/llvm-project/pull/207444
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
