Handle 'device_type(nohost)' in OpenMP target regions
OpenMP target regions now support the `device_type(nohost)` clause for offloading, enabling code execution only on devices.
This commit adds support for the device_type(nohost) clause in OpenMP target regions. It modifies the middle-end to generate code that uses the omp_is_initial_device intrinsic, which is evaluated at compile time. This ensures that only the device-specific code is compiled for nohost regions, while host code is omitted, simplifying offloading constructs.
In Details
The gimplify_omp_workshare and lower_omp_target functions in gimplify.cc and omp-low.cc are updated to handle the device_type(nohost) clause. This is achieved by introducing an if-else structure that uses the omp_is_initial_device intrinsic. This intrinsic is then evaluated to a constant during late optimization, effectively dead-code eliminating either the host or device part of the target region based on the nohost specifier.
- OpenMP target regions
- Code blocks marked with
#pragma omp targetdirectives, intended to be offloaded to a separate device (like a GPU or a coprocessor). - device_type(nohost)
- An OpenMP clause that specifies the code within an OpenMP target region should only execute on the target device and not on the host CPU.
- omp_is_initial_device intrinsic
- A built-in function in OpenMP that returns true if the current code is executing on the initial device (usually the host CPU), and false otherwise. It's often evaluated to a compile-time constant for target regions.