Hi!
On 2021-08-20T15:54:34+0200, I wrote:
> On 2021-08-20T15:21:12+0200, Jakub Jelinek <ja...@redhat.com> wrote:
>> On Fri, Aug 20, 2021 at 03:11:45PM +0200, Thomas Schwinge wrote:
>>> > --- libgomp/error.c.jj 2021-08-19 12:53:44.693106618 +0200
>>> > +++ libgomp/error.c 2021-08-19 17:58:55.633203432 +0200
>>>
>>> > +void
>>> > +GOMP_warning (const char *msg, size_t msglen)
>>> > +{
>>> > + if (msg && msglen == (size_t) -1)
>>> > + gomp_error ("error directive encountered: %s", msg);
>>> > + else if (msg)
>>> > + {
>>> > + fputs ("\nlibgomp: error directive encountered: ", stderr);
>>> > + fwrite (msg, 1, msglen, stderr);
>>> > + fputc ('\n', stderr);
>>> > + }
>>> > + else
>>> > + gomp_error ("error directive encountered");
>>> > +}
>>> > +
>>> > +void
>>> > +GOMP_error (const char *msg, size_t msglen)
>>> > +{
>>> > + if (msg && msglen == (size_t) -1)
>>> > + gomp_fatal ("fatal error: error directive encountered: %s", msg);
>>> > + else if (msg)
>>> > + {
>>> > + fputs ("\nlibgomp: fatal error: error directive encountered: ",
>>> > stderr);
>>> > + fwrite (msg, 1, msglen, stderr);
>>> > + fputc ('\n', stderr);
>>> > + exit (EXIT_FAILURE);
>>> > + }
>>> > + else
>>> > + gomp_fatal ("fatal error: error directive encountered");
>>> > +}
>>>
>>> At least for nvptx offloading, and at least given the newlib sources I'm
>>> using, the 'fputs'/'fwrite' calls here drag in 'isatty', which isn't
>>> provided by my nvptx newlib at present, so we get, for example:
>>
>> fputs/fputc/vfprintf/exit/stderr have been in use by error.c already before,
>> so this must be the fwrite call.
>
> ACK.
>
>> The above is for Fortran which doesn't have zero terminated strings.
>> Initially I wanted to use instead ... encountered: %.*s", (int) msglen,
>> stderr);
>> which doesn't handle > 2GB messages, but with offloading who cares, nobody
>> sane would be trying to print > 2GB messages from offloading regions.
>
> (... likewise from the host...) ;-)
>
>> The question is if it should be achieved through copy of error.c in
>> config/nvptx/, or just include_next there with say fwrite redefined as a
>> macro that does fprintf ("%.*s", (int) msglen, msg, file)?
>
> (Right, that was also my plan.)
>
> | Ah, I just re-discovered 'libgomp/config/nvptx/error.c' -- I'll cook
> | something up.
>
> So, guess what this newlib 'printf ("%.*s", [...]);' prints?
> Yes: literal '%.*s'... Next try: a 'fputc' loop?
Did that; "works". But actually, I think that's good enough for the
intended purpose: there's not much point in optimizing the OpenMP 'error'
directive as long as we still have more than enough real
correctness/performance tasks to be worked on.
Tobias suggested using 'fputc_unlocked', "avoiding repreated locks and
locking for a single stderr char is also pointless", but it's not clear
to me if that's safe to do given that a ton of threads may be hammering
on this in parallel; it's not clear to me if there isn't any
newlib-internal state that needs to be accessed in a serialized way (even
if no actual 'FILE *stream' is involved here)?
>>> permissible to use the 'error' directive also inside 'target' regions, as
>>> far as I can tell?
>>
>> !$omp error at(execution) message('whatever')
>> can be used in offloading regions.
(Also should add test cases for OpenMP 'error' with 'at (execution)' from
deep inside parallelized loop nests, etc., offloaded and non-offloaded?)
> Yes, generally works, but at least for Fortran, 'severity (fatal)' seems
> to cause a hang, so another thing to be looked into...
We thus additionally acquired in 'libgomp/config/nvptx/error.c':
+/* 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 ()
... which is another thing to be resolved incrementally. (Plus adding
corresponding test cases for OpenMP 'error' with 'at (execution)' and
'severity (fatal)' inside OpenMP 'target'.)
Is the attached "Make the OpenMP 'error' directive work for nvptx
offloading" OK to push for now?
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
>From 0762945a17c3ff1a0268edc76f87c0063714a0fc Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Fri, 20 Aug 2021 15:12:56 +0200
Subject: [PATCH] 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 | 32 +++++++++++++++++--
.../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)" }
--
2.25.1
