From patchwork Fri Mar 18 14:47:18 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 52099 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D83263888C6D for ; Fri, 18 Mar 2022 14:47:53 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org D83263888C6D DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1647614873; bh=TOFKGqHYDW0JqU+a5aWUlnul2T96aowV7nrcKRyFnDg=; h=Date:To:Subject:List-Id:List-Unsubscribe:List-Archive:List-Post: List-Help:List-Subscribe:From:Reply-To:Cc:From; b=HFDNda+N6L4WSdr4Iao1OxVG/Sx2cpUjf32BsuqO+xtl0P09zevMHdBcMmu953jwZ 2Sakqhei0qB5WwHCsioJl9wjJMP7qChIxhsK+eaZZb/lTPQvaCYmknjaOtWa3zIgM5 oBWywGw10FmxPtHgcpmdvd3XW9/58TjggZCMREiM= X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from smtp-out1.suse.de (smtp-out1.suse.de [195.135.220.28]) by sourceware.org (Postfix) with ESMTPS id C830A3888C65 for ; Fri, 18 Mar 2022 14:47:20 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org C830A3888C65 Received: from imap2.suse-dmz.suse.de (imap2.suse-dmz.suse.de [192.168.254.74]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-521) server-digest SHA512) (No client certificate requested) by smtp-out1.suse.de (Postfix) with ESMTPS id E8CB3210C2; Fri, 18 Mar 2022 14:47:19 +0000 (UTC) Received: from imap2.suse-dmz.suse.de (imap2.suse-dmz.suse.de [192.168.254.74]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-521) server-digest SHA512) (No client certificate requested) by imap2.suse-dmz.suse.de (Postfix) with ESMTPS id CF4E6133D1; Fri, 18 Mar 2022 14:47:19 +0000 (UTC) Received: from dovecot-director2.suse.de ([192.168.254.65]) by imap2.suse-dmz.suse.de with ESMTPSA id ef95MXebNGJFUAAAMHmgww (envelope-from ); Fri, 18 Mar 2022 14:47:19 +0000 Date: Fri, 18 Mar 2022 15:47:18 +0100 To: gcc-patches@gcc.gnu.org Subject: [committed][openmp] Fix SIMT reduction using TRUTH_{AND,OR}IF_EXPR Message-ID: <20220318144716.GA30949@delia.home> MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.10.1 (2018-07-13) X-Spam-Status: No, score=-12.6 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) 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: , X-Patchwork-Original-From: Tom de Vries via Gcc-patches From: Tom de Vries Reply-To: Tom de Vries Cc: Jakub Jelinek Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi, 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 = 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, while it needs to be executed unconditionally. Fix this by translating instead as: ... D.tmp_bfly = .GOMP_SIMT_XCHG_BFLY (D.2163, D.2164) D.2163 = D.2163 || D.tmp_bfly ... Tested on x86_64-linux with nvptx accelerator. Committed to trunk. Thanks, - Tom [openmp] Fix SIMT reduction using TRUTH_{AND,OR}IF_EXPR gcc/ChangeLog: 2022-03-17 Tom de Vries 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 PR target/104952 * testsuite/libgomp.c/pr104952-1.c: New test. * testsuite/libgomp.c/pr104952-2.c: New test. --- gcc/omp-low.cc | 5 ++++- libgomp/testsuite/libgomp.c/pr104952-1.c | 24 ++++++++++++++++++++++++ libgomp/testsuite/libgomp.c/pr104952-2.c | 22 ++++++++++++++++++++++ 3 files changed, 50 insertions(+), 1 deletion(-) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index cfc63d6a104..392bb18bc5d 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -6743,7 +6743,10 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, x = build_call_expr_internal_loc (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_BFLY, TREE_TYPE (ivar), 2, ivar, simt_lane); - x = build2 (code, TREE_TYPE (ivar), ivar, x); + /* Make sure x is evaluated unconditionally. */ + tree bfly_var = create_tmp_var (TREE_TYPE (ivar)); + gimplify_assign (bfly_var, x, &llist[2]); + x = build2 (code, TREE_TYPE (ivar), ivar, bfly_var); gimplify_assign (ivar, x, &llist[2]); } tree ivar2 = ivar; diff --git a/libgomp/testsuite/libgomp.c/pr104952-1.c b/libgomp/testsuite/libgomp.c/pr104952-1.c new file mode 100644 index 00000000000..a3bfb1e77df --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr104952-1.c @@ -0,0 +1,24 @@ +#define N 32 + +static char arr[N]; + +int +main (void) +{ + unsigned int result = 0; + + for (unsigned int i = 0; i < N; ++i) + arr[i] = 0; + + arr[5] = 42; + +#pragma omp target map(tofrom:result) map(to:arr) +#pragma omp simd reduction(||: result) + for (unsigned int i = 0; i < N; ++i) + result = result || arr[i]; + + if (result != 1) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/pr104952-2.c b/libgomp/testsuite/libgomp.c/pr104952-2.c new file mode 100644 index 00000000000..7ab4bcdb8af --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr104952-2.c @@ -0,0 +1,22 @@ +#define N 32 + +static char arr[N]; + +int +main (void) +{ + unsigned int result = 2; + + for (unsigned int i = 0; i < N; ++i) + arr[i] = i + 1; + +#pragma omp target map(tofrom:result) map(to:arr) +#pragma omp simd reduction(&&: result) + for (unsigned int i = 0; i < N; ++i) + result = result && arr[i]; + + if (result != 1) + __builtin_abort (); + + return 0; +}