| Message ID | 87sf0s6e9l.fsf@euler.schwinge.ddns.net |
|---|---|
| State | New |
| Headers | show |
| Series | OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing (was: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends) | expand |
Hi all, hi Thomas & Chung-Lin, Thomas Schwinge wrote: > But I realized another thing: don't we have to handle the 'readonly' > modifier also in Fortran module files, that is, next to the OpenACC > 'declare' 'copyin' handling in 'gcc/fortran/module.cc': > 'AB_OACC_DECLARE_COPYIN' etc.? I bet so; it is not as bad as with the others as it is "only" an optimization hint, but it makes sense to make it available. Note that when you place the 'module' in the same file as the module users ('use'), the compiler might know things because they are in the same translation unit / file not because it is in the module ... > Chung-Lin, please check, via test cases. > 'gfortran.dg/goacc/routine-module*', for example, should provide some > guidance of how to achieve actual module file use, and then do the same > 'scan-tree-dump' as in the current 'readonly' modifier test cases. ... > By means of only emitting a tag > in the module file if the 'readonly' modifier is specified, we should > maintain compatibility with the current 'MOD_VERSION'. That was the idea: If only new information gets added (if used), older compilers still work. This has huge limitations and does not work as well as imagined but here it should work: Older .mod will work with new compilers, even though the reverse might not be true. Tobias
Hi all, hi Thomas & Chung-Lin, Thomas Schwinge wrote: > But I realized another thing: don't we have to handle the 'readonly' > modifier also in Fortran module files, that is, next to the OpenACC > 'declare' 'copyin' handling in 'gcc/fortran/module.cc': > 'AB_OACC_DECLARE_COPYIN' etc.? I bet so; it is not as bad as with the others as it is "only" an optimization hint, but it makes sense to make it available. Note that when you place the 'module' in the same file as the module users ('use'), the compiler might know things because they are in the same translation unit / file not because it is in the module ... > Chung-Lin, please check, via test cases. > 'gfortran.dg/goacc/routine-module*', for example, should provide some > guidance of how to achieve actual module file use, and then do the same > 'scan-tree-dump' as in the current 'readonly' modifier test cases. ... > By means of only emitting a tag > in the module file if the 'readonly' modifier is specified, we should > maintain compatibility with the current 'MOD_VERSION'. That was the idea: If only new information gets added (if used), older compilers still work. This has huge limitations and does not work as well as imagined but here it should work: Older .mod will work with new compilers, even though the reverse might not be true. Tobias
From 38958ac987dc3e6162e2ddaba3c7e7f41381e079 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <tschwinge@baylibre.com> Date: Thu, 14 Mar 2024 15:01:01 +0100 Subject: [PATCH] OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing ... to complement commit ddf852dac2abaca317c10b8323f338123b0585c8 "OpenACC 2.7: front-end support for readonly modifier". gcc/testsuite/ * c-c++-common/goacc/readonly-1.c: Add basic OpenACC 'declare' testing. * gfortran.dg/goacc/readonly-1.f90: Likewise. --- gcc/testsuite/c-c++-common/goacc/readonly-1.c | 5 +++++ gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 | 6 ++++++ 2 files changed, 11 insertions(+) diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c index 34fc92c24d5..300464c92e3 100644 --- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c @@ -8,12 +8,15 @@ struct S int a[32], b[32]; #pragma acc declare copyin(readonly: a) copyin(b) +/* Not visible in 'original' dump; handled via 'offload_vars'. */ int main (void) { int x[32], y[32]; struct S s = {x, 0}; + #pragma acc declare copyin(readonly: x/*[:32]*/, s/*.ptr[:16]*/) copyin(y/*[:32]*/) + #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) { #pragma acc cache (readonly: x[:32]) @@ -43,6 +46,8 @@ int main (void) return 0; } +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc declare map\\(to:y\\) map\\(readonly,to:s\\) map\\(readonly,to:x\\)" 1 "original" } } */ + /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 index 696ebd08321..fc1e2719e67 100644 --- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 @@ -3,6 +3,9 @@ subroutine foo (a, n) integer :: n, a(:) integer :: i, b(n), c(n) + !!$acc declare copyin(readonly: a(:), b(:n)) copyin(c(:)) + !$acc declare copyin(readonly: b) copyin(c) + !$acc parallel copyin(readonly: a(:), b(:n)) copyin(c(:)) do i = 1,32 !$acc cache (readonly: a(:), b(:n)) @@ -74,6 +77,9 @@ program main end program main +! The front end turns OpenACC 'declare' into OpenACC 'data'. +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*b\\) map\\(alloc:b.+ map\\(to:\\*c\\) map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:g\\) map\\(to:h\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } } ! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } -- 2.34.1