Re: [PATCH] error on missing LTO symbols

[Tom de Vries <tdevries at suse dot de>]

On 13-12-18 14:44, Jakub Jelinek wrote:
> On Thu, Dec 13, 2018 at 02:31:45PM +0100, Tom de Vries wrote:
>>> 2015-07-24  Cesar Philippidis  <cesar@codesourcery.com>
> 
> Please use current date ;)
> 
>>>
>>> 	gcc/
>>> 	* lto-cgraph.c (input_overwrite_node): Error instead of assert
>>> 	on missing cgraph partitions.
>>> 	(input_varpool_node): Likewise.
>>>
>>>
>>> diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c
>>> index d70537d..7e2fc80 100644
>>> --- a/gcc/lto-cgraph.c
>>> +++ b/gcc/lto-cgraph.c
>>> @@ -1218,9 +1218,11 @@ input_overwrite_node (struct lto_file_decl_data
>>> *file_data,
>>>  				     LDPR_NUM_KNOWN);
>>>    node->instrumentation_clone = bp_unpack_value (bp, 1);
>>>    node->split_part = bp_unpack_value (bp, 1);
>>> -  gcc_assert (flag_ltrans
>>> -	      || (!node->in_other_partition
>>> -		  && !node->used_from_other_partition));
>>> +
>>> +  int success = flag_ltrans || (!node->in_other_partition
>>> +				&& !node->used_from_other_partition);
> 
> I'd use internal_error (or internal_error_no_backtrace ?) here if it isn't
> an offloading compiler (so #ifndef ACCEL_COMPILER), or just gcc_assert in
> that case?  Richard/Honza, your thoughts on that?

Left it as gcc_assert for now.

> And error if ACCEL_COMPILER is defined.
> 

Done.

>>> +  if (!success)
>>> +    error ("Missing %<%s%>", node->name ());
> 
> Diagnostics shouldn't start with capital letters.  %<%s%> should be %qs.
> 

Done.

> The diagnostics for the non-accel case if any can be less verbose, but
> perhaps should at least say whether it is a function (above case) or
> variable.
> 

I've left the non-accel case alone for now.

> For the ACCEL_COMPILER case, I think we should be more verbose, say that
> certain function or variable has been referenced in offloaded code but
> corresponding function or variable definition hasn't been marked to be
> included in the offloaded code.

Done.

> Would be nice to have a location where it has been referenced if
> that can be dug from somewhere.

It's not obvious to me how to do that, so I've used the location of the
declaration of the symbol instead.

> Or if e.g. from attribute we can figure out
> whether it was OpenMP or OpenACC offloading and even suggest how to fix it
> in the corresponding language extension.

I left it language-agnostic for now.

OK for trunk?

Thanks,
- Tom

[offloading] Error on missing symbols

When compiling an OpenMP or OpenACC program containing a reference in the
offloaded code to a symbol that has not been included in the offloaded code,
the offloading compiler may ICE in lto1.

Fix this by erroring out instead, mentioning the problematic symbol:
...
error: variable 'results' has been referenced in offloaded code but hasn't
  been marked to be included in the offloaded code
lto1: fatal error: errors during merging of translation units
compilation terminated.
...

Build and reg-tested on x86_64 with nvptx accelerator.

2018-12-13  Tom de Vries  <tdevries@suse.de>

	* lto-cgraph.c (verify_node_partition): New function.
	(input_overwrite_node, input_varpool_node): Use verify_node_partition.

	* testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test.
	* testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test.

---
 gcc/lto-cgraph.c                                   | 42 ++++++++++++++++++----
 .../libgomp.c-c++-common/function-not-offloaded.c  | 21 +++++++++++
 .../libgomp.c-c++-common/variable-not-offloaded.c  | 21 +++++++++++
 3 files changed, 77 insertions(+), 7 deletions(-)

diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c
index 99998cc3c75..f10a96bd2a4 100644
--- a/gcc/lto-cgraph.c
+++ b/gcc/lto-cgraph.c
@@ -1091,6 +1091,37 @@ output_offload_tables (void)
     }
 }
 
+/* Verify the partitioning of a varpool_node or cgraph_node with DECL and NAME,
+   as specified by IN_OTHER_PARTITION and USED_FROM_OTHER_PARTITION.  */
+
+static void
+verify_node_partition (tree decl ATTRIBUTE_UNUSED,
+		       const char *name ATTRIBUTE_UNUSED,
+		       bool in_other_partition,
+		       bool used_from_other_partition ATTRIBUTE_UNUSED)
+{
+  if (flag_ltrans)
+    return;
+
+#ifdef ACCEL_COMPILER
+  if (in_other_partition)
+    {
+      if (TREE_CODE (decl) == FUNCTION_DECL
+	  || TREE_CODE (decl) == VAR_DECL)
+	error_at (DECL_SOURCE_LOCATION (decl),
+		  "%s %qs has been referenced in offloaded code but"
+		  " hasn't been marked to be included in the offloaded code",
+		  TREE_CODE (decl) == FUNCTION_DECL ? "function" : "variable",
+		  name);
+      else
+	gcc_unreachable ();
+    }
+#else
+  gcc_assert (!in_other_partition
+	      && !used_from_other_partition);
+#endif
+}
+
 /* Overwrite the information in NODE based on FILE_DATA, TAG, FLAGS,
    STACK_SIZE, SELF_TIME and SELF_SIZE.  This is called either to initialize
    NODE or to replace the values in it, for instance because the first
@@ -1153,9 +1184,8 @@ input_overwrite_node (struct lto_file_decl_data *file_data,
   node->resolution = bp_unpack_enum (bp, ld_plugin_symbol_resolution,
 				     LDPR_NUM_KNOWN);
   node->split_part = bp_unpack_value (bp, 1);
-  gcc_assert (flag_ltrans
-	      || (!node->in_other_partition
-		  && !node->used_from_other_partition));
+  verify_node_partition (node->decl, node->name (), node->in_other_partition,
+			 node->used_from_other_partition);
 }
 
 /* Return string alias is alias of.  */
@@ -1366,10 +1396,8 @@ input_varpool_node (struct lto_file_decl_data *file_data,
     node->set_section_for_node (section);
   node->resolution = streamer_read_enum (ib, ld_plugin_symbol_resolution,
 					        LDPR_NUM_KNOWN);
-  gcc_assert (flag_ltrans
-	      || (!node->in_other_partition
-		  && !node->used_from_other_partition));
-
+  verify_node_partition (node->decl, node->name (), node->in_other_partition,
+			 node->used_from_other_partition);
   return node;
 }
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c
new file mode 100644
index 00000000000..38b68cc28a6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c
@@ -0,0 +1,21 @@
+/* { dg-do link } */
+/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" } */
+
+#pragma omp declare target
+int results[2000];
+#pragma omp end declare target
+
+void __attribute__((noinline, noclone))
+baz (int i) /* { dg-error "function 'baz' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" } */
+{
+  results[i]++;
+}
+
+int
+main ()
+{
+#pragma omp target
+#pragma omp for
+  for (int i = 0; i < 2000; i++)
+    baz (i);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c b/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c
new file mode 100644
index 00000000000..c2e1d57adea
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c
@@ -0,0 +1,21 @@
+/* { dg-do link } */
+/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" } */
+
+int results[2000]; /* { dg-error "variable 'results' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" } */
+
+#pragma omp declare target
+void  __attribute__((noinline, noclone))
+baz (int i)
+{
+  results[i]++;
+}
+#pragma omp end declare target
+
+int
+main ()
+{
+#pragma omp target
+#pragma omp for
+  for (int i = 0; i < 2000; i++)
+    baz (i);
+}