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 <a...@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
