mirrors/gcc
0
0
Fork 0
mirror of https://gcc.gnu.org/git/gcc.git synced 2024-11-27 13:54:19 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors

Browse Source 
Add support for gfx803 as an alias for fiji.
Add test cases for all supported 'isa' values.

gcc/ChangeLog:

	* config/gcn/gcn.cc (gcn_omp_device_kind_arch_isa): Add gfx803.
	* config/gcn/t-omp-device: Add gfx803.

libgomp/ChangeLog:

	* testsuite/libgomp.c/declare-variant-4-fiji.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx803.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx900.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx906.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx908.c: New test.
	* testsuite/libgomp.c/declare-variant-4-gfx90a.c: New test.
	* testsuite/libgomp.c/declare-variant-4.h: New header file.
This commit is contained in:
Paul-Antoine Arras 2022-11-29 16:22:07 +01:00
parent a1b5cdf381
commit 1fd508744e
9 changed files with 108 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
gcc/config/gcn/gcn.cc
View File

@ -2985,7 +2985,7 @@ gcn_omp_device_kind_arch_isa (enum omp_device_kind_arch_isa trait,
case omp_device_arch:
return strcmp (name, "amdgcn") == 0 || strcmp (name, "gcn") == 0;
case omp_device_isa:
if (strcmp (name, "fiji") == 0)
if (strcmp (name, "fiji") == 0 || strcmp (name, "gfx803") == 0)
return gcn_arch == PROCESSOR_FIJI;
if (strcmp (name, "gfx900") == 0)
return gcn_arch == PROCESSOR_VEGA10;

2
gcc/config/gcn/t-omp-device
View File

@ -1,4 +1,4 @@
omp-device-properties-gcn: $(srcdir)/config/gcn/gcn.cc
echo kind: gpu > $@
echo arch: amdgcn gcn >> $@
echo isa: fiji gfx900 gfx906 gfx908 gfx90a >> $@
echo isa: fiji gfx803 gfx900 gfx906 gfx908 gfx90a >> $@

8
libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c Normal file
View File

@ -0,0 +1,8 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#define USE_FIJI_FOR_GFX803
#include "declare-variant-4.h"
/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */

7
libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c Normal file
View File

@ -0,0 +1,7 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
/* { dg-final { scan-offload-tree-dump "= gfx803 \\(\\);" "optimized" } } */

7
libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c Normal file
View File

@ -0,0 +1,7 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
/* { dg-final { scan-offload-tree-dump "= gfx900 \\(\\);" "optimized" } } */

7
libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c Normal file
View File

@ -0,0 +1,7 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
/* { dg-final { scan-offload-tree-dump "= gfx906 \\(\\);" "optimized" } } */

7
libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c Normal file
View File

@ -0,0 +1,7 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
/* { dg-final { scan-offload-tree-dump "= gfx908 \\(\\);" "optimized" } } */

7
libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c Normal file
View File

@ -0,0 +1,7 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"
/* { dg-final { scan-offload-tree-dump "= gfx90a \\(\\);" "optimized" } } */

63
libgomp/testsuite/libgomp.c/declare-variant-4.h Normal file
View File

@ -0,0 +1,63 @@
#pragma omp declare target
int
gfx803 (void)
{
return 0x803;
}
int
gfx900 (void)
{
return 0x900;
}
int
gfx906 (void)
{
return 0x906;
}
int
gfx908 (void)
{
return 0x908;
}
int
gfx90a (void)
{
return 0x90a;
}
#ifdef USE_FIJI_FOR_GFX803
#pragma omp declare variant(gfx803) match(device = {isa("fiji")})
#else
#pragma omp declare variant(gfx803) match(device = {isa("gfx803")})
#endif
#pragma omp declare variant(gfx900) match(device = {isa("gfx900")})
#pragma omp declare variant(gfx906) match(device = {isa("gfx906")})
#pragma omp declare variant(gfx908) match(device = {isa("gfx908")})
#pragma omp declare variant(gfx90a) match(device = {isa("gfx90a")})
int
f (void)
{
return 0;
}
#pragma omp end declare target
int
main (void)
{
int v = 0;
#pragma omp target map(from : v)
v = f ();
if (v == 0)
__builtin_abort ();
__builtin_printf ("AMDGCN accelerator: gfx%x\n", v);
return 0;
}