On 05/09/2018 03:50 AM, Thomas Schwinge wrote:
>> In addition to XPASS'ing devicetpr-1.f90, this patch [...]
>
> Apart from one remaining XFAIL for "-Os" (see PR80995), I now too see the
> following XPASSes on my main development machine:
>
> PASS: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test for excess errors)
> PASS: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test
> PASS: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 (test for excess errors)
> PASS: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 execution test
> [-XFAIL:-]{+XPASS:+} libgomp.oacc-fortran/deviceptr-1.f90
> -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
> (test for excess errors)
> PASS: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 execution test
> [-XFAIL:-]{+XPASS:+} libgomp.oacc-fortran/deviceptr-1.f90
> -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O3
> -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions
> (test for excess errors)
> PASS: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -O3 -fomit-frame-pointer
> -funroll-loops -fpeel-loops -ftracer -finline-functions execution test
> [-XFAIL:-]{+XPASS:+} libgomp.oacc-fortran/deviceptr-1.f90
> -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O3 -g
> (test for excess errors)
> PASS: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -O3 -g execution test
> XFAIL: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -Os (test for excess errors)
> PASS: libgomp.oacc-fortran/deviceptr-1.f90 -DACC_DEVICE_TYPE_nvidia=1
> -DACC_MEM_SHARED=0 -foffload=nvptx-none -Os execution test
>
>> I've applied this patch to og7 [...]. It was tempting to remove the
>> XFAIL from deviceptr-1.f90, but the test case still fails on at least
>> one legacy driver.
>
> That's surprising. These XFAILs were because "OpenACC kernels construct
> will be executed sequentially", so shouldn't have any relationship to
> Nvidia driver versions. If you identified such a problem (which versions
> and hardware exactly?), that's a separate problam and needs to be filed
> as a new issue, and the reference in the test case file updated. So
> please verify that, and/or alternatively remove the non-"-Os" XFAILs.
You're correct. On further inspection, only -Os fails. The attached
patch removes the xfails for -O2 and -O3.
> Also please verify and resolve the following regression introduced by
> your patch:
>
> PASS: c-c++-common/goacc/deviceptr-4.c (test for excess errors)
> [-PASS:-]{+FAIL:+} c-c++-common/goacc/deviceptr-4.c scan-tree-dump-times
> gimple "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1
>
> [-PASS:-]{+FAIL:+} c-c++-common/goacc/deviceptr-4.c -std=c++11
> scan-tree-dump-times gimple "#pragma omp target
> oacc_parallel.*map\\(tofrom:a" 1
> PASS: c-c++-common/goacc/deviceptr-4.c -std=c++11 (test for excess
> errors)
> [-PASS:-]{+FAIL:+} c-c++-common/goacc/deviceptr-4.c -std=c++14
> scan-tree-dump-times gimple "#pragma omp target
> oacc_parallel.*map\\(tofrom:a" 1
> PASS: c-c++-common/goacc/deviceptr-4.c -std=c++14 (test for excess
> errors)
> [-PASS:-]{+FAIL:+} c-c++-common/goacc/deviceptr-4.c -std=c++98
> scan-tree-dump-times gimple "#pragma omp target
> oacc_parallel.*map\\(tofrom:a" 1
> PASS: c-c++-common/goacc/deviceptr-4.c -std=c++98 (test for excess
> errors)
I forgot to update the expected data mapping in devicetpr-4.c. Now,
instead of implicitly adding a 'copy' clause for know deviceptr
variables, the gimplifier will assign a force_deviceptr clause.
I've applied the attached patch to og7 to fix both of the issues you've
identified.
Cesar
2018-05-09 Cesar Philippidis <ce...@codesourcery.com>
gcc/testsuite/
* c-c++-common/goacc/deviceptr-4.c: Update expected data mapping.
libgomp/
* libgomp.oacc-fortran/deviceptr-1.f90: Remove xfail for -O2 and -O3.
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
index db1b91633a6..79a51620db9 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
@@ -8,4 +8,4 @@ subr (int *a)
a[0] += 1.0;
}
-/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_deviceptr:a" 1 "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
index 610d071393c..7c8b063b220 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deviceptr-1.f90
@@ -7,7 +7,7 @@
! regressed with the "Partially enable GOMP_MAP_FIRSTPRIVATE_POINTER in
! gfortran" changes.
! warning: OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty
-! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" "-O2" "-O3" } { "" } }
+! { dg-xfail-if "TODO" { openacc_nvidia_accel_selected } { "-Os" } { "" } }
subroutine subr1 (a, b)
implicit none
