Hi Jakub!
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?
See attached "[WIP] Make the OpenMP 'error' directive work for nvptx
offloading".
>> Now, there are many ways of addressing this... The most simple one:
>> conditionalize these 'GOMP_warning'/'GOMP_error' definitions on
>> '#ifndef LIBGOMP_OFFLOADED_ONLY' is not possible here, because it's
>> 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.
Yes, generally works, but at least for Fortran, 'severity (fatal)' seems
to cause a hang, so another thing to be looked into...
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 d524726ff5f319658ef317afaf0077f43b8eccf8 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Fri, 20 Aug 2021 15:12:56 +0200
Subject: [PATCH] [WIP] Make the OpenMP 'error' directive work for nvptx
offloading
The 'fputs'/'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 | 2 ++
libgomp/error.c | 6 ++++++
libgomp/testsuite/libgomp.c-c++-common/error-1.c | 5 +++++
libgomp/testsuite/libgomp.fortran/error-1.f90 | 6 ++++++
4 files changed, 19 insertions(+)
diff --git a/libgomp/config/nvptx/error.c b/libgomp/config/nvptx/error.c
index dfa75da354f..40c907e0c9e 100644
--- a/libgomp/config/nvptx/error.c
+++ b/libgomp/config/nvptx/error.c
@@ -34,9 +34,11 @@
#undef vfprintf
#undef fputs
#undef fputc
+#undef fwrite
#define vfprintf(stream, fmt, list) vprintf (fmt, list)
#define fputs(s, stream) printf ("%s", s)
#define fputc(c, stream) printf ("%c", c)
+#define fwrite(ptr, size, nmemb, stream) printf ("%.*s", (int) (size * nmemb), ptr)
#include "../../error.c"
diff --git a/libgomp/error.c b/libgomp/error.c
index 9b69a4b33fe..35fc823abf9 100644
--- a/libgomp/error.c
+++ b/libgomp/error.c
@@ -90,6 +90,12 @@ gomp_fatal (const char *fmt, ...)
va_end (list);
}
+//TODO just for testing
+#if 0
+#undef fwrite
+#define fwrite(ptr, size, nmemb, stream) fprintf (stream, "%.*s", (int) (size * nmemb), ptr)
+#endif
+
void
GOMP_warning (const char *msg, size_t msglen)
{
diff --git a/libgomp/testsuite/libgomp.c-c++-common/error-1.c b/libgomp/testsuite/libgomp.c-c++-common/error-1.c
index 5f454c1adaa..dd3e8b135dd 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/error-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/error-1.c
@@ -39,11 +39,16 @@ main ()
#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;
}
/* { dg-output "libgomp: error directive encountered(\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: error directive encountered: message(\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..496512aa6a7 100644
--- a/libgomp/testsuite/libgomp.fortran/error-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/error-1.f90
@@ -47,9 +47,14 @@ 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))
+ !TODO $omp target
!$omp error at (execution) severity (fatal) message (msg)
+ !TODO $omp end target
! unreachable due to 'fatal'---------^
!$omp error at (execution) severity (warning) message ("foobar")
contains
@@ -73,6 +78,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
