From aae57a9e19ba5d2016fa8a88db6d7ff15a40ca35 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 21 Nov 2023 16:36:48 +0100 Subject: [PATCH] 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 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d "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. --- libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c | 5 +++-- libgomp/testsuite/libgomp.c/declare-variant-4.c | 8 ++++++++ libgomp/testsuite/libgomp.c/declare-variant-4.h | 5 +++++ 8 files changed, 31 insertions(+), 12 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c/declare-variant-4.c diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c index 8a4e0f4728c..a138fb092f8 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-fiji.c @@ -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 diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c index 050d7c9dd79..03dffddac49 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx803.c @@ -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" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c index 2eeb4a248c1..f3f5244c7bc 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c @@ -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" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c index 73a670dcc2a..ac43388a59f 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c @@ -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" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c index ead330f9f2c..f60741f202f 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c @@ -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" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c index a9b2d62a49d..832d174e4a5 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c @@ -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" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.c b/libgomp/testsuite/libgomp.c/declare-variant-4.c new file mode 100644 index 00000000000..3c8802a6081 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4.c @@ -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 } } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h index 47517b75ee7..a70352430c2 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4.h +++ b/libgomp/testsuite/libgomp.c/declare-variant-4.h @@ -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; } -- 2.34.1