[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 'var' 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 x86_64 with nvptx accelerator and reg-tested libgomp.

Build x86_64 and reg-tested libgomp.

2018-12-14  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-aux.c: New test.
	* testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test.
	* testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c: New test.

From-SVN: r267134

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index e5c226498aec75210bbd946ddafd5697db0e0ca9..afbdf75c62c96f7aeb7ae49a1b1bf4bff6a2e219 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,8 @@
+2018-12-14  Tom de Vries  <tdevries@suse.de>
+
+       * lto-cgraph.c (verify_node_partition): New function.
+       (input_overwrite_node, input_varpool_node): Use verify_node_partition.
+
 2018-12-14  H.J. Lu  <hongjiu.lu@intel.com>
 
        PR target/88483

diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c
index 99998cc3c754c1fd3720213271c93808d2affeb0..536e73c9ecd0c3bc668c67a9af7bb093b6edc0b8 100644 (file)
--- a/gcc/lto-cgraph.c
+++ b/gcc/lto-cgraph.c
@@ -1091,6 +1091,36 @@ output_offload_tables (void)
     }
 }
 
+/* Verify the partitioning of NODE.  */
+
+static inline void
+verify_node_partition (symtab_node *node)
+{
+  if (flag_ltrans)
+    return;
+
+#ifdef ACCEL_COMPILER
+  if (node->in_other_partition)
+    {
+      if (TREE_CODE (node->decl) == FUNCTION_DECL)
+       error_at (DECL_SOURCE_LOCATION (node->decl),
+                 "function %qs has been referenced in offloaded code but"
+                 " hasn%'t been marked to be included in the offloaded code",
+                 node->name ());
+      else if (VAR_P (node->decl))
+       error_at (DECL_SOURCE_LOCATION (node->decl),
+                 "variable %qs has been referenced in offloaded code but"
+                 " hasn%'t been marked to be included in the offloaded code",
+                 node->name ());
+      else
+       gcc_unreachable ();
+    }
+#else
+  gcc_assert (!node->in_other_partition
+             && !node->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 +1183,7 @@ 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);
 }
 
 /* Return string alias is alias of.  */
@@ -1366,10 +1394,7 @@ 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);
   return node;
 }
 

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 7b42e87f50a6f8f794cbcf8b0d7a2a696df00193..4c66021c367daa63d2d3fae4d7dab8179b584b6b 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,11 @@
+2018-12-14  Tom de Vries  <tdevries@suse.de>
+
+       * testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c: New test.
+       * testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test.
+       * testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test.
+       * testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: New test.
+       * testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c: New test.
+
 2018-12-13  Tom de Vries  <tdevries@suse.de>
 
        * affinity-fmt.c (gomp_print_string): New function, factored out of ...

diff --git a/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c
new file mode 100644 (file)
index 0000000..b8aa3da
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c
@@ -0,0 +1,12 @@
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp declare target
+extern int var;
+#pragma omp end declare target
+
+void __attribute__((noinline, noclone))
+foo (void)
+{
+  var++;
+}
+

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 (file)
index 0000000..9e59ef8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c
@@ -0,0 +1,16 @@
+/* { dg-do link } */
+/* { dg-excess-errors "unresolved symbol foo, lto1, mkoffload and lto-wrapper fatal errors" { target offload_device_nonshared_as } } */
+/* { dg-additional-sources "function-not-offloaded-aux.c" } */
+
+#pragma omp declare target
+int var;
+#pragma omp end declare target
+
+extern void foo ();
+
+int
+main ()
+{
+#pragma omp target
+  foo ();
+}

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 (file)
index 0000000..bc4b916
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c
@@ -0,0 +1,19 @@
+/* { dg-do link } */
+/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target offload_device_nonshared_as } } */
+
+int var; /* { dg-error "variable 'var' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target offload_device_nonshared_as } } */
+
+#pragma omp declare target
+void  __attribute__((noinline, noclone))
+foo (void)
+{
+  var++;
+}
+#pragma omp end declare target
+
+int
+main ()
+{
+#pragma omp target
+  foo ();
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
new file mode 100644 (file)
index 0000000..c94f268
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
@@ -0,0 +1,18 @@
+/* { dg-do link } */
+/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target openacc_nvidia_accel_configured } } */
+
+int var;
+#pragma acc declare create (var)
+
+void __attribute__((noinline, noclone))
+foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target openacc_nvidia_accel_configured } } */
+{
+  var++;
+}
+
+int
+main ()
+{
+#pragma acc parallel
+  foo ();
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c
new file mode 100644 (file)
index 0000000..8e10271
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c
@@ -0,0 +1,17 @@
+/* { dg-do link } */
+
+int var; /* { dg-error "'var' requires a 'declare' directive for use in a 'routine' function" } */
+
+#pragma acc routine
+void  __attribute__((noinline, noclone))
+foo (void)
+{
+  var++;
+}
+
+int
+main ()
+{
+#pragma acc parallel
+  foo ();
+}