From mboxrd@z Thu Jan  1 00:00:00 1970
Return-Path: <gcc-bugzilla@gcc.gnu.org>
Received: by sourceware.org (Postfix, from userid 48)
 id 9FAC73888C41; Fri, 18 Mar 2022 14:46:56 +0000 (GMT)
DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 9FAC73888C41
From: "cvs-commit at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug target/104952] [nvptx][OpenMP] wrong code with OR / AND
 reduction ('reduction(||:' and '&&') with SIMT
Date: Fri, 18 Mar 2022 14:46:56 +0000
X-Bugzilla-Reason: CC
X-Bugzilla-Type: changed
X-Bugzilla-Watch-Reason: None
X-Bugzilla-Product: gcc
X-Bugzilla-Component: target
X-Bugzilla-Version: 12.0
X-Bugzilla-Keywords: openmp, wrong-code
X-Bugzilla-Severity: normal
X-Bugzilla-Who: cvs-commit at gcc dot gnu.org
X-Bugzilla-Status: UNCONFIRMED
X-Bugzilla-Resolution: 
X-Bugzilla-Priority: P3
X-Bugzilla-Assigned-To: unassigned at gcc dot gnu.org
X-Bugzilla-Target-Milestone: ---
X-Bugzilla-Flags: 
X-Bugzilla-Changed-Fields: 
Message-ID: <bug-104952-4-dyJM1YoG1u@http.gcc.gnu.org/bugzilla/>
In-Reply-To: <bug-104952-4@http.gcc.gnu.org/bugzilla/>
References: <bug-104952-4@http.gcc.gnu.org/bugzilla/>
Content-Type: text/plain; charset="UTF-8"
Content-Transfer-Encoding: quoted-printable
X-Bugzilla-URL: http://gcc.gnu.org/bugzilla/
Auto-Submitted: auto-generated
MIME-Version: 1.0
X-BeenThere: gcc-bugs@gcc.gnu.org
X-Mailman-Version: 2.1.29
Precedence: list
List-Id: Gcc-bugs mailing list <gcc-bugs.gcc.gnu.org>
List-Unsubscribe: <https://gcc.gnu.org/mailman/options/gcc-bugs>,
 <mailto:gcc-bugs-request@gcc.gnu.org?subject=unsubscribe>
List-Archive: <https://gcc.gnu.org/pipermail/gcc-bugs/>
List-Post: <mailto:gcc-bugs@gcc.gnu.org>
List-Help: <mailto:gcc-bugs-request@gcc.gnu.org?subject=help>
List-Subscribe: <https://gcc.gnu.org/mailman/listinfo/gcc-bugs>,
 <mailto:gcc-bugs-request@gcc.gnu.org?subject=subscribe>
X-List-Received-Date: Fri, 18 Mar 2022 14:46:56 -0000

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=3D104952

--- Comment #11 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tom de Vries <vries@gcc.gnu.org>:

https://gcc.gnu.org/g:093cdadbce30ce2d36846a05d979b8afc2eff618

commit r12-7702-g093cdadbce30ce2d36846a05d979b8afc2eff618
Author: Tom de Vries <tdevries@suse.de>
Date:   Thu Mar 17 14:37:28 2022 +0100

    [openmp] Fix SIMT reduction using TRUTH_{AND,OR}IF_EXPR

    Consider test-case pr104952-1.c, included in this commit, containing:
    ...
      #pragma omp target map(tofrom:result) map(to:arr)
      #pragma omp simd reduction(||: result)
    ...

    When run on x86_64 with nvptx accelerator, the test-case either aborts =
or
    hangs.

    The reduction clause is translated by the SIMT code (active for nvptx) =
as a
    butterfly reduction loop with this butterfly shuffle / update pair:
    ...
      D.2163 =3D D.2163 || .GOMP_SIMT_XCHG_BFLY (D.2163, D.2164)
    ...
    in the loop body.

    The problem is that the butterfly shuffle is possibly not executed, whi=
le
it
    needs to be executed unconditionally.

    Fix this by translating instead as:
    ...
      D.tmp_bfly =3D .GOMP_SIMT_XCHG_BFLY (D.2163, D.2164)
      D.2163 =3D D.2163 || D.tmp_bfly
    ...

    Tested on x86_64-linux with nvptx accelerator.

    gcc/ChangeLog:

    2022-03-17  Tom de Vries  <tdevries@suse.de>

            PR target/104952
            * omp-low.cc (lower_rec_input_clauses): Make sure
GOMP_SIMT_XCHG_BFLY
            is executed unconditionally.

    libgomp/ChangeLog:

    2022-03-17  Tom de Vries  <tdevries@suse.de>

            PR target/104952
            * testsuite/libgomp.c/pr104952-1.c: New test.
            * testsuite/libgomp.c/pr104952-2.c: New test.=