Re: [PATCH] error on missing LTO symbols

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

[ cc-ing HSAIL maintainer ]

On 14-12-18 10:54, Jakub Jelinek wrote:
> 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  <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.
> 
>> +      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.
> 

Done.

>> @@ -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.
> 

Done.

>> --- /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).
>

Done, but that ends up not triggering the introduced error condition
anymore. Instead we generate an "unresolved symbol foo".

So I've added the corresponding OpenACC test-cases as well (and the
introduced error triggers only for OpenACC functions and OpenMP variables).

>> 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.
> 

Done.

[ quote-pasted from follow-up email ]
> Ah, one more thing, the testcase doesn't really fail when offloading isn't
> configured, so it would need some effective target or something that
> it actually does the offloading.  And not really sure about shared memory
> offloading like hsail, that doesn't really need the variables declared
> either, just the functions.

Done, using offload_device_nonshared_as for
libgomp.c-c++-common/variable-not-offloaded.c and
openacc_nvidia_accel_configured for
libgomp.oacc-c-c++-common/function-not-offloaded.c.

> Otherwise LGTM.

Updated patch OK?

Thanks,
- Tom