Hello world, hi Chung-Lin,
first, can you add 'PR middle-end/122205' to your changelog?
that's about what you implemented ( https://gcc.gnu.org/PR122205 )
Chung-Lin Tang wrote:
this implements 'link' clause behavior for 'enter' maps, when
detected that 'requires self_maps' is in effect.
We should to do the same with 'requires unified_shared_memory'.
For 'self_maps', the implementation is required that the device
and the host variable access the same memory. For 'unified_shared_memory',
it is left open whether explicit map clauses have an effect or not.
However, we ignore all map operations (aka as self mapping) with
either requirement, which leads to wrong code for:
#pragma omp requires unified_shared_memory // or 'self_maps' before your patch
int var;
#pragma omp declare target enter(var)
int main() {
var = 5;
#pragma omp target map(always, tofrom: var)
var++;
__builtin_printf ("%d\n", var); // should print 6, might print 1 or garbage
}
This basically does it in the front-ends, changing the tagged
attribute to "target link" under the supposed right conditions.
This fails for:
#pragma omp requires self_maps
#pragma omp begin declare target
void f() { }
#pragma omp declare target enter(f) device_type(nohost)
#pragma omp end declare target
with the error message:
foo.c:5:33: error: ‘f’ specified both in declare target ‘link’ and ‘enter’
clauses
5 | #pragma omp declare target enter(f) device_type(nohost)
* * *
Actually: This one is odd in multiple ways, because for *functions* it
shouldn't the clause to 'link'!
* * *
On the other hand, 'begin declare target' doesn't seem to work with variables
either - while it compiles ...
#pragma omp requires self_maps
#pragma omp begin declare target
static int val;
int f() { return val; }
#pragma omp end declare target
int main() {
int res = 0;
val = 41;
#pragma omp target map(from: res)
res = f();
__builtin_printf ("val = %d\n", res);
}
... it gives here 0 instead of the expected 41 as result.
* * *
The following also fails (device-side variable not created) - but
it is unrelated to your patch → https://gcc.gnu.org/PR126255
program main
implicit none
!$omp requires self_maps
integer :: var, res
var = 5
call test
!$omp target
res = f()
!$omp end target
contains
subroutine test
!$omp declare target enter(var)
end
integer function f()
f = var
end
end
* * *
I noticed that the following code is currently rejected with
foo.c:3:12: error: ‘init’ specified both in declare target ‘link’ and
implicitly in ‘to’ clauses
but it is no longer with you patch:
#pragma omp requires self_maps
static int init;
#pragma omp declare target link(init)
static int *val = &init;
#pragma omp declare target enter(val)
#pragma omp begin declare target
int f() { init++; return *val; }
#pragma omp end declare target
int main() {
int res = 0;
init = 42;
#pragma omp target map(from: res)
res = f();
__builtin_printf ("val = %d\n", res);
}
Expected: As 'init' is implicitly 'enter', it should be rejected
as there is an explicit 'link' clause as well.
* * *
The reverse is invalid as well, but not diagnosed, either
→ https://gcc.gnu.org/PR126256
but that's unrelated to your patch.
* * *
The following code works with manually using 'link' but
fails with 'enter':
implicit none
!$omp requires self_maps
integer :: x = 7
common /com/ x
! !$omp declare target link(/com/) ! -- works: prints 8
!$omp declare target enter(/com/) ! -- fails: prints 7
x = 8
!$omp target
print *, f()
!$omp end target
contains
integer function f()
f = x
end
end
* * *
Side remark: For the following code, we should not change
'enter' to 'link' - as the variable is supposed to only
exist on the device (likewise for 'host' but there is does
not really matter):
#pragma omp requires self_maps
static int val = 5;
#pragma omp declare target enter(val) device_type(nohost)
#pragma omp begin declare target device_type(nohost)
int f() { return omp_is_initial_device() ? 99 : val; }
#pragma omp end declare target
int main() {
int res = 0;
#pragma omp target map(from: res) // device_type(nohost)
res = f();
__builtin_printf ("val = %d\n", res);
}
This is currently not run-time detectable but the performance
is definitely better when accessing the device variable and
not the host variable.
(It starts to matter when we stop outputting nohost variables
and functions on the host.)
Note: 'device_type([no]host)' on target is not yet implemented;
the example works around this by using the 'omp_is_initial_device'
builtin that is compile-time evaluated.
* * *
Tobias,
who still has to read through the actual patch.
