[libgomp] Add tests for print from offload target

Message ID 763f86ad-c8b0-de34-e752-f591a87ea851@codesourcery.com
State New
Headers show
Series
  • [libgomp] Add tests for print from offload target
Related show

Commit Message

Andrew Stubbs Nov. 14, 2019, 4:47 p.m.
Hi,

This patch adds new libgomp tests to ensure that C "printf" and Fortran 
"write" work correctly within offload kernels. Both should work for 
amdgcn, but nvptx uses the libgfortran "minimal" mode which lacks 
"write" support.

Obviously, printing from offload kernels is not recommended in 
production, but can be useful in development.

OK to commit?

Thanks

Andrew

Comments

Jakub Jelinek Nov. 14, 2019, 5:05 p.m. | #1
On Thu, Nov 14, 2019 at 04:47:49PM +0000, Andrew Stubbs wrote:
> This patch adds new libgomp tests to ensure that C "printf" and Fortran

> "write" work correctly within offload kernels. Both should work for amdgcn,

> but nvptx uses the libgfortran "minimal" mode which lacks "write" support.


So, do those *.f90 testcases now FAIL with nvptx offloading?
If yes, perhaps there should be effective target check that it is not nvptx
offloading.
Once the declare variant support is finished, at least in OpenMP it could be
handled through that, but Fortran support for that will definitely not make
GCC 10.

	Jakub
Tobias Burnus Nov. 14, 2019, 5:18 p.m. | #2
On 11/14/19 5:47 PM, Andrew Stubbs wrote:
> This patch adds new libgomp tests to ensure that C "printf" and 

> Fortran "write" work correctly within offload kernels. Both should 

> work for amdgcn, but nvptx uses the libgfortran "minimal" mode which 

> lacks "write" support.


Can't you add something like:

! { dg-do run { target { ! { openacc_nvidia_accel_selected } } } }
! For openacc_nvidia_accel_selected, there is no I/O support.

To avoid FAILs?

Cheers,

Tobias
Andrew Stubbs Nov. 14, 2019, 5:18 p.m. | #3
On 14/11/2019 17:05, Jakub Jelinek wrote:
> On Thu, Nov 14, 2019 at 04:47:49PM +0000, Andrew Stubbs wrote:

>> This patch adds new libgomp tests to ensure that C "printf" and Fortran

>> "write" work correctly within offload kernels. Both should work for amdgcn,

>> but nvptx uses the libgfortran "minimal" mode which lacks "write" support.

> 

> So, do those *.f90 testcases now FAIL with nvptx offloading?

> If yes, perhaps there should be effective target check that it is not nvptx

> offloading.

> Once the declare variant support is finished, at least in OpenMP it could be

> handled through that, but Fortran support for that will definitely not make

> GCC 10.


Oops, I forgot to regenerate the patch before posting it.

Here's the version with the nvptx xfails.

OK?

Andrew
Add tests for print from offload target.

2019-11-14  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* testsuite/libgomp.c/target-print-1.c: New file.
	* testsuite/libgomp.fortran/target-print-1.f90: New file.
	* testsuite/libgomp.oacc-c/print-1.c: New file.
	* testsuite/libgomp.oacc-fortran/print-1.f90: New file.

diff --git a/libgomp/testsuite/libgomp.c/target-print-1.c b/libgomp/testsuite/libgomp.c/target-print-1.c
new file mode 100644
index 00000000000..5857b875ced
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-print-1.c
@@ -0,0 +1,17 @@
+/* Ensure that printf on the offload device works.  */
+
+/* { dg-do run } */
+/* { dg-output "The answer is 42(\n|\r\n|\r)+" } */
+
+#include <stdio.h>
+
+int var = 42;
+
+int
+main ()
+{
+#pragma omp target
+    {
+      printf ("The answer is %d\n", var);
+    }
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-print-1.f90 b/libgomp/testsuite/libgomp.fortran/target-print-1.f90
new file mode 100644
index 00000000000..c71a0952483
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-print-1.f90
@@ -0,0 +1,15 @@
+! Ensure that printf on the offload device works.
+
+! { dg-do run }
+! { dg-output "The answer is 42(\n|\r\n|\r)+" }
+! { dg-xfail-if "no write for nvidia" { openacc_nvidia_accel_selected } }                                                                                                                                                                
+
+program main
+  implicit none
+  integer :: var = 42
+
+!$omp target 
+  write (0, '("The answer is ", I2)') var
+!$omp end target
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-c/print-1.c b/libgomp/testsuite/libgomp.oacc-c/print-1.c
new file mode 100644
index 00000000000..593885b5c2c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c/print-1.c
@@ -0,0 +1,17 @@
+/* Ensure that printf on the offload device works.  */
+
+/* { dg-do run } */
+/* { dg-output "The answer is 42(\n|\r\n|\r)+" } */
+
+#include <stdio.h>
+
+int var = 42;
+
+int
+main ()
+{
+#pragma acc parallel
+    {
+      printf ("The answer is %d\n", var);
+    }
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
new file mode 100644
index 00000000000..a83280d1edb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
@@ -0,0 +1,15 @@
+! Ensure that printf on the offload device works.
+
+! { dg-do run }
+! { dg-output "The answer is 42(\n|\r\n|\r)+" }
+! { dg-xfail-if "no write for nvidia" { openacc_nvidia_accel_selected } }                                                                                                                                                                
+
+program main
+  implicit none
+  integer :: var = 42
+
+!$acc parallel
+  write (0, '("The answer is ", I2)') var
+!$acc end parallel
+
+end program main
Jakub Jelinek Nov. 14, 2019, 5:22 p.m. | #4
On Thu, Nov 14, 2019 at 05:18:41PM +0000, Andrew Stubbs wrote:
> On 14/11/2019 17:05, Jakub Jelinek wrote:

> > On Thu, Nov 14, 2019 at 04:47:49PM +0000, Andrew Stubbs wrote:

> > > This patch adds new libgomp tests to ensure that C "printf" and Fortran

> > > "write" work correctly within offload kernels. Both should work for amdgcn,

> > > but nvptx uses the libgfortran "minimal" mode which lacks "write" support.

> > 

> > So, do those *.f90 testcases now FAIL with nvptx offloading?

> > If yes, perhaps there should be effective target check that it is not nvptx

> > offloading.

> > Once the declare variant support is finished, at least in OpenMP it could be

> > handled through that, but Fortran support for that will definitely not make

> > GCC 10.

> 

> Oops, I forgot to regenerate the patch before posting it.

> 

> Here's the version with the nvptx xfails.

> 

> OK?


Ok.

> Add tests for print from offload target.

> 

> 2019-11-14  Andrew Stubbs  <ams@codesourcery.com>

> 

> 	libgomp/

> 	* testsuite/libgomp.c/target-print-1.c: New file.

> 	* testsuite/libgomp.fortran/target-print-1.f90: New file.

> 	* testsuite/libgomp.oacc-c/print-1.c: New file.

> 	* testsuite/libgomp.oacc-fortran/print-1.f90: New file.


> +int

> +main ()

> +{

> +#pragma omp target

> +    {

> +      printf ("The answer is %d\n", var);

> +    }


Just a nit,
#pragma omp target
  printf ("The answer is %d\n", var);
would be valid too, but no need to change the testcase.

	Jakub
Thomas Schwinge Nov. 27, 2019, 5:54 p.m. | #5
Hi!

On 2019-11-14T18:22:39+0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Thu, Nov 14, 2019 at 05:18:41PM +0000, Andrew Stubbs wrote:

>> On 14/11/2019 17:05, Jakub Jelinek wrote:

>> > On Thu, Nov 14, 2019 at 04:47:49PM +0000, Andrew Stubbs wrote:

>> > > This patch adds new libgomp tests to ensure that C "printf" and Fortran

>> > > "write" work correctly within offload kernels.


Thanks.

>> > > Both should work for amdgcn,

>> > > but nvptx uses the libgfortran "minimal" mode which lacks "write" support.

>> > 

>> > So, do those *.f90 testcases now FAIL with nvptx offloading?


Note that for libgomp OpenMP testing that means FAIL as soon as nvptx
plus possibly any other offload targets are configured, as we will always
compile for all offload targets, as no specific '-foffload' gets passed
-- in contrast to the libgomp OpenACC testing.  (Changing that would be a
separate discussion.)

>> > If yes, perhaps there should be effective target check that it is not nvptx

>> > offloading.


>> Here's the version with the nvptx xfails.

>> 

>> OK?

>

> Ok.


... but that doesn't work as expected:

    {+ERROR: libgomp.fortran/target-print-1.f90   -O0 : can't read "openacc_device_type": no such variable for " dg-xfail-if 5 "no write for nvidia" { openacc_nvidia_accel_selected } "+}
    [Etc.]

..., and:

    {+XFAIL: libgomp.oacc-fortran/print-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none  -O0  (test for excess errors)+}
    {+UNRESOLVED: libgomp.oacc-fortran/print-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none  -O0  compilation failed to produce executable+}
    [Etc.]

To fix the ERRORs, I'm adding
'check_effective_target_offload_target_nvptx' to generally test whether
"compiling for offload target nvptx", and to fix the UNRESOLVEDs, I'm
adding separate "nvptx" variants, as I couldn't find a way to express the
'dg-*' directives in one combined file.  I want to actually see the nvptx
XFAILs, so that we'll notice XPASSes once/if nvptx switches away from the
minimal libgfortran.  (Changing that would be a separate discussion.)

Committed "Fix 'libgomp.fortran/target-print-1.f90',
'libgomp.oacc-fortran/print-1.f90' for offload target nvptx" to trunk in
r278779, see attached.


Grüße
 Thomas
From 21d18541ada7c3bdc70f92d9c79ced33a2f5f98a Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>

Date: Wed, 27 Nov 2019 17:50:55 +0000
Subject: [PATCH] Fix 'libgomp.fortran/target-print-1.f90',
 'libgomp.oacc-fortran/print-1.f90' for offload target nvptx

	libgomp/
	* testsuite/lib/libgomp.exp
	(check_effective_target_offload_target_nvptx): New proc.
	* testsuite/libgomp.fortran/target-print-1.f90: Use it with
	'dg-skip-if'.
	* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
	* testsuite/libgomp.fortran/target-print-1-nvptx.f90: New file.
	* testsuite/libgomp.oacc-fortran/print-1-nvptx.f90: Likewise.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@278779 138bc75d-0d04-0410-961f-82ee72b054a4
---
 libgomp/ChangeLog                             | 10 +++++++++
 libgomp/testsuite/lib/libgomp.exp             | 21 +++++++++++++++++++
 .../libgomp.fortran/target-print-1-nvptx.f90  | 11 ++++++++++
 .../libgomp.fortran/target-print-1.f90        |  6 ++++--
 .../libgomp.oacc-fortran/print-1-nvptx.f90    | 11 ++++++++++
 .../libgomp.oacc-fortran/print-1.f90          |  6 ++++--
 6 files changed, 61 insertions(+), 4 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.fortran/target-print-1-nvptx.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/print-1-nvptx.f90

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 65478302ba8..d8d76ba3ce8 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,13 @@
+2019-11-27  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/lib/libgomp.exp
+	(check_effective_target_offload_target_nvptx): New proc.
+	* testsuite/libgomp.fortran/target-print-1.f90: Use it with
+	'dg-skip-if'.
+	* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
+	* testsuite/libgomp.fortran/target-print-1-nvptx.f90: New file.
+	* testsuite/libgomp.oacc-fortran/print-1-nvptx.f90: Likewise.
+
 2019-11-21  Rainer Orth  <ro@CeBiTec.Uni-Bielefeld.DE>
 
 	* testsuite/libgomp.c/pr39591-1.c: Rename err to e.
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 74d032623c9..06e3186d966 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -336,6 +336,27 @@ proc offload_target_to_openacc_device_type { offload_target } {
     }
 }
 
+# Return 1 if compiling for offload target nvptx.
+proc check_effective_target_offload_target_nvptx { } {
+    # Consider all actual options, including the flags passed to
+    # 'gcc-dg-runtest', or 'gfortran-dg-runtest' (see the 'libgomp.*/*.exp'
+    # files; in particular, '-foffload', 'libgomp.oacc-*/*.exp'), which don't
+    # get passed on to 'check_effective_target_*' functions.  (Not caching the
+    # result due to that.)
+    set options [current_compiler_flags]
+    # Instead of inspecting command-line options, look what the compiler driver
+    # decides.  This is somewhat modelled after
+    # 'gcc/testsuite/lib/target-supports.exp:check_configured_with'.
+    set gcc_output [libgomp_target_compile "-v $options" "" "none" ""]
+    if [regexp "(?n)^OFFLOAD_TARGET_NAMES=(.*)" $gcc_output dummy offload_targets] {
+	verbose "compiling for offload targets: $offload_targets"
+	return [string match "*:nvptx*:*" ":$offload_targets:"]
+    }
+
+    verbose "not compiling for any offload targets"
+    return 0
+}
+
 # Return 1 if offload device is available.
 proc check_effective_target_offload_device { } {
     return [check_runtime_nocache offload_device_available_ {
diff --git a/libgomp/testsuite/libgomp.fortran/target-print-1-nvptx.f90 b/libgomp/testsuite/libgomp.fortran/target-print-1-nvptx.f90
new file mode 100644
index 00000000000..a89c9c33484
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-print-1-nvptx.f90
@@ -0,0 +1,11 @@
+! Ensure that write on the offload device works, nvptx offloading variant.
+
+! This doesn't compile: for nvptx offloading we're using a minimal libgfortran
+! configuration.
+! { dg-do link } ! ..., but still apply 'dg-do run' options.
+! { dg-xfail-if "minimal libgfortran" { offload_target_nvptx } }
+
+! Skip duplicated testing.
+! { dg-skip-if "separate file" { ! offload_target_nvptx } }
+
+include 'target-print-1.f90'
diff --git a/libgomp/testsuite/libgomp.fortran/target-print-1.f90 b/libgomp/testsuite/libgomp.fortran/target-print-1.f90
index c71a0952483..327bb22cb6d 100644
--- a/libgomp/testsuite/libgomp.fortran/target-print-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-print-1.f90
@@ -1,8 +1,10 @@
-! Ensure that printf on the offload device works.
+! Ensure that write on the offload device works.
 
 ! { dg-do run }
 ! { dg-output "The answer is 42(\n|\r\n|\r)+" }
-! { dg-xfail-if "no write for nvidia" { openacc_nvidia_accel_selected } }                                                                                                                                                                
+
+! Separate file 'target-print-1-nvptx.f90' for nvptx offloading.
+! { dg-skip-if "separate file" { offload_target_nvptx } }
 
 program main
   implicit none
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1-nvptx.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1-nvptx.f90
new file mode 100644
index 00000000000..866c8654355
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1-nvptx.f90
@@ -0,0 +1,11 @@
+! Ensure that write on the offload device works, nvptx offloading variant.
+
+! This doesn't compile: for nvptx offloading we're using a minimal libgfortran
+! configuration.
+! { dg-do link } ! ..., but still apply 'dg-do run' options.
+! { dg-xfail-if "minimal libgfortran" { offload_target_nvptx } }
+
+! Skip duplicated testing.
+! { dg-skip-if "separate file" { ! offload_target_nvptx } }
+
+include 'print-1.f90'
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
index a83280d1edb..7b7f73741fe 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
@@ -1,8 +1,10 @@
-! Ensure that printf on the offload device works.
+! Ensure that write on the offload device works.
 
 ! { dg-do run }
 ! { dg-output "The answer is 42(\n|\r\n|\r)+" }
-! { dg-xfail-if "no write for nvidia" { openacc_nvidia_accel_selected } }                                                                                                                                                                
+
+! Separate file 'print-1-nvptx.f90' for nvptx offloading.
+! { dg-skip-if "separate file" { offload_target_nvptx } }
 
 program main
   implicit none
-- 
2.17.1

Patch

Add tests for print from offload target.

2019-11-14  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* testsuite/libgomp.c/target-print-1.c: New file.
	* testsuite/libgomp.fortran/target-print-1.f90: New file.
	* testsuite/libgomp.oacc-c/print-1.c: New file.
	* testsuite/libgomp.oacc-fortran/print-1.f90: New file.

diff --git a/libgomp/testsuite/libgomp.c/target-print-1.c b/libgomp/testsuite/libgomp.c/target-print-1.c
new file mode 100644
index 00000000000..5857b875ced
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-print-1.c
@@ -0,0 +1,17 @@ 
+/* Ensure that printf on the offload device works.  */
+
+/* { dg-do run } */
+/* { dg-output "The answer is 42(\n|\r\n|\r)+" } */
+
+#include <stdio.h>
+
+int var = 42;
+
+int
+main ()
+{
+#pragma omp target
+    {
+      printf ("The answer is %d\n", var);
+    }
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-print-1.f90 b/libgomp/testsuite/libgomp.fortran/target-print-1.f90
new file mode 100644
index 00000000000..73ee09a2b79
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-print-1.f90
@@ -0,0 +1,14 @@ 
+! Ensure that printf on the offload device works.
+
+! { dg-do run }
+! { dg-output "The answer is 42(\n|\r\n|\r)+" }
+
+program main
+  implicit none
+  integer :: var = 42
+
+!$omp target 
+  write (0, '("The answer is ", I2)') var
+!$omp end target
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-c/print-1.c b/libgomp/testsuite/libgomp.oacc-c/print-1.c
new file mode 100644
index 00000000000..593885b5c2c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c/print-1.c
@@ -0,0 +1,17 @@ 
+/* Ensure that printf on the offload device works.  */
+
+/* { dg-do run } */
+/* { dg-output "The answer is 42(\n|\r\n|\r)+" } */
+
+#include <stdio.h>
+
+int var = 42;
+
+int
+main ()
+{
+#pragma acc parallel
+    {
+      printf ("The answer is %d\n", var);
+    }
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
new file mode 100644
index 00000000000..bef664df4fa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
@@ -0,0 +1,14 @@ 
+! Ensure that printf on the offload device works.
+
+! { dg-do run }
+! { dg-output "The answer is 42(\n|\r\n|\r)+" }
+
+program main
+  implicit none
+  integer :: var = 42
+
+!$acc parallel
+  write (0, '("The answer is ", I2)') var
+!$acc end parallel
+
+end program main