From patchwork Mon May 3 10:24:10 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: 1473124 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=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (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 4FYfGT6gk3z9rx6 for ; Mon, 3 May 2021 20:24:17 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 7A93F389443C; Mon, 3 May 2021 10:24:14 +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 AC441385782A for ; Mon, 3 May 2021 10:24:12 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org AC441385782A 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 C8F3FAC6A; Mon, 3 May 2021 10:24:11 +0000 (UTC) Date: Mon, 3 May 2021 12:24:10 +0200 From: Tom de Vries To: gcc-patches@gcc.gnu.org Subject: [PATCH][openmp, simt] Error out for user-defined reduction Message-ID: <20210503102409.GA20090@delia> 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, The test-case included in this patch contains this target region: ... for (int i0 = 0 ; i0 < N0 ; i0++ ) counter_N0.i += 1; ... When running with nvptx accelerator, the counter variable is expected to be N0 after the region, but instead is N0 / 32. The problem is that rather than getting the result for all warp lanes, we get it for just one lane. This is caused by the implementation of SIMT being incomplete. It handles regular reductions, but appearantly not user-defined reductions. For now, make this explicit by erroring out for nvptx, like this: ... target-44.c: In function 'main': target-44.c:20:9: error: SIMT reduction not fully implemented ... Tested libgomp on x86_64-linux with and without nvptx accelerator. Any comments? Thanks, - Tom [openmp, simt] Error out for user-defined reduction gcc/ChangeLog: 2021-05-03 Tom de Vries PR target/100321 * omp-low.c (lower_rec_input_clauses): Error out for user-defined reduction for SIMT. libgomp/ChangeLog: 2021-05-03 Tom de Vries PR target/100321 * testsuite/libgomp.c/target-44.c: New test. --- gcc/omp-low.c | 2 ++ libgomp/testsuite/libgomp.c/target-44.c | 28 ++++++++++++++++++++++++++++ 2 files changed, 30 insertions(+) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 7b122059c6e..0f122857a3a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -6005,6 +6005,8 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, tree placeholder = OMP_CLAUSE_REDUCTION_PLACEHOLDER (c); gimple *tseq; tree ptype = TREE_TYPE (placeholder); + if (sctx.is_simt) + error ("SIMT reduction not fully implemented"); if (cond) { x = error_mark_node; diff --git a/libgomp/testsuite/libgomp.c/target-44.c b/libgomp/testsuite/libgomp.c/target-44.c new file mode 100644 index 00000000000..497931cd14c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-44.c @@ -0,0 +1,28 @@ +/* { dg-do link { target { offload_target_nvptx } } } */ +/* { dg-additional-options "-foffload=-latomic" { target { offload_target_nvptx } } } */ +/* { dg-error "SIMT reduction not fully implemented" "" { target { offload_target_nvptx } } 0 } */ +#include + +struct s +{ + int i; +}; + +#pragma omp declare reduction(+: struct s: omp_out.i += omp_in.i) + +int +main (void) +{ + const int N0 = 32768; + + struct s counter_N0 = { 0 }; +#pragma omp target +#pragma omp for simd reduction(+: counter_N0) + for (int i0 = 0 ; i0 < N0 ; i0++ ) + counter_N0.i += 1; + + if (counter_N0.i != N0) + abort (); + + return 0; +}