Patch Detail
get:
Show a patch.
patch:
Update a patch.
put:
Update a patch.
GET /api/1.2/patches/2230140/?format=api
{ "id": 2230140, "url": "http://patchwork.ozlabs.org/api/1.2/patches/2230140/?format=api", "web_url": "http://patchwork.ozlabs.org/project/gcc/patch/fb7dc55f-b9f3-454a-9f6b-a0f295f6e570@baylibre.com/", "project": { "id": 17, "url": "http://patchwork.ozlabs.org/api/1.2/projects/17/?format=api", "name": "GNU Compiler Collection", "link_name": "gcc", "list_id": "gcc-patches.gcc.gnu.org", "list_email": "gcc-patches@gcc.gnu.org", "web_url": null, "scm_url": null, "webscm_url": null, "list_archive_url": "", "list_archive_url_format": "", "commit_url_format": "" }, "msgid": "<fb7dc55f-b9f3-454a-9f6b-a0f295f6e570@baylibre.com>", "list_archive_url": null, "date": "2026-04-29T10:29:06", "name": "[committed] OpenMP: Expand \"declare mapper\" mappers for target {enter,exit,} data directives", "commit_ref": null, "pull_url": null, "state": "new", "archived": false, "hash": "a13f67ef80becf18148d5a8e9f4676d3ed5ea55d", "submitter": { "id": 87873, "url": "http://patchwork.ozlabs.org/api/1.2/people/87873/?format=api", "name": "Tobias Burnus", "email": "tburnus@baylibre.com" }, "delegate": null, "mbox": "http://patchwork.ozlabs.org/project/gcc/patch/fb7dc55f-b9f3-454a-9f6b-a0f295f6e570@baylibre.com/mbox/", "series": [ { "id": 502036, "url": "http://patchwork.ozlabs.org/api/1.2/series/502036/?format=api", "web_url": "http://patchwork.ozlabs.org/project/gcc/list/?series=502036", "date": "2026-04-29T10:29:06", "name": "[committed] OpenMP: Expand \"declare mapper\" mappers for target {enter,exit,} data directives", "version": 1, "mbox": "http://patchwork.ozlabs.org/series/502036/mbox/" } ], "comments": "http://patchwork.ozlabs.org/api/patches/2230140/comments/", "check": "pending", "checks": "http://patchwork.ozlabs.org/api/patches/2230140/checks/", "tags": {}, "related": [], "headers": { "Return-Path": "<gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org>", "X-Original-To": [ "incoming@patchwork.ozlabs.org", "gcc-patches@gcc.gnu.org" ], "Delivered-To": [ "patchwork-incoming@legolas.ozlabs.org", "gcc-patches@gcc.gnu.org" ], "Authentication-Results": [ "legolas.ozlabs.org;\n\tdkim=fail reason=\"signature verification failed\" (2048-bit key;\n unprotected) header.d=baylibre-com.20251104.gappssmtp.com\n header.i=@baylibre-com.20251104.gappssmtp.com header.a=rsa-sha256\n header.s=20251104 header.b=XVNC83+R;\n\tdkim-atps=neutral", "legolas.ozlabs.org;\n spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org\n (client-ip=2620:52:6:3111::32; helo=vm01.sourceware.org;\n envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org;\n receiver=patchwork.ozlabs.org)", "sourceware.org;\n\tdkim=fail reason=\"signature verification failed\" (2048-bit key,\n unprotected) header.d=baylibre-com.20251104.gappssmtp.com\n header.i=@baylibre-com.20251104.gappssmtp.com header.a=rsa-sha256\n header.s=20251104 header.b=XVNC83+R", "sourceware.org;\n dmarc=none (p=none dis=none) header.from=baylibre.com", "sourceware.org; spf=pass smtp.mailfrom=baylibre.com", "server2.sourceware.org;\n arc=none smtp.remote-ip=209.85.128.48" ], "Received": [ "from vm01.sourceware.org (vm01.sourceware.org\n [IPv6:2620:52:6:3111::32])\n\t(using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits)\n\t key-exchange x25519 server-signature ECDSA (secp384r1) server-digest SHA384)\n\t(No client certificate requested)\n\tby legolas.ozlabs.org (Postfix) with ESMTPS id 4g5D8F4HZkz1yHZ\n\tfor <incoming@patchwork.ozlabs.org>; Wed, 29 Apr 2026 20:29:44 +1000 (AEST)", "from vm01.sourceware.org (localhost [127.0.0.1])\n\tby sourceware.org (Postfix) with ESMTP id EDCCA4BB1C1D\n\tfor <incoming@patchwork.ozlabs.org>; Wed, 29 Apr 2026 10:29:41 +0000 (GMT)", "from mail-wm1-f48.google.com (mail-wm1-f48.google.com\n [209.85.128.48])\n by sourceware.org (Postfix) with ESMTPS id B26294BB1C0D\n for <gcc-patches@gcc.gnu.org>; Wed, 29 Apr 2026 10:29:09 +0000 (GMT)", "by mail-wm1-f48.google.com with SMTP id\n 5b1f17b1804b1-48896199cbaso122232935e9.1\n for <gcc-patches@gcc.gnu.org>; Wed, 29 Apr 2026 03:29:09 -0700 (PDT)", "from ?IPV6:2001:16b8:26c8:8300:2e73:e549:e815:505e?\n (200116b826c883002e73e549e815505e.dip.versatel-1u1.de.\n [2001:16b8:26c8:8300:2e73:e549:e815:505e])\n by smtp.gmail.com with ESMTPSA id\n 5b1f17b1804b1-48a7c33f51asm40518755e9.0.2026.04.29.03.29.07\n (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128);\n Wed, 29 Apr 2026 03:29:07 -0700 (PDT)" ], "DKIM-Filter": [ "OpenDKIM Filter v2.11.0 sourceware.org EDCCA4BB1C1D", "OpenDKIM Filter v2.11.0 sourceware.org B26294BB1C0D" ], "DMARC-Filter": "OpenDMARC Filter v1.4.2 sourceware.org B26294BB1C0D", "ARC-Filter": "OpenARC Filter v1.0.0 sourceware.org B26294BB1C0D", "ARC-Seal": "i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1777458550; cv=none;\n b=ivgL8C6mkyUs+ZLhszhm8S+XDXru/Vb3yYWTMfbszMCBdX8z3YI69pAWGaiPxfVaYaFBtHGlp5geWY5zkICtHX4qPROvucePeZFt0b6FdMMl6ckLt//1Ib9v9T/+Iy2o+cf8TB7POq9cSXzilQmslteNmwUjpFycBauZXRahaxY=", "ARC-Message-Signature": "i=1; a=rsa-sha256; d=sourceware.org; s=key;\n t=1777458550; c=relaxed/simple;\n bh=H0vt4dYlDfhIzL949DJIrFMpMfJIxcof24Uiz+aRSpA=;\n h=DKIM-Signature:Message-ID:Date:MIME-Version:To:From:Subject;\n b=I5Nir23kWXfR+DM/ZNBlk0MsGSGp79NyQuq1mubJndbNrg3pIsU8Ojqo3mq6H5So7PsY3w9FhOu6ZVE3nwu7Cqj/yf1m5xlVqdONcfpNXzVAm7wP8KfTuq61KfFLP0zonBm+cv7AGY4uaFNtFreQVQm7JYJoGRB6Dwm+jNNXSdc=", "ARC-Authentication-Results": "i=1; server2.sourceware.org", "DKIM-Signature": "v=1; a=rsa-sha256; c=relaxed/relaxed;\n d=baylibre-com.20251104.gappssmtp.com; s=20251104; t=1777458548;\n x=1778063348;\n darn=gcc.gnu.org;\n h=subject:from:to:content-language:user-agent:mime-version:date\n :message-id:from:to:cc:subject:date:message-id:reply-to;\n bh=j6mr+DKF6Irj96vhLpw8P5kkbFnXT+peKMN/oGav/go=;\n b=XVNC83+RADBuNQNABwFZf5BeMpa5pTMcGogVP74Yxa+g8Iy+RoF+ddq/hVtuiWUx96\n qq7eeiB5ot6t+nssruDjXKzHsDY5NKbhZttF1wknUSIjsOzTUrkIFQu8lk9nnFleSLHM\n Cs8CtKphhs+eoPPmhV09CLQMyBfiKHZAGE9Jw2WTrRBV4FDkK6BI9/IPrsk6PrHTC1gM\n 1yNYg9OgEMlm95ePodmqj4ukoppXcZ48L6BTjK88CiChYoFHOrxaYDbYZUVcPnZVA9Bp\n MON1ZKBPBbU91/LVxavn//qIBEfR8QwZfRxLdLcEdCsyRoQfn2TUg0+GHHXZAxKzv4Vq\n Ksvg==", "X-Google-DKIM-Signature": "v=1; a=rsa-sha256; c=relaxed/relaxed;\n d=1e100.net; s=20251104; t=1777458548; x=1778063348;\n h=subject:from:to:content-language:user-agent:mime-version:date\n :message-id:x-gm-gg:x-gm-message-state:from:to:cc:subject:date\n :message-id:reply-to;\n bh=j6mr+DKF6Irj96vhLpw8P5kkbFnXT+peKMN/oGav/go=;\n b=RbHnRJCmoJvUZmOnWpKsG3AYPAu9KxOiGdt9WFUQ0KSAY9XfTETI5nxkbDaqHhjULM\n udNQjq/GmnKTXWVDuQd1i1jiEMDtSgApraZDgkA7vPwVg/0BV/q1hhFdf4sqdQ+sEuG2\n uPwhMyWPniN07kd9PKKwyk5+GSh6P33oBCn19kV4zUZnuuHPD+dudU2CZMLB6CukVTXu\n G+Bwxcch39qZR8r2O4RRKr24QFxFQoHpMzJkG1i7frPsOEqD+x7ZJ+yKr4ok/gB1eec+\n dwIJ9oHyewRN5wRrSW2rCR3cZGfwRJtVt/ckGH7A9WgwcaCsp3/UKY9MqcCdNdYQfmvc\n /n7A==", "X-Gm-Message-State": "AOJu0YwuKPebEzOAf2B8E1TMSZi9Mfkan64OyXge2TacRCnw5j0RCuqG\n zC6DnYPkiLdpqiNGp6MuJJQdwY1nzu9jkP0qL3zhzFCCd/XEtkHbbk45NtC6DWrfjWgjZmuQKQe\n AFhO2", "X-Gm-Gg": "AeBDiesNFrXzm/YcMO+qR9ul+zpVIhh96GImrLNTOEcUKP+wjKnyFPvYDZ/5HZKymIE\n z1LJs6PFS3it3joAvaaFrL1dA46UPMGNlH9aXBRgAWdc2BOd9ehJ8zfXyXsJJwNVzQ+2kR4Ez4r\n YJdlmB8Lg0luvnwj6WCxnLMKO+Uh/wuYTp4zrI1LjXEHaRsr+z8NL6aT/P0omZL8TrUJOLUmzw+\n +mjVUo2jvNpdOv8f8hhfk536El1KNHE2TU1YZeF7cIZ4OVDVtSu/VqzELIUj/eqFBrxNZjQ6+cw\n V98Us2fttgUNs0ejHb35xJxqI283cGTGMO6v6WAyW7T9bdYu75d9vuyFX0oFqdjm/ef1CIK92tg\n wLiWqVoN9SsLBLWBS1XmxuIx1EM+ue/Szx49hg+ByjxtYuRRj4fAslpKHyCzFP91FVydjULp5QH\n GaHrfe1ji0nNUX+9bqcDX75eRRC+jUNKF6Y5K1lQrNIpeeLDASVADrpssHCyXHgJTTKSCUZ5Snl\n 4dPsanDBugz8j0ABBGUvUboSZbCTw4dzF0T5mYJ0EmeuHjs8y7RDBmGyE5luuich5Eqy4y/wiC+\n pqg+U99cKg==", "X-Received": "by 2002:a05:600c:a311:b0:488:b749:8482 with SMTP id\n 5b1f17b1804b1-48a77ad59e2mr85393635e9.4.1777458548369;\n Wed, 29 Apr 2026 03:29:08 -0700 (PDT)", "Content-Type": "multipart/mixed; boundary=\"------------0pCtPgsI9OdKZfau7lfryabR\"", "Message-ID": "<fb7dc55f-b9f3-454a-9f6b-a0f295f6e570@baylibre.com>", "Date": "Wed, 29 Apr 2026 12:29:06 +0200", "MIME-Version": "1.0", "User-Agent": "Mozilla Thunderbird", "Content-Language": "en-US", "To": "gcc-patches <gcc-patches@gcc.gnu.org>", "From": "Tobias Burnus <tburnus@baylibre.com>", "Subject": "[committed] OpenMP: Expand \"declare mapper\" mappers for target\n {enter,exit,} data directives", "X-BeenThere": "gcc-patches@gcc.gnu.org", "X-Mailman-Version": "2.1.30", "Precedence": "list", "List-Id": "Gcc-patches mailing list <gcc-patches.gcc.gnu.org>", "List-Unsubscribe": "<https://gcc.gnu.org/mailman/options/gcc-patches>,\n <mailto:gcc-patches-request@gcc.gnu.org?subject=unsubscribe>", "List-Archive": "<https://gcc.gnu.org/pipermail/gcc-patches/>", "List-Post": "<mailto:gcc-patches@gcc.gnu.org>", "List-Help": "<mailto:gcc-patches-request@gcc.gnu.org?subject=help>", "List-Subscribe": "<https://gcc.gnu.org/mailman/listinfo/gcc-patches>,\n <mailto:gcc-patches-request@gcc.gnu.org?subject=subscribe>", "Errors-To": "gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org" }, "content": "This is the long overdue patch by Julian that extends the\nOpenMP mapper to support 'target ... data'.\n\nThis is based on his (rebased) commit to OG14 (branch: devel/omp/gcc-14)\ncommit 1188d7e329d329d0e92534fcc8c739450d3d6eae\nhttps://gcc.gnu.org/g:1188d7e329d\n\nThe original commit was for OG13 (2427905b40c) with\nAuthorDate: Mon Jul 3 19:07:11 2023 +0000\nCommitDate: Thu Jul 6 15:59:55 2023 +0000\n\nNOTE: Contrary to the original patch, I excluded the Fortran part\nas the Fortran support for 'declare mapper' still has to land.\nOtherwise, it is unchanged and looked good to me.\n\nCommitted as Rev. r17-202-g3f8c7483112807\n\nTobias\n\nPS: Due to the long delays, other changes and patch submissions\nvs. OG... branch commits + rebasing, the patch history it a bit\ncomplex. :-/", "diff": "commit 3f8c7483112807f8f2c48852f9cbad9a4d5b1da4\nAuthor: Julian Brown <julian@codesourcery.com>\nAuthorDate: Wed Apr 29 12:12:13 2026 +0200\nCommit: Tobias Burnus <tburnus@baylibre.com>\nCommitDate: Wed Apr 29 12:12:13 2026 +0200\n\n OpenMP: Expand \"declare mapper\" mappers for target {enter,exit,} data directives\n \n This patch allows 'declare mapper' mappers to be used on 'omp target\n data', 'omp target enter data' and 'omp target exit data' directives.\n For each of these, only explicit mappings are supported, unlike for\n 'omp target' directives where implicit uses of variables inside an\n offload region might trigger mappers also.\n \n Add support for C and C++.\n \n The patch also adjusts 'map kind decay' to match OpenMP 5.2 semantics,\n which is particularly important with regard to 'exit data' operations.\n \n gcc/c-family/\n * c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA,\n C_ORT_OMP_EXIT_DATA.\n (c_omp_instantiate_mappers): Add region type parameter.\n * c-omp.cc (omp_split_map_kind, omp_join_map_kind,\n omp_map_decayed_kind): New functions.\n (omp_instantiate_mapper): Add ORT parameter. Implement map kind decay\n for instantiated mapper clauses.\n (c_omp_instantiate_mappers): Add ORT parameter, pass to\n omp_instantiate_mapper.\n \n gcc/c/\n * c-parser.cc (c_parser_omp_target_data): Instantiate mappers for\n 'omp target data'.\n (c_parser_omp_target_enter_data): Instantiate mappers for 'omp target\n enter data'.\n (c_parser_omp_target_exit_data): Instantiate mappers for 'omp target\n exit data'.\n (c_parser_omp_target): Add c_omp_region_type argument to\n c_omp_instantiate_mappers call.\n * c-tree.h (c_omp_instantiate_mappers): Remove spurious prototype.\n \n gcc/cp/\n * parser.cc (cp_parser_omp_target_data): Instantiate mappers for 'omp\n target data'.\n (cp_parser_omp_target_enter_data): Instantiate mappers for 'omp target\n enter data'.\n (cp_parser_omp_target_exit_data): Instantiate mappers for 'omp target\n exit data'.\n (cp_parser_omp_target): Add c_omp_region_type argument to\n c_omp_instantiate_mappers call.\n * pt.cc (tsubst_omp_clauses): Instantiate mappers for OMP regions other\n than just C_ORT_OMP_TARGET.\n (tsubst_expr): Update call to tsubst_omp_clauses for OMP_TARGET_UPDATE,\n OMP_TARGET_ENTER_DATA, OMP_TARGET_EXIT_DATA stanza.\n * semantics.cc (cxx_omp_map_array_section): Avoid calling\n build_array_ref for non-array/non-pointer bases (error reported\n already).\n \n gcc/testsuite/\n * c-c++-common/gomp/declare-mapper-15.c: New test.\n * c-c++-common/gomp/declare-mapper-16.c: New test.\n * g++.dg/gomp/declare-mapper-1.C: Adjust expected scan output.\n---\n gcc/c-family/c-common.h | 2 +-\n gcc/c-family/c-omp.cc | 193 ++++++++++++++++++++-\n gcc/c/c-parser.cc | 11 +-\n gcc/c/c-tree.h | 1 -\n gcc/cp/parser.cc | 15 +-\n gcc/cp/pt.cc | 8 +-\n gcc/cp/semantics.cc | 5 +-\n .../c-c++-common/gomp/declare-mapper-15.c | 59 +++++++\n .../c-c++-common/gomp/declare-mapper-16.c | 39 +++++\n gcc/testsuite/g++.dg/gomp/declare-mapper-1.C | 2 +-\n 10 files changed, 313 insertions(+), 22 deletions(-)\n\ndiff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h\nindex 89517e2a80c..a4cd6786109 100644\n--- a/gcc/c-family/c-common.h\n+++ b/gcc/c-family/c-common.h\n@@ -1357,7 +1357,7 @@ extern void c_omp_mark_declare_variant (location_t, tree, tree);\n extern void c_omp_adjust_map_clauses (tree, bool);\n template<typename T> struct omp_mapper_list;\n extern void c_omp_find_nested_mappers (struct omp_mapper_list<tree> *, tree);\n-extern tree c_omp_instantiate_mappers (tree);\n+extern tree c_omp_instantiate_mappers (tree, enum c_omp_region_type);\n \n namespace omp_addr_tokenizer { struct omp_addr_token; }\n typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;\ndiff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc\nindex bb121ede06b..f75f0788ac6 100644\n--- a/gcc/c-family/c-omp.cc\n+++ b/gcc/c-family/c-omp.cc\n@@ -4384,13 +4384,189 @@ remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)\n return NULL_TREE;\n }\n \n+static enum gomp_map_kind\n+omp_split_map_kind (enum gomp_map_kind op, bool *force_p, bool *always_p,\n+\t\t bool *present_p)\n+{\n+ *force_p = *always_p = *present_p = false;\n+\n+ switch (op)\n+ {\n+ case GOMP_MAP_FORCE_ALLOC:\n+ case GOMP_MAP_FORCE_TO:\n+ case GOMP_MAP_FORCE_FROM:\n+ case GOMP_MAP_FORCE_TOFROM:\n+ case GOMP_MAP_FORCE_PRESENT:\n+ *force_p = true;\n+ break;\n+ case GOMP_MAP_ALWAYS_TO:\n+ case GOMP_MAP_ALWAYS_FROM:\n+ case GOMP_MAP_ALWAYS_TOFROM:\n+ *always_p = true;\n+ break;\n+ case GOMP_MAP_ALWAYS_PRESENT_TO:\n+ case GOMP_MAP_ALWAYS_PRESENT_FROM:\n+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:\n+ *always_p = true;\n+ /* Fallthrough. */\n+ case GOMP_MAP_PRESENT_ALLOC:\n+ case GOMP_MAP_PRESENT_TO:\n+ case GOMP_MAP_PRESENT_FROM:\n+ case GOMP_MAP_PRESENT_TOFROM:\n+ *present_p = true;\n+ break;\n+ default:\n+ ;\n+ }\n+\n+ switch (op)\n+ {\n+ case GOMP_MAP_ALLOC:\n+ case GOMP_MAP_FORCE_ALLOC:\n+ case GOMP_MAP_PRESENT_ALLOC:\n+ return GOMP_MAP_ALLOC;\n+ case GOMP_MAP_TO:\n+ case GOMP_MAP_FORCE_TO:\n+ case GOMP_MAP_ALWAYS_TO:\n+ case GOMP_MAP_PRESENT_TO:\n+ case GOMP_MAP_ALWAYS_PRESENT_TO:\n+ return GOMP_MAP_TO;\n+ case GOMP_MAP_FROM:\n+ case GOMP_MAP_FORCE_FROM:\n+ case GOMP_MAP_ALWAYS_FROM:\n+ case GOMP_MAP_PRESENT_FROM:\n+ case GOMP_MAP_ALWAYS_PRESENT_FROM:\n+ return GOMP_MAP_FROM;\n+ case GOMP_MAP_TOFROM:\n+ case GOMP_MAP_FORCE_TOFROM:\n+ case GOMP_MAP_ALWAYS_TOFROM:\n+ case GOMP_MAP_PRESENT_TOFROM:\n+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:\n+ return GOMP_MAP_TOFROM;\n+ default:\n+ ;\n+ }\n+\n+ return op;\n+}\n+\n+static enum gomp_map_kind\n+omp_join_map_kind (enum gomp_map_kind op, bool force_p, bool always_p,\n+\t\t bool present_p)\n+{\n+ gcc_assert (!force_p || !(always_p || present_p));\n+\n+ switch (op)\n+ {\n+ case GOMP_MAP_ALLOC:\n+ if (force_p)\n+\treturn GOMP_MAP_FORCE_ALLOC;\n+ else if (present_p)\n+\treturn GOMP_MAP_PRESENT_ALLOC;\n+ break;\n+\n+ case GOMP_MAP_TO:\n+ if (force_p)\n+\treturn GOMP_MAP_FORCE_TO;\n+ else if (always_p && present_p)\n+\treturn GOMP_MAP_ALWAYS_PRESENT_TO;\n+ else if (always_p)\n+\treturn GOMP_MAP_ALWAYS_TO;\n+ else if (present_p)\n+\treturn GOMP_MAP_PRESENT_TO;\n+ break;\n+\n+ case GOMP_MAP_FROM:\n+ if (force_p)\n+\treturn GOMP_MAP_FORCE_FROM;\n+ else if (always_p && present_p)\n+\treturn GOMP_MAP_ALWAYS_PRESENT_FROM;\n+ else if (always_p)\n+\treturn GOMP_MAP_ALWAYS_FROM;\n+ else if (present_p)\n+\treturn GOMP_MAP_PRESENT_FROM;\n+ break;\n+\n+ case GOMP_MAP_TOFROM:\n+ if (force_p)\n+\treturn GOMP_MAP_FORCE_TOFROM;\n+ else if (always_p && present_p)\n+\treturn GOMP_MAP_ALWAYS_PRESENT_TOFROM;\n+ else if (always_p)\n+\treturn GOMP_MAP_ALWAYS_TOFROM;\n+ else if (present_p)\n+\treturn GOMP_MAP_PRESENT_TOFROM;\n+ break;\n+\n+ default:\n+ ;\n+ }\n+\n+ return op;\n+}\n+\n+/* Map kind decay (OpenMP 5.2, 5.8.8 \"declare mapper Directive\"). Return the\n+ map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS\n+ specified on the clause that invokes the mapper. See also\n+ fortran/trans-openmp.cc:omp_map_decayed_kind. */\n+\n+static enum gomp_map_kind\n+omp_map_decayed_kind (enum gomp_map_kind mapper_kind,\n+\t\t enum gomp_map_kind invoked_as, bool exit_p)\n+{\n+ if (invoked_as == GOMP_MAP_RELEASE || invoked_as == GOMP_MAP_DELETE)\n+ return invoked_as;\n+\n+ bool force_p, always_p, present_p;\n+\n+ invoked_as = omp_split_map_kind (invoked_as, &force_p, &always_p, &present_p);\n+ gomp_map_kind decay_to;\n+\n+ switch (mapper_kind)\n+ {\n+ case GOMP_MAP_ALLOC:\n+ if (exit_p && invoked_as == GOMP_MAP_FROM)\n+\tdecay_to = GOMP_MAP_RELEASE;\n+ else\n+\tdecay_to = GOMP_MAP_ALLOC;\n+ break;\n+\n+ case GOMP_MAP_TO:\n+ if (invoked_as == GOMP_MAP_FROM)\n+\tdecay_to = exit_p ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;\n+ else if (invoked_as == GOMP_MAP_ALLOC)\n+\tdecay_to = GOMP_MAP_ALLOC;\n+ else\n+\tdecay_to = GOMP_MAP_TO;\n+ break;\n+\n+ case GOMP_MAP_FROM:\n+ if (invoked_as == GOMP_MAP_ALLOC || invoked_as == GOMP_MAP_TO)\n+\tdecay_to = GOMP_MAP_ALLOC;\n+ else\n+\tdecay_to = GOMP_MAP_FROM;\n+ break;\n+\n+ case GOMP_MAP_TOFROM:\n+ case GOMP_MAP_UNSET:\n+ decay_to = invoked_as;\n+ break;\n+\n+ default:\n+ gcc_unreachable ();\n+ }\n+\n+ return omp_join_map_kind (decay_to, force_p, always_p, present_p);\n+}\n+\n /* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to\n OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in\n the mapper declaration. */\n \n static tree *\n omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,\n-\t\t\tenum gomp_map_kind outer_kind)\n+\t\t\tenum gomp_map_kind outer_kind,\n+\t\t\tenum c_omp_region_type ort)\n {\n tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);\n tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);\n@@ -4447,8 +4623,10 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,\n \n walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);\n \n- if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)\n-\tOMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);\n+ enum gomp_map_kind decayed_kind\n+\t= omp_map_decayed_kind (clause_kind, outer_kind,\n+\t\t\t\t(ort & C_ORT_EXIT_DATA) != 0);\n+ OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_kind);\n \n type = TYPE_MAIN_VARIANT (type);\n \n@@ -4465,11 +4643,8 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,\n \t = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);\n \t if (nested_mapper != mapper)\n \t {\n-\t if (clause_kind == GOMP_MAP_UNSET)\n-\t\tclause_kind = outer_kind;\n-\n \t outlist = omp_instantiate_mapper (outlist, nested_mapper,\n-\t\t\t\t\t\tt, clause_kind);\n+\t\t\t\t\t\tt, outer_kind, ort);\n \t continue;\n \t }\n \t}\n@@ -4491,7 +4666,7 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,\n visible in the current parsing context. */\n \n tree\n-c_omp_instantiate_mappers (tree clauses)\n+c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort)\n {\n tree c, *pc, mapper_name = NULL_TREE;\n \n@@ -4564,7 +4739,7 @@ c_omp_instantiate_mappers (tree clauses)\n \t {\n \t\ttree mapper\n \t\t = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);\n-\t\tpc = omp_instantiate_mapper (pc, mapper, t, kind);\n+\t\tpc = omp_instantiate_mapper (pc, mapper, t, kind, ort);\n \t\tusing_mapper = true;\n \t }\n \t else if (mapper_name)\ndiff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc\nindex 0b8b2387109..9ddb449d710 100644\n--- a/gcc/c/c-parser.cc\n+++ b/gcc/c/c-parser.cc\n@@ -26928,7 +26928,9 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)\n \n tree clauses\n = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,\n-\t\t\t\t\"#pragma omp target data\");\n+\t\t\t\t\"#pragma omp target data\", false);\n+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);\n+ clauses = c_finish_omp_clauses (clauses, C_ORT_OMP);\n c_omp_adjust_map_clauses (clauses, false);\n int map_seen = 0;\n for (tree *pc = &clauses; *pc;)\n@@ -27086,7 +27088,9 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,\n \n tree clauses\n = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,\n-\t\t\t\t\"#pragma omp target enter data\");\n+\t\t\t\t\"#pragma omp target enter data\", false);\n+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);\n+ clauses = c_finish_omp_clauses (clauses, C_ORT_OMP);\n c_omp_adjust_map_clauses (clauses, false);\n int map_seen = 0;\n for (tree *pc = &clauses; *pc;)\n@@ -27197,6 +27201,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,\n tree clauses\n = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,\n \t\t\t\t\"#pragma omp target exit data\", false);\n+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA);\n clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);\n c_omp_adjust_map_clauses (clauses, false);\n int map_seen = 0;\n@@ -27456,7 +27461,7 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)\n \tOMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);\n \tOMP_CLAUSE_CHAIN (c) = nc;\n }\n- clauses = c_omp_instantiate_mappers (clauses);\n+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_TARGET);\n clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_TARGET);\n c_omp_adjust_map_clauses (clauses, true);\n \ndiff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h\nindex e151095c3a1..0b863cb7d1d 100644\n--- a/gcc/c/c-tree.h\n+++ b/gcc/c/c-tree.h\n@@ -1009,7 +1009,6 @@ extern tree c_check_omp_declare_reduction_r (tree *, int *, void *);\n extern tree c_omp_mapper_id (tree);\n extern tree c_omp_mapper_decl (tree);\n extern void c_omp_scan_mapper_bindings (location_t, tree *, tree);\n-extern tree c_omp_instantiate_mappers (tree);\n extern bool c_check_in_current_scope (tree);\n extern void c_pushtag (location_t, tree, tree);\n extern void c_bind (location_t, tree, bool);\ndiff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc\nindex dc67cfd9f7c..3723aed1e69 100644\n--- a/gcc/cp/parser.cc\n+++ b/gcc/cp/parser.cc\n@@ -51556,7 +51556,10 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)\n \n tree clauses\n = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,\n-\t\t\t\t \"#pragma omp target data\", pragma_tok);\n+\t\t\t\t \"#pragma omp target data\", pragma_tok, false);\n+ if (!processing_template_decl)\n+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);\n+ clauses = finish_omp_clauses (clauses, C_ORT_OMP);\n c_omp_adjust_map_clauses (clauses, false);\n int map_seen = 0;\n for (tree *pc = &clauses; *pc;)\n@@ -51671,7 +51674,11 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,\n \n tree clauses\n = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,\n-\t\t\t\t \"#pragma omp target enter data\", pragma_tok);\n+\t\t\t\t \"#pragma omp target enter data\", pragma_tok,\n+\t\t\t\t false);\n+ if (!processing_template_decl)\n+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP);\n+ clauses = finish_omp_clauses (clauses, C_ORT_OMP);\n c_omp_adjust_map_clauses (clauses, false);\n int map_seen = 0;\n for (tree *pc = &clauses; *pc;)\n@@ -51788,6 +51795,8 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,\n = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,\n \t\t\t\t \"#pragma omp target exit data\", pragma_tok,\n \t\t\t\t false);\n+ if (!processing_template_decl)\n+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA);\n clauses = finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);\n c_omp_adjust_map_clauses (clauses, false);\n int map_seen = 0;\n@@ -52084,7 +52093,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,\n \tOMP_CLAUSE_CHAIN (c) = nc;\n }\n if (!processing_template_decl)\n- clauses = c_omp_instantiate_mappers (clauses);\n+ clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_TARGET);\n clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);\n \n c_omp_adjust_map_clauses (clauses, true);\ndiff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc\nindex 05ecda2a0d8..9a5eb7ef9b0 100644\n--- a/gcc/cp/pt.cc\n+++ b/gcc/cp/pt.cc\n@@ -18936,8 +18936,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,\n new_clauses = nreverse (new_clauses);\n if (ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_OMP_DECLARE_MAPPER)\n {\n- if (ort == C_ORT_OMP_TARGET)\n-\tnew_clauses = c_omp_instantiate_mappers (new_clauses);\n+ if (ort & C_ORT_OMP)\n+\tnew_clauses = c_omp_instantiate_mappers (new_clauses, ort);\n new_clauses = finish_omp_clauses (new_clauses, ort);\n if (linear_no_step)\n \tfor (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))\n@@ -20534,7 +20534,9 @@ tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl)\n case OMP_TARGET_UPDATE:\n case OMP_TARGET_ENTER_DATA:\n case OMP_TARGET_EXIT_DATA:\n- tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), C_ORT_OMP, args,\n+ tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t),\n+\t\t\t\t(TREE_CODE (t) == OMP_TARGET_EXIT_DATA\n+\t\t\t\t ? C_ORT_OMP_EXIT_DATA : C_ORT_OMP), args,\n \t\t\t\tcomplain, in_decl);\n t = copy_node (t);\n OMP_STANDALONE_CLAUSES (t) = tmp;\ndiff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc\nindex 92ec4836a6f..1cb2b509fa2 100644\n--- a/gcc/cp/semantics.cc\n+++ b/gcc/cp/semantics.cc\n@@ -6945,7 +6945,10 @@ cxx_omp_map_array_section (location_t loc, tree t)\n if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)\n \tt = convert_from_reference (t);\n \n- t = build_array_ref (loc, t, low);\n+ if (TYPE_PTR_P (TREE_TYPE (t)))\n+\tt = build_array_ref (loc, t, low);\n+ else\n+\tt = error_mark_node;\n }\n \n return t;\ndiff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c\nnew file mode 100644\nindex 00000000000..ecda2e5ebd1\n--- /dev/null\n+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c\n@@ -0,0 +1,59 @@\n+/* { dg-do compile } */\n+/* { dg-options \"-fopenmp -fdump-tree-gimple\" } */\n+\n+typedef struct {\n+ int a, b, c, d;\n+} S;\n+\n+int main ()\n+{\n+ S s;\n+ #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \\\n+\t\t\t\t map(from: x.c) map(tofrom: x.d)\n+\n+ #pragma omp target enter data map(to: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 4\\]\\) map\\(alloc:s\\.a \\[len: [0-9]+\\]\\) map\\(to:s\\.b \\[len: [0-9]+\\]\\) map\\(alloc:s\\.c \\[len: [0-9]+\\]\\) map\\(to:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+ #pragma omp target exit data map(from: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(release:s\\.a \\[len: 4\\]\\) map\\(release:s\\.b \\[len: [0-9]+\\]\\) map\\(from:s\\.c \\[len: [0-9]+\\]\\) map\\(from:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+\n+ #pragma omp target enter data map(alloc: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 4\\]\\) map\\(alloc:s\\.a \\[len: [0-9]+\\]\\) map\\(alloc:s\\.b \\[len: [0-9]+\\]\\) map\\(alloc:s\\.c \\[len: [0-9]+\\]\\) map\\(alloc:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+ #pragma omp target exit data map(release: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(release:s\\.a \\[len: [0-9]+\\]\\) map\\(release:s\\.b \\[len: [0-9]+\\]\\) map\\(release:s\\.c \\[len: [0-9]+\\]\\) map\\(release:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+\n+ #pragma omp target enter data map(present, to: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 4\\]\\) map\\(force_present:s\\.a \\[len: [0-9]+\\]\\) map\\(force_present:s\\.b \\[len: [0-9]+\\]\\) map\\(force_present:s\\.c \\[len: [0-9]+\\]\\) map\\(force_present:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+ #pragma omp target exit data map(present, from: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(release:s\\.a \\[len: [0-9]+\\]\\) map\\(release:s\\.b \\[len: [0-9]+\\]\\) map\\(force_present:s\\.c \\[len: [0-9]+\\]\\) map\\(force_present:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+\n+ #pragma omp target enter data map(always, to: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 4\\]\\) map\\(alloc:s\\.a \\[len: [0-9]+\\]\\) map\\(always,to:s\\.b \\[len: [0-9]+\\]\\) map\\(alloc:s\\.c \\[len: [0-9]+\\]\\) map\\(always,to:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+ #pragma omp target exit data map(always, from: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(release:s\\.a \\[len: [0-9]+\\]\\) map\\(release:s\\.b \\[len: [0-9]+\\]\\) map\\(always,from:s\\.c \\[len: [0-9]+\\]\\) map\\(always,from:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+\n+ #pragma omp target enter data map(always, present, to: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 4\\]\\) map\\(force_present:s\\.a \\[len: [0-9]+\\]\\) map\\(always,present,to:s\\.b \\[len: [0-9]+\\]\\) map\\(force_present:s\\.c \\[len: [0-9]+\\]\\) map\\(always,present,to:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+ #pragma omp target exit data map(always, present, from: s)\n+\n+ /* { dg-final { scan-tree-dump-times {map\\(release:s\\.a \\[len: [0-9]+\\]\\) map\\(release:s\\.b \\[len: [0-9]+\\]\\) map\\(always,present,from:s\\.c \\[len: [0-9]+\\]\\) map\\(always,present,from:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+\n+ return 0;\n+}\ndiff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c\nnew file mode 100644\nindex 00000000000..20383cc2d69\n--- /dev/null\n+++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c\n@@ -0,0 +1,39 @@\n+/* { dg-do compile } */\n+/* { dg-options \"-fopenmp -fdump-tree-gimple\" } */\n+\n+typedef struct {\n+ int a, b, c, d;\n+} S;\n+\n+int main ()\n+{\n+ S s = { 0, 0, 0, 0 };\n+ #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \\\n+\t\t\t\t map(from: x.c) map(tofrom: x.d)\n+\n+ #pragma omp target data map(s)\n+ /* { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 4\\]\\) map\\(alloc:s\\.a \\[len: [0-9]+\\]\\) map\\(to:s\\.b \\[len: [0-9]+\\]\\) map\\(from:s\\.c \\[len: [0-9]+\\]\\) map\\(tofrom:s\\.d \\[len: [0-9]+\\]\\)} 3 \"gimple\" } } */\n+ {\n+ #pragma omp target\n+ {\n+ s.a++;\n+ s.b++;\n+ s.c++;\n+ s.d++;\n+ }\n+ }\n+\n+ #pragma omp target data map(alloc: s)\n+ /* { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 4\\]\\) map\\(alloc:s\\.a \\[len: [0-9]+\\]\\) map\\(alloc:s\\.b \\[len: [0-9]+\\]\\) map\\(alloc:s\\.c \\[len: [0-9]+\\]\\) map\\(alloc:s\\.d \\[len: [0-9]+\\]\\)} 1 \"gimple\" } } */\n+ {\n+ #pragma omp target\n+ {\n+ s.a++;\n+ s.b++;\n+ s.c++;\n+ s.d++;\n+ }\n+ }\n+\n+ return 0;\n+}\ndiff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C\nindex 2523d02b4a6..2f2dd219bcb 100644\n--- a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C\n+++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C\n@@ -55,4 +55,4 @@ int main (int argc, char *argv[])\n }\n \n // { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 2\\]\\) map\\(alloc:s\\.ptr \\[len: [0-9]+\\]\\) map\\(tofrom:s\\.size \\[len: [0-9]+\\]\\) map\\(tofrom:\\*_[0-9]+ \\[len: _[0-9]+\\]\\) map\\(attach:s\\.ptr \\[bias: 0\\]\\)} 4 \"gimple\" } }\n-// { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 2\\]\\) map\\(alloc:s\\.ptr \\[len: [0-9]+\\]\\) map\\(to:s\\.size \\[len: [0-9]+\\]\\) map\\(alloc:\\*_[0-9]+ \\[len: _[0-9]+\\]\\) map\\(attach:s\\.ptr \\[bias: 0\\]\\)} 1 \"gimple\" } }\n+// { dg-final { scan-tree-dump-times {map\\(struct:s \\[len: 2\\]\\) map\\(alloc:s\\.ptr \\[len: [0-9]+\\]\\) map\\(alloc:s\\.size \\[len: [0-9]+\\]\\) map\\(alloc:\\*_[0-9]+ \\[len: _[0-9]+\\]\\) map\\(attach:s\\.ptr \\[bias: 0\\]\\)} 1 \"gimple\" } }\n", "prefixes": [ "committed" ] }