From 81476bc4f4a20bcf3af7ac2548c2322d48499402 Mon Sep 17 00:00:00 2001 From: Marcel Vollweiler Date: Tue, 6 Dec 2022 05:42:46 -0800 Subject: OpenMP: omp_get_max_teams, omp_set_num_teams, and omp_{gs}et_teams_thread_limit on offload devices This patch adds support for omp_get_max_teams, omp_set_num_teams, and omp_{gs}et_teams_thread_limit on offload devices. That includes the usage of device-specific ICV values (specified as environment variables or changed on a device). In order to reuse device-specific ICV values, a copy back mechanism is implemented that copies ICV values back from device to the host. Additionally, a limitation of the number of teams on gcn offload devices is implemented. The number of teams is limited by twice the number of compute units (one team is executed on one compute unit). This avoids queueing unnessecary many teams and a corresponding allocation of large amounts of memory. Without that limitation the memory allocation for a large number of user-specified teams can result in an "memory access fault". A limitation of the number of teams is already also implemented for nvptx devices (see nvptx_adjust_launch_bounds in libgomp/plugin/plugin-nvptx.c). gcc/ChangeLog: * gimplify.cc (optimize_target_teams): Set initial num_teams_upper to "-2" instead of "1" for non-existing num_teams clause in order to disambiguate from the case of an existing num_teams clause with value 1. libgomp/ChangeLog: * config/gcn/icv-device.c (omp_get_teams_thread_limit): Added to allow processing of device-specific values. (omp_set_teams_thread_limit): Likewise. (ialias): Likewise. * config/nvptx/icv-device.c (omp_get_teams_thread_limit): Likewise. (omp_set_teams_thread_limit): Likewise. (ialias): Likewise. * icv-device.c (omp_get_teams_thread_limit): Likewise. (ialias): Likewise. (omp_set_teams_thread_limit): Likewise. * icv.c (omp_set_teams_thread_limit): Removed. (omp_get_teams_thread_limit): Likewise. (ialias): Likewise. * libgomp.texi: Updated documentation for nvptx and gcn corresponding to the limitation of the number of teams. * plugin/plugin-gcn.c (limit_teams): New helper function that limits the number of teams by twice the number of compute units. (parse_target_attributes): Limit the number of teams on gcn offload devices. * target.c (get_gomp_offload_icvs): Added teams_thread_limit_var handling. (gomp_load_image_to_device): Added a size check for the ICVs struct variable. (gomp_copy_back_icvs): New function that is used in GOMP_target_ext to copy back the ICV values from device to host. (GOMP_target_ext): Update the number of teams and threads in the kernel args also considering device-specific values. * testsuite/libgomp.c-c++-common/icv-4.c: Fixed an error in the reading of OMP_TEAMS_THREAD_LIMIT from the environment. * testsuite/libgomp.c-c++-common/icv-5.c: Extended. * testsuite/libgomp.c-c++-common/icv-6.c: Extended. * testsuite/libgomp.c-c++-common/icv-7.c: Extended. * testsuite/libgomp.c-c++-common/icv-9.c: New test. * testsuite/libgomp.fortran/icv-5.f90: New test. * testsuite/libgomp.fortran/icv-6.f90: New test. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-teams-1.c: Adapt expected values for num_teams from "1" to "-2" in cases without num_teams clause. * g++.dg/gomp/target-teams-1.C: Likewise. * gfortran.dg/gomp/defaultmap-4.f90: Likewise. * gfortran.dg/gomp/defaultmap-5.f90: Likewise. * gfortran.dg/gomp/defaultmap-6.f90: Likewise. --- libgomp/testsuite/libgomp.c-c++-common/icv-4.c | 2 +- libgomp/testsuite/libgomp.c-c++-common/icv-5.c | 198 ++++++++++++++++++++-- libgomp/testsuite/libgomp.c-c++-common/icv-6.c | 61 ++++++- libgomp/testsuite/libgomp.c-c++-common/icv-7.c | 83 ++++++++- libgomp/testsuite/libgomp.c-c++-common/icv-9.c | 72 ++++++++ libgomp/testsuite/libgomp.fortran/icv-5.f90 | 226 +++++++++++++++++++++++++ libgomp/testsuite/libgomp.fortran/icv-6.f90 | 140 +++++++++++++++ 7 files changed, 757 insertions(+), 25 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c-c++-common/icv-9.c create mode 100644 libgomp/testsuite/libgomp.fortran/icv-5.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/icv-6.f90 (limited to 'libgomp/testsuite') diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-4.c b/libgomp/testsuite/libgomp.c-c++-common/icv-4.c index b987a33..9da0d63 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/icv-4.c +++ b/libgomp/testsuite/libgomp.c-c++-common/icv-4.c @@ -16,7 +16,7 @@ main () } else omp_set_num_teams (6); - if (getenv ("OMP_TEAMS_THREAD_LIMIT") == NULL + if (getenv ("OMP_TEAMS_THREAD_LIMIT") != NULL && strcmp (getenv ("OMP_TEAMS_THREAD_LIMIT"), "12") == 0) { if (omp_get_teams_thread_limit () != 12) diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-5.c b/libgomp/testsuite/libgomp.c-c++-common/icv-5.c index 431cfc7..72d7af6 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/icv-5.c +++ b/libgomp/testsuite/libgomp.c-c++-common/icv-5.c @@ -1,25 +1,203 @@ /* { dg-do run } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_0 "42" } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1 "43" } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_2 "44" } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "45" } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "46" } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS "47" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "3" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "4" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS "5" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_0 "6" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1 "7" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_2 "8" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "2" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV "3" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT "4" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_0 "5" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_1 "6" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_2 "7" } */ #include #include +#include int main () { - if (omp_get_max_teams () != 47) + if (omp_get_max_teams () != 5 + || omp_get_teams_thread_limit () != 4) abort (); + #pragma omp teams + { + if (omp_get_num_teams () > 5 + || omp_get_team_num () >= 5) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 4 + || omp_get_thread_num () >= 4) + abort (); + } + + omp_set_num_teams (4); + omp_set_teams_thread_limit (3); + if (omp_get_max_teams () != 4 + || omp_get_teams_thread_limit () != 3) + abort (); + + #pragma omp teams + { + if (omp_get_num_teams () > 4 + || omp_get_team_num () >= 4) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 3 + || omp_get_thread_num () >= 3) + abort (); + } + + #pragma omp teams num_teams(3) thread_limit(2) + { + if (omp_get_num_teams () != 3 + || omp_get_team_num () >= 3) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 2 + || omp_get_thread_num () >= 2) + abort (); + } + + #pragma omp teams num_teams(5) thread_limit(4) + { + if (omp_get_num_teams () != 5 + || omp_get_team_num () >= 5) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 4 + || omp_get_thread_num () >= 4) + abort (); + } + int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices (); - for (int i=0; i < num_devices; i++) - #pragma omp target device (i) - if (omp_get_max_teams () != 42 + i) + + for (int i = 0; i < num_devices; i++) + { + #pragma omp target device (i) + if (omp_get_max_teams () != 6 + i + || omp_get_teams_thread_limit () != 5 + i) + abort (); + + #pragma omp target device (i) + #pragma omp teams + #pragma omp parallel + if (omp_get_thread_limit () > 5 + i + || omp_get_thread_num () >= 5 + i) + abort (); + + #pragma omp target device (i) + { + omp_set_num_teams (5 + i); + omp_set_teams_thread_limit (4 + i); + if (omp_get_max_teams () != 5 + i + || omp_get_teams_thread_limit () != 4 + i) + abort (); + } + + /* omp_set_num_teams and omp_set_teams_thread_limit above set the value + of nteams-var and teams-thread-limit-var ICVs on device 'i', which has + scope 'device' and should be avaible in subsequent target regions. */ + #pragma omp target device (i) + if (omp_get_max_teams () != 5 + i + || omp_get_teams_thread_limit () != 4 + i) + abort (); + + #pragma omp target device (i) + #pragma omp teams + { + if (omp_get_num_teams () > 5 + i + || omp_get_team_num () >= 5 + i) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 4 + i + || omp_get_thread_num () >= 4 + i) + abort (); + } + + #pragma omp target device (i) + #pragma omp teams num_teams(6 + i) thread_limit(5 + i) + { + if (omp_get_num_teams () > 6 + i + || omp_get_team_num () >= 6 + i) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 5 + i + || omp_get_thread_num () >= 5 + i + || omp_get_num_teams () > 6 + i + || omp_get_team_num () >= 6 + i) + abort (); + } + + #pragma omp target device (i) + #pragma omp teams num_teams(4 + i) thread_limit(3 + i) + { + if (omp_get_num_teams () > 4 + i + || omp_get_team_num () >= 4 + i) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 3 + i + || omp_get_thread_num () >= 3 + i + || omp_get_num_teams () > 4 + i + || omp_get_team_num () >= 4 + i) + abort (); + } + + #pragma omp target device (i) + #pragma omp teams thread_limit(3 + i) num_teams(4 + i) + { + if (omp_get_num_teams () > 4 + i + || omp_get_team_num () >= 4 + i) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 3 + i + || omp_get_thread_num () >= 3 + i + || omp_get_num_teams () > 4 + i + || omp_get_team_num () >= 4 + i) + abort (); + } + + /* The NUM_TEAMS and THREAD_LIMIT clauses should not change the values + of the corresponding ICVs. */ + #pragma omp target device (i) + if (omp_get_max_teams () != 5 + i + || omp_get_teams_thread_limit () != 4 + i) + abort (); + + /* This tests a large number of teams and threads. If it is larger than + 2^15+1 then the according argument in the kernels arguments list + is encoded with two items instead of one. */ + intptr_t large_num_teams = 66000; + intptr_t large_threads_limit = 67000; + #pragma omp target device (i) + { + omp_set_num_teams (large_num_teams + i); + omp_set_teams_thread_limit (large_threads_limit + i); + if (omp_get_max_teams () != large_num_teams + i + || omp_get_teams_thread_limit () != large_threads_limit + i) + abort (); + } + + #pragma omp target device (i) + if (omp_get_max_teams () != large_num_teams + i + || omp_get_teams_thread_limit () != large_threads_limit + i) abort (); + #pragma omp target device (i) + #pragma omp teams + { + if (omp_get_num_teams () > large_num_teams + i + || omp_get_team_num () >= large_num_teams + i) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > large_threads_limit + i + || omp_get_thread_num () >= large_threads_limit + i) + abort (); + } + } + return 0; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-6.c b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c index e199a18..1b17f2d 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/icv-6.c +++ b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c @@ -1,9 +1,10 @@ /* { dg-do run } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "42" } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "43" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "3" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV "4" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "2" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV "3" } */ /* { dg-set-target-env-var OMP_SCHEDULE_ALL "guided,4" } */ /* { dg-set-target-env-var OMP_DYNAMIC_ALL "true" } */ -/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "44" } */ /* { dg-set-target-env-var OMP_THREAD_LIMIT_ALL "45" } */ /* { dg-set-target-env-var OMP_NUM_THREADS_ALL "46,3,2" } */ /* { dg-set-target-env-var OMP_MAX_ACTIVE_LEVELS_ALL "47" } */ @@ -12,8 +13,8 @@ /* This tests the hierarchical usage of ICVs on the device, i.e. if OMP_NUM_TEAMS_DEV_ is not configured, then the value of - OMP_NUM_TEAMS_DEV should be used. And if there is no environment variable - without suffix, then the corresponding _ALL variant should be used. */ + OMP_NUM_TEAMS_DEV should be used. And if OMP_NUM_TEAMS (without suffix) is + not defined, then OMP_NUM_TEAMS_ALL should be used for the host. */ #include #include @@ -26,10 +27,10 @@ main () int chunk_size; omp_get_schedule(&kind, &chunk_size); - if ((!getenv ("OMP_NUM_TEAMS") && omp_get_max_teams () != 42) + if ((!getenv ("OMP_NUM_TEAMS") && omp_get_max_teams () != 3) || (!getenv ("OMP_DYNAMIC") && !omp_get_dynamic ()) || (!getenv ("OMP_SCHEDULE") && (kind != 3 || chunk_size != 4)) - || (!getenv ("OMP_TEAMS_THREAD_LIMIT") && omp_get_teams_thread_limit () != 44) + || (!getenv ("OMP_TEAMS_THREAD_LIMIT") && omp_get_teams_thread_limit () != 2) || (!getenv ("OMP_THREAD_LIMIT") && omp_get_thread_limit () != 45) || (!getenv ("OMP_NUM_THREADS") && omp_get_max_threads () != 46) || (!getenv ("OMP_PROC_BIND") && omp_get_proc_bind () != omp_proc_bind_spread) @@ -44,9 +45,51 @@ main () name[sizeof ("OMP_NUM_TEAMS_DEV_1") - 2] = '0' + i; if (getenv (name)) continue; - #pragma omp target device (i) - if (omp_get_max_teams () != 43) + + #pragma omp target device (i) + if (omp_get_max_teams () != 4 + || omp_get_teams_thread_limit () != 3) abort (); + #pragma omp target device (i) + #pragma omp teams + { + if (omp_get_num_teams () > 4 + || omp_get_team_num () >= 4) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 3 + || omp_get_thread_num () >= 3) + abort (); + } + + #pragma omp target device (i) + { + omp_set_num_teams (3 + i); + omp_set_teams_thread_limit (2 + i); + if (omp_get_max_teams () != 3 + i + || omp_get_teams_thread_limit () != 2 + i) + abort (); + } + + /* omp_set_num_teams above set the value of nteams-var ICV on device 'i', + which has scope 'device' and should be avaible in subsequent target + regions. */ + #pragma omp target device (i) + if (omp_get_max_teams () != 3 + i + || omp_get_teams_thread_limit () != 2 + i) + abort (); + + #pragma omp target device (i) + #pragma omp teams + { + if (omp_get_num_teams () > 3 + i + || omp_get_team_num () >= 3 + i) + abort (); + #pragma omp parallel + if (omp_get_thread_limit () > 2 + i + || omp_get_thread_num () >= 2 + i) + abort (); + } } return 0; diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-7.c b/libgomp/testsuite/libgomp.c-c++-common/icv-7.c index 70a716d..bbbd6df 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/icv-7.c +++ b/libgomp/testsuite/libgomp.c-c++-common/icv-7.c @@ -1,5 +1,6 @@ /* { dg-do run } */ -/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "42" } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL "7" } */ +/* { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "2" } */ /* This tests the hierarchical usage of ICVs on the host and on devices, i.e. if OMP_NUM_TEAMS_DEV_, OMP_NUM_TEAMS_DEV, and @@ -9,18 +10,90 @@ #include #include +#include int main () { - if (omp_get_max_teams () != 42) + if ((!getenv ("OMP_NUM_TEAMS") && omp_get_max_teams () != 7) + || (!getenv ("OMP_TEAMS_THREAD_LIMIT") && omp_get_teams_thread_limit () != 2)) abort (); + #pragma omp teams + if ((!getenv ("OMP_NUM_TEAMS")) + && (omp_get_num_teams () > 7 || omp_get_team_num () >= 7)) + abort (); + + omp_set_num_teams (9); + omp_set_teams_thread_limit (3); + if (omp_get_max_teams () != 9 + || omp_get_teams_thread_limit () != 3) + abort (); + + #pragma omp teams + if (omp_get_num_teams () > 9 + || omp_get_team_num () >= 9) + abort (); + + #pragma omp teams num_teams(5) + if (omp_get_num_teams () > 5 + || omp_get_team_num () >= 5) + abort (); + + if (getenv ("OMP_NUM_TEAMS_DEV") || getenv ("OMP_TEAMS_THREAD_LIMIT_DEV")) + return 0; + int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices (); - for (int i=0; i < num_devices; i++) - #pragma omp target device (i) - if (omp_get_max_teams () != 42) + for (int i = 0; i < num_devices; i++) + { + char nteams[sizeof ("OMP_NUM_TEAMS_DEV_1")]; + strcpy (nteams, "OMP_NUM_TEAMS_DEV_1"); + nteams[sizeof ("OMP_NUM_TEAMS_DEV_1") - 2] = '0' + i; + char teams_thread_limit[sizeof ("OMP_TEAMS_THREAD_LIMIT_DEV_1")]; + strcpy (teams_thread_limit, "OMP_TEAMS_THREAD_LIMIT_DEV_1"); + teams_thread_limit[sizeof ("OMP_TEAMS_THREAD_LIMIT_DEV_1") - 2] = '0' + i; + if (getenv (nteams) || getenv (teams_thread_limit)) + continue; + + #pragma omp target device (i) + if (omp_get_max_teams () != 7 + || omp_get_teams_thread_limit () != 2) + abort (); + + #pragma omp target device (i) + #pragma omp teams + if (omp_get_num_teams () > 7 + || omp_get_team_num () >= 7) + abort (); + + #pragma omp target device (i) + { + omp_set_num_teams (8 + i); + omp_set_teams_thread_limit (4 + i); + if (omp_get_max_teams () != 8 + i + || omp_get_teams_thread_limit () != 4 + i) + abort (); + } + + /* omp_set_num_teams above set the value of nteams-var ICV on device 'i', + which has scope 'device' and should be avaible in subsequent target + regions. */ + #pragma omp target device (i) + if (omp_get_max_teams () != 8 + i + || omp_get_teams_thread_limit () != 4 + i) + abort (); + + #pragma omp target device (i) + #pragma omp teams + if (omp_get_num_teams () > 8 + i + || omp_get_team_num () >= 8 + i) + abort (); + + #pragma omp target device (i) + #pragma omp teams num_teams(5 + i) + if (omp_get_num_teams () != 5 + i) abort (); + } return 0; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-9.c b/libgomp/testsuite/libgomp.c-c++-common/icv-9.c new file mode 100644 index 0000000..21b874f --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/icv-9.c @@ -0,0 +1,72 @@ +/* { dg-do run } */ + +/* This tests usage of ICVs on the host and on devices if no corresponding + environment variables are configured. */ + +#include +#include + +int +main () +{ + if (omp_get_max_teams () != 0 + || omp_get_teams_thread_limit () != 0) + abort (); + + omp_set_num_teams (9); + omp_set_teams_thread_limit (2); + if (omp_get_max_teams () != 9 + || omp_get_teams_thread_limit () != 2) + abort (); + + #pragma omp teams + if (omp_get_num_teams () > 9 + || omp_get_team_num () >= 9) + abort (); + + #pragma omp teams num_teams(5) + if (omp_get_num_teams () > 5 + || omp_get_team_num () >= 5) + abort (); + + int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices (); + for (int i = 0; i < num_devices; i++) + { + #pragma omp target device (i) + if (omp_get_max_teams () != 0 + || omp_get_teams_thread_limit () != 0) + abort (); + + #pragma omp target device (i) + { + omp_set_num_teams (8 + i); + omp_set_teams_thread_limit (3 + i); + if (omp_get_max_teams () != 8 + i + || omp_get_teams_thread_limit () != 3 + i) + abort (); + } + + /* omp_set_num_teams above set the value of nteams-var ICV on device 'i', + which has scope 'device' and should be avaible in subsequent target + regions. */ + #pragma omp target device (i) + if (omp_get_max_teams () != 8 + i + || omp_get_teams_thread_limit () != 3 + i) + abort (); + + #pragma omp target device (i) + #pragma omp teams + if (omp_get_num_teams () > 8 + i + || omp_get_team_num () >= 8 + i) + abort (); + + /* NUM_TEAMS clause has priority over previously set NUM_TEAMS value. */ + #pragma omp target device (i) + #pragma omp teams num_teams(5 + i) + if (omp_get_num_teams () > 5 + i + || omp_get_team_num () >= 5 + i) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/icv-5.f90 b/libgomp/testsuite/libgomp.fortran/icv-5.f90 new file mode 100644 index 0000000..05a35fa --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/icv-5.f90 @@ -0,0 +1,226 @@ +! { dg-set-target-env-var OMP_NUM_TEAMS_ALL "3" } +! { dg-set-target-env-var OMP_NUM_TEAMS_DEV "4" } +! { dg-set-target-env-var OMP_NUM_TEAMS "5" } +! { dg-set-target-env-var OMP_NUM_TEAMS_DEV_0 "6" } +! { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1 "7" } +! { dg-set-target-env-var OMP_NUM_TEAMS_DEV_2 "8" } +! { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "2" } +! { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV "3" } +! { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT "4" } +! { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_0 "5" } +! { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_1 "6" } +! { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV_2 "7" } + +use omp_lib +implicit none (type, external) + integer :: num_devices, i, large_num_teams, large_threads_limit + logical :: err + + if (omp_get_num_devices () > 3) then + num_devices = 3 + else + num_devices = omp_get_num_devices () + end if + + do i=0,num_devices-1 + + ! Testing NUM_TEAMS. + if (env_is_set_dev ("OMP_NUM_TEAMS_DEV_", i, 6 + i)) then + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_max_teams () /= 6 + i) err = .true. + !$omp end target + if (err) stop 1 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + if (omp_get_num_teams () > 6 + i .or. omp_get_team_num () >= 6 + i) & + err = .true. + !$omp end teams + !$omp end target + if (err) stop 2 + + err = .false. + !$omp target device(i) map(tofrom: err) + call omp_set_num_teams (5 + i) + if (omp_get_max_teams () /= 5 + i) err = .true. + !$omp end target + if (err) stop 3 + + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_max_teams () /= 5 + i) err = .true. + !$omp end target + if (err) stop 4 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + if (omp_get_num_teams () > 5 + i .or. omp_get_team_num () >= 5 + i) & + err = .true. + !$omp end teams + !$omp end target + if (err) stop 5 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams num_teams(6 + i) + if (omp_get_num_teams () > 6 + i .or. omp_get_team_num () >= 6 + i) & + err = .true. + !$omp end teams + !$omp end target + if (err) stop 6 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams num_teams(4 + i) + if (omp_get_num_teams () > 4 + i .or. omp_get_team_num () >= 4 + i) & + err = .true. + !$omp end teams + !$omp end target + if (err) stop 7 + + large_num_teams = 66000 + err = .false. + !$omp target device(i) map(tofrom: err) + call omp_set_num_teams (large_num_teams + i) + if (omp_get_max_teams () /= large_num_teams + i) err = .true. + !$omp end target + if (err) stop 8 + + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_max_teams () /= large_num_teams + i) err = .true. + !$omp end target + if (err) stop 9 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + if (omp_get_num_teams () > large_num_teams + i & + .or. omp_get_team_num () >= large_num_teams + i) err = .true. + !$omp end teams + !$omp end target + if (err) stop 10 + end if + + ! Testing TEAMS-THREAD-LIMIT + if (env_is_set_dev ("OMP_TEAMS_THREAD_LIMIT_DEV_", i, 5 + i)) then + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_teams_thread_limit () /= 5 + i) err = .true. + !$omp end target + if (err) stop 11 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + !$omp parallel + if (omp_get_thread_limit () > 5 + i .or. omp_get_thread_num () >= 5 + i) & + err = .true. + !$omp end parallel + !$omp end teams + !$omp end target + if (err) stop 12 + + err = .false. + !$omp target device(i) map(tofrom: err) + call omp_set_teams_thread_limit (4 + i) + if (omp_get_teams_thread_limit () /= 4 + i) err = .true. + !$omp end target + if (err) stop 13 + + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_teams_thread_limit () /= 4 + i) err = .true. + !$omp end target + if (err) stop 14 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + !$omp parallel + if (omp_get_thread_limit () > 4 + i .or. omp_get_thread_num () >= 4 + i) & + err = .true. + !$omp end parallel + !$omp end teams + !$omp end target + if (err) stop 15 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams thread_limit(5 + i) + !$omp parallel + if (omp_get_thread_limit () > 5 + i .or. omp_get_thread_num () >= 5 + i) & + err = .true. + !$omp end parallel + !$omp end teams + !$omp end target + if (err) stop 16 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams thread_limit(3 + i) + !$omp parallel + if (omp_get_thread_limit () > 3 + i .or. omp_get_thread_num () >= 3 + i) & + err = .true. + !$omp end parallel + !$omp end teams + !$omp end target + if (err) stop 17 + + large_threads_limit = 67000 + err = .false. + !$omp target device(i) map(tofrom: err) + call omp_set_teams_thread_limit (large_threads_limit + i) + if (omp_get_teams_thread_limit () /= large_threads_limit + i) err = .true. + !$omp end target + if (err) stop 18 + + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_teams_thread_limit () /= large_threads_limit + i) err = .true. + !$omp end target + if (err) stop 19 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + !$omp parallel + if (omp_get_thread_limit () > large_threads_limit + i & + .or. omp_get_thread_num () >= large_threads_limit + i) err = .true. + !$omp end parallel + !$omp end teams + !$omp end target + if (err) stop 20 + end if + + end do + +contains + logical function env_is_set (name, val) + character(len=*) :: name, val + character(len=40) :: val2 + integer :: stat + call get_environment_variable (name, val2, status=stat) + if (stat == 0) then + if (val == val2) then + env_is_set = .true. + return + end if + else if (stat /= 1) then + error stop 30 + endif + env_is_set = .false. + end + logical function env_is_set_dev (name, dev_num, val) + character(len=*) :: name + integer :: dev_num, val + character(len=64) :: dev_num_str, env_var, val_str + dev_num_str = ADJUSTL(dev_num_str) + env_var = name // dev_num_str + val_str = ADJUSTL(val_str) + env_is_set_dev = env_is_set (TRIM(env_var), TRIM(val_str)) + end +end diff --git a/libgomp/testsuite/libgomp.fortran/icv-6.f90 b/libgomp/testsuite/libgomp.fortran/icv-6.f90 new file mode 100644 index 0000000..c8e6a0d --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/icv-6.f90 @@ -0,0 +1,140 @@ +! { dg-set-target-env-var OMP_NUM_TEAMS_ALL "3" } +! { dg-set-target-env-var OMP_NUM_TEAMS_DEV "4" } +! { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_ALL "2" } +! { dg-set-target-env-var OMP_TEAMS_THREAD_LIMIT_DEV "3" } + +! This test considers the hierarchical usage of ICVs on the device, i.e. if +! e.g. OMP_NUM_TEAMS_DEV_ is not configured, then the value of +! OMP_NUM_TEAMS_DEV should be used for the targets. + +use omp_lib +implicit none (type, external) + integer :: num_devices, i, stat, tmp + logical :: err + character(len=40) :: val + + ! The following environment variables should not be set. + call get_environment_variable ("OMP_NUM_TEAMS_DEV_0", val, status=stat) + if (stat /= 1) return + call get_environment_variable ("OMP_NUM_TEAMS_DEV_1", val, status=stat) + if (stat /= 1) return + call get_environment_variable ("OMP_NUM_TEAMS_DEV_2", val, status=stat) + if (stat /= 1) return + call get_environment_variable ("OMP_TEAMS_THREAD_LIMIT_DEV_0", val, status=stat) + if (stat /= 1) return + call get_environment_variable ("OMP_TEAMS_THREAD_LIMIT_DEV_1", val, status=stat) + if (stat /= 1) return + call get_environment_variable ("OMP_TEAMS_THREAD_LIMIT_DEV_2", val, status=stat) + if (stat /= 1) return + + if (omp_get_num_devices () > 3) then + num_devices = 3 + else + num_devices = omp_get_num_devices () + end if + + do i=0,num_devices-1 + + ! Testing NUM_TEAMS. + if (env_is_set ("OMP_NUM_TEAMS_DEV", "4")) then + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_max_teams () /= 4) err = .true. + !$omp end target + if (err) stop 1 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + if (omp_get_num_teams () > 4 .or. omp_get_team_num () >= 4) & + err = .true. + !$omp end teams + !$omp end target + if (err) stop 2 + + err = .false. + !$omp target device(i) map(tofrom: err) + call omp_set_num_teams (3 + i) + if (omp_get_max_teams () /= 3 + i) err = .true. + !$omp end target + if (err) stop 3 + + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_max_teams () /= 3 + i) err = .true. + !$omp end target + if (err) stop 4 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + if (omp_get_num_teams () > 3 + i .or. omp_get_team_num () >= 3 + i) & + err = .true. + !$omp end teams + !$omp end target + if (err) stop 5 + end if + + ! Testing TEAMS-THREAD-LIMIT + if (env_is_set ("OMP_TEAMS_THREAD_LIMIT_DEV", "3")) then + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_teams_thread_limit () /= 3) err = .true. + !$omp end target + if (err) stop 6 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + !$omp parallel + if (omp_get_thread_limit () > 3 .or. omp_get_thread_num () >= 3) & + err = .true. + !$omp end parallel + !$omp end teams + !$omp end target + if (err) stop 7 + + err = .false. + !$omp target device(i) map(tofrom: err) + call omp_set_teams_thread_limit (2 + i) + if (omp_get_teams_thread_limit () /= 2 + i) err = .true. + !$omp end target + if (err) stop 8 + + err = .false. + !$omp target device(i) map(tofrom: err) + if (omp_get_teams_thread_limit () /= 2 + i) err = .true. + !$omp end target + if (err) stop 9 + + err = .false. + !$omp target device(i) map(tofrom: err) + !$omp teams + !$omp parallel + if (omp_get_thread_limit () > 2 + i .or. omp_get_thread_num () >= 2 + i) & + err = .true. + !$omp end parallel + !$omp end teams + !$omp end target + if (err) stop 10 + end if + + end do + +contains + logical function env_is_set (name, val) + character(len=*) :: name, val + character(len=40) :: val2 + integer :: stat + call get_environment_variable (name, val2, status=stat) + if (stat == 0) then + if (val == val2) then + env_is_set = .true. + return + end if + else if (stat /= 1) then + error stop 10 + endif + env_is_set = .false. + end +end -- cgit v1.1