[gcc r12-3058] Make the OpenMP 'error' directive work for nvptx offloading

[Thomas Schwinge <tschwinge@gcc.gnu.org>]

https://gcc.gnu.org/g:a5416bf369419428fb139c432bcd88f6f8ee4910

commit r12-3058-ga5416bf369419428fb139c432bcd88f6f8ee4910
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Fri Aug 20 15:12:56 2021 +0200

    Make the OpenMP 'error' directive work for nvptx offloading
    
    ... and add a minimum amount of offloading testing.
    
    (Leaving aside that 'fwrite' to 'stderr' probably wouldn't work anyway) the
    'fwrite' calls in 'libgomp/error.c:GOMP_warning', 'libgomp/error.c:GOMP_error'
    drag in 'isatty', which isn't provided by my nvptx newlib build at present, so
    we get, for example:
    
        [...]
        FAIL: libgomp.c/../libgomp.c-c++-common/declare_target-1.c (test for excess errors)
        Excess errors:
        unresolved symbol isatty
        mkoffload: fatal error: [...]/build-gcc/./gcc/x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status
        [...]
    
    ..., and many more.
    
    Fix up for recent commit 0d973c0a0d90a0a302e7eda1a4d9709be3c5b102
    "openmp: Implement the error directive".
    
            libgomp/
            * config/nvptx/error.c (fwrite, exit): Override, too.
            * testsuite/libgomp.c-c++-common/error-1.c: Add a minimum amount
            of offloading testing.
            * testsuite/libgomp.fortran/error-1.f90: Likewise.

Diff:
---
 libgomp/config/nvptx/error.c                     | 32 +++++++++++++++++++++---
 libgomp/testsuite/libgomp.c-c++-common/error-1.c | 10 ++++++++
 libgomp/testsuite/libgomp.fortran/error-1.f90    |  9 +++++++
 3 files changed, 48 insertions(+), 3 deletions(-)

diff --git a/libgomp/config/nvptx/error.c b/libgomp/config/nvptx/error.c
index dfa75da354f..c55791e34b4 100644
--- a/libgomp/config/nvptx/error.c
+++ b/libgomp/config/nvptx/error.c
@@ -31,12 +31,38 @@
 #include <stdio.h>
 #include <stdlib.h>
 
-#undef vfprintf
-#undef fputs
-#undef fputc
 
+/* No 'FILE *stream's, just basic 'vprintf' etc.  */
+
+#undef vfprintf
 #define vfprintf(stream, fmt, list) vprintf (fmt, list)
+
+#undef fputs
 #define fputs(s, stream) printf ("%s", s)
+
+#undef fputc
 #define fputc(c, stream) printf ("%c", c)
 
+#undef fwrite
+#if 0
+# define fwrite(ptr, size, nmemb, stream) \
+  printf ("%.*s", (int) (size * nmemb), (int) (size * nmemb), ptr)
+/* ... prints literal '%.*s'.  */
+#else
+# define fwrite(ptr, size, nmemb, stream) \
+  do { \
+    /* Yuck!  */ \
+    for (size_t i = 0; i < size * nmemb; ++i) \
+      printf ("%c", ptr[i]); \
+  } while (0)
+#endif
+
+
+/* The 'exit (EXIT_FAILURE);' of an Fortran (only, huh?) OpenMP 'error'
+   directive with 'severity (fatal)' causes a hang, so 'abort' instead of
+   'exit'.  */
+#undef exit
+#define exit(status) abort ()
+
+
 #include "../../error.c"
diff --git a/libgomp/testsuite/libgomp.c-c++-common/error-1.c b/libgomp/testsuite/libgomp.c-c++-common/error-1.c
index 5f454c1adaa..04c0519bf63 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/error-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/error-1.c
@@ -34,11 +34,20 @@ foo (int i, int x)
 int
 main ()
 {
+  /* Initialize offloading early, so that any output this may produce doesn't
+     disturb the 'dg-output' scanning below.  */
+  #pragma omp target
+  ;
+
   if (foo (5, 0) != 13 || foo (6, 1) != 17)
     abort ();
   #pragma omp error at (execution) severity (warning)
   const char *msg = "my message" + 2;
   #pragma omp error at (execution) severity (warning) message (msg + 1)
+  #pragma omp target
+  {
+    #pragma omp error at (execution) severity (warning) message ("hello from a distance")
+  }
   #pragma omp error at (execution) severity (fatal) message (msg - 2)
   #pragma omp error at (execution) severity (warning) message ("foobar")
   return 0;
@@ -46,4 +55,5 @@ main ()
 
 /* { dg-output "libgomp: error directive encountered(\n|\r|\n\r)(\n|\r|\n\r)" } */
 /* { dg-output "libgomp: error directive encountered: message(\n|\r|\n\r)(\n|\r|\n\r)" } */
+/* { dg-output "libgomp: error directive encountered: hello from a distance(\n|\r|\n\r)(\n|\r|\n\r)" } */
 /* { dg-output "libgomp: fatal error: error directive encountered: my message" } */
diff --git a/libgomp/testsuite/libgomp.fortran/error-1.f90 b/libgomp/testsuite/libgomp.fortran/error-1.f90
index 92c246cfcaf..7c497fd002e 100644
--- a/libgomp/testsuite/libgomp.fortran/error-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/error-1.f90
@@ -37,6 +37,11 @@ program main
   character(len=13) :: msg
   character(len=:), allocatable :: msg2, msg3
 
+  ! Initialize offloading early, so that any output this may produce doesn't
+  ! disturb the 'dg-output' scanning below.
+  !$omp target
+  !$omp end target
+
   msg = "my message"
   if (foo (5, 0) /= 15 .or. foo (7, 1) /= 16) &
     stop 1
@@ -47,6 +52,9 @@ program main
   !$omp error at (execution) severity (warning)
   !$omp error at (execution) severity (warning) message(trim(msg(4:)))
   !$omp error at (execution) severity (warning) message ("Farewell")
+  !$omp target
+  !$omp error at (execution) severity (warning) message ("ffrom a distanceee"(2:16))
+  !$omp end target
   !$omp error at (execution) severity (warning) message (msg2)
   !$omp error at (execution) severity (warning) message (msg(4:6))
   !$omp error at (execution) severity (fatal) message (msg)
@@ -73,6 +81,7 @@ end
 ! { dg-output "libgomp: error directive encountered(\n|\r|\n\r)(\n|\r|\n\r)" }
 ! { dg-output "libgomp: error directive encountered: message(\n|\r|\n\r)(\n|\r|\n\r)" }
 ! { dg-output "libgomp: error directive encountered: Farewell(\n|\r|\n\r)(\n|\r|\n\r)" }
+! { dg-output "libgomp: error directive encountered: from a distance(\n|\r|\n\r)(\n|\r|\n\r)" }
 ! { dg-output "libgomp: error directive encountered: Hello World(\n|\r|\n\r)(\n|\r|\n\r)" }
 ! { dg-output "libgomp: error directive encountered: mes(\n|\r|\n\r)(\n|\r|\n\r)" }
 ! { dg-output "libgomp: fatal error: error directive encountered: my message   (\n|\r|\n\r)" }