diff options
author | Martin Liska <mliska@suse.cz> | 2022-11-07 08:24:48 +0100 |
---|---|---|
committer | Martin Liska <mliska@suse.cz> | 2022-11-07 08:24:48 +0100 |
commit | 1b09b78ee61bd921ae78ebd0f7905b95b9e1c903 (patch) | |
tree | 9c04b59cdd2cd460f0727501d15402d31ffcf5a4 /libgomp | |
parent | 1eb021edb27e26f95cda63df121f6bc951647599 (diff) | |
parent | c4f8f8afd07680f9e718de1331cd09607bdd9ac8 (diff) | |
download | gcc-1b09b78ee61bd921ae78ebd0f7905b95b9e1c903.zip gcc-1b09b78ee61bd921ae78ebd0f7905b95b9e1c903.tar.gz gcc-1b09b78ee61bd921ae78ebd0f7905b95b9e1c903.tar.bz2 |
Merge branch 'master' into devel/sphinx
Diffstat (limited to 'libgomp')
23 files changed, 2114 insertions, 122 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index d61a806..b39144e 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,74 @@ +2022-11-04 Thomas Schwinge <thomas@codesourcery.com> + + * libgomp-plugin.h (OFFLOAD_TARGET_TYPE_INTEL_MIC): Remove. + * libgomp.texi (OpenMP Context Selectors): Remove Intel MIC + documentation. + * plugin/configfrag.ac <enable_offload_targets> + [*-intelmic-* | *-intelmicemul-*]: Remove. + * configure: Regenerate. + * testsuite/lib/libgomp.exp (libgomp_init): Remove 'liboffloadmic' + handling. + (offload_target_to_openacc_device_type) + [$offload_target = *-intelmic*]: Remove. + (check_effective_target_offload_device_intel_mic) + (check_effective_target_offload_device_any_intel_mic): Remove. + * testsuite/libgomp.c-c++-common/on_device_arch.h + (device_arch_intel_mic, on_device_arch_intel_mic, any_device_arch) + (any_device_arch_intel_mic): Remove. + * testsuite/libgomp.c-c++-common/target-45.c: Remove + 'offload_device_any_intel_mic' XFAIL. + * testsuite/libgomp.fortran/target10.f90: Likewise. + +2022-11-03 Tobias Burnus <tobias@codesourcery.com> + + * testsuite/libgomp.fortran/target-11.f90: New test. + * testsuite/libgomp.fortran/target-13.f90: New test. + +2022-11-02 Thomas Schwinge <thomas@codesourcery.com> + + PR libgomp/106643 + PR fortran/96668 + * oacc-mem.c (goacc_enter_data_internal): Support + OpenACC 'declare create' with Fortran allocatable arrays, part II. + * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90: + Adjust. + * testsuite/libgomp.oacc-fortran/pr106643-1.f90: New. + +2022-11-02 Thomas Schwinge <thomas@codesourcery.com> + + PR libgomp/106643 + * oacc-mem.c (goacc_enter_data_internal): Support + OpenACC 'declare create' with Fortran allocatable arrays, part I. + * testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90: + New. + * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90: + New. + +2022-11-02 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90: + New. + +2022-11-02 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90: + New. + +2022-11-02 Cesar Philippidis <cesar@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90: New. + +2022-10-28 Julian Brown <julian@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + + PR middle-end/90115 + * testsuite/libgomp.oacc-fortran/declare-1.f90: Adjust scan output. + * testsuite/libgomp.oacc-fortran/host_data-5.F90: Likewise. + * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise. + * testsuite/libgomp.oacc-fortran/print-1.f90: Likewise. + * testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise. + 2022-10-24 Thomas Schwinge <thomas@codesourcery.com> * plugin/plugin-nvptx.c (nvptx_open_device): Initialize diff --git a/libgomp/configure b/libgomp/configure index 35424d2..45a769e 100755 --- a/libgomp/configure +++ b/libgomp/configure @@ -15200,9 +15200,6 @@ if test x"$enable_offload_targets" != x; then tgt=`echo $tgt | sed 's/=.*//'` tgt_plugin= case $tgt in - *-intelmic-* | *-intelmicemul-*) - tgt_plugin=intelmic - ;; nvptx*) case "${target}" in aarch64*-*-* | powerpc64le-*-* | x86_64-*-*) diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h index 875f967..ac38782 100644 --- a/libgomp/libgomp-plugin.h +++ b/libgomp/libgomp-plugin.h @@ -49,7 +49,6 @@ enum offload_target_type OFFLOAD_TARGET_TYPE_HOST = 2, /* OFFLOAD_TARGET_TYPE_HOST_NONSHM = 3 removed. */ OFFLOAD_TARGET_TYPE_NVIDIA_PTX = 5, - OFFLOAD_TARGET_TYPE_INTEL_MIC = 6, OFFLOAD_TARGET_TYPE_HSA = 7, OFFLOAD_TARGET_TYPE_GCN = 8 }; diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 4ff1fdb..3e8340b 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -4303,7 +4303,7 @@ offloading devices (it's not clear if they should be): @multitable @columnfractions .60 .10 .25 @headitem @code{arch} @tab @code{kind} @tab @code{isa} -@item @code{intel_mic}, @code{x86}, @code{x86_64}, @code{i386}, @code{i486}, +@item @code{x86}, @code{x86_64}, @code{i386}, @code{i486}, @code{i586}, @code{i686}, @code{ia32} @tab @code{host} @tab See @code{-m...} flags in ``x86 Options'' (without @code{-m}) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 73b2710..233fe0e 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1150,8 +1150,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, } else if (n && groupnum > 1) { - assert (n->refcount != REFCOUNT_INFINITY - && n->refcount != REFCOUNT_LINK); + assert (n->refcount != REFCOUNT_LINK); for (size_t j = i + 1; j <= group_last; j++) if ((kinds[j] & 0xff) == GOMP_MAP_ATTACH) @@ -1166,6 +1165,44 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, bool processed = false; struct target_mem_desc *tgt = n->tgt; + + /* Minimal OpenACC variant corresponding to PR96668 + "[OpenMP] Re-mapping allocated but previously unallocated + allocatable does not work" 'libgomp/target.c' changes, so that + OpenACC 'declare' code à la PR106643 + "[gfortran + OpenACC] Allocate in module causes refcount error" + has a chance to work. */ + if ((kinds[i] & 0xff) == GOMP_MAP_TO_PSET + && tgt->list_count == 0) + { + /* 'declare target'. */ + assert (n->refcount == REFCOUNT_INFINITY); + + for (size_t k = 1; k < groupnum; k++) + { + /* The only thing we expect to see here. */ + assert ((kinds[i + k] & 0xff) == GOMP_MAP_POINTER); + } + + /* Let 'goacc_map_vars' -> 'gomp_map_vars_internal' handle + this. */ + gomp_mutex_unlock (&acc_dev->lock); + struct target_mem_desc *tgt_ + = goacc_map_vars (acc_dev, aq, groupnum, &hostaddrs[i], NULL, + &sizes[i], &kinds[i], true, + GOMP_MAP_VARS_ENTER_DATA); + assert (tgt_ == NULL); + gomp_mutex_lock (&acc_dev->lock); + + /* Given that 'goacc_exit_data_internal'/'goacc_exit_datum_1' + will always see 'n->refcount == REFCOUNT_INFINITY', + there's no need to adjust 'n->dynamic_refcount' here. */ + + processed = true; + } + else + assert (n->refcount != REFCOUNT_INFINITY); + for (size_t j = 0; j < tgt->list_count; j++) if (tgt->list[j].key == n) { diff --git a/libgomp/plugin/configfrag.ac b/libgomp/plugin/configfrag.ac index ab03f94..d3b2589 100644 --- a/libgomp/plugin/configfrag.ac +++ b/libgomp/plugin/configfrag.ac @@ -59,9 +59,6 @@ if test x"$enable_offload_targets" != x; then tgt=`echo $tgt | sed 's/=.*//'` tgt_plugin= case $tgt in - *-intelmic-* | *-intelmicemul-*) - tgt_plugin=intelmic - ;; nvptx*) case "${target}" in aarch64*-*-* | powerpc64le-*-* | x86_64-*-*) diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index 4b8c64d..1801fdc 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -119,18 +119,6 @@ proc libgomp_init { args } { # Compute what needs to be put into LD_LIBRARY_PATH set always_ld_library_path ".:${blddir}/.libs" - # Add liboffloadmic build directory in LD_LIBRARY_PATH to support - # Intel MIC offloading testing. - global offload_plugins - if { [string match "*,intelmic,*" ",$offload_plugins,"] } { - append always_ld_library_path ":${blddir}/../liboffloadmic/.libs" - append always_ld_library_path ":${blddir}/../liboffloadmic/plugin/.libs" - # libstdc++ is required by liboffloadmic - append always_ld_library_path ":${blddir}/../libstdc++-v3/src/.libs" - # libgcc_s is required by libstdc++ - append always_ld_library_path ":${blddir}/../libgcc" - } - global offload_additional_lib_paths if { $offload_additional_lib_paths != "" } { append always_ld_library_path "${offload_additional_lib_paths}" @@ -313,9 +301,6 @@ proc offload_target_to_openacc_device_type { offload_target } { disable { return "host" } - *-intelmic* { - return "" - } nvptx* { return "nvidia" } @@ -449,28 +434,6 @@ proc check_effective_target_openacc_nvidia_accel_selected { } { return [string match "nvidia" $openacc_device_type] } -# Return 1 if using Intel MIC offload device. -proc check_effective_target_offload_device_intel_mic { } { - return [check_runtime_nocache offload_device_intel_mic { - #include "testsuite/libgomp.c-c++-common/on_device_arch.h" - int main () - { - return !on_device_arch_intel_mic (); - } - } ] -} - -# Return 1 if any Intel MIC offload device is available. -proc check_effective_target_offload_device_any_intel_mic { } { - return [check_runtime_nocache offload_device_any_intel_mic { - #include "testsuite/libgomp.c-c++-common/on_device_arch.h" - int main () - { - return !any_device_arch_intel_mic (); - } - } ] -} - # Return 1 if the OpenACC 'host' device type is selected. proc check_effective_target_openacc_host_selected { } { diff --git a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h index 6f66dbd..3fb5021 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h +++ b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h @@ -13,15 +13,8 @@ device_arch_gcn (void) return GOMP_DEVICE_GCN; } -/* static */ int -device_arch_intel_mic (void) -{ - return GOMP_DEVICE_INTEL_MIC; -} - #pragma omp declare variant (device_arch_nvptx) match(construct={target},device={arch(nvptx)}) #pragma omp declare variant (device_arch_gcn) match(construct={target},device={arch(gcn)}) -#pragma omp declare variant (device_arch_intel_mic) match(construct={target},device={arch(intel_mic)}) /* static */ int device_arch (void) { @@ -49,31 +42,3 @@ on_device_arch_gcn () { return on_device_arch (GOMP_DEVICE_GCN); } - -int -on_device_arch_intel_mic () -{ - return on_device_arch (GOMP_DEVICE_INTEL_MIC); -} - -static int -any_device_arch (int d) -{ - int nd = omp_get_num_devices (); - for (int i = 0; i < nd; ++i) - { - int d_cur; - #pragma omp target device(i) map(from:d_cur) - d_cur = device_arch (); - if (d_cur == d) - return 1; - } - - return 0; -} - -int -any_device_arch_intel_mic () -{ - return any_device_arch (GOMP_DEVICE_INTEL_MIC); -} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-45.c b/libgomp/testsuite/libgomp.c-c++-common/target-45.c index 27bbedd..73c105d 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-45.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-45.c @@ -1,5 +1,3 @@ -/* { dg-xfail-run-if TODO { offload_device_any_intel_mic } } */ - #include <omp.h> #include <stdlib.h> diff --git a/libgomp/testsuite/libgomp.fortran/target-11.f90 b/libgomp/testsuite/libgomp.fortran/target-11.f90 new file mode 100644 index 0000000..b0faa2e --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-11.f90 @@ -0,0 +1,75 @@ +! Based on libgomp.c/target-23.c + +! { dg-additional-options "-fdump-tree-original" } +! { dg-final { scan-tree-dump "omp target update to\\(xxs\\\[3\\\] \\\[len: 2\\\]\\)" "original" } } +! { dg-final { scan-tree-dump "omp target update to\\(s\\.s \\\[len: 4\\\]\\)" "original" } } +! { dg-final { scan-tree-dump "omp target update from\\(s\\.s \\\[len: 4\\\]\\)" "original" } } + +module m + implicit none + type S_type + integer s + integer, pointer :: u(:) => null() + integer :: v(0:4) + end type S_type + integer, volatile :: z +end module m + +program main + use m + implicit none + integer, target :: u(0:9) = [0, 1, 2, 3, 4, 5, 6, 7, 8, 9] + logical :: err + type (S_type) :: s + integer, pointer :: v(:) + integer(kind=2) :: xxs(5) + err = .false. + s = S_type(9, v=[10, 11, 12, 13, 14]) + s%u(0:) => u(3:) + v(-4+3:) => u(3:) + xxs = [-1,-2,-3,-4,-5] + !$omp target enter data map (to: s%s, s%u, s%u(0:5)) map (alloc: s%v(1:4), xxs(3:5)) + s%s = s%s + 1 + u(3) = u(3) + 1 + s%v(1) = s%v(1) + 1 + xxs(3) = -33 + xxs(4) = -44 + xxs(5) = -55 + !$omp target update to (xxs(4)) + !$omp target update to (s%s) to (s%u(0:2), s%v(1:4)) + + !$omp target map (alloc: s%s, s%v(1:4)) map (from: err) + err = .false. + if (s%s /= 10 .or. s%v(1) /= 12 .or. s%v(2) /= 12 .or. s%v(3) /= 13) & + err = .true. + if (v(-1) /= 4 .or. v(0) /= 4 .or. v(1) /= 5 .or. v(2) /= 6 .or. v(3) /= 7) & + err = .true. + if (xxs(4) /= -44) & + err = .true. + s%s = s%s + 1 + s%v(2) = s%v(2) + 2 + v(-1) = 5 + v(3) = 9 + !$omp end target + + if (err) & + error stop + + !$omp target map (alloc: s%u(0:5)) + err = .false. + if (s%u(0) /= 5 .or. s%u(1) /= 4 .or. s%u(2) /= 5 .or. s%u(3) /= 6 .or. s%u(4) /= 9) & + err = .true. + s%u(1) = 12 + !$omp end target + + !$omp target update from (s%s, s%u(0:5)) from (s%v(1:4)) + if (err .or. s%s /= 11 .or. u(0) /= 0 .or. u(1) /= 1 .or. u(2) /= 2 .or. u(3) /= 5 & + .or. u(4) /= 12 .or. u(5) /= 5 .or. u(6) /= 6 .or. u(7) /= 9 .or. u(8) /= 8 & + .or. u(9) /= 9 .or. s%v(0) /= 10 .or. s%v(1) /= 12 .or. s%v(2) /= 14 & + .or. s%v(3) /= 13 .or. s%v(4) /= 14) & + error stop + ! !$omp target exit data map (release: s%s) + ! !$omp target exit data map (release: s%u(0:5)) + ! !$omp target exit data map (delete: s%v(1:4)) + ! !$omp target exit data map (release: s%s) +end diff --git a/libgomp/testsuite/libgomp.fortran/target-13.f90 b/libgomp/testsuite/libgomp.fortran/target-13.f90 new file mode 100644 index 0000000..6aacc77 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-13.f90 @@ -0,0 +1,159 @@ +module m + implicit none + type t + integer :: s, a(5) + end type t + + type t2 + integer :: s, a(5) + type(t) :: st, at(2:3) + end type t2 + + interface operator(/=) + procedure ne_compare_t + procedure ne_compare_t2 + end interface + +contains + + logical pure elemental function ne_compare_t (a, b) result(res) + type(t), intent(in) :: a, b + res = (a%s /= b%s) .or. any(a%a /= b%a) + end function + + logical pure elemental function ne_compare_t2 (a, b) result(res) + type(t2), intent(in) :: a, b + res = (a%s /= b%s) .or. any(a%a /= b%a) & + .or. (a%st /= b%st) .or. any(a%at /= b%at) + end function +end module m + +program p +use m +implicit none + +type(t2) :: var1, var2(5), var3(:) +type(t2) :: var1a, var2a(5), var3a(:) +allocatable :: var3, var3a +logical :: shared_memory = .false. + +!$omp target map(to: shared_memory) + shared_memory = .true. +!$omp end target + +var1 = T2(1, [1,2,3,4,5], T(11, [11,22,33,44,55]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]) + +var2 = [T2(101, [201,202,203,204,205], T(2011, [2011,2022,2033,2044,2055]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]), & + T2(111, [211,212,213,214,215], T(2111, [2111,2122,2133,2144,2155]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]), & + T2(121, [221,222,223,224,225], T(2211, [2211,2222,2233,2244,2255]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]), & + T2(131, [231,232,233,234,235], T(2311, [2311,2322,2333,2344,2355]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]), & + T2(141, [241,242,243,244,245], T(2411, [2411,2422,2433,2444,2455]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])])] + +var3 = [T2(301, [401,402,403,404,405], T(4011, [4011,4022,4033,4044,4055]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]), & + T2(311, [411,412,413,414,415], T(4111, [4111,4122,4133,4144,4155]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]), & + T2(321, [421,422,423,424,425], T(4211, [4211,4222,4233,4244,4255]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]), & + T2(331, [431,432,433,434,435], T(4311, [4311,4322,4333,4344,4355]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]), & + T2(341, [441,442,443,444,445], T(4411, [4411,4422,4433,4444,4455]), & + [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])])] + +var1a = var1 +var2a = var2 +var3a = var3 + +!$omp target enter data map(to:var1) +!$omp target enter data map(to:var2) +!$omp target enter data map(to:var3) + +! --------------- + +!$omp target update from(var1%at(2:3)) + +if (var1a /= var1) error stop +if (any (var2a /= var2)) error stop +if (any (var3a /= var3)) error stop + +! --------------- + +!$omp target + var1%st%s = 1243 + var2(3)%at(2) = T(123, [345,64,356,39,13]) + var2(3)%at(3) = T(48, [74,162,572,357,3]) +!$omp end target + +if (.not. shared_memory) then + if (var1 /= var1) error stop + if (any (var2a /= var2)) error stop + if (any (var3a /= var3)) error stop +endif + +!$omp target update from(var1%st) from(var2(3)%at(2:3)) + +var1a%st%s = 1243 +var2a(3)%at(2) = T(123, [345,64,356,39,13]) +var2a(3)%at(3) = T(48, [74,162,572,357,3]) +if (var1 /= var1) error stop +if (any (var2a /= var2)) error stop +if (any (var3a /= var3)) error stop + +! --------------- + +var3(1) = var2(1) +var1%at(2)%a = var2(1)%a +var1%at(3)%a = var2(2)%a + +var1a = var1 +var2a = var2 +var3a = var3 + +!$omp target update to(var3) to(var1%at(2:3)) + +!$omp target + var3(1)%s = var3(1)%s + 123 + var1%at(2)%a = var1%at(2)%a * 7 + var1%at(3)%s = var1%at(3)%s * (-3) +!$omp end target + +if (.not. shared_memory) then + if (var1 /= var1) error stop + if (any (var2a /= var2)) error stop + if (any (var3a /= var3)) error stop +endif + +var3a(1)%s = var3a(1)%s + 123 +var1a%at(2)%a = var1a%at(2)%a * 7 +var1a%at(3)%s = var1a%at(3)%s * (-3) + +block + integer, volatile :: i1,i2,i3,i4 + i1 = 1 + i2 = 2 + i3 = 1 + i4 = 2 + !$omp target update from(var3(i1:i2)) from(var1%at(i3:i4)) + i1 = 3 + i2 = 3 + i3 = 1 + i4 = 5 + !$omp target update from(var1%at(i1)%s) from(var1%at(i2)%a(i3:i4)) +end block + +if (var1 /= var1) error stop +if (any (var2a /= var2)) error stop +if (any (var3a /= var3)) error stop + +! --------------- + +!$omp target exit data map(from:var1) +!$omp target exit data map(from:var2) +!$omp target exit data map(from:var3) +end diff --git a/libgomp/testsuite/libgomp.fortran/target10.f90 b/libgomp/testsuite/libgomp.fortran/target10.f90 index 3145255..4876cce 100644 --- a/libgomp/testsuite/libgomp.fortran/target10.f90 +++ b/libgomp/testsuite/libgomp.fortran/target10.f90 @@ -1,5 +1,4 @@ ! { dg-do run } -! { dg-xfail-run-if TODO { offload_device_any_intel_mic } } program main use omp_lib diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 index 51776a1..89bd4a2 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 @@ -215,7 +215,7 @@ program main ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-1 } ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 } ! { dg-note {variable 'S\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 } - ! { dg-note {variable 'desc\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-4 } + ! { dg-note {variable 'desc\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-4 } use vars use openacc implicit none diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 new file mode 100644 index 0000000..759873b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 @@ -0,0 +1,278 @@ +! Test OpenACC 'declare create' with allocatable arrays. + +! { dg-do run } + +!TODO-OpenACC-declare-allocate +! Missing support for OpenACC "Changes from Version 2.0 to 2.5": +! "The 'declare create' directive with a Fortran 'allocatable' has new behavior". +! Thus, after 'allocate'/before 'deallocate', do +! '!$acc enter data create'/'!$acc exit data delete' manually. + +!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'. + +! { dg-additional-options -fopt-info-all-omp } +! { dg-additional-options -foffload=-fopt-info-all-omp } + +! { dg-additional-options --param=openacc-privatization=noisy } +! { dg-additional-options -foffload=--param=openacc-privatization=noisy } +! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): +! { dg-prune-output {note: variable '[Di]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } + +! { dg-additional-options -Wopenacc-parallelism } + +! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' +! passed to 'incr' may be unset, and in that case, it will be set to [...]", +! so to maintain compatibility with earlier Tcl releases, we manually +! initialize counter variables: +! { dg-line l_dummy[variable c 0] } +! { dg-message dummy {} { target iN-VAl-Id } l_dummy } to avoid +! "WARNING: dg-line var l_dummy defined, but not used". + + +module vars + implicit none + integer, parameter :: n = 100 + real*8, allocatable :: b(:) + !$acc declare create (b) +end module vars + +program test + use vars + use openacc + implicit none + real*8 :: a + integer :: i + + interface + subroutine sub1 + !$acc routine gang + end subroutine sub1 + + subroutine sub2 + end subroutine sub2 + + real*8 function fun1 (ix) + integer ix + !$acc routine seq + end function fun1 + + real*8 function fun2 (ix) + integer ix + !$acc routine seq + end function fun2 + end interface + + if (allocated (b)) error stop + + ! Test local usage of an allocated declared array. + + allocate (b(n)) + !$acc enter data create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + a = 2.0 + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = i * a + end do + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i*a) error stop + end do + + !$acc exit data delete (b) + deallocate (b) + + ! Test the usage of an allocated declared array inside an acc + ! routine subroutine. + + allocate (b(n)) + !$acc enter data create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel + call sub1 ! { dg-line l[incr c] } + ! { dg-optimized {assigned OpenACC gang worker vector loop parallelism} {} { target *-*-* } l$c } + !$acc end parallel + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i*2) error stop + end do + + !$acc exit data delete (b) + deallocate (b) + + ! Test the usage of an allocated declared array inside a host + ! subroutine. + + call sub2 + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= 1.0) error stop + end do + + !$acc exit data delete (b) + deallocate (b) + + if (allocated (b)) error stop + + ! Test the usage of an allocated declared array inside an acc + ! routine function. + + allocate (b(n)) + !$acc enter data create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = fun1 (i) ! { dg-line l[incr c] } + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l$c } + end do + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i) error stop + end do + + !$acc exit data delete (b) + deallocate (b) + + ! Test the usage of an allocated declared array inside a host + ! function. + + allocate (b(n)) + !$acc enter data create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do + + !$acc update host(b) + + do i = 1, n + b(i) = fun2 (i) + end do + + if (.not.acc_is_present (b)) error stop + + do i = 1, n + if (b(i) /= i*i) error stop + end do + + !$acc exit data delete (b) + deallocate (b) +end program test ! { dg-line l[incr c] } +! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c } + +! Set each element in array 'b' at index i to i*2. + +subroutine sub1 ! { dg-line subroutine_sub1 } + use vars + implicit none + integer i + !$acc routine gang + ! { dg-bogus {[Ww]arning: region is worker partitioned but does not contain worker partitioned code} {TODO default 'gang' 'vector'} { xfail *-*-* } subroutine_sub1 } + + !$acc loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = i*2 + end do +end subroutine sub1 + +! Allocate array 'b', and set it to all 1.0. + +subroutine sub2 + use vars + use openacc + implicit none + integer i + + allocate (b(n)) + !$acc enter data create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do +end subroutine sub2 + +! Return b(i) * i; + +real*8 function fun1 (i) + use vars + implicit none + integer i + !$acc routine seq + + fun1 = b(i) * i +end function fun1 + +! Return b(i) * i * i; + +real*8 function fun2 (i) + use vars + implicit none + integer i + + fun2 = b(i) * i * i +end function fun2 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 new file mode 100644 index 0000000..e4cb9c3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 @@ -0,0 +1,278 @@ +! Test OpenACC 'declare create' with allocatable arrays. + +! { dg-do run } + +!TODO-OpenACC-declare-allocate +! Missing support for OpenACC "Changes from Version 2.0 to 2.5": +! "The 'declare create' directive with a Fortran 'allocatable' has new behavior". +! Thus, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete' +! manually. + +!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'. + +! { dg-additional-options -fopt-info-all-omp } +! { dg-additional-options -foffload=-fopt-info-all-omp } + +! { dg-additional-options --param=openacc-privatization=noisy } +! { dg-additional-options -foffload=--param=openacc-privatization=noisy } +! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): +! { dg-prune-output {note: variable '[Di]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } + +! { dg-additional-options -Wopenacc-parallelism } + +! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' +! passed to 'incr' may be unset, and in that case, it will be set to [...]", +! so to maintain compatibility with earlier Tcl releases, we manually +! initialize counter variables: +! { dg-line l_dummy[variable c 0] } +! { dg-message dummy {} { target iN-VAl-Id } l_dummy } to avoid +! "WARNING: dg-line var l_dummy defined, but not used". + + +module vars + implicit none + integer, parameter :: n = 100 + real*8, allocatable :: b(:) + !$acc declare create (b) +end module vars + +program test + use vars + use openacc + implicit none + real*8 :: a + integer :: i + + interface + subroutine sub1 + !$acc routine gang + end subroutine sub1 + + subroutine sub2 + end subroutine sub2 + + real*8 function fun1 (ix) + integer ix + !$acc routine seq + end function fun1 + + real*8 function fun2 (ix) + integer ix + !$acc routine seq + end function fun2 + end interface + + if (allocated (b)) error stop + + ! Test local usage of an allocated declared array. + + allocate (b(n)) + call acc_create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + a = 2.0 + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = i * a + end do + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i*a) error stop + end do + + call acc_delete (b) + deallocate (b) + + ! Test the usage of an allocated declared array inside an acc + ! routine subroutine. + + allocate (b(n)) + call acc_create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel + call sub1 ! { dg-line l[incr c] } + ! { dg-optimized {assigned OpenACC gang worker vector loop parallelism} {} { target *-*-* } l$c } + !$acc end parallel + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i*2) error stop + end do + + call acc_delete (b) + deallocate (b) + + ! Test the usage of an allocated declared array inside a host + ! subroutine. + + call sub2 + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= 1.0) error stop + end do + + call acc_delete (b) + deallocate (b) + + if (allocated (b)) error stop + + ! Test the usage of an allocated declared array inside an acc + ! routine function. + + allocate (b(n)) + call acc_create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = fun1 (i) ! { dg-line l[incr c] } + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l$c } + end do + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i) error stop + end do + + call acc_delete (b) + deallocate (b) + + ! Test the usage of an allocated declared array inside a host + ! function. + + allocate (b(n)) + call acc_create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do + + !$acc update host(b) + + do i = 1, n + b(i) = fun2 (i) + end do + + if (.not.acc_is_present (b)) error stop + + do i = 1, n + if (b(i) /= i*i) error stop + end do + + call acc_delete (b) + deallocate (b) +end program test ! { dg-line l[incr c] } +! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c } + +! Set each element in array 'b' at index i to i*2. + +subroutine sub1 ! { dg-line subroutine_sub1 } + use vars + implicit none + integer i + !$acc routine gang + ! { dg-bogus {[Ww]arning: region is worker partitioned but does not contain worker partitioned code} {TODO default 'gang' 'vector'} { xfail *-*-* } subroutine_sub1 } + + !$acc loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = i*2 + end do +end subroutine sub1 + +! Allocate array 'b', and set it to all 1.0. + +subroutine sub2 + use vars + use openacc + implicit none + integer i + + allocate (b(n)) + call acc_create (b) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do +end subroutine sub2 + +! Return b(i) * i; + +real*8 function fun1 (i) + use vars + implicit none + integer i + !$acc routine seq + + fun1 = b(i) * i +end function fun1 + +! Return b(i) * i * i; + +real*8 function fun2 (i) + use vars + implicit none + integer i + + fun2 = b(i) * i * i +end function fun2 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 new file mode 100644 index 0000000..1c8ccd9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 @@ -0,0 +1,268 @@ +! Test OpenACC 'declare create' with allocatable arrays. + +! { dg-do run } + +!TODO-OpenACC-declare-allocate +! Not currently implementing correct '-DACC_MEM_SHARED=0' behavior: +! Missing support for OpenACC "Changes from Version 2.0 to 2.5": +! "The 'declare create' directive with a Fortran 'allocatable' has new behavior". +! { dg-xfail-run-if TODO { *-*-* } { -DACC_MEM_SHARED=0 } } + +!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'. + +! { dg-additional-options -fopt-info-all-omp } +! { dg-additional-options -foffload=-fopt-info-all-omp } + +! { dg-additional-options --param=openacc-privatization=noisy } +! { dg-additional-options -foffload=--param=openacc-privatization=noisy } +! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): +! { dg-prune-output {note: variable '[Di]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } + +! { dg-additional-options -Wopenacc-parallelism } + +! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' +! passed to 'incr' may be unset, and in that case, it will be set to [...]", +! so to maintain compatibility with earlier Tcl releases, we manually +! initialize counter variables: +! { dg-line l_dummy[variable c 0] } +! { dg-message dummy {} { target iN-VAl-Id } l_dummy } to avoid +! "WARNING: dg-line var l_dummy defined, but not used". + + +module vars + implicit none + integer, parameter :: n = 100 + real*8, allocatable :: b(:) + !$acc declare create (b) +end module vars + +program test + use vars + use openacc + implicit none + real*8 :: a + integer :: i + + interface + subroutine sub1 + !$acc routine gang + end subroutine sub1 + + subroutine sub2 + end subroutine sub2 + + real*8 function fun1 (ix) + integer ix + !$acc routine seq + end function fun1 + + real*8 function fun2 (ix) + integer ix + !$acc routine seq + end function fun2 + end interface + + if (allocated (b)) error stop + + ! Test local usage of an allocated declared array. + + allocate (b(n)) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + a = 2.0 + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = i * a + end do + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i*a) error stop + end do + + deallocate (b) + + ! Test the usage of an allocated declared array inside an acc + ! routine subroutine. + + allocate (b(n)) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel + call sub1 ! { dg-line l[incr c] } + ! { dg-optimized {assigned OpenACC gang worker vector loop parallelism} {} { target *-*-* } l$c } + !$acc end parallel + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i*2) error stop + end do + + deallocate (b) + + ! Test the usage of an allocated declared array inside a host + ! subroutine. + + call sub2 + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= 1.0) error stop + end do + + deallocate (b) + + if (allocated (b)) error stop + + ! Test the usage of an allocated declared array inside an acc + ! routine function. + + allocate (b(n)) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = fun1 (i) ! { dg-line l[incr c] } + ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l$c } + end do + + if (.not.acc_is_present (b)) error stop + + !$acc update host(b) + + do i = 1, n + if (b(i) /= i) error stop + end do + + deallocate (b) + + ! Test the usage of an allocated declared array inside a host + ! function. + + allocate (b(n)) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c } + ! { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c } + ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do + + !$acc update host(b) + + do i = 1, n + b(i) = fun2 (i) + end do + + if (.not.acc_is_present (b)) error stop + + do i = 1, n + if (b(i) /= i*i) error stop + end do + + deallocate (b) +end program test ! { dg-line l[incr c] } +! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c } +! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c } + +! Set each element in array 'b' at index i to i*2. + +subroutine sub1 ! { dg-line subroutine_sub1 } + use vars + implicit none + integer i + !$acc routine gang + ! { dg-bogus {[Ww]arning: region is worker partitioned but does not contain worker partitioned code} {TODO default 'gang' 'vector'} { xfail *-*-* } subroutine_sub1 } + + !$acc loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = i*2 + end do +end subroutine sub1 + +! Allocate array 'b', and set it to all 1.0. + +subroutine sub2 + use vars + use openacc + implicit none + integer i + + allocate (b(n)) + + if (.not.allocated (b)) error stop + if (.not.acc_is_present (b)) error stop + + !$acc parallel loop ! { dg-line l[incr c] } + ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c } + ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c } + do i = 1, n + b(i) = 1.0 + end do +end subroutine sub2 + +! Return b(i) * i; + +real*8 function fun1 (i) + use vars + implicit none + integer i + !$acc routine seq + + fun1 = b(i) * i +end function fun1 + +! Return b(i) * i * i; + +real*8 function fun2 (i) + use vars + implicit none + integer i + + fun2 = b(i) * i * i +end function fun2 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90 new file mode 100644 index 0000000..6604f72 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90 @@ -0,0 +1,438 @@ +! Test OpenACC 'declare create' with allocatable arrays. + +! { dg-do run } + +! Note that we're not testing OpenACC semantics here, but rather documenting +! current GCC behavior, specifically, behavior concerning updating of +! host/device array descriptors. +! { dg-skip-if n/a { *-*-* } { -DACC_MEM_SHARED=1 } } + +!TODO-OpenACC-declare-allocate +! Missing support for OpenACC "Changes from Version 2.0 to 2.5": +! "The 'declare create' directive with a Fortran 'allocatable' has new behavior". +! Thus, after 'allocate'/before 'deallocate', do +! '!$acc enter data create'/'!$acc exit data delete' manually. + + +!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'. + + +!TODO OpenACC 'serial' vs. GCC/nvptx: +!TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} } + + +! { dg-additional-options -fdump-tree-original } +! { dg-additional-options -fdump-tree-gimple } + + +module vars + implicit none + integer, parameter :: n1_lb = -3 + integer, parameter :: n1_ub = 6 + integer, parameter :: n2_lb = -9999 + integer, parameter :: n2_ub = 22222 + + integer, allocatable :: b(:) + !$acc declare create (b) + +end module vars + +program test + use vars + use openacc + implicit none + integer :: i + + ! Identifiers for purposes of reliable '-fdump-tree-[...]' scanning. + integer :: id1_1, id1_2 + + interface + + subroutine verify_initial + implicit none + !$acc routine seq + end subroutine verify_initial + + subroutine verify_n1_allocated + implicit none + !$acc routine seq + end subroutine verify_n1_allocated + + subroutine verify_n1_values (addend) + implicit none + !$acc routine gang + integer, value :: addend + end subroutine verify_n1_values + + subroutine verify_n1_deallocated (expect_allocated) + implicit none + !$acc routine seq + logical, value :: expect_allocated + end subroutine verify_n1_deallocated + + subroutine verify_n2_allocated + implicit none + !$acc routine seq + end subroutine verify_n2_allocated + + subroutine verify_n2_values (addend) + implicit none + !$acc routine gang + integer, value :: addend + end subroutine verify_n2_values + + subroutine verify_n2_deallocated (expect_allocated) + implicit none + !$acc routine seq + logical, value :: expect_allocated + end subroutine verify_n2_deallocated + + end interface + + call acc_create (id1_1) + call acc_create (id1_2) + + call verify_initial + ! It is important here (and similarly, following) that there is no data + ! clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'. + !$acc serial + call verify_initial + !$acc end serial + + allocate (b(n1_lb:n1_ub)) + call verify_n1_allocated + if (acc_is_present (b)) error stop + !$acc enter data create (b) + ! This is now OpenACC "present": + if (.not.acc_is_present (b)) error stop + ! ..., and got the actual array descriptor installed: + !$acc serial + call verify_n1_allocated + !$acc end serial + + do i = n1_lb, n1_ub + b(i) = i - 1 + end do + + ! In 'declare-allocatable-array_descriptor-1-runtime.f90', this does "verify + ! that host-to-device copy doesn't touch the device-side (still initial) + ! array descriptor (but it does copy the array data"). This is here not + ! applicable anymore, as we've already gotten the actual array descriptor + ! installed. Thus now verify that it does copy the array data. + call acc_update_device (b) + !$acc serial + call verify_n1_allocated + !$acc end serial + + b = 40 + + !$acc parallel copyout (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'. + call verify_n1_values (-1) + id1_1 = 0 + !$acc end parallel + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(from:id1_1\)$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } } + + !$acc parallel copy (b) copyout (id1_2) + ! As already present, 'copy (b)' doesn't copy; addend is still '-1'. + call verify_n1_values (-1) + id1_2 = 0 + !$acc end parallel + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } } + !TODO ..., but without an actual use of 'b', the gimplifier removes the + !TODO 'GOMP_MAP_TO_PSET': + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + + ! In 'declare-allocatable-array_descriptor-1-runtime.f90', this does "verify + ! that device-to-host copy doesn't touch the host-side array descriptor, + ! doesn't copy out the device-side (still initial) array descriptor (but it + ! does copy the array data)". This is here not applicable anymore, as we've + ! already gotten the actual array descriptor installed. Thus now verify that + ! it does copy the array data. + call acc_update_self (b) + call verify_n1_allocated + + do i = n1_lb, n1_ub + if (b(i) /= i - 1) error stop + b(i) = b(i) + 2 + end do + + ! The same using the OpenACC 'update' directive. + + !$acc update device (b) self (id1_1) + ! We do have 'GOMP_MAP_TO_PSET' here: + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_to:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1\);$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_to:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } } + ! ..., but it's silently skipped in 'GOACC_update'. + !$acc serial + call verify_n1_allocated + !$acc end serial + + b = 41 + + !$acc parallel + call verify_n1_values (1) + !$acc end parallel + + !$acc parallel copy (b) + call verify_n1_values (1) + !$acc end parallel + + !$acc update self (b) self (id1_2) + ! We do have 'GOMP_MAP_TO_PSET' here: + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_from:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2\);$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_from:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + ! ..., but it's silently skipped in 'GOACC_update'. + call verify_n1_allocated + + do i = n1_lb, n1_ub + if (b(i) /= i + 1) error stop + b(i) = b(i) + 2 + end do + + ! Now test that (potentially re-)installing the actual array descriptor is a + ! no-op, via a data clause for 'b' (explicit or implicit): must get a + ! 'GOMP_MAP_TO_PSET'. + !$acc serial present (b) copyin (id1_2) + call verify_n1_allocated + !TODO Use of 'b': + id1_2 = ubound (b, 1) + !$acc end serial + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + + !$acc parallel copyin (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'. + call verify_n1_values (1) + id1_1 = 0 + !$acc end parallel + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(to:id1_1\)$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } } + + !$acc parallel copy (b) copyin (id1_2) + ! As already present, 'copy (b)' doesn't copy; addend is still '1'. + call verify_n1_values (1) + id1_2 = 0 + !$acc end parallel + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } } + !TODO ..., but without an actual use of 'b', the gimplifier removes the + !TODO 'GOMP_MAP_TO_PSET': + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + + call verify_n1_allocated + if (.not.acc_is_present (b)) error stop + + !$acc exit data delete (b) + if (.not.allocated (b)) error stop + if (acc_is_present (b)) error stop + ! The device-side array descriptor doesn't get updated, so 'b' still appears + ! as "allocated": + !$acc serial + call verify_n1_allocated + !$acc end serial + + deallocate (b) + call verify_n1_deallocated (.false.) + ! The device-side array descriptor doesn't get updated, so 'b' still appears + ! as "allocated": + !$acc serial + call verify_n1_allocated + !$acc end serial + + ! Now try to install the actual array descriptor, via a data clause for 'b' + ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in + ! 'gomp_map_vars_internal' is handled as 'declare target', but because of + ! '*(void **) hostaddrs[i] == NULL', we've got 'has_always_ptrset == false', + ! 'always_to_cnt == 0', and therefore 'gomp_map_vars_existing' doesn't update + ! the 'GOMP_MAP_TO_PSET'. + ! The device-side array descriptor doesn't get updated, so 'b' still appears + ! as "allocated": + !TODO Why does 'present (b)' still work here? + !$acc serial present (b) copyout (id1_2) + call verify_n1_deallocated (.true.) + !TODO Use of 'b'. + id1_2 = ubound (b, 1) + !$acc end serial + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + + + ! Restart the procedure, with different array dimensions. + + allocate (b(n2_lb:n2_ub)) + call verify_n2_allocated + if (acc_is_present (b)) error stop + !$acc enter data create (b) + if (.not.acc_is_present (b)) error stop + ! ..., and got the actual array descriptor installed: + !$acc serial + call verify_n2_allocated + !$acc end serial + + do i = n2_lb, n2_ub + b(i) = i + 20 + end do + + call acc_update_device (b) + !$acc serial + call verify_n2_allocated + !$acc end serial + + b = -40 + + !$acc parallel + call verify_n2_values (20) + !$acc end parallel + + !$acc parallel copy (b) + call verify_n2_values (20) + !$acc end parallel + + call acc_update_self (b) + call verify_n2_allocated + + do i = n2_lb, n2_ub + if (b(i) /= i + 20) error stop + b(i) = b(i) - 40 + end do + + !$acc update device (b) + !$acc serial + call verify_n2_allocated + !$acc end serial + + b = -41 + + !$acc parallel + call verify_n2_values (-20) + !$acc end parallel + + !$acc parallel copy (b) + call verify_n2_values (-20) + !$acc end parallel + + !$acc update self (b) + call verify_n2_allocated + + do i = n2_lb, n2_ub + if (b(i) /= i - 20) error stop + b(i) = b(i) + 10 + end do + + !$acc serial present (b) copy (id1_2) + call verify_n2_allocated + !TODO Use of 'b': + id1_2 = ubound (b, 1) + !$acc end serial + + !$acc parallel + call verify_n2_values (-20) + !$acc end parallel + + !$acc parallel copy (b) + call verify_n2_values (-20) + !$acc end parallel + + call verify_n2_allocated + if (.not.acc_is_present (b)) error stop + + !$acc exit data delete (b) + if (.not.allocated (b)) error stop + if (acc_is_present (b)) error stop + !$acc serial + call verify_n2_allocated + !$acc end serial + + deallocate (b) + call verify_n2_deallocated (.false.) + !$acc serial + call verify_n2_allocated + !$acc end serial + + !$acc serial present (b) copy (id1_2) + call verify_n2_deallocated (.true.) + !TODO Use of 'b': + id1_2 = ubound (b, 1) + !$acc end serial + +end program test + + +subroutine verify_initial + use vars + implicit none + !$acc routine seq + + if (allocated (b)) error stop "verify_initial allocated" + if (any (lbound (b) /= [0])) error stop "verify_initial lbound" + if (any (ubound (b) /= [0])) error stop "verify_initial ubound" +end subroutine verify_initial + +subroutine verify_n1_allocated + use vars + implicit none + !$acc routine seq + + if (.not.allocated (b)) error stop "verify_n1_allocated allocated" + if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_allocated lbound" + if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_allocated ubound" +end subroutine verify_n1_allocated + +subroutine verify_n1_values (addend) + use vars + implicit none + !$acc routine gang + integer, value :: addend + integer :: i + + !$acc loop + do i = n1_lb, n1_ub + if (b(i) /= i + addend) error stop + end do +end subroutine verify_n1_values + +subroutine verify_n1_deallocated (expect_allocated) + use vars + implicit none + !$acc routine seq + logical, value :: expect_allocated + + if (allocated(b) .neqv. expect_allocated) error stop "verify_n1_deallocated allocated" + ! Apparently 'deallocate'ing doesn't unset the bounds. + if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_deallocated lbound" + if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_deallocated ubound" +end subroutine verify_n1_deallocated + +subroutine verify_n2_allocated + use vars + implicit none + !$acc routine seq + + if (.not.allocated(b)) error stop "verify_n2_allocated allocated" + if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_allocated lbound" + if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_allocated ubound" +end subroutine verify_n2_allocated + +subroutine verify_n2_values (addend) + use vars + implicit none + !$acc routine gang + integer, value :: addend + integer :: i + + !$acc loop + do i = n2_lb, n2_ub + if (b(i) /= i + addend) error stop + end do +end subroutine verify_n2_values + +subroutine verify_n2_deallocated (expect_allocated) + use vars + implicit none + !$acc routine seq + logical, value :: expect_allocated + + if (allocated(b) .neqv. expect_allocated) error stop "verify_n2_deallocated allocated" + ! Apparently 'deallocate'ing doesn't unset the bounds. + if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_deallocated lbound" + if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_deallocated ubound" +end subroutine verify_n2_deallocated diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90 new file mode 100644 index 0000000..b27f312 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90 @@ -0,0 +1,402 @@ +! Test OpenACC 'declare create' with allocatable arrays. + +! { dg-do run } + +! Note that we're not testing OpenACC semantics here, but rather documenting +! current GCC behavior, specifically, behavior concerning updating of +! host/device array descriptors. +! { dg-skip-if n/a { *-*-* } { -DACC_MEM_SHARED=1 } } + +!TODO-OpenACC-declare-allocate +! Missing support for OpenACC "Changes from Version 2.0 to 2.5": +! "The 'declare create' directive with a Fortran 'allocatable' has new behavior". +! Thus, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete' +! manually. + + +!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'. + + +!TODO OpenACC 'serial' vs. GCC/nvptx: +!TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} } + + +! { dg-additional-options -fdump-tree-original } +! { dg-additional-options -fdump-tree-gimple } + + +module vars + implicit none + integer, parameter :: n1_lb = -3 + integer, parameter :: n1_ub = 6 + integer, parameter :: n2_lb = -9999 + integer, parameter :: n2_ub = 22222 + + integer, allocatable :: b(:) + !$acc declare create (b) + +end module vars + +program test + use vars + use openacc + implicit none + integer :: i + + ! Identifiers for purposes of reliable '-fdump-tree-[...]' scanning. + integer :: id1_1, id1_2 + + interface + + subroutine verify_initial + implicit none + !$acc routine seq + end subroutine verify_initial + + subroutine verify_n1_allocated + implicit none + !$acc routine seq + end subroutine verify_n1_allocated + + subroutine verify_n1_values (addend) + implicit none + !$acc routine gang + integer, value :: addend + end subroutine verify_n1_values + + subroutine verify_n1_deallocated (expect_allocated) + implicit none + !$acc routine seq + logical, value :: expect_allocated + end subroutine verify_n1_deallocated + + subroutine verify_n2_allocated + implicit none + !$acc routine seq + end subroutine verify_n2_allocated + + subroutine verify_n2_values (addend) + implicit none + !$acc routine gang + integer, value :: addend + end subroutine verify_n2_values + + subroutine verify_n2_deallocated (expect_allocated) + implicit none + !$acc routine seq + logical, value :: expect_allocated + end subroutine verify_n2_deallocated + + end interface + + call acc_create (id1_1) + call acc_create (id1_2) + + call verify_initial + ! It is important here (and similarly, following) that there is no data + ! clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'. + !$acc serial + call verify_initial + !$acc end serial + + allocate (b(n1_lb:n1_ub)) + call verify_n1_allocated + if (acc_is_present (b)) error stop + call acc_create (b) + ! This is now OpenACC "present": + if (.not.acc_is_present (b)) error stop + ! This still has the initial array descriptor: + !$acc serial + call verify_initial + !$acc end serial + + do i = n1_lb, n1_ub + b(i) = i - 1 + end do + + ! Verify that host-to-device copy doesn't touch the device-side (still + ! initial) array descriptor (but it does copy the array data). + call acc_update_device (b) + !$acc serial + call verify_initial + !$acc end serial + + b = 40 + + ! Verify that device-to-host copy doesn't touch the host-side array + ! descriptor, doesn't copy out the device-side (still initial) array + ! descriptor (but it does copy the array data). + call acc_update_self (b) + call verify_n1_allocated + + do i = n1_lb, n1_ub + if (b(i) /= i - 1) error stop + b(i) = b(i) + 2 + end do + + ! The same using the OpenACC 'update' directive. + + !$acc update device (b) self (id1_1) + ! We do have 'GOMP_MAP_TO_PSET' here: + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_to:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1\);$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_to:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } } + ! ..., but it's silently skipped in 'GOACC_update'. + !$acc serial + call verify_initial + !$acc end serial + + b = 41 + + !$acc update self (b) self (id1_2) + ! We do have 'GOMP_MAP_TO_PSET' here: + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_from:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2\);$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_from:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + ! ..., but it's silently skipped in 'GOACC_update'. + call verify_n1_allocated + + do i = n1_lb, n1_ub + if (b(i) /= i + 1) error stop + b(i) = b(i) + 2 + end do + + ! Now install the actual array descriptor, via a data clause for 'b' + ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in + ! 'gomp_map_vars_internal' is handled as 'declare target', and because of + ! '*(void **) hostaddrs[i] != NULL', we've got 'has_always_ptrset == true', + ! 'always_to_cnt == 1', and therefore 'gomp_map_vars_existing' does update + ! the 'GOMP_MAP_TO_PSET'. + !$acc serial present (b) copyin (id1_1) + call verify_initial + id1_1 = 0 + !$acc end serial + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_1\)$} 1 original } } + !TODO ..., but without an actual use of 'b', the gimplifier removes the + !TODO 'GOMP_MAP_TO_PSET': + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } } + !$acc serial present (b) copyin (id1_2) + call verify_n1_allocated + !TODO Use of 'b': + id1_2 = ubound (b, 1) + !$acc end serial + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + + !$acc parallel copyin (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'. + call verify_n1_values (1) + id1_1 = 0 + !$acc end parallel + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(to:id1_1\)$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } } + + !$acc parallel copy (b) copyin (id1_2) + ! As already present, 'copy (b)' doesn't copy; addend is still '1'. + call verify_n1_values (1) + id1_2 = 0 + !$acc end parallel + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } } + !TODO ..., but without an actual use of 'b', the gimplifier removes the + !TODO 'GOMP_MAP_TO_PSET': + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + + call verify_n1_allocated + if (.not.acc_is_present (b)) error stop + + call acc_delete (b) + if (.not.allocated (b)) error stop + if (acc_is_present (b)) error stop + ! The device-side array descriptor doesn't get updated, so 'b' still appears + ! as "allocated": + !$acc serial + call verify_n1_allocated + !$acc end serial + + deallocate (b) + call verify_n1_deallocated (.false.) + ! The device-side array descriptor doesn't get updated, so 'b' still appears + ! as "allocated": + !$acc serial + call verify_n1_allocated + !$acc end serial + + ! Now try to install the actual array descriptor, via a data clause for 'b' + ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in + ! 'gomp_map_vars_internal' is handled as 'declare target', but because of + ! '*(void **) hostaddrs[i] == NULL', we've got 'has_always_ptrset == false', + ! 'always_to_cnt == 0', and therefore 'gomp_map_vars_existing' doesn't update + ! the 'GOMP_MAP_TO_PSET'. + ! The device-side array descriptor doesn't get updated, so 'b' still appears + ! as "allocated": + !TODO Why does 'present (b)' still work here? + !$acc serial present (b) copyout (id1_2) + call verify_n1_deallocated (.true.) + !TODO Use of 'b'. + id1_2 = ubound (b, 1) + !$acc end serial + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } } + ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } } + + + ! Restart the procedure, with different array dimensions. + + allocate (b(n2_lb:n2_ub)) + call verify_n2_allocated + if (acc_is_present (b)) error stop + call acc_create (b) + if (.not.acc_is_present (b)) error stop + ! This still has the previous (n1) array descriptor: + !$acc serial + call verify_n1_deallocated (.true.) + !$acc end serial + + do i = n2_lb, n2_ub + b(i) = i + 20 + end do + + call acc_update_device (b) + !$acc serial + call verify_n1_deallocated (.true.) + !$acc end serial + + b = -40 + + call acc_update_self (b) + call verify_n2_allocated + + do i = n2_lb, n2_ub + if (b(i) /= i + 20) error stop + b(i) = b(i) - 40 + end do + + !$acc update device (b) + !$acc serial + call verify_n1_deallocated (.true.) + !$acc end serial + + b = -41 + + !$acc update self (b) + call verify_n2_allocated + + do i = n2_lb, n2_ub + if (b(i) /= i - 20) error stop + b(i) = b(i) + 10 + end do + + !$acc serial present (b) copy (id1_2) + call verify_n2_allocated + !TODO Use of 'b': + id1_2 = ubound (b, 1) + !$acc end serial + + !$acc parallel + call verify_n2_values (-20) + !$acc end parallel + + !$acc parallel copy (b) + call verify_n2_values (-20) + !$acc end parallel + + call verify_n2_allocated + if (.not.acc_is_present (b)) error stop + + call acc_delete (b) + if (.not.allocated (b)) error stop + if (acc_is_present (b)) error stop + !$acc serial + call verify_n2_allocated + !$acc end serial + + deallocate (b) + call verify_n2_deallocated (.false.) + !$acc serial + call verify_n2_allocated + !$acc end serial + + !$acc serial present (b) copy (id1_2) + call verify_n2_deallocated (.true.) + !TODO Use of 'b': + id1_2 = ubound (b, 1) + !$acc end serial + +end program test + + +subroutine verify_initial + use vars + implicit none + !$acc routine seq + + if (allocated (b)) error stop "verify_initial allocated" + if (any (lbound (b) /= [0])) error stop "verify_initial lbound" + if (any (ubound (b) /= [0])) error stop "verify_initial ubound" +end subroutine verify_initial + +subroutine verify_n1_allocated + use vars + implicit none + !$acc routine seq + + if (.not.allocated (b)) error stop "verify_n1_allocated allocated" + if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_allocated lbound" + if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_allocated ubound" +end subroutine verify_n1_allocated + +subroutine verify_n1_values (addend) + use vars + implicit none + !$acc routine gang + integer, value :: addend + integer :: i + + !$acc loop + do i = n1_lb, n1_ub + if (b(i) /= i + addend) error stop + end do +end subroutine verify_n1_values + +subroutine verify_n1_deallocated (expect_allocated) + use vars + implicit none + !$acc routine seq + logical, value :: expect_allocated + + if (allocated(b) .neqv. expect_allocated) error stop "verify_n1_deallocated allocated" + ! Apparently 'deallocate'ing doesn't unset the bounds. + if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_deallocated lbound" + if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_deallocated ubound" +end subroutine verify_n1_deallocated + +subroutine verify_n2_allocated + use vars + implicit none + !$acc routine seq + + if (.not.allocated(b)) error stop "verify_n2_allocated allocated" + if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_allocated lbound" + if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_allocated ubound" +end subroutine verify_n2_allocated + +subroutine verify_n2_values (addend) + use vars + implicit none + !$acc routine gang + integer, value :: addend + integer :: i + + !$acc loop + do i = n2_lb, n2_ub + if (b(i) /= i + addend) error stop + end do +end subroutine verify_n2_values + +subroutine verify_n2_deallocated (expect_allocated) + use vars + implicit none + !$acc routine seq + logical, value :: expect_allocated + + if (allocated(b) .neqv. expect_allocated) error stop "verify_n2_deallocated allocated" + ! Apparently 'deallocate'ing doesn't unset the bounds. + if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_deallocated lbound" + if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_deallocated ubound" +end subroutine verify_n2_deallocated diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 index 93e9ee0..c3453a5 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 @@ -43,7 +43,7 @@ subroutine foo (p2, parr, host_p, host_parr, cond) ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target { ! openacc_host_selected } } .-2 } ! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 } ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 } - ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target { ! openacc_host_selected } } .-5 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-5 } #if !ACC_MEM_SHARED if (acc_is_present(p, c_sizeof(p))) stop 5 if (acc_is_present(parr, 1)) stop 6 @@ -54,8 +54,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond) ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 } ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 } ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } - ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 } - ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 } + ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 } + ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 } ! not mapped yet, so it will be equal to the host pointer. if (transfer(c_loc(p), host_p) /= host_p) stop 7 @@ -74,9 +74,9 @@ subroutine foo (p2, parr, host_p, host_parr, cond) ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 } ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-7 } - ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 } - ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-9 } - ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-10 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 } + ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-9 } + ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-10 } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-11 } if (.not. acc_is_present(p, c_sizeof(p))) stop 11 if (.not. acc_is_present(parr, 1)) stop 12 @@ -90,8 +90,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond) ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 } ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 } ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } - ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 } - ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 } + ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 } + ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 } #if ACC_MEM_SHARED if (transfer(c_loc(p), host_p) /= host_p) stop 15 @@ -110,8 +110,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond) ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 } ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 } - ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 } - ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 } + ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 } + ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 } #if ACC_MEM_SHARED if (transfer(c_loc(p), host_p) /= host_p) stop 19 if (transfer(c_loc(parr), host_parr) /= host_parr) stop 20 @@ -129,8 +129,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond) ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 } ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 } - ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 } - ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 } + ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 } + ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 } #if ACC_MEM_SHARED if (transfer(c_loc(p), host_p) /= host_p) stop 23 if (transfer(c_loc(parr), host_parr) /= host_parr) stop 24 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 index c6d6764..e0cfd91 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 @@ -382,7 +382,7 @@ program main b(:) = 1.0 !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1) - ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 } #if !ACC_MEM_SHARED if (acc_is_present (a) .eqv. .TRUE.) STOP 21 @@ -396,7 +396,7 @@ program main !$acc data copyin (a(1:N)) if (1 == 1) ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } #if !ACC_MEM_SHARED if (acc_is_present (a) .eqv. .FALSE.) STOP 23 @@ -404,7 +404,7 @@ program main !$acc data copyout (b(1:N)) if (0 == 1) ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } #if !ACC_MEM_SHARED if (acc_is_present (b) .eqv. .TRUE.) STOP 24 #endif @@ -877,7 +877,7 @@ program main b(:) = 1.0 !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1) - ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 } #if !ACC_MEM_SHARED if (acc_is_present (a) .eqv. .TRUE.) STOP 56 @@ -891,7 +891,7 @@ program main !$acc data copyin (a(1:N)) if (1 == 1) ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } #if !ACC_MEM_SHARED if (acc_is_present (a) .eqv. .FALSE.) STOP 58 @@ -899,7 +899,7 @@ program main !$acc data copyout (b(1:N)) if (0 == 1) ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } #if !ACC_MEM_SHARED if (acc_is_present (b) .eqv. .TRUE.) STOP 59 #endif diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr106643-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr106643-1.f90 new file mode 100644 index 0000000..a9c969e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr106643-1.f90 @@ -0,0 +1,83 @@ +! { dg-do run } +! { dg-additional-options -cpp } + + +!TODO OpenACC 'serial' vs. GCC/nvptx: +!TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} } + + +module m_macron + + implicit none + + real(kind(0d0)), allocatable, dimension(:) :: valls + !$acc declare create(valls) + +contains + + subroutine s_macron_compute(size) + + integer :: size + + !$acc routine seq + +#if ACC_MEM_SHARED + if (valls(size) /= 1) error stop +#else + if (valls(size) /= size - 2) error stop +#endif + + valls(size) = size + 2 + + end subroutine s_macron_compute + + subroutine s_macron_init(size) + + integer :: size + + print*, "size=", size + + print*, "allocate(valls(1:size))" + allocate(valls(1:size)) + + print*, "acc enter data create(valls(1:size))" + !$acc enter data create(valls(1:size)) + + print*, "!$acc update device(valls(1:size))" + valls(size) = size - 2 + !$acc update device(valls(1:size)) + + valls(size) = 1 + + !$acc serial + call s_macron_compute(size) + !$acc end serial + + valls(size) = -1 + + !$acc update host(valls(1:size)) +#if ACC_MEM_SHARED + if (valls(size) /= -1) error stop +#else + if (valls(size) /= size + 2) error stop +#endif + + print*, valls(1:size) + + print*, "acc exit data delete(valls)" + !$acc exit data delete(valls) + + end subroutine s_macron_init + +end module m_macron + + +program p_main + + use m_macron + + implicit none + + call s_macron_init(10) + +end program p_main diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 index 42a8538..d2f89d9 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 @@ -6,15 +6,6 @@ ! Separate file 'print-1-nvptx.f90' for nvptx offloading. ! { dg-skip-if "separate file" { offload_target_nvptx } } -! For GCN offloading compilation, when gang-privatizing 'dt_parm.N' -! (see below), we run into an 'gang-private data-share memory exhausted' -! error: the default '-mgang-private-size' is too small. Per -! 'gcc/fortran/trans-io.cc'/'libgfortran/io/io.h', that one is -! 'struct st_parameter_dt', which indeed is rather big. Instead of -! working out its exact size (which may vary per GCC configuration), -! raise '-mgang-private-size' to an arbitrary high value. -! { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mgang-private-size=13579" { target openacc_radeon_accel_selected } } - ! { dg-additional-options "-fopt-info-note-omp" } ! { dg-additional-options "-foffload=-fopt-info-note-omp" } @@ -36,9 +27,7 @@ program main integer :: var = 42 !$acc parallel ! { dg-line l_compute[incr c_compute] } - ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } - ! { dg-note {variable 'dt_parm\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_compute$c_compute } - ! { dg-note {variable 'dt_parm\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_compute$c_compute } + ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {} { target *-*-* } l_compute$c_compute } write (0, '("The answer is ", I2)') var !$acc end parallel diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 index b31f406..498ef70 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 @@ -122,9 +122,7 @@ contains ! { dg-note {variable 'str' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'str' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'str' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop } - ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop } + ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop } ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop } do i = 1, 10 str(i:i) = achar(ichar('A') + i) @@ -167,9 +165,7 @@ contains ! { dg-note {variable 'scalar' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'scalar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop } ! { dg-note {variable 'scalar' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop } - ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop } - ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop } + ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop } do i = 1, 15 scalar(i:i) = achar(ichar('A') + i) end do |