[LTO] PR 86416 – improve lto1 diagnostic if a mode does not exist (esp. for offloading targets)
diff mbox series

Message ID ea04f184-d7cb-32b6-032d-8ede3d604d0a@codesourcery.com
State New
Headers show
Series
  • [LTO] PR 86416 – improve lto1 diagnostic if a mode does not exist (esp. for offloading targets)
Related show

Commit Message

Tobias Burnus Dec. 13, 2019, 2:34 p.m. UTC
As long as one compiles for a single target, the message is unlikely to appear.

However, when compiling for offloading, the modes supported on the target
'host' and on the target 'device' can be different. In particular,
'long double' (when larger than double) and '__float128' might not be
available on the device.

This gives currently errors like the following (see PR, comment 0):

> lto1: fatal error: unsupported mode TF  > > compilation terminated. > mkoffload: fatal error: 
x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status

While the device target is hidden in 'x86_64-pc-linux-gnu-accel-nvptx-none-gcc', it might make more sense to 
add it more prominently. Additionally, an average user does
not understand what 'TF' or 'XF' means.

Solutions:
(A) Add the target to the output
(B) Add a better description for the error

I did both in the attached patch, giving:
lto1: fatal error: nvptx-none - 80-bit floating-point numbers unsupported (mode 'XF')
lto1: fatal error: nvptx-none - 128-bit floating-point numbers unsupported (mode 'TF')

* (A) should be fine, I think.

* But I am not 100% happy with (B). I think as interim solution,
it is acceptable as XF/TF are well defined and probably the most
common problem. — Alternatively, one could only commit (A) or
solve it more properly (how?).

* If, e.g., 'long long' or 'integer(kind=16)' are used, only a
generic message is printed.  A message such as "'__float128' not
supported" or "'real(kind=10)' not supported" is more helpful and
supporting all modes and not cherry picking those two would be
useful as well.

The question is how to pass this information to lto-streamer-in.c;
it is available as TYPE_NAME – and, with debugging turned on, this
also gets passed on, but is also not be readily available in
lto_input_mode_table. – Suggestions?


Build on x86-64-gnu-linux and tested without offloading and with nvptx
offloading.

Tobias

Comments

Tobias Burnus Dec. 18, 2019, 11:45 a.m. UTC | #1
*ping*

On 12/13/19 3:34 PM, Tobias Burnus wrote:
> As long as one compiles for a single target, the message is unlikely 
> to appear.
>
> However, when compiling for offloading, the modes supported on the target
> 'host' and on the target 'device' can be different. In particular,
> 'long double' (when larger than double) and '__float128' might not be
> available on the device.
>
> This gives currently errors like the following (see PR, comment 0):
>
>> lto1: fatal error: unsupported mode TF
>> compilation terminated.
>> mkoffload: fatal error:
>> x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status 
>
> While the device target is hidden in 
> 'x86_64-pc-linux-gnu-accel-nvptx-none-gcc', it might make more sense 
> to add it more prominently. Additionally, an average user does
> not understand what 'TF' or 'XF' means.
>
> Solutions:
> (A) Add the target to the output
> (B) Add a better description for the error
>
> I did both in the attached patch, giving:
> lto1: fatal error: nvptx-none - 80-bit floating-point numbers 
> unsupported (mode 'XF')
> lto1: fatal error: nvptx-none - 128-bit floating-point numbers 
> unsupported (mode 'TF')
>
> * (A) should be fine, I think.
>
> * But I am not 100% happy with (B). I think as interim solution,
> it is acceptable as XF/TF are well defined and probably the most
> common problem. — Alternatively, one could only commit (A) or
> solve it more properly (how?).
>
> * If, e.g., 'long long' or 'integer(kind=16)' are used, only a
> generic message is printed.  A message such as "'__float128' not
> supported" or "'real(kind=10)' not supported" is more helpful and
> supporting all modes and not cherry picking those two would be
> useful as well.
>
> The question is how to pass this information to lto-streamer-in.c;
> it is available as TYPE_NAME – and, with debugging turned on, this
> also gets passed on, but is also not be readily available in
> lto_input_mode_table. – Suggestions?
>
>
> Build on x86-64-gnu-linux and tested without offloading and with nvptx
> offloading.
>
> Tobias
>
Jakub Jelinek Dec. 18, 2019, 12:41 p.m. UTC | #2
On Fri, Dec 13, 2019 at 03:34:56PM +0100, Tobias Burnus wrote:
> --- a/gcc/lto-streamer-in.c
> +++ b/gcc/lto-streamer-in.c
> @@ -1700,7 +1700,19 @@ lto_input_mode_table (struct lto_file_decl_data *file_data)
>  		}
>  	      /* FALLTHRU */
>  	    default:
> -	      fatal_error (UNKNOWN_LOCATION, "unsupported mode %qs", mname);
> +	      /* For offloading-target compilions, this is a user-facing
> +		 message.  See also target.def and machmode.def.  */
> +	      if (strcmp (mname, "XF") == 0)
> +		fatal_error (UNKNOWN_LOCATION,
> +			     "%s - 80-bit floating-point numbers unsupported "
> +			     "(mode %qs)", TARGET_MACHINE, mname);
> +	      else if (strcmp (mname, "TF") == 0)
> +		fatal_error (UNKNOWN_LOCATION,
> +			     "%s - 128-bit floating-point numbers unsupported "
> +			     "(mode %qs)", TARGET_MACHINE, mname);

I don't like the above, it looks like a gross hack for two modes that are
often used on x86_64, but as soon as you e.g. use powerpc64 KFmode instead
or many others, you'll be in trouble.
I'd say let lto_output_mode_table stream some string with a description of
the mode (perhaps for offloading streaming only), whether that would be by
looking e.g. through standard lang specific types, finding if they have such
mode and printing the name of the type, or trying to describe the mode in
words somehow, and then print that string during lto_input_mode_table.

	Jakub
Tobias Burnus Dec. 18, 2019, 4:39 p.m. UTC | #3
Hi Jakub,

thanks for the pointers; I was also not happy about part (B).

The mode gets written at lto-streamer-out.c's lto_write_mode_table, 
which is already called with 'if (lto_stream_offload_p)'. — Looking at 
places where the name gets constructed, I realized that required 
information is already written to the mode table: The precision (for 
float, to distinguish between 80bit and 128bit floats) and the size.

One now has a readable error message for float/complex/decimal-float and 
int(eger) modes – and the generic one for the rest; as vector get 
converted to their base type, if not found, this give a readable message 
for all common cases.

Hence, one can write a much cleaner patch which is about as long as the 
hack but not hackish :-)

Bootstrapped + regtested on x86-64-gnu-linux w/o offloading and with 
nvptx offloading.
OK?

Thanks,

Tobias
Jakub Jelinek Dec. 18, 2019, 4:43 p.m. UTC | #4
On Wed, Dec 18, 2019 at 05:39:51PM +0100, Tobias Burnus wrote:
> Hence, one can write a much cleaner patch which is about as long as the hack
> but not hackish :-)
> 
> Bootstrapped + regtested on x86-64-gnu-linux w/o offloading and with nvptx
> offloading.
> OK?

LGTM.

> 2019-12-18  Tobias Burnus  <tobias@codesourcery.com>
> 
> 	PR middle-end/86416
> 	*  Makefile.in (CFLAGS-lto-streamer-in.o): Pass target_noncanonical on.
> 	* lto-streamer-in.c (lto_input_mode_table): Improve unsupported-mode
> 	diagnostic.
> 
> 	PR middle-end/86416
> 	* testsuite/libgomp.c/pr86416-1.c: New.
> 	* testsuite/libgomp.c/pr86416-2.c: New.

	Jakub

Patch
diff mbox series

	PR middle-end/86416
	*  Makefile.in (CFLAGS-lto-streamer-in.o): Pass target_noncanonical on.
	* lto-streamer-in.c (lto_input_mode_table): Use it; add special error
	diagnostic for missing XF and TF modes.

	PR middle-end/86416
	* testsuite/libgomp.c/pr86416-1.c: New.
	* testsuite/libgomp.c/pr86416-2.c: New.

diff --git a/gcc/Makefile.in b/gcc/Makefile.in
index 6b857bd75de..657488d416b 100644
--- a/gcc/Makefile.in
+++ b/gcc/Makefile.in
@@ -2244,6 +2244,8 @@  version.o: $(REVISION) $(DATESTAMP) $(BASEVER) $(DEVPHASE)
 # lto-compress.o needs $(ZLIBINC) added to the include flags.
 CFLAGS-lto-compress.o += $(ZLIBINC)
 
+CFLAGS-lto-streamer-in.o += -DTARGET_MACHINE=\"$(target_noncanonical)\"
+
 bversion.h: s-bversion; @true
 s-bversion: BASE-VER
 	echo "#define BUILDING_GCC_MAJOR `echo $(BASEVER_c) | sed -e 's/^\([0-9]*\).*$$/\1/'`" > bversion.h
diff --git a/gcc/lto-streamer-in.c b/gcc/lto-streamer-in.c
index 128d7640726..f59e8c5363d 100644
--- a/gcc/lto-streamer-in.c
+++ b/gcc/lto-streamer-in.c
@@ -1700,7 +1700,19 @@  lto_input_mode_table (struct lto_file_decl_data *file_data)
 		}
 	      /* FALLTHRU */
 	    default:
-	      fatal_error (UNKNOWN_LOCATION, "unsupported mode %qs", mname);
+	      /* For offloading-target compilions, this is a user-facing
+		 message.  See also target.def and machmode.def.  */
+	      if (strcmp (mname, "XF") == 0)
+		fatal_error (UNKNOWN_LOCATION,
+			     "%s - 80-bit floating-point numbers unsupported "
+			     "(mode %qs)", TARGET_MACHINE, mname);
+	      else if (strcmp (mname, "TF") == 0)
+		fatal_error (UNKNOWN_LOCATION,
+			     "%s - 128-bit floating-point numbers unsupported "
+			     "(mode %qs)", TARGET_MACHINE, mname);
+	      else
+		fatal_error (UNKNOWN_LOCATION, "%s - unsupported mode %qs",
+			     TARGET_MACHINE, mname);
 	      break;
 	    }
 	}
diff --git a/libgomp/testsuite/libgomp.c/pr86416-1.c b/libgomp/testsuite/libgomp.c/pr86416-1.c
new file mode 100644
index 00000000000..c7a162d2f41
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr86416-1.c
@@ -0,0 +1,22 @@ 
+/* { dg-do link } */
+/* { dg-require-effective-target large_long_double } */
+
+/* PR middle-end/86416  */
+/* { dg-error "80-bit floating-point numbers unsupported .mode 'XF'." "" { target offload_device } 0 }  */
+/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target offload_device } }  */
+
+#include <stdlib.h>  /* For abort. */
+
+long double foo (long double x)
+{
+  #pragma omp target map(tofrom:x)
+    x *= 2.0;
+  return x;
+}
+
+int main()
+{
+  long double v = foo (10.0q) - 20.0q;
+  if (v > 1.0e-5 || v < -1.0e-5) abort();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr86416-2.c b/libgomp/testsuite/libgomp.c/pr86416-2.c
new file mode 100644
index 00000000000..75e9cd773a3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr86416-2.c
@@ -0,0 +1,22 @@ 
+/* { dg-do link { target __float128 } } */
+/* { dg-add-options __float128 } */
+
+/* PR middle-end/86416  */
+/* { dg-error "128-bit floating-point numbers unsupported .mode 'TF'." "" { target offload_device } 0 } */
+/* { dg-excess-errors "Follow-up errors from mkoffload and lto-wrapper" { target offload_device } }  */
+
+#include <stdlib.h>  /* For abort. */
+
+__float128 foo(__float128 y)
+{
+  #pragma omp target map(tofrom: y)
+    y *= 4.0;
+  return y;
+}
+
+int main()
+{
+  __float128 v = foo (5.0L) - 20.0L;
+  if (v > 1.0e-5 || v < -1.0e-5) abort();
+  return 0;
+}