Re: [og7] Update deviceptr handling in Fortran

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi Cesar!

On Mon, 7 May 2018 08:49:26 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch teaches both the Fortran FE and the gimplifier how to only
> utilize one data mapping for OpenACC deviceptr clauses.  [...]

Thanks!  (I didn't verify your code changes.)


> 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.


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)


Grüße
 Thomas