From patchwork Thu Aug 10 13:33:02 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 133999 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:b824:0:b0:3f2:4152:657d with SMTP id z4csp424698vqi; Thu, 10 Aug 2023 06:35:23 -0700 (PDT) X-Google-Smtp-Source: AGHT+IHHnezSy//Q3tzqVEE8cY6hwQrJeCje63sTwdt8B3lKAn03MlP4rF5MhIcG2/AE1ovL7Z4l X-Received: by 2002:aa7:d812:0:b0:521:728f:d84e with SMTP id v18-20020aa7d812000000b00521728fd84emr2752542edq.0.1691674523442; Thu, 10 Aug 2023 06:35:23 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1691674523; cv=none; d=google.com; s=arc-20160816; b=pbvhb0t3TzT+4Z28QRPV514Be+P++vf5T0MAQusNDu+urdJp1Qm56nfFVxL9+7njqE ES5PLVKdlacioBAszItk3V2WDcQO4Nas+CO8PVFbGaqMWBb53fwkNAc1SQwOgvheWE3q sIZInqR0zl1XtGe4g0xHuDDaKcPHYw/vynvC0HmG7v3RZko3NQDwhltJGr4cwqSYwAu1 OEEEAl7YWqEX44R+2050sDvCGeVcmmGsU8xgEqFtfS7N+EzjFW3PmjC7yz1ASi2liuqz d2vLlLwaHQN0TpVtOM4+7llnUCuiUI1xXIvq1dmNe9JtJLIymxea7KpDRokQOfNjVfdw gggQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:references:in-reply-to:message-id:date:subject:cc:to :from:ironport-sdr:dmarc-filter:delivered-to; bh=OSWrWPMBIyFGEZDarUY+OZiOr7oXZy9xqK/Q8Mr0daY=; fh=6entJIlBX3Lqynes5VYRRY4zgfJfSDVLFnUjulEf2Vo=; b=DufbJY2P5Hz9t4TIIxjKbaTkspBALtLtPn+Y74W06IledFEmDDKcvhVE+11SkIovsx KJsjLZlpLtODFZF1BAc3AjnP35/du7qCDJ8YJCMm7uqvQsg6INZ+Rrd/4NS2WWCW8j4R uJkYRfBfYUZDIev1VhVTZmv/aOPyx+dE6W/AEucr1/ny6RSH6w4toTKhs4sb0BNaOaEj OAd3klCC+XiC3RoWVDQiMUVSZBT7QERDi6/poItdgT9I9DiCRc4anfxe41FEnvWR8i2g KmvK6bPZkgu45cknB9L8Fi+SfOtQbRD7bO2o16DP7riHGNSuEQnRKZM00bfAlCZEFppu hYHA== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id r6-20020aa7c146000000b0051e1a616e32si1459208edp.89.2023.08.10.06.35.23 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 10 Aug 2023 06:35:23 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c 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 62936385E006 for ; Thu, 10 Aug 2023 13:34:04 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id A85F43857706; Thu, 10 Aug 2023 13:33:25 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org A85F43857706 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,162,1684828800"; d="scan'208";a="14111544" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 10 Aug 2023 05:33:24 -0800 IronPort-SDR: tYWPsq6yA9qEQ7mPPzk3DmptdqbbE/JGZsP2KJIeBRSRQvMgPGPplZS7IxM8yF6C68P+P+Y1ik UUU3v85zJz6e7eUNIN0e3nqPi5R+iaLQqJ69jLiMfJDwZ+nMsgteR8/hA2llAjswDOU/gjcS0h pyayp8oHXD0EiBKYBGsRsy8W6H0JTU5YJeMs2q87HwWkQb8jAy4HhH3IX22vXATpAWUQDsKsz8 zTKZArOfkqM+XPS/7y8dyDAs3CbneWUHju266udroifgAbIjdP8aupcYhDRRk8z6dfkqmSwzgo hMU= From: Julian Brown To: CC: , , Subject: [PATCH 1/5] OpenMP: Move Fortran 'declare mapper' instantiation code Date: Thu, 10 Aug 2023 13:33:02 +0000 Message-ID: X-Mailer: git-send-email 2.25.1 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham 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.29 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 Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1773849304799409370 X-GMAIL-MSGID: 1773849304799409370 This patch moves the code for explicit 'declare mapper' directive instantiation in the Fortran front-end to openmp.cc from trans-openmp.cc. The transformation takes place entirely in the front end's own representation and doesn't involve middle-end trees at all. Also, having the code in openmp.cc is more convenient for the following patch that introduces the 'resolve_omp_mapper_clauses' function. 2023-08-10 Julian Brown gcc/fortran/ * gfortran.h (toc_directive): Move here. (gfc_omp_instantiate_mappers, gfc_get_location): Add prototypes. * openmp.cc (omp_split_map_op, omp_join_map_op, omp_map_decayed_kind, omp_basic_map_kind_name, gfc_subst_replace, gfc_subst_prepend_ref, gfc_subst_in_expr_1, gfc_subst_in_expr, gfc_subst_mapper_var): Move here. (gfc_omp_instantiate_mapper, gfc_omp_instantiate_mappers): Move here and rename. * trans-openmp.cc (toc_directive, omp_split_map_op, omp_join_map_op, omp_map_decayed_kind, gfc_subst_replace, gfc_subst_prepend_ref, gfc_subst_in_expr_1, gfc_subst_in_expr, gfc_subst_mapper_var, gfc_trans_omp_instantiate_mapper, gfc_trans_omp_instantiate_mappers): Remove from here. (gfc_trans_omp_target, gfc_trans_omp_target_data, gfc_trans_omp_target_enter_data, gfc_trans_omp_target_exit_data): Rename calls to gfc_omp_instantiate_mappers. --- gcc/fortran/gfortran.h | 16 ++ gcc/fortran/openmp.cc | 435 ++++++++++++++++++++++++++++++++++++ gcc/fortran/trans-openmp.cc | 388 +------------------------------- 3 files changed, 456 insertions(+), 383 deletions(-) diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 0e7e80e4bf1..788b3797893 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -3246,6 +3246,18 @@ typedef struct gfc_finalizer gfc_finalizer; #define gfc_get_finalizer() XCNEW (gfc_finalizer) +/* Control clause translation per-directive for gfc_trans_omp_clauses. Also + used for gfc_omp_instantiate_mappers. */ + +enum toc_directive +{ + TOC_OPENMP, + TOC_OPENMP_DECLARE_SIMD, + TOC_OPENMP_DECLARE_MAPPER, + TOC_OPENMP_EXIT_DATA, + TOC_OPENACC, + TOC_OPENACC_DECLARE +}; /************************ Function prototypes *************************/ @@ -3707,6 +3719,9 @@ void gfc_resolve_omp_do_blocks (gfc_code *, gfc_namespace *); void gfc_resolve_omp_declare_simd (gfc_namespace *); void gfc_resolve_omp_udrs (gfc_symtree *); void gfc_resolve_omp_udms (gfc_symtree *); +void gfc_omp_instantiate_mappers (gfc_code *, gfc_omp_clauses *, + toc_directive = TOC_OPENMP, + int = OMP_LIST_MAP); void gfc_omp_save_and_clear_state (struct gfc_omp_saved_state *); void gfc_omp_restore_state (struct gfc_omp_saved_state *); void gfc_free_expr_list (gfc_expr_list *); @@ -3956,6 +3971,7 @@ bool gfc_convert_to_structure_constructor (gfc_expr *, gfc_symbol *, /* trans.cc */ void gfc_generate_code (gfc_namespace *); void gfc_generate_module_code (gfc_namespace *); +location_t gfc_get_location (locus *); /* trans-intrinsic.cc */ bool gfc_inline_intrinsic_function_p (gfc_expr *); diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index deccb14a525..0f715a6f997 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -12584,6 +12584,441 @@ gfc_resolve_omp_udrs (gfc_symtree *st) gfc_resolve_omp_udr (omp_udr); } +static enum gfc_omp_map_op +omp_split_map_op (enum gfc_omp_map_op op, bool *force_p, bool *always_p, + bool *present_p) +{ + *force_p = *always_p = *present_p = false; + + switch (op) + { + case OMP_MAP_FORCE_ALLOC: + case OMP_MAP_FORCE_TO: + case OMP_MAP_FORCE_FROM: + case OMP_MAP_FORCE_TOFROM: + case OMP_MAP_FORCE_PRESENT: + *force_p = true; + break; + case OMP_MAP_ALWAYS_TO: + case OMP_MAP_ALWAYS_FROM: + case OMP_MAP_ALWAYS_TOFROM: + *always_p = true; + break; + case OMP_MAP_ALWAYS_PRESENT_TO: + case OMP_MAP_ALWAYS_PRESENT_FROM: + case OMP_MAP_ALWAYS_PRESENT_TOFROM: + *always_p = true; + /* Fallthrough. */ + case OMP_MAP_PRESENT_ALLOC: + case OMP_MAP_PRESENT_TO: + case OMP_MAP_PRESENT_FROM: + case OMP_MAP_PRESENT_TOFROM: + *present_p = true; + break; + default: + ; + } + + switch (op) + { + case OMP_MAP_ALLOC: + case OMP_MAP_FORCE_ALLOC: + case OMP_MAP_PRESENT_ALLOC: + return OMP_MAP_ALLOC; + case OMP_MAP_TO: + case OMP_MAP_FORCE_TO: + case OMP_MAP_ALWAYS_TO: + case OMP_MAP_PRESENT_TO: + case OMP_MAP_ALWAYS_PRESENT_TO: + return OMP_MAP_TO; + case OMP_MAP_FROM: + case OMP_MAP_FORCE_FROM: + case OMP_MAP_ALWAYS_FROM: + case OMP_MAP_PRESENT_FROM: + case OMP_MAP_ALWAYS_PRESENT_FROM: + return OMP_MAP_FROM; + case OMP_MAP_TOFROM: + case OMP_MAP_FORCE_TOFROM: + case OMP_MAP_ALWAYS_TOFROM: + case OMP_MAP_PRESENT_TOFROM: + case OMP_MAP_ALWAYS_PRESENT_TOFROM: + return OMP_MAP_TOFROM; + default: + ; + } + return op; +} + +static enum gfc_omp_map_op +omp_join_map_op (enum gfc_omp_map_op op, bool force_p, bool always_p, + bool present_p) +{ + gcc_assert (!force_p || !(always_p || present_p)); + + switch (op) + { + case OMP_MAP_ALLOC: + if (force_p) + return OMP_MAP_FORCE_ALLOC; + else if (present_p) + return OMP_MAP_PRESENT_ALLOC; + break; + + case OMP_MAP_TO: + if (force_p) + return OMP_MAP_FORCE_TO; + else if (always_p && present_p) + return OMP_MAP_ALWAYS_PRESENT_TO; + else if (always_p) + return OMP_MAP_ALWAYS_TO; + else if (present_p) + return OMP_MAP_PRESENT_TO; + break; + + case OMP_MAP_FROM: + if (force_p) + return OMP_MAP_FORCE_FROM; + else if (always_p && present_p) + return OMP_MAP_ALWAYS_PRESENT_FROM; + else if (always_p) + return OMP_MAP_ALWAYS_FROM; + else if (present_p) + return OMP_MAP_PRESENT_FROM; + break; + + case OMP_MAP_TOFROM: + if (force_p) + return OMP_MAP_FORCE_TOFROM; + else if (always_p && present_p) + return OMP_MAP_ALWAYS_PRESENT_TOFROM; + else if (always_p) + return OMP_MAP_ALWAYS_TOFROM; + else if (present_p) + return OMP_MAP_PRESENT_TOFROM; + break; + + default: + ; + } + + return op; +} + +/* Map kind decay (OpenMP 5.2, 5.8.8 "declare mapper Directive"). Return the + map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS + specified on the clause that invokes the mapper. See also + c-family/c-omp.cc:omp_map_decayed_kind. */ + +static enum gfc_omp_map_op +omp_map_decayed_kind (enum gfc_omp_map_op mapper_kind, + enum gfc_omp_map_op invoked_as, bool exit_p) +{ + if (invoked_as == OMP_MAP_RELEASE || invoked_as == OMP_MAP_DELETE) + return invoked_as; + + bool force_p, always_p, present_p; + + invoked_as = omp_split_map_op (invoked_as, &force_p, &always_p, &present_p); + gfc_omp_map_op decay_to; + + switch (mapper_kind) + { + case OMP_MAP_ALLOC: + if (exit_p && invoked_as == OMP_MAP_FROM) + decay_to = OMP_MAP_RELEASE; + else + decay_to = OMP_MAP_ALLOC; + break; + + case OMP_MAP_TO: + if (invoked_as == OMP_MAP_FROM) + decay_to = exit_p ? OMP_MAP_RELEASE : OMP_MAP_ALLOC; + else if (invoked_as == OMP_MAP_ALLOC) + decay_to = OMP_MAP_ALLOC; + else + decay_to = OMP_MAP_TO; + break; + + case OMP_MAP_FROM: + if (invoked_as == OMP_MAP_ALLOC || invoked_as == OMP_MAP_TO) + decay_to = OMP_MAP_ALLOC; + else + decay_to = OMP_MAP_FROM; + break; + + case OMP_MAP_TOFROM: + case OMP_MAP_UNSET: + decay_to = invoked_as; + break; + + default: + gcc_unreachable (); + } + + return omp_join_map_op (decay_to, force_p, always_p, present_p); +} + +static const char * +omp_basic_map_kind_name (enum gfc_omp_map_op op) +{ + switch (op) + { + case OMP_MAP_ALLOC: + return "ALLOC"; + case OMP_MAP_TO: + return "TO"; + case OMP_MAP_FROM: + return "FROM"; + case OMP_MAP_TOFROM: + return "TOFROM"; + case OMP_MAP_RELEASE: + return "RELEASE"; + case OMP_MAP_DELETE: + return "DELETE"; + default: + gcc_unreachable (); + } +} + +static gfc_symtree *gfc_subst_replace; +static gfc_ref *gfc_subst_prepend_ref; + +static bool +gfc_subst_in_expr_1 (gfc_expr *expr, gfc_symbol *search, int *) +{ + /* The base-object for component accesses may be stored in expr->symtree. + If it's the symbol for our "declare mapper" placeholder variable, + substitute it. */ + if (expr->symtree && expr->symtree->n.sym == search) + { + gfc_ref **lastptr = NULL; + expr->symtree = gfc_subst_replace; + + if (!gfc_subst_prepend_ref) + return false; + + gfc_ref *prepend_ref = gfc_copy_ref (gfc_subst_prepend_ref); + + for (gfc_ref *walk = prepend_ref; walk; walk = walk->next) + lastptr = &walk->next; + + *lastptr = expr->ref; + expr->ref = prepend_ref; + } + + return false; +} + +static void +gfc_subst_in_expr (gfc_expr *expr, gfc_symbol *search, gfc_symtree *replace, + gfc_ref *prepend_ref) +{ + gfc_subst_replace = replace; + gfc_subst_prepend_ref = prepend_ref; + gfc_traverse_expr (expr, search, gfc_subst_in_expr_1, 0); +} + +static void +gfc_subst_mapper_var (gfc_symbol **out_sym, gfc_expr **out_expr, + gfc_symbol *orig_sym, gfc_expr *orig_expr, + gfc_symbol *dummy_var, + gfc_symbol *templ_sym, gfc_expr *templ_expr) +{ + gfc_ref *orig_ref = orig_expr ? orig_expr->ref : NULL; + gfc_symtree *orig_st = gfc_find_symtree (orig_sym->ns->sym_root, + orig_sym->name); + + if (dummy_var == templ_sym) + *out_sym = orig_sym; + else + *out_sym = templ_sym; + + if (templ_expr) + { + *out_expr = gfc_copy_expr (templ_expr); + gfc_subst_in_expr (*out_expr, dummy_var, orig_st, orig_ref); + } + else if (orig_expr) + *out_expr = gfc_copy_expr (orig_expr); + else + *out_expr = NULL; +} + +static gfc_omp_namelist ** +gfc_omp_instantiate_mapper (gfc_omp_namelist **outlistp, + gfc_omp_namelist *clause, + gfc_omp_map_op outer_map_op, gfc_omp_udm *udm, + toc_directive cd, int list) +{ + /* Here "sym" and "expr" describe the clause as written, to be substituted + for the dummy variable in the mapper definition. */ + struct gfc_symbol *sym = clause->sym; + struct gfc_expr *expr = clause->expr; + gfc_omp_namelist *mapper_clause = udm->clauses->lists[OMP_LIST_MAP]; + bool pointer_needed_p = false; + + if (expr) + { + gfc_ref *lastref = expr->ref, *lastcomp = NULL; + + for (; lastref->next; lastref = lastref->next) + if (lastref->type == REF_COMPONENT) + lastcomp = lastref; + + if (lastref + && lastref->type == REF_ARRAY + && (lastref->u.ar.type == AR_SECTION + || lastref->u.ar.type == AR_FULL)) + { + mpz_t elems; + bool multiple_elems_p = false; + + if (gfc_array_size (expr, &elems)) + { + HOST_WIDE_INT nelems = gfc_mpz_get_hwi (elems); + if (nelems > 1) + multiple_elems_p = true; + } + else + multiple_elems_p = true; + + if (multiple_elems_p && clause->u2.udm) + { + clause->u2.udm->multiple_elems_p = true; + *outlistp = clause; + return &(*outlistp)->next; + } + } + + if (lastcomp + && lastcomp->type == REF_COMPONENT + && (lastcomp->u.c.component->attr.pointer + || lastcomp->u.c.component->attr.allocatable)) + pointer_needed_p = true; + } + + if (pointer_needed_p) + { + /* If we're instantiating a mapper via a pointer, we need to map that + pointer as well as mapping the entities explicitly listed in the + mapper definition. Create a node for that. */ + gfc_omp_namelist *new_clause = gfc_get_omp_namelist (); + new_clause->sym = sym; + new_clause->expr = gfc_copy_expr (expr); + /* We want the pointer itself: cut off any further accessors after the + last component reference (e.g. array indices). */ + gfc_ref *lastcomp = NULL; + for (gfc_ref *lastref = new_clause->expr->ref; + lastref; + lastref = lastref->next) + if (lastref->type == REF_COMPONENT) + lastcomp = lastref; + gcc_assert (lastcomp != NULL); + lastcomp->next = NULL; + new_clause->u.map_op = OMP_MAP_POINTER_ONLY; + *outlistp = new_clause; + outlistp = &new_clause->next; + } + + for (; mapper_clause; mapper_clause = mapper_clause->next) + { + gfc_omp_namelist *new_clause = gfc_get_omp_namelist (); + + gfc_subst_mapper_var (&new_clause->sym, &new_clause->expr, + sym, expr, udm->var_sym, mapper_clause->sym, + mapper_clause->expr); + + enum gfc_omp_map_op map_clause_op = mapper_clause->u.map_op; + enum gfc_omp_map_op new_kind + = omp_map_decayed_kind (map_clause_op, outer_map_op, + (cd == TOC_OPENMP_EXIT_DATA + || list == OMP_LIST_FROM)); + if (list == OMP_LIST_FROM || list == OMP_LIST_TO) + { + switch (new_kind) + { + case OMP_MAP_PRESENT_FROM: + case OMP_MAP_PRESENT_TO: + new_clause->u.present_modifier = true; + /* Fallthrough. */ + case OMP_MAP_FROM: + case OMP_MAP_TO: + break; + default: + { + bool present_p, force_p, always_p; + gfc_omp_map_op basic_kind + = omp_split_map_op (map_clause_op, &force_p, &always_p, + &present_p); + free (new_clause); + gfc_warning (0, "Dropping incompatible %qs mapper clause at %C", + omp_basic_map_kind_name (basic_kind)); + inform (gfc_get_location (&mapper_clause->where), + "Defined here"); + continue; + } + } + } + else + new_clause->u.map_op = new_kind; + + new_clause->where = clause->where; + + if (mapper_clause->u2.udm + && mapper_clause->u2.udm->udm != udm) + { + gfc_omp_udm *inner_udm = mapper_clause->u2.udm->udm; + outlistp = gfc_omp_instantiate_mapper (outlistp, new_clause, + outer_map_op, inner_udm, cd, + list); + } + else + { + *outlistp = new_clause; + outlistp = &new_clause->next; + } + } + + return outlistp; +} + +void +gfc_omp_instantiate_mappers (gfc_code *code, gfc_omp_clauses *clauses, + toc_directive cd, int list) +{ + gfc_omp_namelist *clause = clauses->lists[list]; + gfc_omp_namelist **clausep = &clauses->lists[list]; + + for (; clause; clause = *clausep) + { + if (clause->u2.udm) + { + gfc_omp_map_op outer_map_op; + + switch (list) + { + case OMP_LIST_TO: + outer_map_op = clause->u.present_modifier ? OMP_MAP_PRESENT_TO + : OMP_MAP_TO; + break; + case OMP_LIST_FROM: + outer_map_op = clause->u.present_modifier ? OMP_MAP_PRESENT_FROM + : OMP_MAP_FROM; + break; + case OMP_LIST_MAP: + outer_map_op = clause->u.map_op; + break; + default: + gcc_unreachable (); + } + clausep = gfc_omp_instantiate_mapper (clausep, clause, outer_map_op, + clause->u2.udm->udm, cd, list); + *clausep = clause->next; + } + else + clausep = &clause->next; + } +} /* The following functions implement automatic recognition and annotation of DO loops in OpenACC kernels regions. Inside a kernels region, a nest of diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 0ef984720d0..170615974b3 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -3859,18 +3859,6 @@ gfc_convert_expr_to_tree (stmtblock_t *block, gfc_expr *expr) static vec *doacross_steps; -/* Control clause translation per-directive for gfc_trans_omp_clauses. */ - -enum toc_directive -{ - TOC_OPENMP, - TOC_OPENMP_DECLARE_SIMD, - TOC_OPENMP_DECLARE_MAPPER, - TOC_OPENMP_EXIT_DATA, - TOC_OPENACC, - TOC_OPENACC_DECLARE -}; - /* Translate an array section or array element. */ static void @@ -10082,372 +10070,6 @@ gfc_trans_omp_teams (gfc_code *code, gfc_omp_clauses *clausesa, return gfc_finish_block (&block); } -static enum gfc_omp_map_op -omp_split_map_op (enum gfc_omp_map_op op, bool *force_p, bool *always_p, - bool *present_p) -{ - *force_p = *always_p = *present_p = false; - - switch (op) - { - case OMP_MAP_FORCE_ALLOC: - case OMP_MAP_FORCE_TO: - case OMP_MAP_FORCE_FROM: - case OMP_MAP_FORCE_TOFROM: - case OMP_MAP_FORCE_PRESENT: - *force_p = true; - break; - case OMP_MAP_ALWAYS_TO: - case OMP_MAP_ALWAYS_FROM: - case OMP_MAP_ALWAYS_TOFROM: - *always_p = true; - break; - case OMP_MAP_ALWAYS_PRESENT_TO: - case OMP_MAP_ALWAYS_PRESENT_FROM: - case OMP_MAP_ALWAYS_PRESENT_TOFROM: - *always_p = true; - /* Fallthrough. */ - case OMP_MAP_PRESENT_ALLOC: - case OMP_MAP_PRESENT_TO: - case OMP_MAP_PRESENT_FROM: - case OMP_MAP_PRESENT_TOFROM: - *present_p = true; - break; - default: - ; - } - - switch (op) - { - case OMP_MAP_ALLOC: - case OMP_MAP_FORCE_ALLOC: - case OMP_MAP_PRESENT_ALLOC: - return OMP_MAP_ALLOC; - case OMP_MAP_TO: - case OMP_MAP_FORCE_TO: - case OMP_MAP_ALWAYS_TO: - case OMP_MAP_PRESENT_TO: - case OMP_MAP_ALWAYS_PRESENT_TO: - return OMP_MAP_TO; - case OMP_MAP_FROM: - case OMP_MAP_FORCE_FROM: - case OMP_MAP_ALWAYS_FROM: - case OMP_MAP_PRESENT_FROM: - case OMP_MAP_ALWAYS_PRESENT_FROM: - return OMP_MAP_FROM; - case OMP_MAP_TOFROM: - case OMP_MAP_FORCE_TOFROM: - case OMP_MAP_ALWAYS_TOFROM: - case OMP_MAP_PRESENT_TOFROM: - case OMP_MAP_ALWAYS_PRESENT_TOFROM: - return OMP_MAP_TOFROM; - default: - ; - } - return op; -} - -static enum gfc_omp_map_op -omp_join_map_op (enum gfc_omp_map_op op, bool force_p, bool always_p, - bool present_p) -{ - gcc_assert (!force_p || !(always_p || present_p)); - - switch (op) - { - case OMP_MAP_ALLOC: - if (force_p) - return OMP_MAP_FORCE_ALLOC; - else if (present_p) - return OMP_MAP_PRESENT_ALLOC; - break; - - case OMP_MAP_TO: - if (force_p) - return OMP_MAP_FORCE_TO; - else if (always_p && present_p) - return OMP_MAP_ALWAYS_PRESENT_TO; - else if (always_p) - return OMP_MAP_ALWAYS_TO; - else if (present_p) - return OMP_MAP_PRESENT_TO; - break; - - case OMP_MAP_FROM: - if (force_p) - return OMP_MAP_FORCE_FROM; - else if (always_p && present_p) - return OMP_MAP_ALWAYS_PRESENT_FROM; - else if (always_p) - return OMP_MAP_ALWAYS_FROM; - else if (present_p) - return OMP_MAP_PRESENT_FROM; - break; - - case OMP_MAP_TOFROM: - if (force_p) - return OMP_MAP_FORCE_TOFROM; - else if (always_p && present_p) - return OMP_MAP_ALWAYS_PRESENT_TOFROM; - else if (always_p) - return OMP_MAP_ALWAYS_TOFROM; - else if (present_p) - return OMP_MAP_PRESENT_TOFROM; - break; - - default: - ; - } - - return op; -} - -/* Map kind decay (OpenMP 5.2, 5.8.8 "declare mapper Directive"). Return the - map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS - specified on the clause that invokes the mapper. See also - c-family/c-omp.cc:omp_map_decayed_kind. */ - -static enum gfc_omp_map_op -omp_map_decayed_kind (enum gfc_omp_map_op mapper_kind, - enum gfc_omp_map_op invoked_as, bool exit_p) -{ - if (invoked_as == OMP_MAP_RELEASE || invoked_as == OMP_MAP_DELETE) - return invoked_as; - - bool force_p, always_p, present_p; - - invoked_as = omp_split_map_op (invoked_as, &force_p, &always_p, &present_p); - gfc_omp_map_op decay_to; - - switch (mapper_kind) - { - case OMP_MAP_ALLOC: - if (exit_p && invoked_as == OMP_MAP_FROM) - decay_to = OMP_MAP_RELEASE; - else - decay_to = OMP_MAP_ALLOC; - break; - - case OMP_MAP_TO: - if (invoked_as == OMP_MAP_FROM) - decay_to = exit_p ? OMP_MAP_RELEASE : OMP_MAP_ALLOC; - else if (invoked_as == OMP_MAP_ALLOC) - decay_to = OMP_MAP_ALLOC; - else - decay_to = OMP_MAP_TO; - break; - - case OMP_MAP_FROM: - if (invoked_as == OMP_MAP_ALLOC || invoked_as == OMP_MAP_TO) - decay_to = OMP_MAP_ALLOC; - else - decay_to = OMP_MAP_FROM; - break; - - case OMP_MAP_TOFROM: - case OMP_MAP_UNSET: - decay_to = invoked_as; - break; - - default: - gcc_unreachable (); - } - - return omp_join_map_op (decay_to, force_p, always_p, present_p); -} - -static gfc_symtree *gfc_subst_replace; -static gfc_ref *gfc_subst_prepend_ref; - -static bool -gfc_subst_in_expr_1 (gfc_expr *expr, gfc_symbol *search, int *) -{ - /* The base-object for component accesses may be stored in expr->symtree. - If it's the symbol for our "declare mapper" placeholder variable, - substitute it. */ - if (expr->symtree && expr->symtree->n.sym == search) - { - gfc_ref **lastptr = NULL; - expr->symtree = gfc_subst_replace; - - if (!gfc_subst_prepend_ref) - return false; - - gfc_ref *prepend_ref = gfc_copy_ref (gfc_subst_prepend_ref); - - for (gfc_ref *walk = prepend_ref; walk; walk = walk->next) - lastptr = &walk->next; - - *lastptr = expr->ref; - expr->ref = prepend_ref; - } - - return false; -} - -static void -gfc_subst_in_expr (gfc_expr *expr, gfc_symbol *search, gfc_symtree *replace, - gfc_ref *prepend_ref) -{ - gfc_subst_replace = replace; - gfc_subst_prepend_ref = prepend_ref; - gfc_traverse_expr (expr, search, gfc_subst_in_expr_1, 0); -} - -static void -gfc_subst_mapper_var (gfc_symbol **out_sym, gfc_expr **out_expr, - gfc_symbol *orig_sym, gfc_expr *orig_expr, - gfc_symbol *dummy_var, - gfc_symbol *templ_sym, gfc_expr *templ_expr) -{ - gfc_ref *orig_ref = orig_expr ? orig_expr->ref : NULL; - gfc_symtree *orig_st = gfc_find_symtree (orig_sym->ns->sym_root, - orig_sym->name); - - if (dummy_var == templ_sym) - *out_sym = orig_sym; - else - *out_sym = templ_sym; - - if (templ_expr) - { - *out_expr = gfc_copy_expr (templ_expr); - gfc_subst_in_expr (*out_expr, dummy_var, orig_st, orig_ref); - } - else if (orig_expr) - *out_expr = gfc_copy_expr (orig_expr); - else - *out_expr = NULL; -} - -static gfc_omp_namelist ** -gfc_trans_omp_instantiate_mapper (gfc_omp_namelist **outlistp, - gfc_omp_namelist *clause, gfc_omp_udm *udm, - toc_directive cd) -{ - /* Here "sym" and "expr" describe the clause as written, to be substituted - for the dummy variable in the mapper definition. */ - struct gfc_symbol *sym = clause->sym; - struct gfc_expr *expr = clause->expr; - gfc_omp_namelist *mapper_clause = udm->clauses->lists[OMP_LIST_MAP]; - gfc_omp_map_op outer_map_op = clause->u.map_op; - bool pointer_needed_p = false; - - if (expr) - { - gfc_ref *lastref = expr->ref, *lastcomp = NULL; - - for (; lastref->next; lastref = lastref->next) - if (lastref->type == REF_COMPONENT) - lastcomp = lastref; - - if (lastref - && lastref->type == REF_ARRAY - && (lastref->u.ar.type == AR_SECTION - || lastref->u.ar.type == AR_FULL)) - { - mpz_t elems; - bool multiple_elems_p = false; - - if (gfc_array_size (expr, &elems)) - { - HOST_WIDE_INT nelems = gfc_mpz_get_hwi (elems); - if (nelems > 1) - multiple_elems_p = true; - } - else - multiple_elems_p = true; - - if (multiple_elems_p && clause->u2.udm) - { - clause->u2.udm->multiple_elems_p = true; - *outlistp = clause; - return &(*outlistp)->next; - } - } - - if (lastcomp - && lastcomp->type == REF_COMPONENT - && (lastcomp->u.c.component->attr.pointer - || lastcomp->u.c.component->attr.allocatable)) - pointer_needed_p = true; - } - - if (pointer_needed_p) - { - /* If we're instantiating a mapper via a pointer, we need to map that - pointer as well as mapping the entities explicitly listed in the - mapper definition. Create a node for that. */ - gfc_omp_namelist *new_clause = gfc_get_omp_namelist (); - new_clause->sym = sym; - new_clause->expr = gfc_copy_expr (expr); - /* We want the pointer itself: cut off any further accessors after the - last component reference (e.g. array indices). */ - gfc_ref *lastcomp = NULL; - for (gfc_ref *lastref = new_clause->expr->ref; - lastref; - lastref = lastref->next) - if (lastref->type == REF_COMPONENT) - lastcomp = lastref; - gcc_assert (lastcomp != NULL); - lastcomp->next = NULL; - new_clause->u.map_op = OMP_MAP_POINTER_ONLY; - *outlistp = new_clause; - outlistp = &new_clause->next; - } - - for (; mapper_clause; mapper_clause = mapper_clause->next) - { - gfc_omp_namelist *new_clause = gfc_get_omp_namelist (); - - gfc_subst_mapper_var (&new_clause->sym, &new_clause->expr, - sym, expr, udm->var_sym, mapper_clause->sym, - mapper_clause->expr); - - enum gfc_omp_map_op map_clause_op = mapper_clause->u.map_op; - new_clause->u.map_op - = omp_map_decayed_kind (map_clause_op, outer_map_op, - (cd == TOC_OPENMP_EXIT_DATA)); - - new_clause->where = clause->where; - - if (mapper_clause->u2.udm - && mapper_clause->u2.udm->udm != udm) - { - gfc_omp_udm *inner_udm = mapper_clause->u2.udm->udm; - outlistp = gfc_trans_omp_instantiate_mapper (outlistp, new_clause, - inner_udm, cd); - } - else - { - *outlistp = new_clause; - outlistp = &new_clause->next; - } - } - - return outlistp; -} - -static void -gfc_trans_omp_instantiate_mappers (gfc_omp_clauses *clauses, - toc_directive cd = TOC_OPENMP) -{ - gfc_omp_namelist *clause = clauses->lists[OMP_LIST_MAP]; - gfc_omp_namelist **clausep = &clauses->lists[OMP_LIST_MAP]; - - for (; clause; clause = *clausep) - { - if (clause->u2.udm) - { - clausep = gfc_trans_omp_instantiate_mapper (clausep, clause, - clause->u2.udm->udm, cd); - *clausep = clause->next; - } - else - clausep = &clause->next; - } -} - /* Code callback for gfc_code_walker. */ static int @@ -10612,7 +10234,7 @@ gfc_trans_omp_target (gfc_code *code) if (flag_openmp) { gfc_omp_clauses *target_clauses = &clausesa[GFC_OMP_SPLIT_TARGET]; - gfc_trans_omp_instantiate_mappers (target_clauses); + gfc_omp_instantiate_mappers (code, target_clauses); omp_clauses = gfc_trans_omp_clauses (&block, target_clauses, code->loc); } @@ -10895,7 +10517,7 @@ gfc_trans_omp_target_data (gfc_code *code) gfc_start_block (&block); gfc_omp_clauses *target_data_clauses = code->ext.omp_clauses; - gfc_trans_omp_instantiate_mappers (target_data_clauses); + gfc_omp_instantiate_mappers (code, target_data_clauses); omp_clauses = gfc_trans_omp_clauses (&block, target_data_clauses, code->loc); stmt = gfc_trans_omp_code (code->block->next, true); stmt = build2_loc (gfc_get_location (&code->loc), OMP_TARGET_DATA, @@ -10912,7 +10534,7 @@ gfc_trans_omp_target_enter_data (gfc_code *code) gfc_start_block (&block); gfc_omp_clauses *target_enter_data_clauses = code->ext.omp_clauses; - gfc_trans_omp_instantiate_mappers (target_enter_data_clauses); + gfc_omp_instantiate_mappers (code, target_enter_data_clauses); omp_clauses = gfc_trans_omp_clauses (&block, target_enter_data_clauses, code->loc); stmt = build1_loc (input_location, OMP_TARGET_ENTER_DATA, void_type_node, @@ -10929,8 +10551,8 @@ gfc_trans_omp_target_exit_data (gfc_code *code) gfc_start_block (&block); gfc_omp_clauses *target_exit_data_clauses = code->ext.omp_clauses; - gfc_trans_omp_instantiate_mappers (target_exit_data_clauses, - TOC_OPENMP_EXIT_DATA); + gfc_omp_instantiate_mappers (code, target_exit_data_clauses, + TOC_OPENMP_EXIT_DATA); omp_clauses = gfc_trans_omp_clauses (&block, target_exit_data_clauses, code->loc, TOC_OPENMP_EXIT_DATA); stmt = build1_loc (input_location, OMP_TARGET_EXIT_DATA, void_type_node, From patchwork Thu Aug 10 13:33:03 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 134002 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:b824:0:b0:3f2:4152:657d with SMTP id z4csp426088vqi; Thu, 10 Aug 2023 06:37:47 -0700 (PDT) X-Google-Smtp-Source: AGHT+IHvk3uRcQmPv/0ql5H7ibrEUoiugQTE14+eHOVX3ZaemtV6Xf+HQljTUFinEeyZrMW/HoSq X-Received: by 2002:aa7:d812:0:b0:521:728f:d84e with SMTP id v18-20020aa7d812000000b00521728fd84emr2760960edq.0.1691674666869; Thu, 10 Aug 2023 06:37:46 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1691674666; cv=none; d=google.com; s=arc-20160816; b=XgKSZ2jsfvYzW7NBHttCwDomZ0Y/6qt4kB7gAb2IZv3uiMZ6IL86ADtkRJCFxUp5M7 nagq2zpBGA7s1nIXDbTcHL71Di4YS1GTGC76vQmftTQ4WMwd/Tu81CSvoOs5w4RxVi00 OB/dttB773RoU18zMGw1S8TL3ILBBmAeqFhUEHsW7CX/wyrOIyyhsfbK1EzqXRRMwoGx A1LxYawj61rieQbM1kb3yS6QqsSK+mIsO0xFl/53vbSPjivOdZGS3p4vK9RyA/LyOVrM uuFEqH55+hLdlE23CzQGrBB8vp9CuQcqVHFhxtI9WB+YK9gP9G2xM+vxFA/kfdRAVc6N Cnsg== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:references:in-reply-to:message-id:date:subject:cc:to :from:ironport-sdr:dmarc-filter:delivered-to; bh=mmOqieH491WsO/DH+HD33ykh6F8yxzGZyrHfmlW/9zI=; fh=6entJIlBX3Lqynes5VYRRY4zgfJfSDVLFnUjulEf2Vo=; b=0TUSo9kahFs2HOYT/k2F5huxknATGiXSCqWrWPMS69sc1rGViRPbA7gvQW6Gcs3zKY hw5SAqwW7XZt2nBGq5Efv8DUuFvvK0sZyNCz6YthwQWiolmmlsGyMR4jFMsDa9XCq9pc Enf0lzIYALMbDZMkMYcpgg1u//WZypd9dvLj4O16L8XQRIOTOruisOwLXKaBhveh6p4U NRrcjI3TbELocX4JUdXgxWQTal2TZ7tPPEqCVrZrSN+h7YLVupBGT//VqeF7LmV9jn+U pkZR0yHL1y5fcYW6zbqB3TaYEsuFSrO4Xv1CeDLwKhnMH+cDQwjiH6FUbWvTUDgAWOgT w3qA== ARC-Authentication-Results: i=1; mx.google.com; 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 (ip-8-43-85-97.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id q25-20020a056402041900b0052228b42848si1340020edv.60.2023.08.10.06.37.46 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 10 Aug 2023 06:37:46 -0700 (PDT) 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; 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 62A9F389364F for ; Thu, 10 Aug 2023 13:34:53 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 23D99385701D; Thu, 10 Aug 2023 13:33:29 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 23D99385701D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,162,1684828800"; d="scan'208";a="14111552" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 10 Aug 2023 05:33:28 -0800 IronPort-SDR: Epz5fqRkr4HvcJKmk3JKuUDeU6tmyV7pmPBEFESxVWa4cpXF8etLlSTm2vmpAjclKW+X3htqHH 3FvmwPbxTIxsovAYGg/xOgaMgqxodsaLKV9pb8h/rN2si5e+3g8dTKZYKxp/9qOg84q2QylDTr 2nQbbaMfFmyE83TYIaffDQrasm1xP+syK/rcUWrvHUulJxyrNuZtu2PVlXOELgRpKbrxWbazf4 h8jAQFidhUDR2oA7JIgT5o1KC351G0URcHxRnbwm7ciJ4EUNVNbSemd/gOE2qse1LKUiTJP0pk H38= From: Julian Brown To: CC: , , Subject: [PATCH 2/5] OpenMP: Reprocess expanded clauses after 'declare mapper' instantiation Date: Thu, 10 Aug 2023 13:33:03 +0000 Message-ID: X-Mailer: git-send-email 2.25.1 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham 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.29 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 Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1773849455386140179 X-GMAIL-MSGID: 1773849455386140179 This patch reprocesses expanded clauses after 'declare mapper' instantiation -- checking things such as duplicated clauses, illegal use of strided accesses, and so forth. Two functions are broken out of the 'resolve_omp_clauses' function and reused in a new function 'resolve_omp_mapper_clauses', called after mapper instantiation. This improves diagnostic output. 2023-08-10 Julian Brown gcc/fortran/ * gfortran.h (gfc_omp_clauses): Add NS field. * openmp.cc (verify_omp_clauses_symbol_dups, omp_verify_map_motion_clauses): New functions, broken out of... (resolve_omp_clauses): Here. Record namespace containing clauses. Call above functions. (resolve_omp_mapper_clauses): New function, using helper functions broken out above. (gfc_resolve_omp_directive): Add NS parameter to resolve_omp_clauses calls. (gfc_omp_instantiate_mappers): Call resolve_omp_mapper_clauses if we instantiate any mappers. gcc/testsuite/ * gfortran.dg/gomp/declare-mapper-26.f90: New test. * gfortran.dg/gomp/declare-mapper-29.f90: New test. --- gcc/fortran/gfortran.h | 1 + gcc/fortran/openmp.cc | 1250 +++++++++-------- .../gfortran.dg/gomp/declare-mapper-26.f90 | 28 + .../gfortran.dg/gomp/declare-mapper-29.f90 | 22 + 4 files changed, 718 insertions(+), 583 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/declare-mapper-29.f90 diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 788b3797893..a98424b3263 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1577,6 +1577,7 @@ typedef struct gfc_omp_clauses struct gfc_omp_assumptions *assume; struct gfc_expr_list *tile_sizes; const char *critical_name; + gfc_namespace *ns; enum gfc_omp_default_sharing default_sharing; enum gfc_omp_atomic_op atomic_op; enum gfc_omp_defaultmap defaultmap[OMP_DEFAULTMAP_CAT_NUM]; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 0f715a6f997..0109df4dfce 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -8123,6 +8123,611 @@ gfc_resolve_omp_assumptions (gfc_omp_assumptions *assume) &el->expr->where); } +/* Check OMP_CLAUSES for duplicate symbols and various other constraints. + Helper function for resolve_omp_clauses and resolve_omp_mapper_clauses. */ + +static void +verify_omp_clauses_symbol_dups (gfc_code *code, gfc_omp_clauses *omp_clauses, + gfc_namespace *ns, bool openacc) +{ + gfc_omp_namelist *n; + int list; + + /* Check that no symbol appears on multiple clauses, except that a symbol + can appear on both firstprivate and lastprivate. */ + for (list = 0; list < OMP_LIST_NUM; list++) + for (n = omp_clauses->lists[list]; n; n = n->next) + { + if (!n->sym) /* omp_all_memory. */ + continue; + n->sym->mark = 0; + n->sym->comp_mark = 0; + n->sym->data_mark = 0; + n->sym->dev_mark = 0; + n->sym->gen_mark = 0; + n->sym->reduc_mark = 0; + if (n->sym->attr.flavor == FL_VARIABLE + || n->sym->attr.proc_pointer + || (!code + && !ns->omp_udm_ns + && (!n->sym->attr.dummy || n->sym->ns != ns))) + { + if (!code + && !ns->omp_udm_ns + && (!n->sym->attr.dummy || n->sym->ns != ns)) + gfc_error ("Variable %qs is not a dummy argument at %L", + n->sym->name, &n->where); + continue; + } + if (n->sym->attr.flavor == FL_PROCEDURE + && n->sym->result == n->sym + && n->sym->attr.function) + { + if (gfc_current_ns->proc_name == n->sym + || (gfc_current_ns->parent + && gfc_current_ns->parent->proc_name == n->sym)) + continue; + if (gfc_current_ns->proc_name->attr.entry_master) + { + gfc_entry_list *el = gfc_current_ns->entries; + for (; el; el = el->next) + if (el->sym == n->sym) + break; + if (el) + continue; + } + if (gfc_current_ns->parent + && gfc_current_ns->parent->proc_name->attr.entry_master) + { + gfc_entry_list *el = gfc_current_ns->parent->entries; + for (; el; el = el->next) + if (el->sym == n->sym) + break; + if (el) + continue; + } + } + if (list == OMP_LIST_MAP + && n->sym->attr.flavor == FL_PARAMETER) + { + if (openacc) + gfc_error ("Object %qs is not a variable at %L; parameters" + " cannot be and need not be copied", n->sym->name, + &n->where); + else + gfc_error ("Object %qs is not a variable at %L; parameters" + " cannot be and need not be mapped", n->sym->name, + &n->where); + } + else + gfc_error ("Object %qs is not a variable at %L", n->sym->name, + &n->where); + } + if (omp_clauses->lists[OMP_LIST_REDUCTION_INSCAN]) + { + locus *loc = &omp_clauses->lists[OMP_LIST_REDUCTION_INSCAN]->where; + if (code->op != EXEC_OMP_DO + && code->op != EXEC_OMP_SIMD + && code->op != EXEC_OMP_DO_SIMD + && code->op != EXEC_OMP_PARALLEL_DO + && code->op != EXEC_OMP_PARALLEL_DO_SIMD) + gfc_error ("% REDUCTION clause on construct other than DO, " + "SIMD, DO SIMD, PARALLEL DO, PARALLEL DO SIMD at %L", loc); + if (omp_clauses->ordered) + gfc_error ("ORDERED clause specified together with % " + "REDUCTION clause at %L", loc); + if (omp_clauses->sched_kind != OMP_SCHED_NONE) + gfc_error ("SCHEDULE clause specified together with % " + "REDUCTION clause at %L", loc); + } + + for (list = 0; list < OMP_LIST_NUM; list++) + if (list != OMP_LIST_FIRSTPRIVATE + && list != OMP_LIST_LASTPRIVATE + && list != OMP_LIST_ALIGNED + && list != OMP_LIST_DEPEND + && list != OMP_LIST_FROM + && list != OMP_LIST_TO + && (list != OMP_LIST_REDUCTION || !openacc) + && list != OMP_LIST_ALLOCATE) + for (n = omp_clauses->lists[list]; n; n = n->next) + { + bool component_ref_p = false; + + /* Allow multiple components of the same (e.g. derived-type) + variable here. Duplicate components are detected elsewhere. */ + if (n->expr && n->expr->expr_type == EXPR_VARIABLE) + for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) + if (ref->type == REF_COMPONENT) + component_ref_p = true; + if ((list == OMP_LIST_IS_DEVICE_PTR + || list == OMP_LIST_HAS_DEVICE_ADDR) + && !component_ref_p) + { + if (n->sym->gen_mark + || n->sym->dev_mark + || n->sym->reduc_mark + || n->sym->mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->dev_mark = 1; + } + else if ((list == OMP_LIST_USE_DEVICE_PTR + || list == OMP_LIST_USE_DEVICE_ADDR + || list == OMP_LIST_PRIVATE + || list == OMP_LIST_SHARED) + && !component_ref_p) + { + if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + { + n->sym->gen_mark = 1; + /* Set both generic and device bits if we have + use_device_*(x) or shared(x). This allows us to diagnose + "map(x) private(x)" below. */ + if (list != OMP_LIST_PRIVATE) + n->sym->dev_mark = 1; + } + } + else if ((list == OMP_LIST_REDUCTION + || list == OMP_LIST_REDUCTION_TASK + || list == OMP_LIST_REDUCTION_INSCAN + || list == OMP_LIST_IN_REDUCTION + || list == OMP_LIST_TASK_REDUCTION) + && !component_ref_p) + { + /* Attempts to mix reduction types are diagnosed below. */ + if (n->sym->gen_mark || n->sym->dev_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + n->sym->reduc_mark = 1; + } + else if ((!component_ref_p && n->sym->comp_mark) + || (component_ref_p && n->sym->mark)) + { + if (openacc) + gfc_error ("Symbol %qs has mixed component and non-component " + "accesses at %L", n->sym->name, &n->where); + } + else if (n->sym->mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + { + if (component_ref_p) + n->sym->comp_mark = 1; + else + n->sym->mark = 1; + } + } + + /* Detect specifically the case where we have "map(x) private(x)" and raise + an error. If we have "...simd" combined directives though, the "private" + applies to the simd part, so this is permitted. */ + for (n = omp_clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next) + if (n->sym->mark + && n->sym->gen_mark + && !n->sym->dev_mark + && !n->sym->reduc_mark + && code->op != EXEC_OMP_TARGET_SIMD + && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD) + gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, + &n->where); + + gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1); + for (list = OMP_LIST_FIRSTPRIVATE; list <= OMP_LIST_LASTPRIVATE; list++) + { + gfc_omp_namelist **pn = &omp_clauses->lists[list]; + while ((n = *pn) != NULL) + { + bool remove = false; + + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) + { + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + n->sym->data_mark = n->sym->gen_mark = n->sym->dev_mark = 0; + } + else if (n->sym->mark + && code->op != EXEC_OMP_TARGET_TEAMS + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE + && code->op != EXEC_OMP_TARGET_TEAMS_LOOP + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO + && code->op != EXEC_OMP_TARGET_PARALLEL + && code->op != EXEC_OMP_TARGET_PARALLEL_DO + && code->op != EXEC_OMP_TARGET_PARALLEL_LOOP + && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD + && (code->op + != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD)) + { + gfc_error ("Symbol %qs present on both data and map clauses " + "at %L", n->sym->name, &n->where); + /* We've already shown an error. Avoid confusing gimplify. */ + remove = true; + } + + if (remove) + *pn = n->next; + else + pn = &n->next; + } + } + + for (n = omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; n; n = n->next) + { + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->data_mark = 1; + } + for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next) + n->sym->data_mark = 0; + + for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next) + { + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->data_mark = 1; + } + + for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next) + n->sym->mark = 0; + + for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next) + { + if (n->sym->mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->mark = 1; + } + + if (omp_clauses->lists[OMP_LIST_ALLOCATE]) + { + for (n = omp_clauses->lists[OMP_LIST_ALLOCATE]; n; n = n->next) + { + if (n->expr && (!gfc_resolve_expr (n->expr) + || n->expr->ts.type != BT_INTEGER + || n->expr->ts.kind != gfc_c_intptr_kind)) + { + gfc_error ("Expected integer expression of the " + "% kind at %L", + &n->expr->where); + break; + } + if (!n->u.align) + continue; + int alignment = 0; + if (!gfc_resolve_expr (n->u.align) + || n->u.align->ts.type != BT_INTEGER + || n->u.align->rank != 0 + || gfc_extract_int (n->u.align, &alignment) + || alignment <= 0 + || !pow2p_hwi (alignment)) + { + gfc_error ("ALIGN modifier requires at %L a scalar positive " + "constant integer alignment expression that is a " + "power of two", &n->u.align->where); + break; + } + } + + /* Check for 2 things here. + 1. There is no duplication of variable in allocate clause. + 2. Variable in allocate clause are also present in some + privatization clase (non-composite case). */ + for (n = omp_clauses->lists[OMP_LIST_ALLOCATE]; n; n = n->next) + n->sym->mark = 0; + + gfc_omp_namelist *prev = NULL; + for (n = omp_clauses->lists[OMP_LIST_ALLOCATE]; n;) + { + if (n->sym->mark == 1) + { + gfc_warning (0, "%qs appears more than once in % " + "clauses at %L" , n->sym->name, &n->where); + /* We have already seen this variable so it is a duplicate. + Remove it. */ + if (prev != NULL && prev->next == n) + { + prev->next = n->next; + n->next = NULL; + gfc_free_omp_namelist (n, OMP_LIST_ALLOCATE); + n = prev->next; + } + continue; + } + n->sym->mark = 1; + prev = n; + n = n->next; + } + + /* Non-composite constructs. */ + if (code && code->op < EXEC_OMP_DO_SIMD) + { + for (list = 0; list < OMP_LIST_NUM; list++) + switch (list) + { + case OMP_LIST_PRIVATE: + case OMP_LIST_FIRSTPRIVATE: + case OMP_LIST_LASTPRIVATE: + case OMP_LIST_REDUCTION: + case OMP_LIST_REDUCTION_INSCAN: + case OMP_LIST_REDUCTION_TASK: + case OMP_LIST_IN_REDUCTION: + case OMP_LIST_TASK_REDUCTION: + case OMP_LIST_LINEAR: + for (n = omp_clauses->lists[list]; n; n = n->next) + n->sym->mark = 0; + break; + default: + break; + } + + for (n = omp_clauses->lists[OMP_LIST_ALLOCATE]; n; n = n->next) + if (n->sym->mark == 1) + gfc_error ("%qs specified in % clause at %L but not " + "in an explicit privatization clause", 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) + if (n->expr == NULL) + n->sym->mark = 1; + for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next) + { + if (n->expr == NULL && n->sym->mark) + gfc_error ("Symbol %qs present on both FROM and TO clauses at %L", + n->sym->name, &n->where); + else + n->sym->mark = 1; + } +} + +/* Check that the parameter of a MAP, TO and FROM clause N meets certain + constraints. Helper function for resolve_omp_clauses and + resolve_omp_mapper_clauses. */ + +static bool +omp_verify_map_motion_clauses (gfc_code *code, int list, const char *name, + gfc_omp_namelist *n, bool openacc) +{ + gfc_ref *lastref = NULL, *lastslice = NULL; + bool resolved = false; + if (n->expr) + { + lastref = n->expr->ref; + resolved = gfc_resolve_expr (n->expr); + + /* Look through component refs to find last array + reference. */ + if (resolved) + { + for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) + if (ref->type == REF_COMPONENT + || ref->type == REF_SUBSTRING + || ref->type == REF_INQUIRY) + lastref = ref; + else if (ref->type == REF_ARRAY) + { + for (int i = 0; i < ref->u.ar.dimen; i++) + if (ref->u.ar.dimen_type[i] == DIMEN_RANGE) + lastslice = ref; + + lastref = ref; + } + + /* The "!$acc cache" directive allows rectangular subarrays to be + specified, with some restrictions on the form of bounds (not + implemented). + Only raise an error here if we're really sure the array isn't + contiguous. An expression such as arr(-n:n,-n:n) could be + contiguous even if it looks like it may not be. Also OpenMP's + 'target update' permits strides for the to/from clause. */ + if (code + && code->op != EXEC_OACC_UPDATE + && code->op != EXEC_OMP_TARGET_UPDATE + && list != OMP_LIST_CACHE + && list != OMP_LIST_DEPEND + && !gfc_is_simply_contiguous (n->expr, false, true) + && gfc_is_not_contiguous (n->expr) + && !(lastslice && (lastslice->next + || lastslice->type != REF_ARRAY))) + gfc_error ("Array is not contiguous at %L", + &n->where); + } + } + if (openacc + && list == OMP_LIST_MAP + && (n->u.map_op == OMP_MAP_ATTACH || n->u.map_op == OMP_MAP_DETACH)) + { + symbol_attribute attr; + if (n->expr) + attr = gfc_expr_attr (n->expr); + else + attr = n->sym->attr; + if (!attr.pointer && !attr.allocatable) + gfc_error ("%qs clause argument must be ALLOCATABLE or a POINTER " + "at %L", + (n->u.map_op == OMP_MAP_ATTACH) ? "attach" : "detach", + &n->where); + } + if (lastref + || (n->expr && (!resolved || n->expr->expr_type != EXPR_VARIABLE))) + { + if (!lastslice && lastref && lastref->type == REF_SUBSTRING) + gfc_error ("Unexpected substring reference in %s clause " + "at %L", name, &n->where); + else if (!lastslice && lastref && lastref->type == REF_INQUIRY) + { + gcc_assert (lastref->u.i == INQUIRY_RE + || lastref->u.i == INQUIRY_IM); + gfc_error ("Unexpected complex-parts designator " + "reference in %s clause at %L", + name, &n->where); + } + else if (!resolved + || n->expr->expr_type != EXPR_VARIABLE + || (lastslice + && (lastslice->next || lastslice->type != REF_ARRAY))) + gfc_error ("%qs in %s clause at %L is not a proper " + "array section", n->sym->name, name, + &n->where); + else if (lastslice) + { + int i; + gfc_array_ref *ar = &lastslice->u.ar; + for (i = 0; i < ar->dimen; i++) + if (ar->stride[i] + && code + && code->op != EXEC_OACC_UPDATE + && code->op != EXEC_OMP_TARGET_UPDATE) + { + gfc_error ("Stride should not be specified for " + "array section in %s clause at %L", + name, &n->where); + return false; + } + else if (ar->dimen_type[i] != DIMEN_ELEMENT + && ar->dimen_type[i] != DIMEN_RANGE) + { + gfc_error ("%qs in %s clause at %L is not a " + "proper array section", + n->sym->name, name, &n->where); + return false; + } + else if ((list == OMP_LIST_DEPEND || list == OMP_LIST_AFFINITY) + && ar->start[i] + && ar->start[i]->expr_type == EXPR_CONSTANT + && ar->end[i] + && ar->end[i]->expr_type == EXPR_CONSTANT + && mpz_cmp (ar->start[i]->value.integer, + ar->end[i]->value.integer) > 0) + { + gfc_error ("%qs in %s clause at %L is a zero size array " + "section", n->sym->name, list == OMP_LIST_DEPEND + ? "DEPEND" : "AFFINITY", &n->where); + return false; + } + } + } + else if (openacc) + { + if (list == OMP_LIST_MAP && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR) + resolve_oacc_deviceptr_clause (n->sym, n->where, name); + else + resolve_oacc_data_clauses (n->sym, n->where, name); + } + else if (list != OMP_LIST_DEPEND + && n->sym->as + && n->sym->as->type == AS_ASSUMED_SIZE) + gfc_error ("Assumed size array %qs in %s clause at %L", + n->sym->name, name, &n->where); + + if (!code || list != OMP_LIST_MAP || openacc) + return true; + + switch (code->op) + { + case EXEC_OMP_TARGET: + case EXEC_OMP_TARGET_DATA: + switch (n->u.map_op) + { + case OMP_MAP_TO: + case OMP_MAP_ALWAYS_TO: + case OMP_MAP_PRESENT_TO: + case OMP_MAP_ALWAYS_PRESENT_TO: + case OMP_MAP_FROM: + case OMP_MAP_ALWAYS_FROM: + case OMP_MAP_PRESENT_FROM: + case OMP_MAP_ALWAYS_PRESENT_FROM: + case OMP_MAP_TOFROM: + case OMP_MAP_ALWAYS_TOFROM: + case OMP_MAP_PRESENT_TOFROM: + case OMP_MAP_ALWAYS_PRESENT_TOFROM: + case OMP_MAP_ALLOC: + case OMP_MAP_PRESENT_ALLOC: + break; + default: + gfc_error ("TARGET%s with map-type other than TO, FROM, TOFROM, or " + "ALLOC on MAP clause at %L", + code->op == EXEC_OMP_TARGET ? "" : " DATA", &n->where); + break; + } + break; + case EXEC_OMP_TARGET_ENTER_DATA: + switch (n->u.map_op) + { + case OMP_MAP_TO: + case OMP_MAP_ALWAYS_TO: + case OMP_MAP_PRESENT_TO: + case OMP_MAP_ALWAYS_PRESENT_TO: + case OMP_MAP_ALLOC: + case OMP_MAP_PRESENT_ALLOC: + break; + case OMP_MAP_TOFROM: + n->u.map_op = OMP_MAP_TO; + break; + case OMP_MAP_ALWAYS_TOFROM: + n->u.map_op = OMP_MAP_ALWAYS_TO; + break; + case OMP_MAP_PRESENT_TOFROM: + n->u.map_op = OMP_MAP_PRESENT_TO; + break; + case OMP_MAP_ALWAYS_PRESENT_TOFROM: + n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO; + break; + default: + gfc_error ("TARGET ENTER DATA with map-type other than TO, TOFROM " + "or ALLOC on MAP clause at %L", &n->where); + break; + } + break; + case EXEC_OMP_TARGET_EXIT_DATA: + switch (n->u.map_op) + { + case OMP_MAP_FROM: + case OMP_MAP_ALWAYS_FROM: + case OMP_MAP_PRESENT_FROM: + case OMP_MAP_ALWAYS_PRESENT_FROM: + case OMP_MAP_RELEASE: + case OMP_MAP_DELETE: + break; + case OMP_MAP_TOFROM: + n->u.map_op = OMP_MAP_FROM; + break; + case OMP_MAP_ALWAYS_TOFROM: + n->u.map_op = OMP_MAP_ALWAYS_FROM; + break; + case OMP_MAP_PRESENT_TOFROM: + n->u.map_op = OMP_MAP_PRESENT_FROM; + break; + case OMP_MAP_ALWAYS_PRESENT_TOFROM: + n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM; + break; + default: + gfc_error ("TARGET EXIT DATA with map-type other than FROM, TOFROM, " + "RELEASE, or DELETE on MAP clause at %L", &n->where); + break; + } + break; + default: + ; + } + + return true; +} /* OpenMP directive resolving routines. */ @@ -8157,6 +8762,11 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, if (omp_clauses->order_concurrent && omp_clauses->ordered) gfc_error ("ORDER clause must not be used together ORDERED at %L", &code->loc); + /* If we're invoking any declared mappers as a result of these clauses, we may + need to know the namespace their directive was originally defined within in + order to resolve clauses again after substitution. Record it here. */ + if (ns) + omp_clauses->ns = ns; if (omp_clauses->if_expr) { gfc_expr *expr = omp_clauses->if_expr; @@ -8349,337 +8959,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, gfc_error ("DEPOBJ in DEPOBJ construct at %L shall be a scalar integer " "of OMP_DEPEND_KIND kind", &omp_clauses->depobj->where); - /* Check that no symbol appears on multiple clauses, except that - a symbol can appear on both firstprivate and lastprivate. */ - for (list = 0; list < OMP_LIST_NUM; list++) - for (n = omp_clauses->lists[list]; n; n = n->next) - { - if (!n->sym) /* omp_all_memory. */ - continue; - n->sym->mark = 0; - n->sym->comp_mark = 0; - n->sym->data_mark = 0; - n->sym->dev_mark = 0; - n->sym->gen_mark = 0; - n->sym->reduc_mark = 0; - if (n->sym->attr.flavor == FL_VARIABLE - || n->sym->attr.proc_pointer - || (!code - && !ns->omp_udm_ns - && (!n->sym->attr.dummy || n->sym->ns != ns))) - { - if (!code - && !ns->omp_udm_ns - && (!n->sym->attr.dummy || n->sym->ns != ns)) - gfc_error ("Variable %qs is not a dummy argument at %L", - n->sym->name, &n->where); - continue; - } - if (n->sym->attr.flavor == FL_PROCEDURE - && n->sym->result == n->sym - && n->sym->attr.function) - { - if (gfc_current_ns->proc_name == n->sym - || (gfc_current_ns->parent - && gfc_current_ns->parent->proc_name == n->sym)) - continue; - if (gfc_current_ns->proc_name->attr.entry_master) - { - gfc_entry_list *el = gfc_current_ns->entries; - for (; el; el = el->next) - if (el->sym == n->sym) - break; - if (el) - continue; - } - if (gfc_current_ns->parent - && gfc_current_ns->parent->proc_name->attr.entry_master) - { - gfc_entry_list *el = gfc_current_ns->parent->entries; - for (; el; el = el->next) - if (el->sym == n->sym) - break; - if (el) - continue; - } - } - if (list == OMP_LIST_MAP - && n->sym->attr.flavor == FL_PARAMETER) - { - if (openacc) - gfc_error ("Object %qs is not a variable at %L; parameters" - " cannot be and need not be copied", n->sym->name, - &n->where); - else - gfc_error ("Object %qs is not a variable at %L; parameters" - " cannot be and need not be mapped", n->sym->name, - &n->where); - } - else - gfc_error ("Object %qs is not a variable at %L", n->sym->name, - &n->where); - } - if (omp_clauses->lists[OMP_LIST_REDUCTION_INSCAN]) - { - locus *loc = &omp_clauses->lists[OMP_LIST_REDUCTION_INSCAN]->where; - if (code->op != EXEC_OMP_DO - && code->op != EXEC_OMP_SIMD - && code->op != EXEC_OMP_DO_SIMD - && code->op != EXEC_OMP_PARALLEL_DO - && code->op != EXEC_OMP_PARALLEL_DO_SIMD) - gfc_error ("% REDUCTION clause on construct other than DO, " - "SIMD, DO SIMD, PARALLEL DO, PARALLEL DO SIMD at %L", - loc); - if (omp_clauses->ordered) - gfc_error ("ORDERED clause specified together with % " - "REDUCTION clause at %L", loc); - if (omp_clauses->sched_kind != OMP_SCHED_NONE) - gfc_error ("SCHEDULE clause specified together with % " - "REDUCTION clause at %L", loc); - } - - for (list = 0; list < OMP_LIST_NUM; list++) - if (list != OMP_LIST_FIRSTPRIVATE - && list != OMP_LIST_LASTPRIVATE - && list != OMP_LIST_ALIGNED - && list != OMP_LIST_DEPEND - && list != OMP_LIST_FROM - && list != OMP_LIST_TO - && (list != OMP_LIST_REDUCTION || !openacc) - && list != OMP_LIST_ALLOCATE) - for (n = omp_clauses->lists[list]; n; n = n->next) - { - bool component_ref_p = false; - - /* Allow multiple components of the same (e.g. derived-type) - variable here. Duplicate components are detected elsewhere. */ - if (n->expr && n->expr->expr_type == EXPR_VARIABLE) - for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) - if (ref->type == REF_COMPONENT) - component_ref_p = true; - if ((list == OMP_LIST_IS_DEVICE_PTR - || list == OMP_LIST_HAS_DEVICE_ADDR) - && !component_ref_p) - { - if (n->sym->gen_mark - || n->sym->dev_mark - || n->sym->reduc_mark - || n->sym->mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - else - n->sym->dev_mark = 1; - } - else if ((list == OMP_LIST_USE_DEVICE_PTR - || list == OMP_LIST_USE_DEVICE_ADDR - || list == OMP_LIST_PRIVATE - || list == OMP_LIST_SHARED) - && !component_ref_p) - { - if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - else - { - n->sym->gen_mark = 1; - /* Set both generic and device bits if we have - use_device_*(x) or shared(x). This allows us to diagnose - "map(x) private(x)" below. */ - if (list != OMP_LIST_PRIVATE) - n->sym->dev_mark = 1; - } - } - else if ((list == OMP_LIST_REDUCTION - || list == OMP_LIST_REDUCTION_TASK - || list == OMP_LIST_REDUCTION_INSCAN - || list == OMP_LIST_IN_REDUCTION - || list == OMP_LIST_TASK_REDUCTION) - && !component_ref_p) - { - /* Attempts to mix reduction types are diagnosed below. */ - if (n->sym->gen_mark || n->sym->dev_mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - n->sym->reduc_mark = 1; - } - else if ((!component_ref_p && n->sym->comp_mark) - || (component_ref_p && n->sym->mark)) - { - if (openacc) - gfc_error ("Symbol %qs has mixed component and non-component " - "accesses at %L", n->sym->name, &n->where); - } - else if (n->sym->mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - else - { - if (component_ref_p) - n->sym->comp_mark = 1; - else - n->sym->mark = 1; - } - } - - /* Detect specifically the case where we have "map(x) private(x)" and raise - an error. If we have "...simd" combined directives though, the "private" - applies to the simd part, so this is permitted though. */ - for (n = omp_clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next) - if (n->sym->mark - && n->sym->gen_mark - && !n->sym->dev_mark - && !n->sym->reduc_mark - && code->op != EXEC_OMP_TARGET_SIMD - && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD - && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD - && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - - gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1); - for (list = OMP_LIST_FIRSTPRIVATE; list <= OMP_LIST_LASTPRIVATE; list++) - for (n = omp_clauses->lists[list]; n; n = n->next) - if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) - { - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - n->sym->data_mark = n->sym->gen_mark = n->sym->dev_mark = 0; - } - else if (n->sym->mark - && code->op != EXEC_OMP_TARGET_TEAMS - && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE - && code->op != EXEC_OMP_TARGET_TEAMS_LOOP - && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD - && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO - && code->op != EXEC_OMP_TARGET_PARALLEL - && code->op != EXEC_OMP_TARGET_PARALLEL_DO - && code->op != EXEC_OMP_TARGET_PARALLEL_LOOP - && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD - && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD) - gfc_error ("Symbol %qs present on both data and map clauses " - "at %L", n->sym->name, &n->where); - - for (n = omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; n; n = n->next) - { - if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - else - n->sym->data_mark = 1; - } - for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next) - n->sym->data_mark = 0; - - for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next) - { - if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - else - n->sym->data_mark = 1; - } - - for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next) - n->sym->mark = 0; - - for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next) - { - if (n->sym->mark) - gfc_error ("Symbol %qs present on multiple clauses at %L", - n->sym->name, &n->where); - else - n->sym->mark = 1; - } - - if (omp_clauses->lists[OMP_LIST_ALLOCATE]) - { - for (n = omp_clauses->lists[OMP_LIST_ALLOCATE]; n; n = n->next) - { - if (n->expr && (!gfc_resolve_expr (n->expr) - || n->expr->ts.type != BT_INTEGER - || n->expr->ts.kind != gfc_c_intptr_kind)) - { - gfc_error ("Expected integer expression of the " - "% kind at %L", - &n->expr->where); - break; - } - if (!n->u.align) - continue; - int alignment = 0; - if (!gfc_resolve_expr (n->u.align) - || n->u.align->ts.type != BT_INTEGER - || n->u.align->rank != 0 - || gfc_extract_int (n->u.align, &alignment) - || alignment <= 0 - || !pow2p_hwi (alignment)) - { - gfc_error ("ALIGN modifier requires at %L a scalar positive " - "constant integer alignment expression that is a " - "power of two", &n->u.align->where); - break; - } - } - - /* Check for 2 things here. - 1. There is no duplication of variable in allocate clause. - 2. Variable in allocate clause are also present in some - privatization clase (non-composite case). */ - for (n = omp_clauses->lists[OMP_LIST_ALLOCATE]; n; n = n->next) - n->sym->mark = 0; - - gfc_omp_namelist *prev = NULL; - for (n = omp_clauses->lists[OMP_LIST_ALLOCATE]; n;) - { - if (n->sym->mark == 1) - { - gfc_warning (0, "%qs appears more than once in % " - "clauses at %L" , n->sym->name, &n->where); - /* We have already seen this variable so it is a duplicate. - Remove it. */ - if (prev != NULL && prev->next == n) - { - prev->next = n->next; - n->next = NULL; - gfc_free_omp_namelist (n, OMP_LIST_ALLOCATE); - n = prev->next; - } - continue; - } - n->sym->mark = 1; - prev = n; - n = n->next; - } - - /* Non-composite constructs. */ - if (code && code->op < EXEC_OMP_DO_SIMD) - { - for (list = 0; list < OMP_LIST_NUM; list++) - switch (list) - { - case OMP_LIST_PRIVATE: - case OMP_LIST_FIRSTPRIVATE: - case OMP_LIST_LASTPRIVATE: - case OMP_LIST_REDUCTION: - case OMP_LIST_REDUCTION_INSCAN: - case OMP_LIST_REDUCTION_TASK: - case OMP_LIST_IN_REDUCTION: - case OMP_LIST_TASK_REDUCTION: - case OMP_LIST_LINEAR: - for (n = omp_clauses->lists[list]; n; n = n->next) - n->sym->mark = 0; - break; - default: - break; - } - - for (n = omp_clauses->lists[OMP_LIST_ALLOCATE]; n; n = n->next) - if (n->sym->mark == 1) - gfc_error ("%qs specified in % clause at %L but not " - "in an explicit privatization clause", - n->sym->name, &n->where); - } - } + verify_omp_clauses_symbol_dups (code, omp_clauses, ns, openacc); /* OpenACC reductions. */ if (openacc) @@ -8702,20 +8982,6 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } } - 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) - if (n->expr == NULL) - n->sym->mark = 1; - for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next) - { - if (n->expr == NULL && n->sym->mark) - gfc_error ("Symbol %qs present on both FROM and TO clauses at %L", - n->sym->name, &n->where); - else - n->sym->mark = 1; - } - bool has_inscan = false, has_notinscan = false; for (list = 0; list < OMP_LIST_NUM; list++) if ((n = omp_clauses->lists[list]) != NULL) @@ -8886,242 +9152,9 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, "type shall be a scalar integer of " "OMP_DEPEND_KIND kind", &n->expr->where); } - gfc_ref *lastref = NULL, *lastslice = NULL; - bool resolved = false; - if (n->expr) - { - lastref = n->expr->ref; - resolved = gfc_resolve_expr (n->expr); - - /* Look through component refs to find last array - reference. */ - if (resolved) - { - for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) - if (ref->type == REF_COMPONENT - || ref->type == REF_SUBSTRING - || ref->type == REF_INQUIRY) - lastref = ref; - else if (ref->type == REF_ARRAY) - { - for (int i = 0; i < ref->u.ar.dimen; i++) - if (ref->u.ar.dimen_type[i] == DIMEN_RANGE) - lastslice = ref; - - lastref = ref; - } - - /* The "!$acc cache" directive allows rectangular - subarrays to be specified, with some restrictions - on the form of bounds (not implemented). - Only raise an error here if we're really sure the - array isn't contiguous. An expression such as - arr(-n:n,-n:n) could be contiguous even if it looks - like it may not be. - And OpenMP's 'target update' permits strides for - the to/from clause. */ - if (code - && code->op != EXEC_OACC_UPDATE - && code->op != EXEC_OMP_TARGET_UPDATE - && list != OMP_LIST_CACHE - && list != OMP_LIST_DEPEND - && !gfc_is_simply_contiguous (n->expr, false, true) - && gfc_is_not_contiguous (n->expr) - && !(lastslice - && (lastslice->next - || lastslice->type != REF_ARRAY))) - gfc_error ("Array is not contiguous at %L", - &n->where); - } - } - if (openacc - && list == OMP_LIST_MAP - && (n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH)) - { - symbol_attribute attr; - if (n->expr) - attr = gfc_expr_attr (n->expr); - else - attr = n->sym->attr; - if (!attr.pointer && !attr.allocatable) - gfc_error ("%qs clause argument must be ALLOCATABLE or " - "a POINTER at %L", - (n->u.map_op == OMP_MAP_ATTACH) ? "attach" - : "detach", &n->where); - } - if (lastref - || (n->expr - && (!resolved || n->expr->expr_type != EXPR_VARIABLE))) - { - if (!lastslice - && lastref - && lastref->type == REF_SUBSTRING) - gfc_error ("Unexpected substring reference in %s clause " - "at %L", name, &n->where); - else if (!lastslice - && lastref - && lastref->type == REF_INQUIRY) - { - gcc_assert (lastref->u.i == INQUIRY_RE - || lastref->u.i == INQUIRY_IM); - gfc_error ("Unexpected complex-parts designator " - "reference in %s clause at %L", - name, &n->where); - } - else if (!resolved - || n->expr->expr_type != EXPR_VARIABLE - || (lastslice - && (lastslice->next - || lastslice->type != REF_ARRAY))) - gfc_error ("%qs in %s clause at %L is not a proper " - "array section", n->sym->name, name, - &n->where); - else if (lastslice) - { - int i; - gfc_array_ref *ar = &lastslice->u.ar; - for (i = 0; i < ar->dimen; i++) - if (ar->stride[i] - && code->op != EXEC_OACC_UPDATE - && code->op != EXEC_OMP_TARGET_UPDATE) - { - gfc_error ("Stride should not be specified for " - "array section in %s clause at %L", - name, &n->where); - break; - } - else if (ar->dimen_type[i] != DIMEN_ELEMENT - && ar->dimen_type[i] != DIMEN_RANGE) - { - gfc_error ("%qs in %s clause at %L is not a " - "proper array section", - n->sym->name, name, &n->where); - break; - } - else if ((list == OMP_LIST_DEPEND - || list == OMP_LIST_AFFINITY) - && ar->start[i] - && ar->start[i]->expr_type == EXPR_CONSTANT - && ar->end[i] - && ar->end[i]->expr_type == EXPR_CONSTANT - && mpz_cmp (ar->start[i]->value.integer, - ar->end[i]->value.integer) > 0) - { - gfc_error ("%qs in %s clause at %L is a " - "zero size array section", - n->sym->name, - list == OMP_LIST_DEPEND - ? "DEPEND" : "AFFINITY", &n->where); - break; - } - } - } - else if (openacc) - { - if (list == OMP_LIST_MAP - && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR) - resolve_oacc_deviceptr_clause (n->sym, n->where, name); - else - resolve_oacc_data_clauses (n->sym, n->where, name); - } - else if (list != OMP_LIST_DEPEND - && n->sym->as - && n->sym->as->type == AS_ASSUMED_SIZE) - gfc_error ("Assumed size array %qs in %s clause at %L", - n->sym->name, name, &n->where); - if (code && list == OMP_LIST_MAP && !openacc) - switch (code->op) - { - case EXEC_OMP_TARGET: - case EXEC_OMP_TARGET_DATA: - switch (n->u.map_op) - { - case OMP_MAP_TO: - case OMP_MAP_ALWAYS_TO: - case OMP_MAP_PRESENT_TO: - case OMP_MAP_ALWAYS_PRESENT_TO: - case OMP_MAP_FROM: - case OMP_MAP_ALWAYS_FROM: - case OMP_MAP_PRESENT_FROM: - case OMP_MAP_ALWAYS_PRESENT_FROM: - case OMP_MAP_TOFROM: - case OMP_MAP_ALWAYS_TOFROM: - case OMP_MAP_PRESENT_TOFROM: - case OMP_MAP_ALWAYS_PRESENT_TOFROM: - case OMP_MAP_ALLOC: - case OMP_MAP_PRESENT_ALLOC: - break; - default: - gfc_error ("TARGET%s with map-type other than TO, " - "FROM, TOFROM, or ALLOC on MAP clause " - "at %L", - code->op == EXEC_OMP_TARGET - ? "" : " DATA", &n->where); - break; - } - break; - case EXEC_OMP_TARGET_ENTER_DATA: - switch (n->u.map_op) - { - case OMP_MAP_TO: - case OMP_MAP_ALWAYS_TO: - case OMP_MAP_PRESENT_TO: - case OMP_MAP_ALWAYS_PRESENT_TO: - case OMP_MAP_ALLOC: - case OMP_MAP_PRESENT_ALLOC: - break; - case OMP_MAP_TOFROM: - n->u.map_op = OMP_MAP_TO; - break; - case OMP_MAP_ALWAYS_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_TO; - break; - case OMP_MAP_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_PRESENT_TO; - break; - case OMP_MAP_ALWAYS_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO; - break; - default: - gfc_error ("TARGET ENTER DATA with map-type other " - "than TO, TOFROM or ALLOC on MAP clause " - "at %L", &n->where); - break; - } - break; - case EXEC_OMP_TARGET_EXIT_DATA: - switch (n->u.map_op) - { - case OMP_MAP_FROM: - case OMP_MAP_ALWAYS_FROM: - case OMP_MAP_PRESENT_FROM: - case OMP_MAP_ALWAYS_PRESENT_FROM: - case OMP_MAP_RELEASE: - case OMP_MAP_DELETE: - break; - case OMP_MAP_TOFROM: - n->u.map_op = OMP_MAP_FROM; - break; - case OMP_MAP_ALWAYS_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_FROM; - break; - case OMP_MAP_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_PRESENT_FROM; - break; - case OMP_MAP_ALWAYS_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM; - break; - default: - gfc_error ("TARGET EXIT DATA with map-type other " - "than FROM, TOFROM, RELEASE, or DELETE on " - "MAP clause at %L", &n->where); - break; - } - break; - default: - break; - } + if (!omp_verify_map_motion_clauses (code, list, name, n, + openacc)) + break; } if (list != OMP_LIST_DEPEND) @@ -9661,6 +9694,46 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, gfc_resolve_omp_assumptions (omp_clauses->assume); } +/* This very simplified version of the above function is for use after mapper + instantiation. It avoids dealing with anything other than basic + verification for map/to/from clauses. */ + +static void +resolve_omp_mapper_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, + gfc_namespace *ns) +{ + gfc_omp_namelist *n; + int list; + + verify_omp_clauses_symbol_dups (code, omp_clauses, ns, false); + + for (list = OMP_LIST_MAP; list <= OMP_LIST_FROM; list++) + if ((n = omp_clauses->lists[list]) != NULL) + { + const char *name = NULL; + switch (list) + { + case OMP_LIST_MAP: + if (name == NULL) + name = "MAP"; + /* Fallthrough. */ + case OMP_LIST_TO: + if (name == NULL) + name = "TO"; + /* Fallthrough. */ + case OMP_LIST_FROM: + if (name == NULL) + name = "FROM"; + for (; n != NULL; n = n->next) + if (!omp_verify_map_motion_clauses (code, list, name, n, false)) + break; + break; + default: + ; + } + } +} + /* Return true if SYM is ever referenced in EXPR except in the SE node. */ @@ -12377,11 +12450,11 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns) case EXEC_OMP_WORKSHARE: case EXEC_OMP_DEPOBJ: if (code->ext.omp_clauses) - resolve_omp_clauses (code, code->ext.omp_clauses, NULL); + resolve_omp_clauses (code, code->ext.omp_clauses, ns); break; case EXEC_OMP_TARGET_UPDATE: if (code->ext.omp_clauses) - resolve_omp_clauses (code, code->ext.omp_clauses, NULL); + resolve_omp_clauses (code, code->ext.omp_clauses, ns); if (code->ext.omp_clauses == NULL || (code->ext.omp_clauses->lists[OMP_LIST_TO] == NULL && code->ext.omp_clauses->lists[OMP_LIST_FROM] == NULL)) @@ -12988,6 +13061,7 @@ gfc_omp_instantiate_mappers (gfc_code *code, gfc_omp_clauses *clauses, { gfc_omp_namelist *clause = clauses->lists[list]; gfc_omp_namelist **clausep = &clauses->lists[list]; + bool invoked_mappers = false; for (; clause; clause = *clausep) { @@ -13014,10 +13088,20 @@ gfc_omp_instantiate_mappers (gfc_code *code, gfc_omp_clauses *clauses, clausep = gfc_omp_instantiate_mapper (clausep, clause, outer_map_op, clause->u2.udm->udm, cd, list); *clausep = clause->next; + invoked_mappers = true; } else clausep = &clause->next; } + + if (invoked_mappers) + { + gfc_namespace *old_ns = gfc_current_ns; + if (clauses->ns) + gfc_current_ns = clauses->ns; + resolve_omp_mapper_clauses (code, clauses, gfc_current_ns); + gfc_current_ns = old_ns; + } } /* The following functions implement automatic recognition and annotation of diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 new file mode 100644 index 00000000000..c408b37f5a9 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 @@ -0,0 +1,28 @@ +! { dg-do compile } + +type t +integer, allocatable :: arr(:) +end type t + +!$omp declare mapper(even: T :: tv) map(tv%arr(2::2)) + +type(t) :: var + +allocate(var%arr(100)) + +var%arr = 0 + +! You can't do this, the mapper specifies a noncontiguous access. +!$omp target enter data map(mapper(even), to: var) +! { dg-error {Stride should not be specified for array section in MAP clause} "" { target *-*-* } .-1 } + +var%arr = 1 + +! But this is fine. (Re-enabled by later patch.) +!!$omp target update to(mapper(even): var) + +! As 'enter data'. +!$omp target exit data map(mapper(even), delete: var) +! { dg-error {Stride should not be specified for array section in MAP clause} "" { target *-*-* } .-1 } + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-29.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-29.f90 new file mode 100644 index 00000000000..e2039e80e57 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-29.f90 @@ -0,0 +1,22 @@ +! { dg-do compile } + +! Check duplicate clause detection after mapper expansion. + +type t +integer :: x +end type t + +real(4) :: unrelated +type(t) :: tvar + +!$omp declare mapper (t :: var) map(unrelated) map(var%x) + +tvar%x = 0 +unrelated = 5 + +!$omp target firstprivate(unrelated) map(tofrom: tvar) +! { dg-error "Symbol .unrelated. present on both data and map clauses" "" { target *-*-* } .-1 } +tvar%x = unrelated +!$omp end target + +end From patchwork Thu Aug 10 13:33:04 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 133998 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:b824:0:b0:3f2:4152:657d with SMTP id z4csp424506vqi; Thu, 10 Aug 2023 06:35:06 -0700 (PDT) X-Google-Smtp-Source: AGHT+IHW7898oI+mfmQllvO+yAzUklsnuX3Wyog+YlImJ4f0xc+Mo3ehyUEmrnXOEMZPSniQiP5U X-Received: by 2002:a17:907:a06a:b0:99c:6671:66d7 with SMTP id ia10-20020a170907a06a00b0099c667166d7mr2226478ejc.39.1691674506431; Thu, 10 Aug 2023 06:35:06 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1691674506; cv=none; d=google.com; s=arc-20160816; b=Me/J3ZVMD75KnesDoNoh+yQ8xm2U+PyG0svyzAC6udZzK4TN33xSA8aDrf1oksloFE dPCVsi+o0x0PRqu2/s2MUe3ExmjCYFyIxspFHUF4QH0aJk02M+AGWlcchCKH16M0GFsf cGva5lE0I7ZHtDrIGp7K8OXmMMZmyG05OXliX9UCphwk8C394ueLELU+vapMX3X60Nf4 xGNz3KA4ohCXOsX4lYyaEXTI6juZOvrKH8ef9sI9UMQpoEo/s1BiXjR7hevJ0mb+2vWj u86mu5lNz8xff8Pgms3gobZB6IiU89dh4f6NgwANNNDFeVloBATumux4wEUsD0c3H/M1 c98A== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:references:in-reply-to:message-id:date:subject:cc:to :from:ironport-sdr:dmarc-filter:delivered-to; bh=iwPwmbuU9CMLJs1W7htjZWz49AKVKJsHkvHFSAYkkAo=; fh=6entJIlBX3Lqynes5VYRRY4zgfJfSDVLFnUjulEf2Vo=; b=VDyh5B7gIBq/Wu7HM6xBTPVnJ0zGSms5gR/L83dBSvbpBV2hw9Etn8jjK4pktvbyjG LAyN9kZOGPx4WU7WUtVESWFh0a0D/hBayOIO2FbxMi06T5Cevn2Ni82eZVpob0HURUeJ N1g0TIRsD6my0ObdA1y9GyYn6hbdj0LybUJ/B0RIy/2U4KaOL9FQVFbnDC4V1vUhMEpk RhgrjX7BSl4XZJlTdEh29qtx/80eo3uXV1R4j1SFDg8JhNsZ4U7EnfmvduiWWIuA01Ay Xn7aW4tuXQLnO5PBdiXbSanT2XmuNAKmZJfk4a/rcaWdaNHpz/Aruf8x5I+2fUDg3lkx qzvg== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id va3-20020a17090711c300b00992c3b85acbsi1487221ejb.128.2023.08.10.06.35.05 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 10 Aug 2023 06:35:06 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c 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 D8603383138F for ; Thu, 10 Aug 2023 13:33:56 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id A821D385703C; Thu, 10 Aug 2023 13:33:30 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org A821D385703C Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,162,1684828800"; d="scan'208";a="14111557" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 10 Aug 2023 05:33:30 -0800 IronPort-SDR: 1U/FUC4xHVMCOvFC0BKc3jKRXXJHp2ECtwn+fJ1ApijNXoLlkV5dWXFZd++pGi/Fmy/U32i0MU O+4CkogiTH/clqGyH7n7yvts5W8Q1TrbStlx4T/oB7vhCwRNyPewpWRKudKlvhjoCVrnAEFEUq MtAGrKkeFUVKYfKjpzkEBsU7jLmlgI8NI2kG7CN4C7/T+JlU9OmfXmQTNT2s8DhOdWwr3uI41l SPOaeDtgYpJ2lL0NvdkQX4WPyIIx9AFnivC/p6mFvru8xCBYgcu/35u7jJISCKyjaJ72SA7jR4 +FE= From: Julian Brown To: CC: , , Subject: [PATCH 3/5] OpenMP: Introduce C_ORT_{, OMP_}DECLARE_MAPPER c_omp_region_type types Date: Thu, 10 Aug 2023 13:33:04 +0000 Message-ID: <7f150aec826622376459cdecd2c870bc2e80d07b.1691672603.git.julian@codesourcery.com> X-Mailer: git-send-email 2.25.1 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham 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.29 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 Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1773849287429257118 X-GMAIL-MSGID: 1773849287429257118 This patch adds C_ORT_DECLARE_MAPPER and C_ORT_OMP_DECLARE_MAPPER region types to the c_omp_region_type enum, and uses them in cp/pt.cc. Previously the C_ORT_DECLARE_SIMD code was being abused to inhibit calling finish_omp_clauses within mapper definitions, but this patch uses one of the new enumeration values for that purpose instead. This shouldn't result in any behaviour change, but improves self-documentation. 2023-08-10 Julian Brown gcc/c-family/ * c-common.h (c_omp_region_type): Add C_ORT_DECLARE_MAPPER and C_ORT_OMP_DECLARE_MAPPER codes. gcc/cp/ * pt.cc (tsubst_omp_clauses): Use C_ORT_OMP_DECLARE_MAPPER. (tsubst_expr): Likewise. --- gcc/c-family/c-common.h | 2 ++ gcc/cp/pt.cc | 4 ++-- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index c805c8b2f7e..079d1eaafaa 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1271,9 +1271,11 @@ enum c_omp_region_type C_ORT_DECLARE_SIMD = 1 << 2, C_ORT_TARGET = 1 << 3, C_ORT_EXIT_DATA = 1 << 4, + C_ORT_DECLARE_MAPPER = 1 << 6, C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD, C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET, C_ORT_OMP_EXIT_DATA = C_ORT_OMP | C_ORT_EXIT_DATA, + C_ORT_OMP_DECLARE_MAPPER = C_ORT_OMP | C_ORT_DECLARE_MAPPER, C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET }; diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index fb50c5ac48d..2794c0ebecb 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -18328,7 +18328,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, } new_clauses = nreverse (new_clauses); - if (ort != C_ORT_OMP_DECLARE_SIMD) + if (ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_OMP_DECLARE_MAPPER) { if (ort & C_ORT_OMP) new_clauses = c_omp_instantiate_mappers (new_clauses, ort); @@ -19905,7 +19905,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl) decl = tsubst (decl, args, complain, in_decl); tree type = tsubst (TREE_TYPE (t), args, complain, in_decl); tree clauses = OMP_DECLARE_MAPPER_CLAUSES (t); - clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_SIMD, args, + clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_MAPPER, args, complain, in_decl); TREE_TYPE (t) = type; OMP_DECLARE_MAPPER_DECL (t) = decl; From patchwork Thu Aug 10 13:33:05 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 134001 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:b824:0:b0:3f2:4152:657d with SMTP id z4csp425679vqi; Thu, 10 Aug 2023 06:37:07 -0700 (PDT) X-Google-Smtp-Source: AGHT+IEIGECb74iGvMvDZBiwNTqqMzCCMaKW6kyakXUCrK9INqutHbD5+2n4Q3m53A641M6oMX5v X-Received: by 2002:a19:3858:0:b0:4fe:179a:18d2 with SMTP id d24-20020a193858000000b004fe179a18d2mr1477415lfj.21.1691674626770; Thu, 10 Aug 2023 06:37:06 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1691674626; cv=none; d=google.com; s=arc-20160816; b=mgdG+g+QpvHgm+HIHJ26Ec9bfjBw16gTvG+E8iGRbxYRfTolHHsPvm3gLijCzEarGC 3onzHXJFgCglZRt+NQGmyrdVd7h6n3kGlQuzVRXfdzRA8C0XG3zgiKcxmDhIrAd+IdXm u4Bs6WQ6fGCQdYgx3rv+B1BaUI1oueo2XxjehXhHG2/kWwBRlFgE7T4AwDcymbevZub7 K7uteM7giodZZW0uUIM6xSAH7tahSyYxjGP8/QGoR82tZYaVx9REg3q3+b18e3XVaUDR s2F+bdwu1kH5XsPXHPYj8UhnKV3Nhoz4l8Si2Mv1/kY+qPyumT6hLDpCsP4wP+/bl9S6 lNig== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:references:in-reply-to:message-id:date:subject:cc:to :from:ironport-sdr:dmarc-filter:delivered-to; bh=Hf5oGhy0Xu0/7JWalkAL5xgBwl3a8t0plT+dn0Oxdvo=; fh=6entJIlBX3Lqynes5VYRRY4zgfJfSDVLFnUjulEf2Vo=; b=ZbLiK6/xEpXsGSeB9RLSf3rTijb3PcvlE3QwqjwkKjs6EAx4q3cQN+LLVkksOrHtvM cBITM0AT851eLHBe9KEy6EFt9tHEXGpC2o0OBmQBIZNyPA1CBDE3pc3PITj+iINpGTA6 naDKHwdUF1W2CErHXY61EftHLDIMQw7nAsu5ngWhSP17J3QsA004tMWJaz7+H5LVETHw lHx6Gwk+GYAKdUqpPCD82cfFoLj1/zURknhrt8vzG7TaF7pPHyqLJBLQTNQqKN5Wc1PL ABF6HeJWfs6dOVPJI9VPZrQEao6lWYGgF9ueEUhROMzcmaTOnFbVW2Dc7TlceJTFWwo1 qtKg== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id pw4-20020a17090720a400b00992671918a3si1432126ejb.270.2023.08.10.06.37.06 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 10 Aug 2023 06:37:06 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c 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 444C93889E3F for ; Thu, 10 Aug 2023 13:34:38 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id C28B33855586; Thu, 10 Aug 2023 13:33:34 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org C28B33855586 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,162,1684828800"; d="scan'208";a="14111564" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 10 Aug 2023 05:33:34 -0800 IronPort-SDR: ehdZlECWTn4ysks+VP8cvMuLnO11zYhT0DInmoxb4XeNoV6DvJfhNsrVvaO4lunY39mqasBveP HuAnrjCaQzT21MfkI/k+i6M8dgLxRhYWtbobqhzBUF8rDijRcnGir/MgzLamIwghzAzjYuHs/V LTddT564wUBGq3gGw3wNOSh7nQM+Kw9ElAkBYh46h59JVrSCZ0fRS4Z8TJSOqVVouyZckDJe12 Z5O41/lqDzDiYVoFl4WWZqDZ5i/6lptrrqaJX/ExhJUCLY1L/SXa9MLieYgzBDXYSgCwRooP9O ou8= From: Julian Brown To: CC: , , Subject: [PATCH 4/5] OpenMP: Look up 'declare mapper' definitions at resolution time not parse time Date: Thu, 10 Aug 2023 13:33:05 +0000 Message-ID: X-Mailer: git-send-email 2.25.1 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham 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.29 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 Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1773849413450184237 X-GMAIL-MSGID: 1773849413450184237 This patch moves 'declare mapper' lookup for OpenMP clauses from parse time to resolution time for Fortran, and adds diagnostics for missing named mappers. This changes clause lookup in a particular case -- where several 'declare mapper's are defined in a context, mappers declared earlier may now instantiate mappers declared later, whereas previously they would not. I think the new behaviour makes more sense -- at an invocation site, all mappers are visible no matter the declaration order in some particular block. I've adjusted tests to account for this. I think the new arrangement better matches the Fortran FE's usual way of doing things -- mapper lookup is a semantic concept, not a syntactical one, so shouldn't be handled in the syntax-handling code. The patch also fixes a case where the user explicitly writes 'default' as the name on the mapper modifier for a clause. 2023-08-10 Julian Brown gcc/fortran/ * gfortran.h (gfc_omp_namelist_udm): Add MAPPER_ID field to store the mapper name to use for lookup during resolution. * match.cc (gfc_free_omp_namelist): Handle OMP_LIST_TO and OMP_LIST_FROM when freeing mapper references. * module.cc (load_omp_udms, write_omp_udm): Handle MAPPER_ID field. * openmp.cc (gfc_match_omp_clauses): Handle explicitly-specified 'default' name. Don't do mapper lookup here, but record mapper name if the user specifies one. (resolve_omp_clauses): Do mapper lookup here instead. Report error for missing named mapper. gcc/testsuite/ * gfortran.dg/gomp/declare-mapper-31.f90: New test. libgomp/ * testsuite/libgomp.fortran/declare-mapper-30.f90: New test. * testsuite/libgomp.fortran/declare-mapper-4.f90: Adjust test for new lookup behaviour. --- gcc/fortran/gfortran.h | 3 ++ gcc/fortran/match.cc | 4 +- gcc/fortran/module.cc | 6 +++ gcc/fortran/openmp.cc | 46 ++++++++++++++----- .../gfortran.dg/gomp/declare-mapper-31.f90 | 34 ++++++++++++++ .../libgomp.fortran/declare-mapper-30.f90 | 24 ++++++++++ .../libgomp.fortran/declare-mapper-4.f90 | 18 +++++--- 7 files changed, 116 insertions(+), 19 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/declare-mapper-31.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/declare-mapper-30.f90 diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index a98424b3263..3b854e14d47 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1784,6 +1784,9 @@ gfc_omp_udm; typedef struct gfc_omp_namelist_udm { + /* Used to store mapper_id before resolution. */ + const char *mapper_id; + bool multiple_elems_p; struct gfc_omp_udm *udm; } diff --git a/gcc/fortran/match.cc b/gcc/fortran/match.cc index 53367ab2a0b..3db8e0f0969 100644 --- a/gcc/fortran/match.cc +++ b/gcc/fortran/match.cc @@ -5537,7 +5537,9 @@ void gfc_free_omp_namelist (gfc_omp_namelist *name, int list) { bool free_ns = (list == OMP_LIST_AFFINITY || list == OMP_LIST_DEPEND); - bool free_mapper = (list == OMP_LIST_MAP); + bool free_mapper = (list == OMP_LIST_MAP + || list == OMP_LIST_TO + || list == OMP_LIST_FROM); bool free_align = (list == OMP_LIST_ALLOCATE); gfc_omp_namelist *n; diff --git a/gcc/fortran/module.cc b/gcc/fortran/module.cc index 5cd52e7729b..acdbfa7924f 100644 --- a/gcc/fortran/module.cc +++ b/gcc/fortran/module.cc @@ -5238,6 +5238,11 @@ load_omp_udms (void) if (peek_atom () != ATOM_RPAREN) { n->u2.udm = gfc_get_omp_namelist_udm (); + mio_pool_string (&n->u2.udm->mapper_id); + + if (n->u2.udm->mapper_id == NULL) + n->u2.udm->mapper_id = gfc_get_string ("%s", ""); + n->u2.udm->multiple_elems_p = mio_name (0, omp_map_cardinality); mio_pointer_ref (&n->u2.udm->udm); } @@ -6314,6 +6319,7 @@ write_omp_udm (gfc_omp_udm *udm) if (n->u2.udm) { + mio_pool_string (&n->u2.udm->mapper_id); mio_name (n->u2.udm->multiple_elems_p, omp_map_cardinality); mio_pointer_ref (&n->u2.udm->udm); } diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 0109df4dfce..ba2a8221b96 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -3615,6 +3615,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, m = gfc_match (" %n ) ", mapper_id); if (m != MATCH_YES) goto error; + if (strcmp (mapper_id, "default") == 0) + mapper_id[0] = '\0'; } else break; @@ -3689,19 +3691,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, for (n = *head; n; n = n->next) { n->u.map_op = map_op; - - gfc_typespec *ts; - if (n->expr) - ts = &n->expr->ts; - else - ts = &n->sym->ts; - - gfc_omp_udm *udm - = gfc_find_omp_udm (gfc_current_ns, mapper_id, ts); - if (udm) + if (mapper_id[0] != '\0') { n->u2.udm = gfc_get_omp_namelist_udm (); - n->u2.udm->udm = udm; + n->u2.udm->mapper_id + = gfc_get_string ("%s", mapper_id); } } continue; @@ -9155,6 +9149,36 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, if (!omp_verify_map_motion_clauses (code, list, name, n, openacc)) break; + if (list == OMP_LIST_MAP + || list == OMP_LIST_TO + || list == OMP_LIST_FROM) + { + gfc_typespec *ts; + + if (n->expr) + ts = &n->expr->ts; + else + ts = &n->sym->ts; + + const char *mapper_id + = n->u2.udm ? n->u2.udm->mapper_id : ""; + + gfc_omp_udm *udm = gfc_find_omp_udm (gfc_current_ns, + mapper_id, ts); + if (mapper_id[0] != '\0' && !udm) + gfc_error ("User-defined mapper %qs not found at %L", + mapper_id, &n->where); + else if (udm) + { + if (!n->u2.udm) + { + n->u2.udm = gfc_get_omp_namelist_udm (); + gcc_assert (mapper_id[0] == '\0'); + n->u2.udm->mapper_id = mapper_id; + } + n->u2.udm->udm = udm; + } + } } if (list != OMP_LIST_DEPEND) diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-31.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-31.f90 new file mode 100644 index 00000000000..bcb0a6c5429 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-31.f90 @@ -0,0 +1,34 @@ +! { dg-do run } + +type t +integer :: x, y +integer, allocatable :: arr(:) +end type t + +type(t) :: var + +allocate(var%arr(1:20)) + +var%arr = 0 + +! If we ask for a named mapper that hasn't been defined, an error should be +! raised. This isn't a *syntax* error, so the !$omp target..!$omp end target +! block should still be parsed correctly. +!$omp target map(mapper(arraymapper), tofrom: var) +! { dg-error "User-defined mapper .arraymapper. not found" "" { target *-*-* } .-1 } +var%arr(5) = 5 +!$omp end target + +! OTOH, this is a syntax error, and the offload block is not recognized. +!$omp target map( +! { dg-error "Syntax error in OpenMP variable list" "" { target *-*-* } .-1 } +var%arr(6) = 6 +!$omp end target +! { dg-error "Unexpected !.OMP END TARGET statement" "" { target *-*-* } .-1 } + +! ...but not for the specific name 'default'. +!$omp target map(mapper(default), tofrom: var) +var%arr(5) = 5 +!$omp end target + +end diff --git a/libgomp/testsuite/libgomp.fortran/declare-mapper-30.f90 b/libgomp/testsuite/libgomp.fortran/declare-mapper-30.f90 new file mode 100644 index 00000000000..bfac28cd45c --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/declare-mapper-30.f90 @@ -0,0 +1,24 @@ +! { dg-do run } + +type t +integer :: x, y +integer, allocatable :: arr(:) +end type t + +!$omp declare mapper (t :: x) map(x%arr) + +type(t) :: var + +allocate(var%arr(1:20)) + +var%arr = 0 + +! The mapper named literally 'default' should be the default mapper, i.e. +! the same as the unnamed mapper defined above. +!$omp target map(mapper(default), tofrom: var) +var%arr(5) = 5 +!$omp end target + +if (var%arr(5).ne.5) stop 1 + +end diff --git a/libgomp/testsuite/libgomp.fortran/declare-mapper-4.f90 b/libgomp/testsuite/libgomp.fortran/declare-mapper-4.f90 index e95dbbd6f96..266845f35c7 100644 --- a/libgomp/testsuite/libgomp.fortran/declare-mapper-4.f90 +++ b/libgomp/testsuite/libgomp.fortran/declare-mapper-4.f90 @@ -3,7 +3,7 @@ program myprog type s integer :: c - integer :: d(99) + integer, allocatable :: d(:) end type s type t @@ -16,21 +16,25 @@ end type u type(u) :: myu -! Here, the mappers are declared out of order, so later ones are not 'seen' by -! earlier ones. Is that right? +! Here, the mappers are declared out of order, but earlier ones can still +! trigger mappers defined later. Implementation-wise, this happens during +! resolution, but from the user perspective it appears to happen at +! instantiation time -- at which point all mappers are visible. I think +! that makes sense. !$omp declare mapper (u :: x) map(tofrom: x%myt) !$omp declare mapper (t :: x) map(tofrom: x%mys) !$omp declare mapper (s :: x) map(tofrom: x%c, x%d(1:x%c)) +allocate(myu%myt%mys%d(1:20)) + myu%myt%mys%c = 1 myu%myt%mys%d = 0 !$omp target map(tofrom: myu) -myu%myt%mys%d(5) = myu%myt%mys%d(5) + 1 +myu%myt%mys%d(1) = myu%myt%mys%d(1) + 1 !$omp end target -! Note: we used the default mapper, not the 's' mapper, so we mapped the -! whole array 'd'. -if (myu%myt%mys%d(5).ne.1) stop 1 +! Note: we only mapped the first element of the array 'd'. +if (myu%myt%mys%d(1).ne.1) stop 1 end program myprog From patchwork Thu Aug 10 13:33:06 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 134003 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:b824:0:b0:3f2:4152:657d with SMTP id z4csp426443vqi; Thu, 10 Aug 2023 06:38:26 -0700 (PDT) X-Google-Smtp-Source: AGHT+IFFzeBiI3ZCKX622xwU/z0HOp+t1s52m/khBxKKWEXs2El0w0CzFxQkfeGVzRmwy76Ezuca X-Received: by 2002:a17:907:7788:b0:99b:bc52:8d2 with SMTP id ky8-20020a170907778800b0099bbc5208d2mr2098906ejc.6.1691674705923; Thu, 10 Aug 2023 06:38:25 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1691674705; cv=none; d=google.com; s=arc-20160816; b=DcYjejqeK9W4iWdJNntA5GGJNDrUqtY7NZhW9tNM9HMEwHsEdI8h7j038UZ6v08Ud5 aHtOgTWdHCTC+Eiq7ha/sDG06JK04Kw8+c9RgvzSXM832Zs+HmIJP9PRgdJelaVxTn7O lYjdmWwX/AiInHkFT1qzpXqoVPenmtCEX74alh+aiMldyCTb/ecExtVjiN+ObJgZbyZC YTy2bwTTvsWo5UIt+d4Gb9LrQR7fRE4uL0Zkg6Q4hkUrffPZQGHiOIHZAl+S+LdYSPGt e8zd+DE7yveG+ecLHU7/6ikbS0bD0HuTJicgNkOGVlXChQE6mWfhKRP6oGjN9e1/vrpz w8sQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:references:in-reply-to:message-id:date:subject:cc:to :from:ironport-sdr:dmarc-filter:delivered-to; bh=RvMBXKbkAvSbK3of911K/ysVCTY92NiklIoeoSz/+c8=; fh=6entJIlBX3Lqynes5VYRRY4zgfJfSDVLFnUjulEf2Vo=; b=tzM/uvi6iSMjRHTs2jlJEASHUa/v1qnrwFRrbioXFypaJS+Spen5BQA50icKCql5HZ FgVpcahM4y8mKj+6s1h/DxSWdEbLHdLzhyDjxVVuP+DZlFvx0TlZyPHFlpAOi8bH9d2U trh8nwWA/HRX0LEFhIQKh9IdhKQJv3tQCv6dIDaC6y6SKES36MTCa/lniaY+EJ2ODfVb 5R3Q0QeHrgxLmAajiipfcVGM4732WqvuXyX+051YyCQUDJOtc3FiihoOKTRqWDSLkmJ+ /OVE9gBUvuHBVzZm0ZHY2MKxUhEPnWlyIClHdr8UIlCELserCMc46vt11ZAMvu74qijg 9C7w== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id z17-20020a1709064e1100b0099cd52f95dfsi1527032eju.332.2023.08.10.06.38.25 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 10 Aug 2023 06:38:25 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c 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 D6CB53830B4F for ; Thu, 10 Aug 2023 13:35:09 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id D6DA43882669; Thu, 10 Aug 2023 13:34:19 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org D6DA43882669 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,162,1684828800"; d="scan'208";a="15953789" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 10 Aug 2023 05:34:18 -0800 IronPort-SDR: wetX+50St6ZdnJJDslC1jBvn+7dL9g4V6PwZK//jjRvT/cxgJ9FN7lJe6VJS1thr2XDDICOtO7 7uVllZQKODj5flLB7JnNaiEvA34Hfjst79QpiWEBSGlvFYVWU+z4IBnuBUByAxKa7MMO2rzNN/ zUfxXBGl0xjp+nby88Bhc6OpOKMggjiLc3lFkiFx00J52Tm9hJAjOhyh3AkbOGJaFCgXOd5mBi RFD+i2FFctETqRlSGuIz8g/dV1lf+OwisMg8L1mwPDamz0NO4wV9jsyGsHEfKXtREYZeWi3CcY x8s= From: Julian Brown To: CC: , , Subject: [PATCH 5/5] OpenMP: Enable 'declare mapper' mappers for 'target update' directives Date: Thu, 10 Aug 2023 13:33:06 +0000 Message-ID: <0b22f24553a75a560f7ed9f264aac0d5a1239191.1691672603.git.julian@codesourcery.com> X-Mailer: git-send-email 2.25.1 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham 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.29 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 Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1773849496172880985 X-GMAIL-MSGID: 1773849496172880985 This patch enables use of 'declare mapper' for 'target update' directives, for each of C, C++ and Fortran. There are some implementation choices here and some "read-between-the-lines" consequences regarding this functionality, as follows: * It is possible to invoke a mapper which contains clauses that don't make sense for a given 'target update' operation. E.g. if a mapper definition specifies a "from:" mapping and the user does "target update to(...)" which triggers that mapper, the resulting map kind (OpenMP 5.2, "Table 5.3: Map-Type Decay of Map Type Combinations") is "alloc" (and for the inverse case "release"). For such cases, an unconditional warning is issued and the map clause in question is dropped from the mapper expansion. (Other choices might be to make this an error, or to do the same thing but silently, or warn only given some special option.) * The array-shaping operator is *permitted* for map clauses within 'declare mapper' definitions. That is because such mappers may be used for 'target update' directives, where the array-shaping operator is permitted. I think that makes sense, depending on the semantic model of how and when substitution is supposed to take place, but I couldn't find such behaviour explicitly mentioned in the spec (as of 5.2). If the mapper is triggered by a different directive ("omp target", "omp target data", etc.), an error will be raised. Support is also added for the "mapper" modifier on to/from clauses for all three base languages. 2023-08-10 Julian Brown gcc/c-family/ * c-common.h (c_omp_region_type): Add C_ORT_UPDATE and C_ORT_OMP_UPDATE codes. * c-omp.cc (omp_basic_map_kind_name): New function. (omp_instantiate_mapper): Add LOC parameter. Add 'target update' support. (c_omp_instantiate_mappers): Add 'target update' support. gcc/c/ * c-parser.cc (c_parser_omp_variable_list): Support array-shaping operator in 'declare mapper' definitions. (c_parser_omp_clause_map): Pass C_ORT_OMP_DECLARE_MAPPER to c_parser_omp_variable_list in mapper definitions. (c_parser_omp_clause_from_to): Add parsing for mapper modifier. (c_parser_omp_target_update): Instantiate mappers. gcc/cp/ * parser.cc (cp_parser_omp_var_list_no_open): Support array-shaping operator in 'declare mapper' definitions. (cp_parser_omp_clause_from_to): Add parsing for mapper modifier. (cp_parser_omp_clause_map): Pass C_ORT_OMP_DECLARE_MAPPER to cp_parser_omp_var_list_no_open in mapper definitions. (cp_parser_omp_target_update): Instantiate mappers. gcc/fortran/ * openmp.cc (gfc_match_motion_var_list): Add parsing for mapper modifier. (gfc_match_omp_clauses): Adjust error handling for changes to gfc_match_motion_var_list. * trans-openmp.cc (gfc_trans_omp_clauses): Use correct ref for update operations. (gfc_trans_omp_target_update): Instantiate mappers. gcc/testsuite/ * c-c++-common/gomp/declare-mapper-17.c: New test. * c-c++-common/gomp/declare-mapper-19.c: New test. * gfortran.dg/gomp/declare-mapper-24.f90: New test. * gfortran.dg/gomp/declare-mapper-26.f90: Uncomment 'target update' part of test. * gfortran.dg/gomp/declare-mapper-27.f90: New test. libgomp/ * testsuite/libgomp.c-c++-common/declare-mapper-18.c: New test. * testsuite/libgomp.fortran/declare-mapper-25.f90: New test. * testsuite/libgomp.fortran/declare-mapper-28.f90: New test. --- gcc/c-family/c-common.h | 2 + gcc/c-family/c-omp.cc | 117 +++++++++++-- gcc/c/c-parser.cc | 152 +++++++++++++++-- gcc/cp/parser.cc | 160 ++++++++++++++++-- gcc/fortran/openmp.cc | 86 ++++++++-- gcc/fortran/trans-openmp.cc | 20 ++- .../c-c++-common/gomp/declare-mapper-17.c | 38 +++++ .../c-c++-common/gomp/declare-mapper-19.c | 40 +++++ .../gfortran.dg/gomp/declare-mapper-24.f90 | 43 +++++ .../gfortran.dg/gomp/declare-mapper-26.f90 | 4 +- .../gfortran.dg/gomp/declare-mapper-27.f90 | 25 +++ .../libgomp.c-c++-common/declare-mapper-18.c | 33 ++++ .../libgomp.fortran/declare-mapper-25.f90 | 44 +++++ .../libgomp.fortran/declare-mapper-28.f90 | 38 +++++ 14 files changed, 746 insertions(+), 56 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-17.c create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c create mode 100644 gcc/testsuite/gfortran.dg/gomp/declare-mapper-24.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/declare-mapper-27.f90 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-mapper-18.c create mode 100644 libgomp/testsuite/libgomp.fortran/declare-mapper-25.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/declare-mapper-28.f90 diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 079d1eaafaa..0998d875e3f 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1271,10 +1271,12 @@ enum c_omp_region_type C_ORT_DECLARE_SIMD = 1 << 2, C_ORT_TARGET = 1 << 3, C_ORT_EXIT_DATA = 1 << 4, + C_ORT_UPDATE = 1 << 5, C_ORT_DECLARE_MAPPER = 1 << 6, C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD, C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET, C_ORT_OMP_EXIT_DATA = C_ORT_OMP | C_ORT_EXIT_DATA, + C_ORT_OMP_UPDATE = C_ORT_OMP | C_ORT_UPDATE, C_ORT_OMP_DECLARE_MAPPER = C_ORT_OMP | C_ORT_DECLARE_MAPPER, C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET }; diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 2ddb30b54f8..9ff09b59bc6 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -5176,12 +5176,37 @@ omp_map_decayed_kind (enum gomp_map_kind mapper_kind, return omp_join_map_kind (decay_to, force_p, always_p, present_p); } +/* Return a name to use for a "basic" map kind, e.g. as output from + omp_split_map_kind above. */ + +static const char * +omp_basic_map_kind_name (enum gomp_map_kind kind) +{ + switch (kind) + { + case GOMP_MAP_ALLOC: + return "alloc"; + case GOMP_MAP_TO: + return "to"; + case GOMP_MAP_FROM: + return "from"; + case GOMP_MAP_TOFROM: + return "tofrom"; + case GOMP_MAP_RELEASE: + return "release"; + case GOMP_MAP_DELETE: + return "delete"; + default: + gcc_unreachable (); + } +} + /* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in the mapper declaration. */ static tree * -omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, +omp_instantiate_mapper (location_t loc, tree *outlist, tree mapper, tree expr, enum gomp_map_kind outer_kind, enum c_omp_region_type ort) { @@ -5214,7 +5239,6 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - location_t loc = OMP_CLAUSE_LOCATION (c); tree t2 = lang_hooks.decls.omp_map_array_section (loc, t); if (t2 == t) @@ -5240,9 +5264,13 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL); + OMP_CLAUSE_LOCATION (unshared) = loc; + enum gomp_map_kind decayed_kind = omp_map_decayed_kind (clause_kind, outer_kind, - (ort & C_ORT_EXIT_DATA) != 0); + (ort & C_ORT_EXIT_DATA) != 0 + || (outer_kind == GOMP_MAP_FROM + && (ort & C_ORT_UPDATE) != 0)); OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_kind); type = TYPE_MAIN_VARIANT (type); @@ -5260,7 +5288,7 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); if (nested_mapper != mapper) { - outlist = omp_instantiate_mapper (outlist, nested_mapper, + outlist = omp_instantiate_mapper (loc, outlist, nested_mapper, t, outer_kind, ort); continue; } @@ -5271,8 +5299,51 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, continue; } - *outlist = unshared; - outlist = &OMP_CLAUSE_CHAIN (unshared); + if (ort & C_ORT_UPDATE) + { + bool force_p, always_p, present_p; + decayed_kind + = omp_split_map_kind (decayed_kind, &force_p, &always_p, + &present_p); + /* We don't expect to see these flags here. */ + gcc_assert (!force_p && !always_p); + /* For a "target update" operation, we want to turn the map node + expanded from the mapper back into a OMP_CLAUSE_TO or + OMP_CLAUSE_FROM node. If we can do neither, emit a warning and + drop the clause. */ + switch (decayed_kind) + { + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + { + tree xfer + = build_omp_clause (loc, (decayed_kind == GOMP_MAP_TO + ? OMP_CLAUSE_TO : OMP_CLAUSE_FROM)); + OMP_CLAUSE_DECL (xfer) = OMP_CLAUSE_DECL (unshared); + OMP_CLAUSE_SIZE (xfer) = OMP_CLAUSE_SIZE (unshared); + /* For FROM/TO clauses, "present" is represented by a flag. + Set it for the expanded clause here. */ + if (present_p) + OMP_CLAUSE_MOTION_PRESENT (xfer) = 1; + *outlist = xfer; + outlist = &OMP_CLAUSE_CHAIN (xfer); + } + break; + default: + clause_kind + = omp_split_map_kind (clause_kind, &force_p, &always_p, + &present_p); + warning_at (loc, 0, "dropping %qs clause during mapper expansion " + "in %<#pragma omp target update%>", + omp_basic_map_kind_name (clause_kind)); + inform (OMP_CLAUSE_LOCATION (c), "for map clause here"); + } + } + else + { + *outlist = unshared; + outlist = &OMP_CLAUSE_CHAIN (unshared); + } } return outlist; @@ -5290,17 +5361,25 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort) for (pc = &clauses, c = clauses; c; c = *pc) { bool using_mapper = false; + bool update_p = false, update_present_p = false; switch (OMP_CLAUSE_CODE (c)) { + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: + update_p = true; + if (OMP_CLAUSE_MOTION_PRESENT (c)) + update_present_p = true; + /* Fallthrough. */ case OMP_CLAUSE_MAP: { tree t = OMP_CLAUSE_DECL (c); tree type = NULL_TREE; bool nonunit_array_with_mapper = false; - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME) + if (!update_p + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)) { if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME) mapper_name = OMP_CLAUSE_DECL (c); @@ -5337,9 +5416,22 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort) continue; } - enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c); - if (kind == GOMP_MAP_UNSET) - kind = GOMP_MAP_TOFROM; + enum gomp_map_kind kind; + if (update_p) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO) + kind = update_present_p ? GOMP_MAP_PRESENT_TO + : GOMP_MAP_TO; + else + kind = update_present_p ? GOMP_MAP_PRESENT_FROM + : GOMP_MAP_FROM; + } + else + { + kind = OMP_CLAUSE_MAP_KIND (c); + if (kind == GOMP_MAP_UNSET) + kind = GOMP_MAP_TOFROM; + } type = TYPE_MAIN_VARIANT (type); @@ -5356,7 +5448,8 @@ c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort) { tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); - pc = omp_instantiate_mapper (pc, mapper, t, kind, ort); + pc = omp_instantiate_mapper (OMP_CLAUSE_LOCATION (c), + pc, mapper, t, kind, ort); using_mapper = true; } else if (mapper_name) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index f1348a7e5f3..b37c2d60593 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -14196,7 +14196,9 @@ c_parser_omp_variable_list (c_parser *parser, bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p; c_omp_array_section_p = true; c_omp_array_shaping_op_p - = (kind == OMP_CLAUSE_TO || kind == OMP_CLAUSE_FROM); + = (kind == OMP_CLAUSE_TO + || kind == OMP_CLAUSE_FROM + || ort == C_ORT_OMP_DECLARE_MAPPER); c_expr expr = c_parser_expr_no_commas (parser, NULL); if (expr.value != error_mark_node) mark_exp_read (expr.value); @@ -18078,7 +18080,9 @@ c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind) } nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list, - C_ORT_OMP, true); + (kind == GOMP_MAP_UNSET + ? C_ORT_OMP_DECLARE_MAPPER + : C_ORT_OMP), true); tree last_new = NULL_TREE; @@ -18355,26 +18359,148 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, if (!parens.require_open (parser)) return list; - bool present = false; - c_token *token = c_parser_peek_token (parser); + int pos = 1, colon_pos = 0; - if (token->type == CPP_NAME - && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0 - && c_parser_peek_2nd_token (parser)->type == CPP_COLON) + while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME) { - present = true; - c_parser_consume_token (parser); - c_parser_consume_token (parser); + if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA) + pos += 2; + else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type + == CPP_OPEN_PAREN) + { + unsigned int npos = pos + 2; + if (c_parser_check_balanced_raw_token_sequence (parser, &npos) + && (c_parser_peek_nth_token_raw (parser, npos)->type + == CPP_CLOSE_PAREN)) + pos = npos + 1; + } + else + pos++; + if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON) + { + colon_pos = pos; + break; + } } + int present_modifier = false; + int mapper_modifier = false; + tree mapper_name = NULL_TREE; + + for (int pos = 1; pos < colon_pos; ++pos) + { + c_token *tok = c_parser_peek_token (parser); + if (tok->type == CPP_COMMA) + { + c_parser_consume_token (parser); + continue; + } + const char *p = IDENTIFIER_POINTER (tok->value); + if (strcmp ("present", p) == 0) + { + if (present_modifier) + { + c_parser_error (parser, "too many % modifiers"); + parens.skip_until_found_close (parser); + return list; + } + present_modifier++; + c_parser_consume_token (parser); + } + else if (strcmp ("mapper", p) == 0) + { + c_parser_consume_token (parser); + + matching_parens mparens; + if (mparens.require_open (parser)) + { + if (mapper_modifier) + { + c_parser_error (parser, "too many % modifiers"); + /* Assume it's a well-formed mapper modifier, even if it + seems to be in the wrong place. */ + c_parser_consume_token (parser); + mparens.require_close (parser); + parens.skip_until_found_close (parser); + return list; + } + + tok = c_parser_peek_token (parser); + + switch (tok->type) + { + case CPP_NAME: + { + mapper_name = tok->value; + c_parser_consume_token (parser); + } + break; + + case CPP_KEYWORD: + if (tok->keyword == RID_DEFAULT) + { + c_parser_consume_token (parser); + break; + } + /* Fallthrough. */ + + default: + error_at (tok->location, + "expected identifier or %"); + return list; + } + + if (!mparens.require_close (parser)) + { + parens.skip_until_found_close (parser); + return list; + } + + mapper_modifier++; + pos += 3; + } + } + else + { + c_parser_error (parser, "% or % clause with modifier " + "other than % or %"); + parens.skip_until_found_close (parser); + return list; + } + } + + if (colon_pos) + c_parser_require (parser, CPP_COLON, "expected %<:%>"); + tree nl = c_parser_omp_variable_list (parser, loc, kind, list, C_ORT_OMP, true); parens.skip_until_found_close (parser); - if (present) + if (present_modifier) for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_MOTION_PRESENT (c) = 1; + if (mapper_name) + { + tree last_new = NULL_TREE; + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + last_new = c; + + tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME); + OMP_CLAUSE_DECL (name) = mapper_name; + OMP_CLAUSE_CHAIN (name) = nl; + nl = name; + + gcc_assert (last_new); + + name = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME); + OMP_CLAUSE_DECL (name) = null_pointer_node; + OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new); + OMP_CLAUSE_CHAIN (last_new) = name; + } + return nl; } @@ -23023,7 +23149,9 @@ c_parser_omp_target_update (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK, - "#pragma omp target update"); + "#pragma omp target update", false); + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_UPDATE); + clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_UPDATE); bool to_clause = false, from_clause = false; for (tree c = clauses; c && !to_clause && !from_clause; diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 4984accb0bd..f7576c5d137 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -37770,7 +37770,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, auto s = make_temp_override (parser->omp_array_section_p, true); auto o = make_temp_override (parser->omp_array_shaping_op_p, (kind == OMP_CLAUSE_TO - || kind == OMP_CLAUSE_FROM)); + || kind == OMP_CLAUSE_FROM + || ort == C_ORT_OMP_DECLARE_MAPPER)); tree reshaped_to = NULL_TREE; token = cp_lexer_peek_token (parser->lexer); location_t loc = token->location; @@ -41308,24 +41309,153 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return list; - bool present = false; - cp_token *token = cp_lexer_peek_token (parser->lexer); + int pos = 1; + int colon_pos = 0; - if (token->type == CPP_NAME - && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0 - && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME) { - present = true; - cp_lexer_consume_token (parser->lexer); - cp_lexer_consume_token (parser->lexer); + if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA) + pos += 2; + else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type + == CPP_OPEN_PAREN) + pos = cp_parser_skip_balanced_tokens (parser, pos + 1); + else + pos++; + if (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_COLON) + { + colon_pos = pos; + break; + } } + bool present_modifier = false; + bool mapper_modifier = false; + tree mapper_name = NULL_TREE; + + for (int pos = 1; pos < colon_pos; ++pos) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + if (tok->type == CPP_COMMA) + { + cp_lexer_consume_token (parser->lexer); + continue; + } + const char *p = IDENTIFIER_POINTER (tok->u.value); + if (strcmp ("present", p) == 0) + { + if (present_modifier) + { + cp_parser_error (parser, "too many % modifiers"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + present_modifier = true; + cp_lexer_consume_token (parser->lexer); + } + else if (strcmp ("mapper", p) == 0) + { + cp_lexer_consume_token (parser->lexer); + matching_parens parens; + if (parens.require_open (parser)) + { + if (mapper_modifier) + { + cp_parser_error (parser, "too many % modifiers"); + /* Assume it's a well-formed mapper modifier, even if it + seems to be in the wrong place. */ + cp_lexer_consume_token (parser->lexer); + parens.require_close (parser); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/ + true); + return list; + } + tok = cp_lexer_peek_token (parser->lexer); + switch (tok->type) + { + case CPP_NAME: + { + cp_expr e = cp_parser_identifier (parser); + if (e != error_mark_node) + mapper_name = e; + else + goto err; + } + break; + case CPP_KEYWORD: + if (tok->keyword == RID_DEFAULT) + { + cp_lexer_consume_token (parser->lexer); + break; + } + /* Fallthrough. */ + default: + err: + cp_parser_error (parser, + "expected identifier or %"); + return list; + } + + if (!parens.require_close (parser)) + { + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/ + true); + return list; + } + mapper_modifier = true; + pos += 3; + } + else + { + cp_parser_error (parser, "% or % clause with " + "modifier other than % or %"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + } + } + + if (colon_pos) + cp_parser_require (parser, CPP_COLON, RT_COLON); + tree nl = cp_parser_omp_var_list_no_open (parser, kind, list, NULL, C_ORT_OMP, true); - if (present) + if (present_modifier) for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_MOTION_PRESENT (c) = 1; + if (mapper_name) + { + tree last_new = NULL_TREE; + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + last_new = c; + + tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME); + OMP_CLAUSE_DECL (name) = mapper_name; + OMP_CLAUSE_CHAIN (name) = nl; + nl = name; + + gcc_assert (last_new); + + name = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME); + OMP_CLAUSE_DECL (name) = null_pointer_node; + OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new); + OMP_CLAUSE_CHAIN (last_new) = name; + } + return nl; } @@ -41559,7 +41689,9 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind) legally. */ begin_scope (sk_omp, NULL); nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, - NULL, C_ORT_OMP, true); + NULL, (kind == GOMP_MAP_UNSET + ? C_ORT_OMP_DECLARE_MAPPER + : C_ORT_OMP), true); finish_scope (); tree last_new = NULL_TREE; @@ -46931,7 +47063,11 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK, - "#pragma omp target update", pragma_tok); + "#pragma omp target update", pragma_tok, + false); + if (!processing_template_decl) + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_UPDATE); + clauses = finish_omp_clauses (clauses, C_ORT_OMP_UPDATE); bool to_clause = false, from_clause = false; for (tree c = clauses; c && !to_clause && !from_clause; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index ba2a8221b96..96e764b90a6 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -1448,16 +1448,68 @@ gfc_match_motion_var_list (const char *str, gfc_omp_namelist **list, if (m != MATCH_YES) return m; - match m_present = gfc_match (" present : "); + locus old_loc = gfc_current_locus; + int present_modifier = 0; + int mapper_modifier = 0; + locus second_mapper_locus = old_loc; + locus second_present_locus = old_loc; + char mapper_id[GFC_MAX_SYMBOL_LEN + 1] = { '\0' }; + + for (;;) + { + locus current_locus = gfc_current_locus; + if (gfc_match ("present ") == MATCH_YES) + { + if (present_modifier++ == 1) + second_present_locus = current_locus; + } + else if (gfc_match ("mapper ( ") == MATCH_YES) + { + if (mapper_modifier++ == 1) + second_mapper_locus = current_locus; + m = gfc_match (" %n ) ", mapper_id); + if (m != MATCH_YES) + return m; + if (strcmp (mapper_id, "default") == 0) + mapper_id[0] = '\0'; + } + else + break; + gfc_match (", "); + } + + if (gfc_match (" : ") != MATCH_YES) + { + gfc_current_locus = old_loc; + present_modifier = 0; + mapper_modifier = 0; + } + + if (present_modifier > 1) + { + gfc_error ("too many % modifiers at %L", &second_present_locus); + return MATCH_ERROR; + } + if (mapper_modifier > 1) + { + gfc_error ("too many % modifiers at %L", &second_mapper_locus); + return MATCH_ERROR; + } m = gfc_match_omp_variable_list ("", list, false, NULL, headp, true, true); if (m != MATCH_YES) return m; - if (m_present == MATCH_YES) + gfc_omp_namelist *n; + for (n = **headp; n; n = n->next) { - gfc_omp_namelist *n; - for (n = **headp; n; n = n->next) + if (present_modifier) n->u.present_modifier = true; + + if (mapper_id[0] != '\0') + { + n->u2.udm = gfc_get_omp_namelist_udm (); + n->u2.udm->mapper_id = gfc_get_string ("%s", mapper_id); + } } return MATCH_YES; } @@ -3215,10 +3267,15 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, &c->lists[OMP_LIST_FIRSTPRIVATE], true) == MATCH_YES) continue; - if ((mask & OMP_CLAUSE_FROM) - && gfc_match_motion_var_list ("from (", &c->lists[OMP_LIST_FROM], - &head) == MATCH_YES) - continue; + if (mask & OMP_CLAUSE_FROM) + { + m = gfc_match_motion_var_list ("from (", &c->lists[OMP_LIST_FROM], + &head); + if (m == MATCH_YES) + continue; + else if (m == MATCH_ERROR) + goto error; + } if ((mask & OMP_CLAUSE_UNROLL_FULL) && (m = gfc_match_dupl_check (!c->unroll_full, "full")) != MATCH_NO) @@ -4240,10 +4297,15 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if (m == MATCH_YES) continue; } - else if ((mask & OMP_CLAUSE_TO) - && gfc_match_motion_var_list ("to (", &c->lists[OMP_LIST_TO], - &head) == MATCH_YES) - continue; + else if (mask & OMP_CLAUSE_TO) + { + m = gfc_match_motion_var_list ("to (", &c->lists[OMP_LIST_TO], + &head); + if (m == MATCH_YES) + continue; + else if (m == MATCH_ERROR) + goto error; + } break; case 'u': if ((mask & OMP_CLAUSE_UNIFORM) diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 170615974b3..746efb75369 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -6171,11 +6171,14 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gcc_unreachable (); } + gfc_ref *lastref = NULL; + if (n->expr) + for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) + if (ref->type == REF_COMPONENT || ref->type == REF_ARRAY) + lastref = ref; + if ((list == OMP_LIST_TO || list == OMP_LIST_FROM) - && (!n->expr - || (n->expr - && n->expr->ref - && n->expr->ref->type == REF_ARRAY)) + && (!n->expr || (lastref && lastref->type == REF_ARRAY)) && !gfc_omp_contiguous_update_p (n)) { int ndims; @@ -6197,7 +6200,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, or other data between elements, e.g. of a derived-type array). */ span = gfc_get_array_span (desc, n->expr); - ndims = n->expr->ref->u.ar.dimen; + ndims = lastref->u.ar.dimen; } else { @@ -10568,7 +10571,12 @@ gfc_trans_omp_target_update (gfc_code *code) tree stmt, omp_clauses; gfc_start_block (&block); - omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses, + gfc_omp_clauses *target_update_clauses = code->ext.omp_clauses; + gfc_omp_instantiate_mappers (code, target_update_clauses, TOC_OPENMP, + OMP_LIST_TO); + gfc_omp_instantiate_mappers (code, target_update_clauses, TOC_OPENMP, + OMP_LIST_FROM); + omp_clauses = gfc_trans_omp_clauses (&block, target_update_clauses, code->loc); stmt = build1_loc (input_location, OMP_TARGET_UPDATE, void_type_node, omp_clauses); diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-17.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-17.c new file mode 100644 index 00000000000..ddbb59e4f7f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-17.c @@ -0,0 +1,38 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +typedef struct { + int a, b, c, d; +} S; + +#pragma omp declare mapper (S s) map(alloc: s.a) map(to: s.b) map(from: s.c) \ + map(tofrom: s.d) +#pragma omp declare mapper (update: S s) map(s.a, s.b, s.c, s.d) + +int main() +{ + S v; +#pragma omp target update to(v) +/* { dg-warning {dropping .from. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-1 } */ +/* { dg-warning {dropping .alloc. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-2 } */ +/* { dg-final { scan-tree-dump-times {(?n)update to\(v\.d\) to\(v\.b\)$} 1 "original" } } */ +#pragma omp target update from(v) +/* { dg-warning {dropping .to. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-1 } */ +/* { dg-warning {dropping .alloc. clause during mapper expansion in .#pragma omp target update.} "" { target *-*-* } .-2 } */ +/* { dg-final { scan-tree-dump-times {(?n)update from\(v\.d\) from\(v\.c\)$} 1 "original" } } */ + +#pragma omp target update to(mapper(update): v) +/* { dg-final { scan-tree-dump-times {(?n)update to\(v\.d\) to\(v\.c\) to\(v\.b\) to\(v\.a\)$} 1 "original" } } */ +#pragma omp target update from(mapper(update): v) +/* { dg-final { scan-tree-dump-times {(?n)update from\(v\.d\) from\(v\.c\) from\(v\.b\) from\(v\.a\)$} 1 "original" } } */ + +#pragma omp target update to(present, mapper(update): v) +/* { dg-final { scan-tree-dump-times {(?n)update to\(present:v\.d\) to\(present:v\.c\) to\(present:v\.b\) to\(present:v\.a\)$} 2 "original" } } */ +#pragma omp target update from(present, mapper(update): v) +/* { dg-final { scan-tree-dump-times {(?n)update from\(present:v\.d\) from\(present:v\.c\) from\(present:v\.b\) from\(present:v\.a\)$} 2 "original" } } */ + +#pragma omp target update to(present: v.a, v.b, v.c, v.d) +#pragma omp target update from(present: v.a, v.b, v.c, v.d) + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c new file mode 100644 index 00000000000..fd40c6a25e8 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-19.c @@ -0,0 +1,40 @@ +/* { dg-do compile } */ + +#include +#include +#include + +typedef struct { + int *ptr; +} S; + +int main(void) +{ +#pragma omp declare mapper(grid: S x) map(([9][11]) x.ptr[3:3:2][1:4:3]) + S q; + q.ptr = (int *) calloc (9 * 11, sizeof (int)); + + /* The 'grid' mapper specifies a noncontiguous region, so it can't be used + for 'map' like this. */ +#pragma omp target enter data map(mapper(grid), to: q) +/* { dg-error {array section is not contiguous in .map. clause} "" { target *-*-* } .-1 } */ +/* { dg-error {.#pragma omp target enter data. must contain at least one .map. clause} "" { target *-*-* } .-2 } */ + +#pragma omp target + for (int i = 0; i < 9*11; i++) + q.ptr[i] = i; + + /* It's OK on a 'target update' directive though. */ +#pragma omp target update from(mapper(grid): q) + + for (int j = 0; j < 9; j++) + for (int i = 0; i < 11; i++) + if (j >= 3 && j <= 7 && ((j - 3) % 2) == 0 + && i >= 1 && i <= 10 && ((i - 1) % 3) == 0) + assert (q.ptr[j * 11 + i] == j * 11 + i); + +#pragma omp target exit data map(mapper(grid), release: q) +/* { dg-error {array section is not contiguous in .map. clause} "" { target *-*-* } .-1 } */ +/* { dg-error {.#pragma omp target exit data. must contain at least one .map. clause} "" { target *-*-* } .-2 } */ + return 0; +} diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-24.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-24.f90 new file mode 100644 index 00000000000..9555a94bada --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-24.f90 @@ -0,0 +1,43 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } + +type t +integer :: a, b, c, d +end type t + +type(t) :: tvar + +!$omp declare mapper (T :: t) map(alloc: t%a) map(to: t%b) map(from: t%c) & +!$omp & map(tofrom: t%d) + +!$omp declare mapper (updatey: T :: t) map(t%a) map(t%b) map(t%c) map(t%d) + +!$omp target update to(tvar) +! { dg-warning "Dropping incompatible .ALLOC. mapper clause" "" { target *-*-* } .-1 } +! { dg-warning "Dropping incompatible .FROM. mapper clause" "" { target *-*-* } .-2 } +! { dg-final { scan-tree-dump-times {(?n)update to\(tvar\.b \[len: [0-9]+\]\) to\(tvar\.d \[len: [0-9]+\]\)$} 1 "original" } } +!$omp target update from(tvar) +! { dg-warning "Dropping incompatible .ALLOC. mapper clause" "" { target *-*-* } .-1 } +! { dg-warning "Dropping incompatible .TO. mapper clause" "" { target *-*-* } .-2 } +! { dg-final { scan-tree-dump-times {(?n)update from\(tvar\.c \[len: [0-9]+\]\) from\(tvar\.d \[len: [0-9]+\]\)$} 1 "original" } } + +!$omp target update to(present: tvar) +! { dg-warning "Dropping incompatible .ALLOC. mapper clause" "" { target *-*-* } .-1 } +! { dg-warning "Dropping incompatible .FROM. mapper clause" "" { target *-*-* } .-2 } +! { dg-final { scan-tree-dump-times {(?n)update to\(present:tvar\.b \[len: [0-9]+\]\) to\(present:tvar\.d \[len: [0-9]+\]\)$} 1 "original" } } +!$omp target update from(present: tvar) +! { dg-warning "Dropping incompatible .ALLOC. mapper clause" "" { target *-*-* } .-1 } +! { dg-warning "Dropping incompatible .TO. mapper clause" "" { target *-*-* } .-2 } +! { dg-final { scan-tree-dump-times {(?n)update from\(present:tvar\.c \[len: [0-9]+\]\) from\(present:tvar\.d \[len: [0-9]+\]\)$} 1 "original" } } + +!$omp target update to(mapper(updatey): tvar) +! { dg-final { scan-tree-dump-times {(?n)update to\(tvar\.a \[len: [0-9]+\]\) to\(tvar\.b \[len: [0-9]+\]\) to\(tvar\.c \[len: [0-9]+\]\) to\(tvar\.d \[len: [0-9]+\]\)$} 1 "original" } } +!$omp target update from(mapper(updatey): tvar) +! { dg-final { scan-tree-dump-times {(?n)update from\(tvar\.a \[len: [0-9]+\]\) from\(tvar\.b \[len: [0-9]+\]\) from\(tvar\.c \[len: [0-9]+\]\) from\(tvar\.d \[len: [0-9]+\]\)$} 1 "original" } } + +!$omp target update to(present, mapper(updatey): tvar) +! { dg-final { scan-tree-dump-times {(?n)update to\(present:tvar\.a \[len: [0-9]+\]\) to\(present:tvar\.b \[len: [0-9]+\]\) to\(present:tvar\.c \[len: [0-9]+\]\) to\(present:tvar\.d \[len: [0-9]+\]\)$} 1 "original" } } +!$omp target update from(present, mapper(updatey): tvar) +! { dg-final { scan-tree-dump-times {(?n)update from\(present:tvar\.a \[len: [0-9]+\]\) from\(present:tvar\.b \[len: [0-9]+\]\) from\(present:tvar\.c \[len: [0-9]+\]\) from\(present:tvar\.d \[len: [0-9]+\]\)$} 1 "original" } } + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 index c408b37f5a9..16afb51f48a 100644 --- a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-26.f90 @@ -18,8 +18,8 @@ var%arr = 0 var%arr = 1 -! But this is fine. (Re-enabled by later patch.) -!!$omp target update to(mapper(even): var) +! But this is fine. +!$omp target update to(mapper(even): var) ! As 'enter data'. !$omp target exit data map(mapper(even), delete: var) diff --git a/gcc/testsuite/gfortran.dg/gomp/declare-mapper-27.f90 b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-27.f90 new file mode 100644 index 00000000000..6b3a181acaa --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/declare-mapper-27.f90 @@ -0,0 +1,25 @@ +! { dg-do compile } + +type t +integer :: x +end type t + +type(t) :: var + +! Error on attempt to use missing named mapper. +!$omp target update to(mapper(boo): var) +! { dg-error {User-defined mapper .boo. not found} "" { target *-*-* } .-1 } + +var%x = 0 + +!$omp target map(mapper(boo), tofrom: var) +! { dg-error {User-defined mapper .boo. not found} "" { target *-*-* } .-1 } +var%x = 5 +!$omp end target + +! These should be fine though... +!$omp target enter data map(mapper(default), to: var) + +!$omp target exit data map(from: var) + +end diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-18.c b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-18.c new file mode 100644 index 00000000000..50f37cba89d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-mapper-18.c @@ -0,0 +1,33 @@ +#include +#include +#include + +typedef struct { + int *ptr; +} S; + +int main(void) +{ +#pragma omp declare mapper(grid: S x) map(([9][11]) x.ptr[3:3:2][1:4:3]) + S q; + q.ptr = (int *) calloc (9 * 11, sizeof (int)); + +#pragma omp target enter data map(to: q.ptr, q.ptr[0:9*11]) + +#pragma omp target + for (int i = 0; i < 9*11; i++) + q.ptr[i] = i; + +#pragma omp target update from(mapper(grid): q) + + for (int j = 0; j < 9; j++) + for (int i = 0; i < 11; i++) + if (j >= 3 && j <= 7 && ((j - 3) % 2) == 0 + && i >= 1 && i <= 10 && ((i - 1) % 3) == 0) + assert (q.ptr[j * 11 + i] == j * 11 + i); + else + assert (q.ptr[j * 11 + i] == 0); + +#pragma omp target exit data map(release: q.ptr, q.ptr[0:9*11]) + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/declare-mapper-25.f90 b/libgomp/testsuite/libgomp.fortran/declare-mapper-25.f90 new file mode 100644 index 00000000000..dc1f5272a26 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/declare-mapper-25.f90 @@ -0,0 +1,44 @@ +! { dg-do run } +! { dg-require-effective-target offload_device_nonshared_as } + +type t +integer, allocatable :: arr(:) +end type t + +!$omp declare mapper(odd: T :: tv) map(tv%arr(1::2)) +!$omp declare mapper(even: T :: tv) map(tv%arr(2::2)) + +type(t) :: var +integer :: i + +allocate(var%arr(100)) + +var%arr = 0 + +!$omp target enter data map(to: var) + +var%arr = 1 + +!$omp target update to(mapper(odd): var) + +!$omp target +do i=1,100 + if (mod(i,2).eq.0.and.var%arr(i).ne.0) stop 1 + if (mod(i,2).eq.1.and.var%arr(i).ne.1) stop 2 +end do +!$omp end target + +var%arr = 2 + +!$omp target update to(mapper(even): var) + +!$omp target +do i=1,100 + if (mod(i,2).eq.0.and.var%arr(i).ne.2) stop 3 + if (mod(i,2).eq.1.and.var%arr(i).ne.1) stop 4 +end do +!$omp end target + +!$omp target exit data map(delete: var) + +end diff --git a/libgomp/testsuite/libgomp.fortran/declare-mapper-28.f90 b/libgomp/testsuite/libgomp.fortran/declare-mapper-28.f90 new file mode 100644 index 00000000000..6561decc49a --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/declare-mapper-28.f90 @@ -0,0 +1,38 @@ +! { dg-do run } + +program p + +type t +integer :: x, y +end type t + +type(t) :: var + +var%x = 0 +var%y = 0 + +var = sub(7) + +contains + +type(t) function sub(arg) +integer :: arg + +!$omp declare mapper (t :: tvar) map(tvar%x, tvar%y) + +!$omp target enter data map(alloc: sub) + +sub%x = 5 +sub%y = arg + +!$omp target update to(sub) + +!$omp target +if (sub%x.ne.5) stop 1 +if (sub%y.ne.7) stop 2 +!$omp end target + +!$omp target exit data map(release: sub) + +end function sub +end program p