| Issue |
181210
|
| Summary |
[OpenMP] Offloading in 'Generic' mode runs with incorrect number of threads
|
| Labels |
new issue
|
| Assignees |
|
| Reporter |
jhuber6
|
The following test case fails the assertion as the outlined parallel region is only executed once. With optimizations the kernel is converted to 'SPMD' mode and no longer fails. This is possible due to the main warp of the state machine not participating in the parallel region properly.
```c
#include <stdio.h>
#include <assert.h>
int g = 0;
#pragma omp declare target to(g)
void foo() {
#pragma omp parallel num_threads(2)
__atomic_fetch_add(&g, 1, __ATOMIC_RELAXED);
}
int main() {
#pragma omp target teams num_teams(1)
foo();
#pragma omp target update from(g)
assert(g == 2 && "Not run with two threads?");
}
```
```console
> clang print.c -fopenmp --offload-arch=gfx1030 -O0
> ./a.out
a.out: print.c:17: int main(): Assertion `g == 2 && "Not run with two threads?"' failed.
[1] 268954 IOT instruction (core dumped) ./a.out
> clang print.c -fopenmp --offload-arch=sm_89 -O0
> ./a.out
a.out: print.c:17: int main(): Assertion `g == 2 && "Not run with two threads?"' failed.
[1] 268997 IOT instruction (core dumped) ./a.out
```
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
