[OpenMP] Device ID `-1` executes accelerator event on default device instead of host device
- Dominant language
- LLVM
- Stars
- 40.5k
- Forks
- 18.7k
- PR merge metrics
- PR metrics pending
Description
Given the following example program, the kernel is executed on the default accelerator instead of the host device:
```c
#include
#include
#include
void foo( int device )
{
bool is_host_device = false;
#pragma omp target device( device ) map(from: is_host_device)
{
is_host_device = omp_is_initial_device();
}
printf("Ran on initial device? %d\n", is_host_device);
}
int main( int argc, char ** argv )
{
foo( -1 );
}
```
Console:
```console
$ clang -fopenmp --offload-arch=gfx1101 test.c
$ ./a.out
Ran on initial device? 0
$ clang --version
clang version 23.0.0git (https://github.com/llvm/llvm-project.git cea56f648ef4c0972c04c5230393f1e2a5d590b9)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/apps/software/Clang/trunk/bin
Build config: +assertions
```
----
Up until OpenMP v5.2, no constant was defined for the host device in OpenMP. Applications were expected to use `omp_get_initial_device()`, which mapped to `omp_get_num_devices()`, hence was always `>=0`.
Version 5.2 introduced the constant `omp_initial_device`, which can be used as an alias. This constant is defined as `-1` (OpenMP v5.2, p. 346, l.14; OpenMP v5.2, p. 347, l.25; OpenMP v6.0, p. 534, l. 5) and is allowed to be passed for a `device` clause.
Similarly `-2` is defined for `omp_invalid_device`. Passing this causes the runtime to (correctly) abort execution, as defined in the OpenMP spec.
----
In [`checkDevice`](https://github.com/llvm/llvm-project/blob/c7f0fd60659e4e3b8c9b472effec8cd7225de825/offload/libomptarget/interface.cpp#L52), we replace the passed device number by the default device, if our passed device matches [`OFFLOAD_DEVICE_DEFAULT`](https://github.com/llvm/llvm-project/blob/c7f0fd60659e4e3b8c9b472effec8cd7225de825/offload/libomptarget/interface.cpp#L58).
This variable is [defined as `-1`](https://github.com/llvm/llvm-project/blob/c7f0fd60659e4e3b8c9b472effec8cd7225de825/offload/include/omptarget.h#L34), clashing with the definition of `omp_initial_device`.
A quick fix would be to re-define the variable to a new unique value.
This change would need to be adapted in other places using `OFFLOAD_DEVICE_DEFAULT`, mainly affecting `llvm::omp::OMP_DEVICEID_UNDEF` used in `mlir` and Clangs `CodeGen` (see definition [here](https://github.com/llvm/llvm-project/blob/5d98f79f52dd56750e826eb6efcd8b5cf41de239/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h#L272))
Contributor guide
Research direction
Start with checkDevice in offload/libomptarget/interface.cpp and compare OFFLOAD_DEVICE_DEFAULT in offload/include/omptarget.h with omp_initial_device. Trace its uses, especially OMP_DEVICEID_UNDEF in llvm/include/llvm/Frontend/OpenMP/OMPConstants.h and the referenced MLIR and Clang CodeGen paths. Validate with the provided C reproducer: device -1 must run on the host while the default-device sentinel remains distinct, and -2 must still abort.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- c
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Mostly clear
- Newbie friendliness
- 48/100