From 38958ac987dc3e6162e2ddaba3c7e7f41381e079 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge 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