From patchwork Thu Feb 8 14:47:13 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 198437 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a05:7300:50ea:b0:106:860b:bbdd with SMTP id r10csp218284dyd; Thu, 8 Feb 2024 06:48:14 -0800 (PST) X-Google-Smtp-Source: AGHT+IGOlxSP3MSafTAMUmVsoCbrGeFUARWrtthu+yBn3R+aOJXsgXNeYPdx8iotbpcwl2aRT+DU X-Received: by 2002:a25:1545:0:b0:dc7:4460:878a with SMTP id 66-20020a251545000000b00dc74460878amr1504909ybv.3.1707403693934; Thu, 08 Feb 2024 06:48:13 -0800 (PST) ARC-Seal: i=2; a=rsa-sha256; t=1707403693; cv=pass; d=google.com; s=arc-20160816; b=HGpev2Prq6CWByoj4epu/pMQdUoBCLWHeEy0BYQaxO1UvSwM0OfjVKypJwXCXe3POq oV0/bL72RLxJ2XxYQJdHkOiykOx8k5c+Ic0/dIW7//lmF/fZJSY8I0YlXqEtQDZkwly7 LEH9ju/OkhhGU9L3yXBSvkJpZXuEWfivfFvvbLV8JIjb6krFATiHZBN7QLoAjbbf1+O0 fLTfEcte0Iv+R69YiV921EP82QbYRibJhJctAfYy4pkxkh41QWhEavk+x9R+HMEXmfrt qbgNw+PrZALj72MESePJcXkhaQxkuu3mGk17W9zjfnDY6eHhuSn7EyzYd6e3BDiYS1G9 xI8Q== ARC-Message-Signature: i=2; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:to:subject:from :content-language:user-agent:mime-version:date:message-id :dkim-signature:arc-filter:dmarc-filter:delivered-to; bh=Ex9s5iUTLtD8Wf/d2e6ZVIOqrl1KqswIHKltvwT4vW0=; fh=A1mIjOnzrxx9pWcqaNrpQ3xupQWMTCha3uszkq2XUNM=; b=GtU+Z9ohrlcw4/ZWpRoxVpVKWm5xVepak61t3Vl0E3YiE5CV5bW/uBP5qKuV633r7a 551rvBsya95Z7mYSdOo9LKeWVmevJmPqgJ6/p0meA2xlbnXXeLxKN8YEW3swUDm7eDkK 4ADcdYIcfBJKmhl9CumcCtXON1ixEnNruRsUSPmrZjSiI3fXYV0vWqZHyHLjQyEbALQO QBmyVr/AWWJ68fVdkWnzAna0hdVrUt2wZ1fdcsHXa4+R8P4Cfnwi4LwvtsG07f403F7b dqHZFHHuH9EqWMOpOfhoBtNpWD+o9JbOLg59tC4f6eG/aquwxxlrNYk5E4p7VBFU2iiJ GnXw==; dara=google.com ARC-Authentication-Results: i=2; mx.google.com; dkim=pass header.i=@baylibre-com.20230601.gappssmtp.com header.s=20230601 header.b=VK4kjCdY; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" X-Forwarded-Encrypted: i=2; AJvYcCXUdniuVhaFaAPkEHNCvkvPhF5K1GENcXoFxdD9FJWIqWuiUdKB3swnJ3VWEmw9Cp/q0qpn4dRfAEMvLOXk1xlFUxdK/Q== Received: from server2.sourceware.org (server2.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id g19-20020ae9e113000000b007856c1a29d1si127245qkm.242.2024.02.08.06.48.13 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 08 Feb 2024 06:48:13 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; dkim=pass header.i=@baylibre-com.20230601.gappssmtp.com header.s=20230601 header.b=VK4kjCdY; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 93F003858C78 for ; Thu, 8 Feb 2024 14:48:13 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-pj1-x102b.google.com (mail-pj1-x102b.google.com [IPv6:2607:f8b0:4864:20::102b]) by sourceware.org (Postfix) with ESMTPS id 17DC53858C60 for ; Thu, 8 Feb 2024 14:47:17 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 17DC53858C60 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 17DC53858C60 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2607:f8b0:4864:20::102b ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1707403643; cv=none; b=UWBzRbDCr3eisGaP/lLA8fjp2bXgMCB0NQHFth6rLuqdrqao8mdSVyfukYKXlZdxJil167mXoj4yyS3KWvH028szDhxpyvBD7I9hkXDlLKIEdZGEjw5/ifhPraJvwmkKoQEtKtybvZJpUqhkPD5CAC/cdZzwcH7fHgSYVW2+2eE= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1707403643; c=relaxed/simple; bh=htI3oGngmYhgTiVXgmmPavzS7055qDUJrEXcl0Yk5g4=; h=DKIM-Signature:Message-ID:Date:MIME-Version:From:Subject:To; b=CEJ0k765IO4bvHBUvUumsceuvSz4wLqnPMzkKNWAhYzkb/KMLn4Bn3/tBJhgB/RuxPP5SEibLgPJrMFIie+DvPTDjmFGGvnBgL3T6UhhYH09meh+eQTnLQoQsf8sFgQt4VyvGKTV/ResbXuE8khb+pOjg33vzczR/CtxEPEpvJg= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-pj1-x102b.google.com with SMTP id 98e67ed59e1d1-296c562ac70so1398739a91.2 for ; Thu, 08 Feb 2024 06:47:17 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1707403636; x=1708008436; darn=gcc.gnu.org; h=to:subject:from:content-language:user-agent:mime-version:date :message-id:from:to:cc:subject:date:message-id:reply-to; bh=Ex9s5iUTLtD8Wf/d2e6ZVIOqrl1KqswIHKltvwT4vW0=; b=VK4kjCdY7VGdnu5c4IcAaYlImz0fTNIig7ifKIZznO1g+X7gxAypmYarQUW3Z2aCmf PoHgKCpQDBnixJ9y/9MLHwibHgCHAI16eAnUXLiVboi63uTzMowu2mwnSt/EiWfut612 QdKw9loRig+FPbB+bHA7CQA/phbe1XS90UYy7rnS9zC7kpgTqJ+LzwAgfb8rGIcZ2EXm 2unK2cEiXHkx+BJAMGpS5kj2GWFzRpxhKUP7fBKn1wG8FmVJ4y6H7r3LfEOVwyRnGEo0 wRzPIH8CvfjEXnmUrJgPUS2L0dVoXanJEf5XSzoQNesqfAGZ5yTfxJkYAqsW01d/nIUf iHrA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1707403636; x=1708008436; h=to:subject:from:content-language:user-agent:mime-version:date :message-id:x-gm-message-state:from:to:cc:subject:date:message-id :reply-to; bh=Ex9s5iUTLtD8Wf/d2e6ZVIOqrl1KqswIHKltvwT4vW0=; b=uQj7Irh/y8isANcngwlwTsEoAS36CGaKtQb3bSdco1goJRbB5Q3y2fZTyShpyCr8Ig Fxi8QE+wUfi4F/W/NP4vp1dg5HGlWReeGEHvH9CdqeIEW+t/zZ902PjWFIIuSX710HwV IpNAzm4skZG69Y5HRgk9RNixGdU19/4USFv8iScrbaV2N8FUVZloxsk5OSdkVTtwPjRj aETbQNNmwMekQ5Ajj0xis27NBxQv/NjNFRPeZ7LdFQ5jCaZ/SIf0M/1OaGpUTwYYQvB/ Nf1A4QL5iJLy3ne1kh8j6G0fiyB8rsLNQLOjcx126TEMozB1hlj0Kvqra3yhDBTUxZL0 hAQg== X-Gm-Message-State: AOJu0YydJ5lvk9U4+rMYXWOhfJfsybn3Lhj8qeiMIKzV5+NvNB+ZHOoS k3agp+MNkFgh+advQvHhD8VAsrasTdJ/isbyZbgz7gGXhqUuTpnYi0aY2CBRh4J8CqN7AcbaIJH bup8= X-Received: by 2002:a17:90b:4f87:b0:296:1d2a:6245 with SMTP id qe7-20020a17090b4f8700b002961d2a6245mr6411943pjb.18.1707403635873; Thu, 08 Feb 2024 06:47:15 -0800 (PST) X-Forwarded-Encrypted: i=1; AJvYcCUW8qGxAjziKAUFA0QOEH3T6vNnIvIaYCccQ9MLbY3G8yu/caKmr4rQLlSyEsXD+tgTvqWeJ5ZOLcy1m3SMr4XccOFzqpY4W9AvzXtn5Izg7okd2LQIHW0taa/9VVv+9DlxJGzZVJczmSZkZaxFq8lUxg== Received: from [192.168.50.226] (112-104-16-194.adsl.dynamic.seed.net.tw. [112.104.16.194]) by smtp.gmail.com with ESMTPSA id br15-20020a17090b0f0f00b00296bae431ecsm1622218pjb.48.2024.02.08.06.47.14 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Thu, 08 Feb 2024 06:47:15 -0800 (PST) Message-ID: <9209bd62-7ca1-4480-8497-d402b2889a72@baylibre.com> Date: Thu, 8 Feb 2024 22:47:13 +0800 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Content-Language: en-US From: Chung-Lin Tang Subject: [PATCH, OpenACC 2.7] struct/array reductions for Fortran To: gcc-patches , gfortran , Tobias Burnus , Thomas Schwinge X-Spam-Status: No, score=-13.7 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=unavailable autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1790342535506858553 X-GMAIL-MSGID: 1790342535506858553 Hi Tobias, Thomas, this patch adds support for Fortran to use arrays and struct(record) types in OpenACC reductions. There is still some shortcomings in the current state, mainly that only explicit-shaped arrays can be used (like its C counterpart). Anything else is currently a bit more complicated in the middle-end, since the existing reduction code creates an "init-op" (literal of initial values) which can't be done when say TYPE_MAX_VALUE (TYPE_DOMAIN (array_type)) is not a tree constant. I think we'll be on the hook to solve this later, but I think the current state is okay to submit. Tested without regressions on mainline (on top of first struct/array reduction patch[1]) Thanks, Chung-Lin [1] https://gcc.gnu.org/pipermail/gcc-patches/2024-January/641669.html 2024-02-08 Chung-Lin Tang gcc/fortran/ChangeLog: * openmp.cc (oacc_reduction_defined_type_p): New function. (resolve_omp_clauses): Adjust OpenACC array reduction error case. Use oacc_reduction_defined_type_p for OpenACC. * trans-openmp.cc (gfc_trans_omp_array_reduction_or_udr): Add 'bool openacc' parameter, adjust part of function to be !openacc only. (gfc_trans_omp_reduction_list): Add 'bool openacc' parameter, pass to calls to gfc_trans_omp_array_reduction_or_udr. (gfc_trans_omp_clauses): Add 'openacc' argument to calls to gfc_trans_omp_reduction_list. (gfc_trans_omp_do): Pass 'op == EXEC_OACC_LOOP' as 'bool openacc' parameter in call to gfc_trans_omp_clauses. gcc/ChangeLog: * omp-low.cc (omp_reduction_init_op): Add checking if reduced array has constant bounds. (lower_oacc_reductions): Add handling of error_mark_node. gcc/testsuite/ChangeLog: * gfortran.dg/goacc/array-reduction.f90: Adjust testcase. * gfortran.dg/goacc/reduction.f95: Likewise. libgomp/ChangeLog: * libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90: New testcase. * libgomp/testsuite/libgomp.oacc-fortran/reduction-10.f90: Likewise. * libgomp/testsuite/libgomp.oacc-fortran/reduction-11.f90: Likewise. * libgomp/testsuite/libgomp.oacc-fortran/reduction-12.f90: Likewise. * libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90: Likewise. diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 0af80d54fad..4bba9e666d6 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -7047,6 +7047,72 @@ oacc_is_loop (gfc_code *code) || code->op == EXEC_OACC_LOOP; } +static bool +oacc_reduction_defined_type_p (enum gfc_omp_reduction_op rop, gfc_typespec *ts) +{ + if (rop == OMP_REDUCTION_USER || rop == OMP_REDUCTION_NONE) + return false; + + if (ts->type == BT_INTEGER) + switch (rop) + { + case OMP_REDUCTION_AND: + case OMP_REDUCTION_OR: + case OMP_REDUCTION_EQV: + case OMP_REDUCTION_NEQV: + return false; + default: + return true; + } + + if (ts->type == BT_LOGICAL) + switch (rop) + { + case OMP_REDUCTION_AND: + case OMP_REDUCTION_OR: + case OMP_REDUCTION_EQV: + case OMP_REDUCTION_NEQV: + return true; + default: + return false; + } + + if (ts->type == BT_REAL || ts->type == BT_COMPLEX) + switch (rop) + { + case OMP_REDUCTION_PLUS: + case OMP_REDUCTION_TIMES: + case OMP_REDUCTION_MINUS: + return true; + + case OMP_REDUCTION_AND: + case OMP_REDUCTION_OR: + case OMP_REDUCTION_EQV: + case OMP_REDUCTION_NEQV: + return false; + + case OMP_REDUCTION_MAX: + case OMP_REDUCTION_MIN: + return ts->type != BT_COMPLEX; + case OMP_REDUCTION_IAND: + case OMP_REDUCTION_IOR: + case OMP_REDUCTION_IEOR: + return false; + default: + gcc_unreachable (); + } + + if (ts->type == BT_DERIVED) + { + for (gfc_component *p = ts->u.derived->components; p; p = p->next) + if (!oacc_reduction_defined_type_p (rop, &p->ts)) + return false; + return true; + } + + return false; +} + static void resolve_scalar_int_expr (gfc_expr *expr, const char *clause) { @@ -8137,13 +8203,15 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, else n->sym->mark = 1; - /* OpenACC does not support reductions on arrays. */ - if (n->sym->as) + /* OpenACC current only supports array reductions on explicit-shape + arrays. */ + if ((n->sym->as && n->sym->as->type != AS_EXPLICIT) + || n->sym->attr.codimension) gfc_error ("Array %qs is not permitted in reduction at %L", n->sym->name, &n->where); } } - + for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next) n->sym->mark = 0; for (n = omp_clauses->lists[OMP_LIST_FROM]; n; n = n->next) @@ -8797,39 +8865,46 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, case OMP_LIST_IN_REDUCTION: case OMP_LIST_TASK_REDUCTION: case OMP_LIST_REDUCTION_INSCAN: - switch (n->u.reduction_op) + if (openacc) { - case OMP_REDUCTION_PLUS: - case OMP_REDUCTION_TIMES: - case OMP_REDUCTION_MINUS: - if (!gfc_numeric_ts (&n->sym->ts)) + if (!oacc_reduction_defined_type_p (n->u.reduction_op, + &n->sym->ts)) bad = true; - break; - case OMP_REDUCTION_AND: - case OMP_REDUCTION_OR: - case OMP_REDUCTION_EQV: - case OMP_REDUCTION_NEQV: - if (n->sym->ts.type != BT_LOGICAL) - bad = true; - break; - case OMP_REDUCTION_MAX: - case OMP_REDUCTION_MIN: - if (n->sym->ts.type != BT_INTEGER - && n->sym->ts.type != BT_REAL) - bad = true; - break; - case OMP_REDUCTION_IAND: - case OMP_REDUCTION_IOR: - case OMP_REDUCTION_IEOR: - if (n->sym->ts.type != BT_INTEGER) - bad = true; - break; - case OMP_REDUCTION_USER: - bad = true; - break; - default: - break; } + else + switch (n->u.reduction_op) + { + case OMP_REDUCTION_PLUS: + case OMP_REDUCTION_TIMES: + case OMP_REDUCTION_MINUS: + if (!gfc_numeric_ts (&n->sym->ts)) + bad = true; + break; + case OMP_REDUCTION_AND: + case OMP_REDUCTION_OR: + case OMP_REDUCTION_EQV: + case OMP_REDUCTION_NEQV: + if (n->sym->ts.type != BT_LOGICAL) + bad = true; + break; + case OMP_REDUCTION_MAX: + case OMP_REDUCTION_MIN: + if (n->sym->ts.type != BT_INTEGER + && n->sym->ts.type != BT_REAL) + bad = true; + break; + case OMP_REDUCTION_IAND: + case OMP_REDUCTION_IOR: + case OMP_REDUCTION_IEOR: + if (n->sym->ts.type != BT_INTEGER) + bad = true; + break; + case OMP_REDUCTION_USER: + bad = true; + break; + default: + break; + } if (!bad) n->u2.udr = NULL; else diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 9599521b97c..29ad880a30c 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -1996,7 +1996,8 @@ omp_udr_find_orig (gfc_expr **e, int *walk_subtrees ATTRIBUTE_UNUSED, } static void -gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where) +gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where, + bool openacc) { gfc_symbol *sym = n->sym; gfc_symtree *root1 = NULL, *root2 = NULL, *root3 = NULL, *root4 = NULL; @@ -2251,21 +2252,24 @@ gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where) poplevel (0, 0); OMP_CLAUSE_REDUCTION_INIT (c) = stmt; - /* Create the merge statement list. */ - pushlevel (); - if (e4) - stmt = gfc_trans_assignment (e3, e4, false, true); - else - stmt = gfc_trans_call (n->u2.udr->combiner, false, - NULL_TREE, NULL_TREE, false); - if (TREE_CODE (stmt) != BIND_EXPR) - stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); - else - poplevel (0, 0); - OMP_CLAUSE_REDUCTION_MERGE (c) = stmt; + if (!openacc) + { + /* Create the merge statement list. */ + pushlevel (); + if (e4) + stmt = gfc_trans_assignment (e3, e4, false, true); + else + stmt = gfc_trans_call (n->u2.udr->combiner, false, + NULL_TREE, NULL_TREE, false); + if (TREE_CODE (stmt) != BIND_EXPR) + stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0)); + else + poplevel (0, 0); + OMP_CLAUSE_REDUCTION_MERGE (c) = stmt; - /* And stick the placeholder VAR_DECL into the clause as well. */ - OMP_CLAUSE_REDUCTION_PLACEHOLDER (c) = outer_decl; + /* And stick the placeholder VAR_DECL into the clause as well. */ + OMP_CLAUSE_REDUCTION_PLACEHOLDER (c) = outer_decl; + } gfc_current_locus = old_loc; @@ -2296,7 +2300,7 @@ gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where) static tree gfc_trans_omp_reduction_list (int kind, gfc_omp_namelist *namelist, tree list, - locus where, bool mark_addressable) + locus where, bool mark_addressable, bool openacc) { omp_clause_code clause = OMP_CLAUSE_REDUCTION; switch (kind) @@ -2376,7 +2380,8 @@ gfc_trans_omp_reduction_list (int kind, gfc_omp_namelist *namelist, tree list, if (namelist->sym->attr.dimension || namelist->u.reduction_op == OMP_REDUCTION_USER || namelist->sym->attr.allocatable) - gfc_trans_omp_array_reduction_or_udr (node, namelist, where); + gfc_trans_omp_array_reduction_or_udr (node, namelist, where, + openacc); list = gfc_trans_add_clause (node, list); } } @@ -2715,7 +2720,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, /* An OpenACC async clause indicates the need to set reduction arguments addressable, to allow asynchronous copy-out. */ omp_clauses = gfc_trans_omp_reduction_list (list, n, omp_clauses, - where, clauses->async); + where, clauses->async, + openacc); break; case OMP_LIST_PRIVATE: clause_code = OMP_CLAUSE_PRIVATE; @@ -5757,7 +5763,8 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock, on the simd construct and DO's clauses are translated elsewhere. */ do_clauses->sched_simd = false; - omp_clauses = gfc_trans_omp_clauses (pblock, do_clauses, code->loc); + omp_clauses = gfc_trans_omp_clauses (pblock, do_clauses, code->loc, false, + op == EXEC_OACC_LOOP); for (i = 0; i < collapse; i++) { diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index f3a056df8f2..4bbf30627c3 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -4426,9 +4426,16 @@ omp_reduction_init_op (location_t loc, enum tree_code op, tree type) { if (TREE_CODE (type) == ARRAY_TYPE) { + tree min_tree = TYPE_MIN_VALUE (TYPE_DOMAIN (type)); + tree max_tree = TYPE_MAX_VALUE (TYPE_DOMAIN (type)); + if (!TREE_CONSTANT (min_tree) || !TREE_CONSTANT (max_tree)) + { + error_at (loc, "array in reduction must be of constant size"); + return error_mark_node; + } vec *v = NULL; - HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (TYPE_DOMAIN (type))); - HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (TYPE_DOMAIN (type))); + HOST_WIDE_INT min = tree_to_shwi (min_tree); + HOST_WIDE_INT max = tree_to_shwi (max_tree); tree t = omp_reduction_init_op (loc, op, TREE_TYPE (type)); for (HOST_WIDE_INT i = min; i <= max; i++) CONSTRUCTOR_APPEND_ELT (v, size_int (i), t); @@ -7559,6 +7566,9 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, has_outer_reduction:; } + if (incoming == error_mark_node) + continue; + if (!ref_to_res) ref_to_res = integer_zero_node; diff --git a/gcc/testsuite/gfortran.dg/goacc/array-reduction.f90 b/gcc/testsuite/gfortran.dg/goacc/array-reduction.f90 index d71c400a5bf..f9a3b43e7f3 100644 --- a/gcc/testsuite/gfortran.dg/goacc/array-reduction.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/array-reduction.f90 @@ -1,74 +1,80 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + program test implicit none integer a(10), i a(:) = 0 - + ! Array reductions. - - !$acc parallel reduction (+:a) ! { dg-error "Array 'a' is not permitted in reduction" } + + !$acc parallel reduction (+:a) do i = 1, 10 a = a + 1 end do !$acc end parallel !$acc parallel - !$acc loop reduction (+:a) ! { dg-error "Array 'a' is not permitted in reduction" } + !$acc loop reduction (+:a) do i = 1, 10 a = a + 1 end do !$acc end parallel !$acc kernels - !$acc loop reduction (+:a) ! { dg-error "Array 'a' is not permitted in reduction" } + !$acc loop reduction (+:a) do i = 1, 10 a = a + 1 end do !$acc end kernels ! Subarray reductions. - - !$acc parallel reduction (+:a(1:5)) ! { dg-error "Array 'a' is not permitted in reduction" } + + !$acc parallel reduction (+:a(1:5)) do i = 1, 10 a = a + 1 end do !$acc end parallel !$acc parallel - !$acc loop reduction (+:a(1:5)) ! { dg-error "Array 'a' is not permitted in reduction" } + !$acc loop reduction (+:a(1:5)) do i = 1, 10 a = a + 1 end do !$acc end parallel !$acc kernels - !$acc loop reduction (+:a(1:5)) ! { dg-error "Array 'a' is not permitted in reduction" } + !$acc loop reduction (+:a(1:5)) do i = 1, 10 a = a + 1 end do !$acc end kernels ! Reductions on array elements. - - !$acc parallel reduction (+:a(1)) ! { dg-error "Array 'a' is not permitted in reduction" } + + !$acc parallel reduction (+:a(1)) do i = 1, 10 a(1) = a(1) + 1 end do !$acc end parallel !$acc parallel - !$acc loop reduction (+:a(1)) ! { dg-error "Array 'a' is not permitted in reduction" } + !$acc loop reduction (+:a(1)) do i = 1, 10 a(1) = a(1) + 1 end do !$acc end parallel !$acc kernels - !$acc loop reduction (+:a(1)) ! { dg-error "Array 'a' is not permitted in reduction" } + !$acc loop reduction (+:a(1)) do i = 1, 10 a(1) = a(1) + 1 end do !$acc end kernels - + print *, a end program test + +! { dg-final { scan-tree-dump-times "(?n)#pragma acc loop private\\(i\\) reduction\\(\\+:a\\)" 6 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_parallel reduction\\(\\+:a\\) map\\(tofrom:a \\\[len: \[0-9\]+\\\]\\)" 3 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction.f95 b/gcc/testsuite/gfortran.dg/goacc/reduction.f95 index a13574b150c..c425f00d87f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/reduction.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/reduction.f95 @@ -72,9 +72,9 @@ common /blk/ i1 !$acc end parallel !$acc parallel reduction (-:a1) ! { dg-error "OMP DECLARE REDUCTION - not found for type CHARACTER" } !$acc end parallel -!$acc parallel reduction (+:t1) ! { dg-error "OMP DECLARE REDUCTION \\+ not found for type TYPE" } +!$acc parallel reduction (+:t1) !$acc end parallel -!$acc parallel reduction (*:ta1) ! { dg-error "OMP DECLARE REDUCTION \\* not found for type TYPE" } +!$acc parallel reduction (*:ta1) !$acc end parallel !$acc parallel reduction (.and.:i3) ! { dg-error "OMP DECLARE REDUCTION \\.and\\. not found for type INTEGER" } !$acc end parallel @@ -108,9 +108,9 @@ common /blk/ i1 !$acc end parallel !$acc parallel reduction (max:a1) ! { dg-error "OMP DECLARE REDUCTION max not found for type CHARACTER" } !$acc end parallel -!$acc parallel reduction (min:t1) ! { dg-error "OMP DECLARE REDUCTION min not found for type TYPE" } +!$acc parallel reduction (min:t1) !$acc end parallel -!$acc parallel reduction (max:ta1) ! { dg-error "OMP DECLARE REDUCTION max not found for type TYPE" } +!$acc parallel reduction (max:ta1) !$acc end parallel !$acc parallel reduction (iand:r1) ! { dg-error "OMP DECLARE REDUCTION iand not found for type REAL" } !$acc end parallel @@ -130,32 +130,12 @@ common /blk/ i1 !$acc end parallel !$acc parallel reduction (ior:a1) ! { dg-error "OMP DECLARE REDUCTION ior not found for type CHARACTER" } !$acc end parallel -!$acc parallel reduction (ieor:t1) ! { dg-error "OMP DECLARE REDUCTION ieor not found for type TYPE" } +!$acc parallel reduction (ieor:t1) !$acc end parallel -!$acc parallel reduction (iand:ta1) ! { dg-error "OMP DECLARE REDUCTION iand not found for type TYPE" } +!$acc parallel reduction (iand:ta1) !$acc end parallel end subroutine -! { dg-error "Array 'ia2' is not permitted in reduction" "" { target "*-*-*" } 27 } -! { dg-error "Array 'ra1' is not permitted in reduction" "" { target "*-*-*" } 29 } -! { dg-error "Array 'ca1' is not permitted in reduction" "" { target "*-*-*" } 31 } -! { dg-error "Array 'da1' is not permitted in reduction" "" { target "*-*-*" } 33 } -! { dg-error "Array 'la1' is not permitted in reduction" "" { target "*-*-*" } 35 } ! { dg-error "Array 'aa1' is not permitted in reduction" "" { target "*-*-*" } 65 } ! { dg-error "Array 'ia1' is not permitted in reduction" "" { target "*-*-*" } 67 } -! { dg-error "Array 'la1' is not permitted in reduction" "" { target "*-*-*" } 71 } -! { dg-error "Array 'ta1' is not permitted in reduction" "" { target "*-*-*" } 77 } -! { dg-error "Array 'ia2' is not permitted in reduction" "" { target "*-*-*" } 81 } -! { dg-error "Array 'ra1' is not permitted in reduction" "" { target "*-*-*" } 85 } -! { dg-error "Array 'da1' is not permitted in reduction" "" { target "*-*-*" } 89 } -! { dg-error "Array 'ca1' is not permitted in reduction" "" { target "*-*-*" } 93 } -! { dg-error "Array 'ta1' is not permitted in reduction" "" { target "*-*-*" } 99 } -! { dg-error "Array 'ca1' is not permitted in reduction" "" { target "*-*-*" } 103 } -! { dg-error "Array 'la1' is not permitted in reduction" "" { target "*-*-*" } 107 } -! { dg-error "Array 'ta1' is not permitted in reduction" "" { target "*-*-*" } 113 } -! { dg-error "Array 'ra1' is not permitted in reduction" "" { target "*-*-*" } 117 } -! { dg-error "Array 'da1' is not permitted in reduction" "" { target "*-*-*" } 121 } -! { dg-error "Array 'ca1' is not permitted in reduction" "" { target "*-*-*" } 125 } -! { dg-error "Array 'la1' is not permitted in reduction" "" { target "*-*-*" } 129 } -! { dg-error "Array 'ta1' is not permitted in reduction" "" { target "*-*-*" } 135 } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-10.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-10.f90 new file mode 100644 index 00000000000..506dfaf29f6 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-10.f90 @@ -0,0 +1,483 @@ +! { dg-do run } + +! real array reductions + +program reduction_10 + implicit none + + integer, parameter :: n = 10, ng = 8, nw = 4, vl = 32 + integer :: i, j + real, dimension (n) :: vresult, rg, rw, rv, rc + logical, dimension (n) :: lrg, lrw, lrv, lrc, lvresult + real, dimension (n) :: array + + do i = 1, n + array(i) = i + end do + + ! + ! '+' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(+:rg) gang + do i = 1, n + do j = 1, n + rg(j) = rg(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(+:rw) worker + do i = 1, n + do j = 1, n + rw(j) = rw(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(+:rv) vector + do i = 1, n + do j = 1, n + rv(j) = rv(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(+:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = rc(j) + array(i) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = vresult(j) + array(i) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 1 + if (count (rw .ne. vresult) .ne. 0) STOP 2 + if (count (rv .ne. vresult) .ne. 0) STOP 3 + if (count (rc .ne. vresult) .ne. 0) STOP 4 + + ! + ! '*' reductions + ! + + rg = 1 + rw = 1 + rv = 1 + rc = 1 + vresult = 1 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(*:rg) gang + do i = 1, n + do j = 1, n + rg(j) = rg(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(*:rw) worker + do i = 1, n + do j = 1, n + rw(j) = rw(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(*:rv) vector + do i = 1, n + do j = 1, n + rv(j) = rv(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(*:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = rc(j) * array(i) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = vresult(j) * array(i) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 5 + if (count (rw .ne. vresult) .ne. 0) STOP 6 + if (count (rv .ne. vresult) .ne. 0) STOP 7 + if (count (rc .ne. vresult) .ne. 0) STOP 8 + + ! + ! 'max' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(max:rg) gang + do i = 1, n + do j = 1, n + rg(j) = max (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(max:rw) worker + do i = 1, n + do j = 1, n + rw(j) = max (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(max:rv) vector + do i = 1, n + do j = 1, n + rv(j) = max (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(max:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = max (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = max (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 9 + if (count (rw .ne. vresult) .ne. 0) STOP 10 + if (count (rv .ne. vresult) .ne. 0) STOP 11 + if (count (rc .ne. vresult) .ne. 0) STOP 12 + + ! + ! 'min' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(min:rg) gang + do i = 1, n + do j = 1, n + rg(j) = min (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(min:rw) worker + do i = 1, n + do j = 1, n + rw(j) = min (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(min:rv) vector + do i = 1, n + do j = 1, n + rv(j) = min (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(min:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = min (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = min (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 13 + if (count (rw .ne. vresult) .ne. 0) STOP 14 + if (count (rv .ne. vresult) .ne. 0) STOP 15 + if (count (rc .ne. vresult) .ne. 0) STOP 16 + + ! + ! '.and.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.and.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.and.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.and.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.and.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .and. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 17 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 18 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 19 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 20 + + ! + ! '.or.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.or.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.or.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.or.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.or.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .or. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 21 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 22 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 23 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 24 + + ! + ! '.eqv.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.eqv.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.eqv.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.eqv.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.eqv.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .eqv. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 25 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 26 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 27 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 28 + + ! + ! '.neqv.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.neqv.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.neqv.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.neqv.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.neqv.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .neqv. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 29 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 30 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 31 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 32 + +end program reduction_10 + diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-11.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-11.f90 new file mode 100644 index 00000000000..4bec1c797cd --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-11.f90 @@ -0,0 +1,483 @@ +! { dg-do run } + +! double precision array reductions + +program reduction_11 + implicit none + + integer, parameter :: n = 10, ng = 8, nw = 4, vl = 32 + integer :: i, j + double precision, dimension (n) :: vresult, rg, rw, rv, rc + logical, dimension (n) :: lrg, lrw, lrv, lrc, lvresult + double precision, dimension (n) :: array + + do i = 1, n + array(i) = i + end do + + ! + ! '+' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(+:rg) gang + do i = 1, n + do j = 1, n + rg(j) = rg(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(+:rw) worker + do i = 1, n + do j = 1, n + rw(j) = rw(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(+:rv) vector + do i = 1, n + do j = 1, n + rv(j) = rv(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(+:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = rc(j) + array(i) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = vresult(j) + array(i) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 1 + if (count (rw .ne. vresult) .ne. 0) STOP 2 + if (count (rv .ne. vresult) .ne. 0) STOP 3 + if (count (rc .ne. vresult) .ne. 0) STOP 4 + + ! + ! '*' reductions + ! + + rg = 1 + rw = 1 + rv = 1 + rc = 1 + vresult = 1 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(*:rg) gang + do i = 1, n + do j = 1, n + rg(j) = rg(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(*:rw) worker + do i = 1, n + do j = 1, n + rw(j) = rw(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(*:rv) vector + do i = 1, n + do j = 1, n + rv(j) = rv(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(*:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = rc(j) * array(i) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = vresult(j) * array(i) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 5 + if (count (rw .ne. vresult) .ne. 0) STOP 6 + if (count (rv .ne. vresult) .ne. 0) STOP 7 + if (count (rc .ne. vresult) .ne. 0) STOP 8 + + ! + ! 'max' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(max:rg) gang + do i = 1, n + do j = 1, n + rg(j) = max (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(max:rw) worker + do i = 1, n + do j = 1, n + rw(j) = max (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(max:rv) vector + do i = 1, n + do j = 1, n + rv(j) = max (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(max:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = max (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = max (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 9 + if (count (rw .ne. vresult) .ne. 0) STOP 10 + if (count (rv .ne. vresult) .ne. 0) STOP 11 + if (count (rc .ne. vresult) .ne. 0) STOP 12 + + ! + ! 'min' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(min:rg) gang + do i = 1, n + do j = 1, n + rg(j) = min (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(min:rw) worker + do i = 1, n + do j = 1, n + rw(j) = min (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(min:rv) vector + do i = 1, n + do j = 1, n + rv(j) = min (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(min:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = min (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = min (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 13 + if (count (rw .ne. vresult) .ne. 0) STOP 14 + if (count (rv .ne. vresult) .ne. 0) STOP 15 + if (count (rc .ne. vresult) .ne. 0) STOP 16 + + ! + ! '.and.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.and.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.and.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.and.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.and.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .and. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 17 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 18 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 19 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 20 + + ! + ! '.or.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.or.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.or.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.or.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.or.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .or. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 21 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 22 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 23 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 24 + + ! + ! '.eqv.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.eqv.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.eqv.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.eqv.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.eqv.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .eqv. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 25 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 26 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 27 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 28 + + ! + ! '.neqv.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.neqv.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.neqv.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.neqv.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.neqv.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .neqv. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 29 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 30 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 31 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 32 + +end program reduction_11 + diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-12.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-12.f90 new file mode 100644 index 00000000000..b609c7a294e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-12.f90 @@ -0,0 +1,135 @@ +! { dg-do run } + +! complex array reductions + +program reduction_12 + implicit none + + integer, parameter :: n = 10, ng = 8, nw = 4, vl = 32 + integer :: i, j + complex, dimension (n) :: vresult, rg, rw, rv, rc + logical, dimension (n) :: lrg, lrw, lrv, lrc, lvresult + complex, dimension (n) :: array + + do i = 1, n + array(i) = i + end do + + ! + ! '+' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(+:rg) gang + do i = 1, n + do j = 1, n + rg(j) = rg(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(+:rw) worker + do i = 1, n + do j = 1, n + rw(j) = rw(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(+:rv) vector + do i = 1, n + do j = 1, n + rv(j) = rv(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(+:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = rc(j) + array(i) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = vresult(j) + array(i) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 1 + if (count (rw .ne. vresult) .ne. 0) STOP 2 + if (count (rv .ne. vresult) .ne. 0) STOP 3 + if (count (rc .ne. vresult) .ne. 0) STOP 4 + + ! + ! '*' reductions + ! + + rg = 1 + rw = 1 + rv = 1 + rc = 1 + vresult = 1 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(*:rg) gang + do i = 1, n + do j = 1, n + rg(j) = rg(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(*:rw) worker + do i = 1, n + do j = 1, n + rw(j) = rw(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(*:rv) vector + do i = 1, n + do j = 1, n + rv(j) = rv(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(*:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = rc(j) * array(i) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = vresult(j) * array(i) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 5 + if (count (rw .ne. vresult) .ne. 0) STOP 6 + if (count (rv .ne. vresult) .ne. 0) STOP 7 + if (count (rc .ne. vresult) .ne. 0) STOP 8 + +end program reduction_12 + diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90 new file mode 100644 index 00000000000..088c5cd3b04 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-13.f90 @@ -0,0 +1,66 @@ +! { dg-do run } + +! record type reductions + +program reduction_13 + implicit none + + type t1 + integer :: i + real :: r + end type t1 + + type t2 + real :: r + integer :: i + double precision :: d + end type t2 + + integer, parameter :: n = 10, ng = 8, nw = 4, vl = 32 + integer :: i + type(t1) :: v1, a1 + type (t2) :: v2, a2 + + v1%i = 0 + v1%r = 0 + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(v1) + !$acc loop reduction (+:v1) + do i = 1, n + v1%i = v1%i + 1 + v1%r = v1%r + 2 + end do + !$acc end parallel + a1%i = 0 + a1%r = 0 + do i = 1, n + a1%i = a1%i + 1 + a1%r = a1%r + 2 + end do + if (v1%i .ne. a1%i) STOP 1 + if (v1%r .ne. a1%r) STOP 2 + + v2%i = 1 + v2%r = 1 + v2%d = 1 + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(v2) + !$acc loop reduction (*:v2) + do i = 1, n + v2%i = v2%i * 2 + v2%r = v2%r * 1.1 + v2%d = v2%d * 1.3 + end do + !$acc end parallel + a2%i = 1 + a2%r = 1 + a2%d = 1 + do i = 1, n + a2%i = a2%i * 2 + a2%r = a2%r * 1.1 + a2%d = a2%d * 1.3 + end do + + if (v2%i .ne. a2%i) STOP 3 + if (v2%r .ne. a2%r) STOP 4 + if (v2%d .ne. a2%d) STOP 5 + +end program reduction_13 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 new file mode 100644 index 00000000000..43ab155aa73 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 @@ -0,0 +1,657 @@ +! { dg-do run } + +! integer array reductions + +program reduction_9 + implicit none + + integer, parameter :: n = 10, ng = 8, nw = 4, vl = 32 + integer :: i, j + integer, dimension (n) :: vresult, rg, rw, rv, rc + logical, dimension (n) :: lrg, lrw, lrv, lrc, lvresult + integer, dimension (n) :: array + + do i = 1, n + array(i) = i + end do + + ! + ! '+' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(+:rg) gang + do i = 1, n + do j = 1, n + rg(j) = rg(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(+:rw) worker + do i = 1, n + do j = 1, n + rw(j) = rw(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(+:rv) vector + do i = 1, n + do j = 1, n + rv(j) = rv(j) + array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(+:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = rc(j) + array(i) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = vresult(j) + array(i) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 1 + if (count (rw .ne. vresult) .ne. 0) STOP 2 + if (count (rv .ne. vresult) .ne. 0) STOP 3 + if (count (rc .ne. vresult) .ne. 0) STOP 4 + + ! + ! '*' reductions + ! + + rg = 1 + rw = 1 + rv = 1 + rc = 1 + vresult = 1 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(*:rg) gang + do i = 1, n + do j = 1, n + rg(j) = rg(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(*:rw) worker + do i = 1, n + do j = 1, n + rw(j) = rw(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(*:rv) vector + do i = 1, n + do j = 1, n + rv(j) = rv(j) * array(i) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(*:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = rc(j) * array(i) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = vresult(j) * array(i) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 5 + if (count (rw .ne. vresult) .ne. 0) STOP 6 + if (count (rv .ne. vresult) .ne. 0) STOP 7 + if (count (rc .ne. vresult) .ne. 0) STOP 8 + + ! + ! 'max' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(max:rg) gang + do i = 1, n + do j = 1, n + rg(j) = max (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(max:rw) worker + do i = 1, n + do j = 1, n + rw(j) = max (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(max:rv) vector + do i = 1, n + do j = 1, n + rv(j) = max (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(max:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = max (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = max (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 9 + if (count (rw .ne. vresult) .ne. 0) STOP 10 + if (count (rv .ne. vresult) .ne. 0) STOP 11 + if (count (rc .ne. vresult) .ne. 0) STOP 12 + + ! + ! 'min' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(min:rg) gang + do i = 1, n + do j = 1, n + rg(j) = min (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(min:rw) worker + do i = 1, n + do j = 1, n + rw(j) = min (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(min:rv) vector + do i = 1, n + do j = 1, n + rv(j) = min (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(min:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = min (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = min (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 13 + if (count (rw .ne. vresult) .ne. 0) STOP 14 + if (count (rv .ne. vresult) .ne. 0) STOP 15 + if (count (rc .ne. vresult) .ne. 0) STOP 16 + + ! + ! 'iand' reductions + ! + + rg = 1 + rw = 1 + rv = 1 + rc = 1 + vresult = 1 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(iand:rg) gang + do i = 1, n + do j = 1, n + rg(j) = iand (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(iand:rw) worker + do i = 1, n + do j = 1, n + rw(j) = iand (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(iand:rv) vector + do i = 1, n + do j = 1, n + rv(j) = iand (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(iand:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = iand (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = iand (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 17 + if (count (rw .ne. vresult) .ne. 0) STOP 18 + if (count (rv .ne. vresult) .ne. 0) STOP 19 + if (count (rc .ne. vresult) .ne. 0) STOP 20 + + ! + ! 'ior' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(ior:rg) gang + do i = 1, n + do j = 1, n + rg(j) = ior (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(ior:rw) worker + do i = 1, n + do j = 1, n + rw(j) = ior (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(ior:rv) vector + do i = 1, n + do j = 1, n + rv(j) = ior (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(ior:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = ior (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = ior (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 21 + if (count (rw .ne. vresult) .ne. 0) STOP 22 + if (count (rv .ne. vresult) .ne. 0) STOP 23 + if (count (rc .ne. vresult) .ne. 0) STOP 24 + + ! + ! 'ieor' reductions + ! + + rg = 0 + rw = 0 + rv = 0 + rc = 0 + vresult = 0 + + !$acc parallel num_gangs(ng) copy(rg) + !$acc loop reduction(ieor:rg) gang + do i = 1, n + do j = 1, n + rg(j) = ieor (rg(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(rw) + !$acc loop reduction(ieor:rw) worker + do i = 1, n + do j = 1, n + rw(j) = ieor (rw(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(rv) + !$acc loop reduction(ieor:rv) vector + do i = 1, n + do j = 1, n + rv(j) = ieor (rv(j), array(i)) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(rc) + !$acc loop reduction(ieor:rc) gang worker vector + do i = 1, n + do j = 1, n + rc(j) = ieor (rc(j), array(i)) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + vresult(j) = ieor (vresult(j), array(i)) + end do + end do + + if (count (rg .ne. vresult) .ne. 0) STOP 25 + if (count (rw .ne. vresult) .ne. 0) STOP 26 + if (count (rv .ne. vresult) .ne. 0) STOP 27 + if (count (rc .ne. vresult) .ne. 0) STOP 28 + + ! + ! '.and.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.and.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.and.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.and.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.and.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .and. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .and. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 29 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 30 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 31 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 32 + + ! + ! '.or.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.or.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.or.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.or.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.or.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .or. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .or. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 33 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 34 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 35 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 36 + + ! + ! '.eqv.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.eqv.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.eqv.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.eqv.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.eqv.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .eqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .eqv. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 37 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 38 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 39 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 40 + + ! + ! '.neqv.' reductions + ! + + lrg = .true. + lrw = .true. + lrv = .true. + lrc = .true. + lvresult = .true. + + !$acc parallel num_gangs(ng) copy(lrg) + !$acc loop reduction(.neqv.:lrg) gang + do i = 1, n + do j = 1, n + lrg(j) = lrg(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_workers(nw) copy(lrw) + !$acc loop reduction(.neqv.:lrw) worker + do i = 1, n + do j = 1, n + lrw(j) = lrw(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel vector_length(vl) copy(lrv) + !$acc loop reduction(.neqv.:lrv) vector + do i = 1, n + do j = 1, n + lrv(j) = lrv(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + !$acc parallel num_gangs(ng) num_workers(nw) vector_length(vl) copy(lrc) + !$acc loop reduction(.neqv.:lrc) gang worker vector + do i = 1, n + do j = 1, n + lrc(j) = lrc(j) .neqv. (array(i) .ge. 5) + end do + end do + !$acc end parallel + + ! Verify the results + do i = 1, n + do j = 1, n + lvresult(j) = lvresult(j) .neqv. (array(i) .ge. 5) + end do + end do + + if (count (lrg .neqv. lvresult) .ne. 0) STOP 41 + if (count (lrw .neqv. lvresult) .ne. 0) STOP 42 + if (count (lrv .neqv. lvresult) .ne. 0) STOP 43 + if (count (lrc .neqv. lvresult) .ne. 0) STOP 44 + +end program reduction_9 +