llvm / llvm/llvm-project

[OpenMP][offload] `allocate(omp_const_mem_alloc:)` on a local variable silently miscompiles

Open
#215,290 2 comments 0 reactions 1 assignee Claimed by @abidh View on GitHub
miscompilation offload openmp
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.