llvm / llvm/llvm-project

[OpenMP] Device ID `-1` executes accelerator event on default device instead of host device

Open
#192,980 5 comments 0 reactions 0 assignees View on GitHub
openmp
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.