From patchwork Fri Oct 2 13:21:04 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1375750 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=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Received: from sourceware.org (server2.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 4C2rGx5ZsGz9sPB for ; Fri, 2 Oct 2020 23:21:12 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id CB134398E405; Fri, 2 Oct 2020 13:21:09 +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 16ED1398C829 for ; Fri, 2 Oct 2020 13:21:07 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 16ED1398C829 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 DD13BB162; Fri, 2 Oct 2020 13:21:05 +0000 (UTC) Date: Fri, 2 Oct 2020 15:21:04 +0200 From: Tom de Vries To: gcc-patches@gcc.gnu.org Subject: [PATCH][omp, simt] Handle alternative IV Message-ID: <20201002132102.GA22373@delia> MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.10.1 (2018-07-13) X-Spam-Status: No, score=-10.8 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, Consider the test-case libgomp.c/pr81778.c added in this commit, with this core loop (note: CANARY_SIZE set to 0 for simplicity): ... int s = 1; #pragma omp target simd for (int i = N - 1; i > -1; i -= s) a[i] = 1; ... which, given that N is 32, sets a[0..31] to 1. After omp-expand, this looks like: ... : simduid.7 = .GOMP_SIMT_ENTER (simduid.7); .omp_simt.8 = .GOMP_SIMT_ENTER_ALLOC (simduid.7); D.3193 = -s; s.9 = s; D.3204 = .GOMP_SIMT_LANE (); D.3205 = -s.9; D.3206 = (int) D.3204; D.3207 = D.3205 * D.3206; i = D.3207 + 31; D.3209 = 0; D.3210 = -s.9; D.3211 = D.3210 - i; D.3210 = -s.9; D.3212 = D.3211 / D.3210; D.3213 = (unsigned int) D.3212; D.3213 = i >= 0 ? D.3213 : 0; : if (D.3209 < D.3213) goto ; [87.50%] else goto ; [12.50%] : a[i] = 1; D.3215 = -s.9; D.3219 = .GOMP_SIMT_VF (); D.3216 = (int) D.3219; D.3220 = D.3215 * D.3216; i = D.3220 + i; D.3209 = D.3209 + 1; goto ; [100.00%] ... On nvptx, the first time bb6 is executed, i is in the 0..31 range (depending on the lane that is executing) at bb entry. So we have the following sequence: - a[0..31] is set to 1 - i is updated to -32..-1 - D.3209 is updated to 1 (being 0 initially) - bb19 is executed, and if condition (D.3209 < D.3213) == (1 < 32) evaluates to true - bb6 is once more executed, which should not happen because all the elements that needed to be handled were already handled. - consequently, elements that should not be written are written - with CANARY_SIZE == 0, we may run into a libgomp error: ... libgomp: cuCtxSynchronize error: an illegal memory access was encountered ... and with CANARY_SIZE unmodified, we run into: ... Expected 0, got 1 at base[-961] Aborted (core dumped) ... The cause of this is as follows: - because the step s is a variable rather than a constant, an alternative IV (D.3209 in our example) is generated in expand_omp_simd, and the loop condition is tested in terms of the alternative IV rather than the original IV (i in our example). - the SIMT code in expand_omp_simd works by modifying step and initial value. - The initial value fd->loop.n1 is loaded into a variable n1, which is modified by the SIMT code and then used there-after. - The step fd->loop.step is loaded into a variable step, which is is modified by the SIMT code, but afterwards there are uses of both step and fd->loop.step. - There are uses of fd->loop.step in the alternative IV handling code, which should use step instead. Fix this by introducing an additional variable orig_step, which is not modified by the SIMT code and replacing all remaining uses of fd->loop.step by either step or orig_step. Build on x86_64-linux with nvptx accelerator, tested libgomp. This fixes for-5.c and for-6.c FAILs I'm currently seeing on a quadro m1200 with driver 450.66. OK for trunk? Thanks, - Tom [omp, simt] Handle alternative IV gcc/ChangeLog: 2020-10-02 Tom de Vries * omp-expand.c (expand_omp_simd): Add step_orig, and replace uses of fd->loop.step by either step or orig_step. libgomp/ChangeLog: 2020-10-02 Tom de Vries * testsuite/libgomp.c/pr81778.c: New test. --- gcc/omp-expand.c | 11 ++++---- libgomp/testsuite/libgomp.c/pr81778.c | 48 +++++++++++++++++++++++++++++++++++ 2 files changed, 54 insertions(+), 5 deletions(-) diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 99cb4f9dda4..80e35ac0294 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -6307,6 +6307,7 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd) n2 = OMP_CLAUSE_DECL (innerc); } tree step = fd->loop.step; + tree orig_step = step; /* May be different from step if is_simt. */ bool is_simt = omp_find_clause (gimple_omp_for_clauses (fd->for_stmt), OMP_CLAUSE__SIMT_); @@ -6455,7 +6456,7 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd) tree altv = NULL_TREE, altn2 = NULL_TREE; if (fd->collapse == 1 && !broken_loop - && TREE_CODE (fd->loops[0].step) != INTEGER_CST) + && TREE_CODE (orig_step) != INTEGER_CST) { /* The vectorizer currently punts on loops with non-constant steps for the main IV (can't compute number of iterations and gives up @@ -6471,7 +6472,7 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd) itype = signed_type_for (itype); t = build_int_cst (itype, (fd->loop.cond_code == LT_EXPR ? -1 : 1)); t = fold_build2 (PLUS_EXPR, itype, - fold_convert (itype, fd->loop.step), t); + fold_convert (itype, step), t); t = fold_build2 (PLUS_EXPR, itype, t, fold_convert (itype, n2)); t = fold_build2 (MINUS_EXPR, itype, t, fold_convert (itype, fd->loop.v)); @@ -6479,10 +6480,10 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd) t = fold_build2 (TRUNC_DIV_EXPR, itype, fold_build1 (NEGATE_EXPR, itype, t), fold_build1 (NEGATE_EXPR, itype, - fold_convert (itype, fd->loop.step))); + fold_convert (itype, step))); else t = fold_build2 (TRUNC_DIV_EXPR, itype, t, - fold_convert (itype, fd->loop.step)); + fold_convert (itype, step)); t = fold_convert (TREE_TYPE (altv), t); altn2 = create_tmp_var (TREE_TYPE (altv)); expand_omp_build_assign (&gsi, altn2, t); @@ -6630,7 +6631,7 @@ expand_omp_simd (struct omp_region *region, struct omp_for_data *fd) if (is_simt) { gsi = gsi_start_bb (l2_bb); - step = fold_build2 (MINUS_EXPR, TREE_TYPE (step), fd->loop.step, step); + step = fold_build2 (MINUS_EXPR, TREE_TYPE (step), orig_step, step); if (POINTER_TYPE_P (type)) t = fold_build_pointer_plus (fd->loop.v, step); else diff --git a/libgomp/testsuite/libgomp.c/pr81778.c b/libgomp/testsuite/libgomp.c/pr81778.c new file mode 100644 index 00000000000..571668eb36a --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr81778.c @@ -0,0 +1,48 @@ +/* Minimized from for-5.c. */ + +#include +#include + +/* Size of array we want to write. */ +#define N 32 + +/* Size of extra space before and after. */ +#define CANARY_SIZE (N * 32) + +/* Start of array we want to write. */ +#define BASE (CANARY_SIZE) + +// Total size to be allocated. +#define ALLOC_SIZE (CANARY_SIZE + N + CANARY_SIZE) + +#pragma omp declare target +int a[ALLOC_SIZE]; +#pragma omp end declare target + +int +main (void) +{ + /* Use variable step in for loop. */ + int s = 1; + +#pragma omp target update to(a) + + /* Write a[BASE] .. a[BASE + N - 1]. */ +#pragma omp target simd + for (int i = N - 1; i > -1; i -= s) + a[BASE + i] = 1; + +#pragma omp target update from(a) + + for (int i = 0; i < ALLOC_SIZE; i++) + { + int expected = (BASE <= i && i < BASE + N) ? 1 : 0; + if (a[i] == expected) + continue; + + printf ("Expected %d, got %d at base[%d]\n", expected, a[i], i - BASE); + abort (); + } + + return 0; +}