Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c'

These test cases being 'dg-skip-if [...] { ! amdgcn-*-* }' meant they were
only ever considered for 'amdgcn-*-*' target -- that is, GCN *target*, not
GCN *offload target*, what is intended to be tested here, but for which they
thus always were UNSUPPORTED.  Use the same style of 'dg-[...]' directives
as for the nvptx offloading test cases, 'libgomp.c/declare-variant-3*.c'.

Fix-up for commit 1fd508744e
"amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors".

	libgomp/
	* testsuite/libgomp.c/declare-variant-4-fiji.c: Adjust.
	* testsuite/libgomp.c/declare-variant-4-gfx803.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx900.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx906.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx908.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4-gfx90a.c: Likewise.
	* testsuite/libgomp.c/declare-variant-4.h: Likewise.
	* testsuite/libgomp.c/declare-variant-4.c: New.
This commit is contained in:
Thomas Schwinge 2023-11-21 16:36:48 +01:00
parent 95e6e32a85
commit aae57a9e19
8 changed files with 31 additions and 12 deletions

View file

@ -1,5 +1,6 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
/* { dg-do link { target { offload_target_amdgcn } } } */
/* { dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-additional-options -foffload=-march=fiji } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#define USE_FIJI_FOR_GFX803

View file

@ -1,5 +1,6 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "fiji/gfx803 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=fiji" } } */
/* { dg-do link { target { offload_target_amdgcn } } } */
/* { dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-additional-options -foffload=-march=fiji } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"

View file

@ -1,5 +1,6 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "gfx900 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx900" } } */
/* { dg-do link { target { offload_target_amdgcn } } } */
/* { dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-additional-options -foffload=-march=gfx900 } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"

View file

@ -1,5 +1,6 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "gfx906 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx906" } } */
/* { dg-do link { target { offload_target_amdgcn } } } */
/* { dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-additional-options -foffload=-march=gfx906 } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"

View file

@ -1,5 +1,6 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "gfx908 only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx908" } } */
/* { dg-do link { target { offload_target_amdgcn } } } */
/* { dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-additional-options -foffload=-march=gfx908 } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"

View file

@ -1,5 +1,6 @@
/* { dg-do run { target { offload_target_amdgcn } } } */
/* { dg-skip-if "gfx90a only" { ! amdgcn-*-* } { "*" } { "-foffload=-march=gfx90a" } } */
/* { dg-do link { target { offload_target_amdgcn } } } */
/* { dg-additional-options -foffload=amdgcn-amdhsa } */
/* { dg-additional-options -foffload=-march=gfx90a } */
/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */
#include "declare-variant-4.h"

View file

@ -0,0 +1,8 @@
/* { dg-additional-options -DOFFLOAD_DEVICE_GCN { target offload_device_gcn } } */
/* { dg-additional-options {-fdump-tree-optimized -foffload-options=-fdump-tree-optimized} } */
#include "declare-variant-4.h"
/* { dg-final { scan-tree-dump "= f \\(\\);" "optimized" } }
{ dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx\[^ \]+ \\(\\);" "optimized" { target offload_target_amdgcn } } }
{ dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f \\(\\);" "optimized" { target offload_target_nvptx } } } */

View file

@ -61,10 +61,15 @@ main (void)
#pragma omp target map(from : v)
v = f ();
#ifdef OFFLOAD_DEVICE_GCN
if (v == 0)
__builtin_abort ();
__builtin_printf ("AMDGCN accelerator: gfx%x\n", v);
#else
if (v != 0)
__builtin_abort ();
#endif
return 0;
}