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