[OpenMP][offload] `allocate(omp_const_mem_alloc:)` on a local variable silently miscompiles
- Dominant language
- LLVM
- Stars
- 40.5k
- Forks
- 18.7k
- PR merge metrics
- PR metrics pending
Description
A variable that is private to a target region and given the `omp_const_mem_alloc` allocator is placed in the device's constant address
space, which the device cannot write. The initialising store is emitted anyway and is then discarded, so everything computed from the variable is lost. There is no diagnostic: the region simply produces a wrong answer.
## Reproducer
```c
#include
#include
int main(void) {
int x = 23;
int got = 0;
#pragma omp target uses_allocators(omp_const_mem_alloc) \
allocate(omp_const_mem_alloc : x) firstprivate(x) map(tofrom : got)
for (int i = 0; i < 1024; ++i)
got += x;
printf("got = %d, expected = %d\n", got, 1024 * 23);
return got != 1024 * 23;
}
```
## How to reproduce
```
$ clang -fopenmp -O1 -g --offload-arch=gfx1100 repro.c -o repro
$ ./repro
got = 0, expected = 23552
```
Expected: `got = 23552`. Observed: `got = 0`, and the program exits non-zero.
The compile is silent, with no warning or error at any point.
Note that using `firstprivate` with `omp_const_mem_alloc` is explicitely allowed by the specs. OpenMP 5.2, section 6.1 (Memory Spaces), on `omp_const_mem_space`, which is the memory space behind `omp_const_mem_alloc`:
> Variables allocated in the `omp_const_mem_space` memory space may be
> initialized through the `firstprivate` clause or with compile time constants
> for static and constant variables. Implementation-defined mechanisms to
> provide the constant value of these variables may also be supported.
The reproducer above is exactly that: a `firstprivate` variable, read but never written inside the region. The one restriction the section imposes, "Variables in the `omp_const_mem_space` memory space may not be written", is not violated by it.
The same section also says the actual storage a memory space maps to is implementation defined, and that selecting one "expresses an intent". The host already uses that latitude: `libomp` reports
```
OMP: Warning #190: Allocator omp_const_mem_alloc is not available, will use default allocator.
```
and the same program prints `23552` on the host at every optimisation level.
## Flag matrix
| flags | result |
| -------- | --------------------- |
| `-O0 -g` | correct (`23552`) |
| `-O1` | correct (`23552`) |
| `-O1 -g` | wrong (`0`) |
| `-O2 -g` | wrong (`0`) |
| `-O3 -g` | wrong (`0`) |
Contributor guide
Assessment
This issue has not been assessed yet.