procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

openmp: Fix up handling of kind(host) and kind(nohost) in ACCEL_COMPILERs [PR103384]

Browse source 
As the testcase shows, we weren't handling kind(host) and kind(nohost) properly
in the ACCEL_COMPILERs, the code written in there is valid for the host
compiler only, where if we are maybe offloaded, we defer resolution after IPA,
otherwise return 0 for kind(nohost) and accept it for kind(host).  Note,
omp_maybe_offloaded is false after IPA.  If ACCEL_COMPILER is defined, it is
the other way around, but also we know we are after IPA.

2021-11-24  Jakub Jelinek  <jakub@redhat.com>

	PR middle-end/103384
gcc/
	* omp-general.c (omp_context_selector_matches): For ACCEL_COMPILER,
	return 0 for kind(host) and continue for kind(nohost).
libgomp/
	* testsuite/libgomp.c/declare-variant-2.c: New test.
This commit is contained in:
Jakub Jelinek 2021-11-24 10:30:32 +01:00
parent 709716b9f4
commit 5bca26742c
2 changed files with 51 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
gcc/omp-general.c
View file

@ -1487,16 +1487,22 @@ omp_context_selector_matches (tree ctx)
continue;
if (!strcmp (prop, "host"))
{
#ifdef ACCEL_COMPILER
return 0;
#else
if (omp_maybe_offloaded ())
ret = -1;
continue;
#endif
}
if (!strcmp (prop, "nohost"))
{
#ifndef ACCEL_COMPILER
if (omp_maybe_offloaded ())
ret = -1;
else
return 0;
#endif
continue;
}
int r = 0;

45
libgomp/testsuite/libgomp.c/declare-variant-2.c Normal file
View file

@ -0,0 +1,45 @@
/* { dg-do run } */
#include <omp.h>
#include <stdlib.h>
void
foo_host (void)
{
if (!omp_is_initial_device ())
abort ();
}
#pragma omp declare variant (foo_host) match (device={kind(host)})
void
foo (void)
{
if (omp_is_initial_device ())
abort ();
}
void
bar_nohost (void)
{
if (omp_is_initial_device ())
abort ();
}
#pragma omp declare variant (bar_nohost) match (device={kind(nohost)})
void
bar (void)
{
if (!omp_is_initial_device ())
abort ();
}
int
main ()
{
#pragma omp target
{
foo ();
bar ();
}
return 0;
}