On Fri, Dec 14, 2018 at 10:21:35AM +0100, Tom de Vries wrote: > Build and reg-tested on x86_64 with nvptx accelerator. > > 2018-12-13 Tom de Vries <tdevr...@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.
> + 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); This is translation unfriendly. Please just do: if (TREE_CODE (decl) == FUNCTION_DECL) error_at (...); else if (VAR_P (decl)) error_at (...); else gcc_unreachable (); (also note VAR_P). And, please use hasn%'t instead of hasn't. > @@ -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); > } Why are you passing all these arguments to that function (especially calling node->name () even when you don't know if it will be needed or not)? Doesn't both cgraph_node and varpool_node inherit from symtab_node, which has all these 3 fields as well as name () method? So, I think if verify_node_partition takes a symtab_node *node argument and the callers just both do verify_node_partition (node); it should work fine. I'd make it inline too. > --- /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); > +} Note, this will be well defined in OpenMP 5.0, just the support isn't there. The spec says: "If a function (C, C++, Fortran) or subroutine (Fortran) is referenced in a target construct then that function or subroutine is treated as if its name had appeared in a to clause on a declare target directive." and "If a function is referenced in a function that appears as a list item in a to clause on a declare target directive then the name of the referenced function is treated as if it had appeared in a to clause on a declare target directive. If a variable with static storage duration or a function (except lambda for C++) is referenced in the initializer expression list of a variable with static storage duration that appears as a list item in a to clause on a declare target directive then the name of the referenced variable or function is treated as if it had appeared in a to clause on a declare target directive." so for functions it should work through implicit propagation of the "omp declare target" attribute. Can't find a restriction I'd expect to see that if a declaration of a function or variable is marked declare target then the definition has to be as well, will talk to omp-lang. In any case, for the above testcase it might be better to split it into two sources with dg-auxiliary-source, declare baz in the source with main and define in another file (where declare instead of define the variable). > 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); > +} Perhaps just simplify the testcases too, the variable can be just called var and be a scalar, baz can be renamed to foo, drop the i argument, drop the #pragma omp for and loop from there too (in both tests). So simply just call a function from target region that increments a variable. Otherwise LGTM. Jakub