On 23.07.21 12:21, Chung-Lin Tang wrote:
Inside offloaded regions, the preferred way to should be that the device
already has this information initialized (once) when the device is
initialized.
And the function merely returns the stored value.
...
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-45.c
@@ -0,0 +1,23 @@
+/* { dg-do run { target { offload_target_nvptx || offload_target_amdgcn } } }
*/
...
+ int device_num;
+ #pragma omp target map(from: device_num)
+ {
+ device_num = omp_get_device_num ();
+ }
+
+ if (host_device_num == device_num)
+ abort ();
I personally prefer having:
int initial_dev;
and inside 'omp target' (with 'map(from:initial_dev)'):
initial_device = omp_is_initial_device();
Then the check would be:
if (initial_device && host_device_num != device_num)
abort();
if (!initial_device && host_device_num == device_num)
abort();
(Likewise for Fortran.)
And instead of restricting the target to nvptx/gcn, we could just add
dg-xfail-run-if for *-intelmic-* and *-intelmicemul-*.
Additionally, offload_target_nvptx/...amdgcn only check whether
compilation support is available not whether a device exists
at run time.
(The device availability is checked by target_offload_device,
using omp_is_initial_device().)
Tobias
PS: For completeness, I want to note that OpenMP 5.1 supports
setting the per-device ICV as via the environment variables,
besides inheriting the generic ICV values, device-specific
settings are possible with:
<ENVIRONMENT VARIABLE>_DEV[_<device>]
Thus, more data will be passed from libgomp to the plugins
in the future.
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
