81476bc4f4
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.
203 lines
5.6 KiB
C
203 lines
5.6 KiB
C
/* { dg-do run } */
|
|
/* { 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 <omp.h>
|
|
#include <stdlib.h>
|
|
#include <unistd.h>
|
|
|
|
int
|
|
main ()
|
|
{
|
|
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 () != 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;
|
|
}
|