Chung-Lin Tang wrote:
as we discussed off-list, here is the updated patch for the omp declare target
local clause.
Thanks for the patch. I am mostly happy with the patch.
* * *
TODO in the future (not this patch):
We should cleanup internal representation mess of using additional attributes
vs. arguments to the 'omp declare target' attribute, unifying OpenMP vs. OpenACC
and Fortran vs. C/C++ handling.
* * *
Regarding Fortran: I think the first change in openmp.cc should be
undone, the second is fine though.
Regarding the first one: For
module m
integer :: x, y
!$omp declare target local(y)
!$omp groupprivate(x, y)
!$omp declare target local(x)
end module m
GCC with the patch compiles the first two directive lines happily (i.e. 'y')
but when encountering 'x' a second time in a directive there is NOW the error:
Error: List item ‘x’ at (1) not appear in the ‘local’ clause as it was
previously specified in a GROUPPRIVATE directive
* * *
First, someone - likely me - forgot a verb. I think it should be:
'List item 'x' at (1)' <VERB> not appear'
with a 'shall', 'must', or 'may' added here.
Undoing the condition, it is no longer touched by your patch, but can
you still add a verb to the error message?
Thanks!
* * *
Second - IMHO the code is valid. OpenMP has:
"9.2 groupprivate Directive"
"For a 'device_type' clause that is specified implicitly or explicitly
on the directive, the behavior is as if the list items appear in a
'local' clause on a 'declare target' directive on which the same
'device_type' clause is specified and at the same program point."
And for there 'device_type' clause, there is:
"If the device_type clause is not specified, the behavior is as if
the device_type clause appears with any specified."
There is also wording related to such variables in "15.8 Declare Target
Directives":
"If a variable with static storage duration that is not a device-local
variable (including that it is not a 'groupprivate' variable) is declared
in a device procedure then the variable is treated as if it had appeared
as a list item in an 'enter' clause on a declare target directive."
Thus 'local' and 'groupprivate' are linked. I do not see a reason
why specifying them in any order should cause an error (if they have
the same device_type). We also accept
omp declare target enter(x)
omp declare target enter(x)
multiple times.
Thus, I think this check (first one in openmp.cc) should be reverted.
* * *
Third: Kudos for catching the thinko in the diagnostic check;
before there was no error when 'link' and 'local' were combined!
* * *
Two small issues with the testcase:
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-local-1.c
@@ -0,0 +1,27 @@
+
+int local_var = 900;
+#pragma omp declare target local(local_var)
+
+#pragma omp begin declare target
+int foo (int x)
+{
+ local_var += x;
+ return local_var;
+}
+#pragma omp end declare target
+
+int main (void)
+{
+ int e;
+ #pragma omp target map(from : e)
+ {
+ e = foo (23);
+ }
+
+ #pragma omp target update from(local_var)
Please add a comment like:
/* As with 'local' no corresponding storage exists,
the 'target update' directive should have no effect
and 'local_var' retains the host value. */
While the code is fine, it is rather confusing and a
comment helps a lot!
* * *
Plus:
+ if (e == local_var)
+ __builtin_abort ();
This will fail if you don't have an offload device.
Hence: Either check for
omp_get_num_device () == 0
and, if true, by check for e != local_var → abort
instead of 'e == local_var'.
Or, alternatively, add, e.g.,
/* { dg-do run { target offload_device } } */
To skip that test if there is no offload device at runtime.
* * *
Likewise for:
+++ b/libgomp/testsuite/libgomp.fortran/declare-target-local-1.f90
...
+program main
+ use m
+ implicit none (type, external)
+ integer :: e
+
+ !$omp target map(from : e)
+ e = incr_local_var (23)
+ !$omp end target
+
+ !$omp target update from(local_var)
Namely: Please add a comment here
And you also need to handle the case that no offload device exists.
* * *
Otherwise: LGTM - Thanks for the patch!
Tobias
