From patchwork Wed Apr 28 14:30:30 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1471213 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (ip-8-43-85-97.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4FVgz42CHQz9sWC for ; Thu, 29 Apr 2021 00:30:39 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id CEC563AA0042; Wed, 28 Apr 2021 14:30:36 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mx2.suse.de (mx2.suse.de [195.135.220.15]) by sourceware.org (Postfix) with ESMTPS id 996F1382E82E for ; Wed, 28 Apr 2021 14:30:33 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 996F1382E82E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=tdevries@suse.de X-Virus-Scanned: by amavisd-new at test-mx.suse.de Received: from relay2.suse.de (unknown [195.135.221.27]) by mx2.suse.de (Postfix) with ESMTP id 7063BAD8A; Wed, 28 Apr 2021 14:30:32 +0000 (UTC) Date: Wed, 28 Apr 2021 16:30:30 +0200 From: Tom de Vries To: gcc-patches@gcc.gnu.org Subject: [PATCH][omp, simt] Fix expand_GOMP_SIMT_* Message-ID: <20210428143029.GA2644@delia.home> MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.10.1 (2018-07-13) X-Spam-Status: No, score=-12.2 required=5.0 tests=BAYES_00, GIT_PATCH_0, KAM_DMARC_STATUS, RCVD_IN_MSPIKE_H3, RCVD_IN_MSPIKE_WL, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Jakub Jelinek , Alexander Monakov Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" Hi, When running the test-case included in this patch using an nvptx accelerator, it fails in execution. The problem is that the expansion of GOMP_SIMT_XCHG_BFLY is optimized away during pass_jump as "trivially dead insns". This is caused by this code in expand_GOMP_SIMT_XCHG_BFLY: ... class expand_operand ops[3]; create_output_operand (&ops[0], target, mode); ... expand_insn (targetm.code_for_omp_simt_xchg_bfly, 3, ops); ... which doesn't guarantee that target is assigned to by the expanded insn. F.i., if target is: ... (gdb) call debug_rtx ( target ) (subreg/s/u:QI (reg:SI 40 [ _61 ]) 0) ... then after expand_insn, we have: ... (gdb) call debug_rtx ( ops[0].value ) (reg:QI 57) ... See commit 3af3bec2e4d "internal-fn: Avoid dropping the lhs of some calls [PR94941]" for a similar problem. Fix this in the same way, by adding: ... if (!rtx_equal_p (target, ops[0].value)) emit_move_insn (target, ops[0].value); ... where applicable in the expand_GOMP_SIMT_* functions. Tested libgomp on x86_64 with nvptx accelerator. Any comments? Thanks, - Tom [omp, simt] Fix expand_GOMP_SIMT_* gcc/ChangeLog: 2021-04-28 Tom de Vries PR target/100232 * internal-fn.c (expand_GOMP_SIMT_ENTER_ALLOC) (expand_GOMP_SIMT_LAST_LANE, expand_GOMP_SIMT_ORDERED_PRED) (expand_GOMP_SIMT_VOTE_ANY, expand_GOMP_SIMT_XCHG_BFLY) (expand_GOMP_SIMT_XCHG_IDX): Ensure target is assigned to. --- gcc/internal-fn.c | 12 ++++++++++++ libgomp/testsuite/libgomp.c/target-43.c | 24 ++++++++++++++++++++++++ 2 files changed, 36 insertions(+) diff --git a/gcc/internal-fn.c b/gcc/internal-fn.c index dd7173126fb..d209a52f823 100644 --- a/gcc/internal-fn.c +++ b/gcc/internal-fn.c @@ -243,6 +243,8 @@ expand_GOMP_SIMT_ENTER_ALLOC (internal_fn, gcall *stmt) create_input_operand (&ops[2], align, Pmode); gcc_assert (targetm.have_omp_simt_enter ()); expand_insn (targetm.code_for_omp_simt_enter, 3, ops); + if (!rtx_equal_p (target, ops[0].value)) + emit_move_insn (target, ops[0].value); } /* Deallocate per-lane storage and leave non-uniform execution region. */ @@ -300,6 +302,8 @@ expand_GOMP_SIMT_LAST_LANE (internal_fn, gcall *stmt) create_input_operand (&ops[1], cond, mode); gcc_assert (targetm.have_omp_simt_last_lane ()); expand_insn (targetm.code_for_omp_simt_last_lane, 2, ops); + if (!rtx_equal_p (target, ops[0].value)) + emit_move_insn (target, ops[0].value); } /* Non-transparent predicate used in SIMT lowering of OpenMP "ordered". */ @@ -319,6 +323,8 @@ expand_GOMP_SIMT_ORDERED_PRED (internal_fn, gcall *stmt) create_input_operand (&ops[1], ctr, mode); gcc_assert (targetm.have_omp_simt_ordered ()); expand_insn (targetm.code_for_omp_simt_ordered, 2, ops); + if (!rtx_equal_p (target, ops[0].value)) + emit_move_insn (target, ops[0].value); } /* "Or" boolean reduction across SIMT lanes: return non-zero in all lanes if @@ -339,6 +345,8 @@ expand_GOMP_SIMT_VOTE_ANY (internal_fn, gcall *stmt) create_input_operand (&ops[1], cond, mode); gcc_assert (targetm.have_omp_simt_vote_any ()); expand_insn (targetm.code_for_omp_simt_vote_any, 2, ops); + if (!rtx_equal_p (target, ops[0].value)) + emit_move_insn (target, ops[0].value); } /* Exchange between SIMT lanes with a "butterfly" pattern: source lane index @@ -361,6 +369,8 @@ expand_GOMP_SIMT_XCHG_BFLY (internal_fn, gcall *stmt) create_input_operand (&ops[2], idx, SImode); gcc_assert (targetm.have_omp_simt_xchg_bfly ()); expand_insn (targetm.code_for_omp_simt_xchg_bfly, 3, ops); + if (!rtx_equal_p (target, ops[0].value)) + emit_move_insn (target, ops[0].value); } /* Exchange between SIMT lanes according to given source lane index. */ @@ -382,6 +392,8 @@ expand_GOMP_SIMT_XCHG_IDX (internal_fn, gcall *stmt) create_input_operand (&ops[2], idx, SImode); gcc_assert (targetm.have_omp_simt_xchg_idx ()); expand_insn (targetm.code_for_omp_simt_xchg_idx, 3, ops); + if (!rtx_equal_p (target, ops[0].value)) + emit_move_insn (target, ops[0].value); } /* This should get expanded in adjust_simduid_builtins. */ diff --git a/libgomp/testsuite/libgomp.c/target-43.c b/libgomp/testsuite/libgomp.c/target-43.c new file mode 100644 index 00000000000..46b1cfc5b20 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-43.c @@ -0,0 +1,24 @@ +/* { dg-do run } */ +#include + +#define N 32 +#define TYPE char + +int +main (void) +{ + TYPE result = 1; + TYPE a[N]; + for (int x = 0; x < N; ++x) + a[x] = 1; + +#pragma omp target map(tofrom: result) map(to:a) +#pragma omp for simd reduction(&&:result) + for (int x = 0; x < N; ++x) + result = result && a[x]; + + if (result != 1) + abort (); + + return 0; +}