Patch Detail

HTTP 200 OK
Allow: GET, PUT, PATCH, HEAD, OPTIONS
Content-Type: application/json
Vary: Accept

{
    "id": 2218664,
    "url": "http://patchwork.ozlabs.org/api/patches/2218664/?format=api",
    "web_url": "http://patchwork.ozlabs.org/project/gcc/patch/4596ca92-75bc-4085-982e-3305002265cf@baylibre.com/",
    "project": {
        "id": 17,
        "url": "http://patchwork.ozlabs.org/api/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": "<4596ca92-75bc-4085-982e-3305002265cf@baylibre.com>",
    "list_archive_url": null,
    "date": "2026-04-01T15:58:47",
    "name": "[OpenMP,v2] uses_allocators support",
    "commit_ref": null,
    "pull_url": null,
    "state": "new",
    "archived": false,
    "hash": "cd7cb79edfa4036083e3632df5d87eded6fd9378",
    "submitter": {
        "id": 87913,
        "url": "http://patchwork.ozlabs.org/api/people/87913/?format=api",
        "name": "Chung-Lin Tang",
        "email": "cltang@baylibre.com"
    },
    "delegate": null,
    "mbox": "http://patchwork.ozlabs.org/project/gcc/patch/4596ca92-75bc-4085-982e-3305002265cf@baylibre.com/mbox/",
    "series": [
        {
            "id": 498355,
            "url": "http://patchwork.ozlabs.org/api/series/498355/?format=api",
            "web_url": "http://patchwork.ozlabs.org/project/gcc/list/?series=498355",
            "date": "2026-04-01T15:58:47",
            "name": "[OpenMP,v2] uses_allocators support",
            "version": 2,
            "mbox": "http://patchwork.ozlabs.org/series/498355/mbox/"
        }
    ],
    "comments": "http://patchwork.ozlabs.org/api/patches/2218664/comments/",
    "check": "pending",
    "checks": "http://patchwork.ozlabs.org/api/patches/2218664/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=pass (2048-bit key;\n unprotected) header.d=baylibre-com.20230601.gappssmtp.com\n header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256\n header.s=20230601 header.b=DeEYcRpm;\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=pass (2048-bit key,\n unprotected) header.d=baylibre-com.20230601.gappssmtp.com\n header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256\n header.s=20230601 header.b=DeEYcRpm",
            "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.215.169"
        ],
        "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 4fm8nK3fsqz1yFv\n\tfor <incoming@patchwork.ozlabs.org>; Thu, 02 Apr 2026 02:59:33 +1100 (AEDT)",
            "from vm01.sourceware.org (localhost [127.0.0.1])\n\tby sourceware.org (Postfix) with ESMTP id 708D04BA23E5\n\tfor <incoming@patchwork.ozlabs.org>; Wed,  1 Apr 2026 15:59:31 +0000 (GMT)",
            "from mail-pg1-f169.google.com (mail-pg1-f169.google.com\n [209.85.215.169])\n by sourceware.org (Postfix) with ESMTPS id 5FE8C4BA543C\n for <gcc-patches@gcc.gnu.org>; Wed,  1 Apr 2026 15:58:53 +0000 (GMT)",
            "by mail-pg1-f169.google.com with SMTP id\n 41be03b00d2f7-c6e2355739dso2986404a12.2\n for <gcc-patches@gcc.gnu.org>; Wed, 01 Apr 2026 08:58:53 -0700 (PDT)",
            "from [192.168.50.226] (112-104-14-82.adsl.dynamic.seed.net.tw.\n [112.104.14.82]) by smtp.gmail.com with ESMTPSA id\n 41be03b00d2f7-c76c657e5f3sm260943a12.23.2026.04.01.08.58.49\n (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128);\n Wed, 01 Apr 2026 08:58:51 -0700 (PDT)"
        ],
        "DKIM-Filter": [
            "OpenDKIM Filter v2.11.0 sourceware.org 708D04BA23E5",
            "OpenDKIM Filter v2.11.0 sourceware.org 5FE8C4BA543C"
        ],
        "DMARC-Filter": "OpenDMARC Filter v1.4.2 sourceware.org 5FE8C4BA543C",
        "ARC-Filter": "OpenARC Filter v1.0.0 sourceware.org 5FE8C4BA543C",
        "ARC-Seal": "i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1775059133; cv=none;\n b=whr2QcTiM0aMGoCwGyhrYdXrIu2/Zn0KpgaTG8KyJlrmMCFHncivjHC+4L3PVTIQJLGfzl5kOZcgi19OzvENUv/Qo4NoOn7Puf+Z02cVDIqNwbp+2AdXWVxD5+msUbzu9UGDFB+BYkav0VSpikl+jgsXgK/OVYm8BW/U0t1aUNc=",
        "ARC-Message-Signature": "i=1; a=rsa-sha256; d=sourceware.org; s=key;\n t=1775059133; c=relaxed/simple;\n bh=WX3xE+2T4il5mxeQPCTIxLRHncxW37pYdhrdLg+5uUk=;\n h=DKIM-Signature:Message-ID:Date:MIME-Version:Subject:From:To;\n b=drr/K2frOzBSwIBGIys4rL7VBXmOZehKOwlkehSDMFmfeGOomloJKnl/8EZmPZcGvJBJg8Yru4lbJUbTm5waScHReUkAEFk9+kqp6vANf+lf44pKzhngbFD8J8G9Dmln2QHpY+yCZvJg6Qbflhhvi+AtC+/0w30LUDOkd0pnrZQ=",
        "ARC-Authentication-Results": "i=1; server2.sourceware.org",
        "DKIM-Signature": "v=1; a=rsa-sha256; c=relaxed/relaxed;\n d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1775059132;\n x=1775663932;\n darn=gcc.gnu.org;\n h=in-reply-to:content-language:references:to:from:subject:user-agent\n :mime-version:date:message-id:from:to:cc:subject:date:message-id\n :reply-to; bh=eLDj8604EdE0u7iomkCtvSpbotZm9XbdvpErIlIAth8=;\n b=DeEYcRpmOOy3vZAMZl58OMXFsDSbT9nY9p+eV5cI9XZ3qtbg8IqQSCdWnngpN3W4sE\n rFdVfHCZvQ1uxY76I8DfPB0vliJzbsJlT0IfyqZ2o+4rZkO0G/1K2GCvaALQCpIdiPNM\n /CKANUoL3AZnQ/aXIFOCC/rA5sxYy9jNIzUQJSOIWrVct5H76AxnXZSCDTGp6rbckWN1\n IWvPve0kDoW3MG+ck5OvgLFCPBixAwKF3V/E+oP1u/NJg384oigR2CYPNxejDK+e86qk\n WDNbY4PtStLuLlxfvdWJSKbNulqP5zKBsOV5lwmOY1WQYkN5bRKXhJY0wD2LdHhcL6DQ\n cXIQ==",
        "X-Google-DKIM-Signature": "v=1; a=rsa-sha256; c=relaxed/relaxed;\n d=1e100.net; s=20251104; t=1775059132; x=1775663932;\n h=in-reply-to:content-language:references:to:from:subject:user-agent\n :mime-version:date:message-id:x-gm-gg:x-gm-message-state:from:to:cc\n :subject:date:message-id:reply-to;\n bh=eLDj8604EdE0u7iomkCtvSpbotZm9XbdvpErIlIAth8=;\n b=rFlgGoLWWIQuvn+Fr6rTW2BCLjL3ZwNpslrtEelr98yK+oEQDrjQRLvINTDG0NEgzC\n bxW5MvxfF2GMJ4xv3YIKe0KTfVBPxadTr4r2zeB0z1IjFeTRJkdvxs3fkXG1/I8/ePVO\n cKUXkOiHY7wRKaVFji+rl0/atzY+LY2DaJc4Ryr3XLOZNLFf/GeY2z9Dy2cG2myjOqQe\n O9EqUcCPP4KXzjUvwnVv01bgh7mm191H5+FS66mspa6BZPLUqrEwmG/Bavsa5KYg+nsc\n AQXDpwN+Bym60NaBD2/ZsGX8wolI2KXzfBoxXmQuU2vXrcv5a5rjU83UgP6nahLbekBl\n ne5Q==",
        "X-Gm-Message-State": "AOJu0Yx1fdmLvPBnRByKimvQep6EEaul5gAo3O63aRaCjBB8o7c7e6JG\n 0sF59ZrQohiWtqZpYgZW+hSKH77Wbyhbqp4pngezYtl2vwhQsHjbOnHhmOsse2z57LubCeW+AuA\n 2vZyLt1vwtA==",
        "X-Gm-Gg": "ATEYQzxUnxBIYHuCAWSkrLsyHZ0l9K8muv+EOoOe1//NG6eQgtHnshUPMPj/Z8fzJtj\n YIkcqIhQ9JzJ56BVgBn5Y9YU1kqaA7UU33jU+qtsbdKrFF0V6b2JSLX1CKFo4w0McdKSLTF5qvX\n S5AUqfyl3v/UbeC0kMCjhM2WkoTeM5qLyRjNATpGiwriwJZ8Uby48l+K35pQ9LooNeLgCkdX1Bf\n od3gFGmdwZFbPTBxywjToNqOnAISy6XdYqlXeg5R8AnSe0EjrRwWgtKhK9v+Yr2UOPexxBxu6l4\n 9syhecjRTjUTDolrjzR4Ty7Vz/6097miB2t+lfs4lkvbmdXIAiRrV1R6c5CO+j7ztIL9UEL7+TH\n FXapzADDHcH30ONts14xQlEUvCzR6yQkLzneM6zkokpXNR8LFMkwonw2FycqaJ+JidFO37idvTu\n TJexMSmDpZ3Rq9ci1NPhbHGtTwqmn+hdcRTQ5erLbqk5sDym+Paj8Vx2mmtRg3BPN1J/5AfYyLA\n 2Tidqn8eIzEMfBOkM0=",
        "X-Received": "by 2002:a05:6a20:918c:b0:39b:c3a0:9f31 with SMTP id\n adf61e73a8af0-39ef73d5cd2mr4648476637.21.1775059131774;\n Wed, 01 Apr 2026 08:58:51 -0700 (PDT)",
        "Content-Type": "multipart/mixed; boundary=\"------------o3FJ0omHD5WL8NlluOzlnVxX\"",
        "Message-ID": "<4596ca92-75bc-4085-982e-3305002265cf@baylibre.com>",
        "Date": "Wed, 1 Apr 2026 23:58:47 +0800",
        "MIME-Version": "1.0",
        "User-Agent": "Mozilla Thunderbird",
        "Subject": "[PATCH, OpenMP, v2] uses_allocators support",
        "From": "Chung-Lin Tang <cltang@baylibre.com>",
        "To": "gcc-patches <gcc-patches@gcc.gnu.org>,\n Tobias Burnus <tburnus@baylibre.com>, Andrew Stubbs <ams@baylibre.com>,\n Thomas Schwinge <tschwinge@baylibre.com>, Jakub Jelinek <jakub@redhat.com>",
        "References": "<0f1188cd-fa04-46f8-a56e-329635902c80@baylibre.com>",
        "Content-Language": "en-US",
        "In-Reply-To": "<0f1188cd-fa04-46f8-a56e-329635902c80@baylibre.com>",
        "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": "On 11/12/25 11:21 PM, Chung-Lin Tang wrote:\n> This is a new updated patch for OpenMP uses_allocators support.\n> The last one was submitted by Tobias:\n> https://gcc.gnu.org/pipermail/gcc-patches/2023-November/637415.html\n> \n> This new version is a combination of all our patches and fixes during this\n> period, now including:\n> \n> 1. C/C++ front-end parts re-written, to be more like established style.\n> \n> 2. The target teams issue has been solved by a host-side implementation:\n>    omp_init_allocator() is called on host side, and mapped to target,\n>    with the allocator passed in as a firstprivate variable.\n>    Some adjustments where made to ensure that host/device side must have\n>    same omp_allocator_data format.\n> \n> 3. Various other fixes, e.g. ntraits now using array_type_nelts, don't crash on VLAs,\n>    omp_null_allocator, etc.\n\nAfter Tobias has committed the front-end parts, here are the middle-end parts.\n\nThe implementation has be revised to use a libgomp map kind (GOMP_MAP_USES_ALLOCATORS)\nto implement the call to omp_init_allocators.\n\nThis should resolve all the issues raised (hiding omp_allocator_data as opaque, nowait, etc.)\n\nThe sole libgomp uses_allocators-7.f90 test (which seemed to raise several issues) has\nbeen moved to a compiler test.\n\nThanks,\nChung-Lin",
    "diff": "diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc\nindex 4452bb32728..4902a0fee98 100644\n--- a/gcc/c/c-typeck.cc\n+++ b/gcc/c/c-typeck.cc\n@@ -17396,7 +17396,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)\n \t  else if (TREE_CODE (t) == CONST_DECL)\n \t    {\n \t      /* omp_null_allocator is ignored and for predefined allocators,\n-\t\t not special handling is required; thus, remove them removed. */\n+\t\t no special handling is required; thus, mark them removed. */\n \t      remove = true;\n \n \t      if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)\ndiff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc\nindex a2da4fe7c7d..477210f2a08 100644\n--- a/gcc/fortran/trans-array.cc\n+++ b/gcc/fortran/trans-array.cc\n@@ -7117,10 +7117,7 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr)\n \t\t\t       &expr->where, flag_max_array_constructor);\n \t      return NULL_TREE;\n \t    }\n-          if (mpz_cmp_si (c->offset, 0) != 0)\n-            index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind);\n-          else\n-            index = NULL_TREE;\n+\t  index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind);\n \n \t  if (mpz_cmp_si (c->repeat, 1) > 0)\n \t    {\n@@ -7191,7 +7188,7 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr)\n \t    CONSTRUCTOR_APPEND_ELT (v, index, se.expr);\n \t  else\n \t    {\n-\t      if (index != NULL_TREE)\n+\t      if (!integer_zerop (index))\n \t\tCONSTRUCTOR_APPEND_ELT (v, index, se.expr);\n \t      CONSTRUCTOR_APPEND_ELT (v, range, se.expr);\n \t    }\ndiff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc\nindex 230395f9d2c..4d819bbe94e 100644\n--- a/gcc/fortran/trans-openmp.cc\n+++ b/gcc/fortran/trans-openmp.cc\n@@ -3943,7 +3943,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,\n \t\t\t      gfc_init_se (&se, NULL);\n \t\t\t      gfc_conv_expr (&se, n->u2.allocator);\n \t\t\t      gfc_add_block_to_block (block, &se.pre);\n-\t\t\t      allocator_ = gfc_evaluate_now (se.expr, block);\n+\t\t\t      t = se.expr;\n+\t\t\t      if (DECL_P (t) && se.post.head == NULL_TREE)\n+\t\t\t\tallocator_ = (POINTER_TYPE_P (TREE_TYPE (t))\n+\t\t\t\t\t      ? build_fold_indirect_ref (t): t);\n+\t\t\t      else\n+\t\t\t\tallocator_ = gfc_evaluate_now (t, block);\n \t\t\t      gfc_add_block_to_block (block, &se.post);\n \t\t\t    }\n \t\t\t  OMP_CLAUSE_ALLOCATE_ALLOCATOR (node) = allocator_;\n@@ -5225,14 +5230,36 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,\n \t    }\n \t  break;\n \tcase OMP_LIST_USES_ALLOCATORS:\n-\t  /* Ignore omp_null_allocator and pre-defined allocators as no\n-\t     special treatment is needed. */\n \t  for (; n != NULL; n = n->next)\n-\t    if (n->sym->attr.flavor == FL_VARIABLE)\n-\t      break;\n-\t  if (n != NULL)\n-\t    sorry_at (input_location, \"%<uses_allocators%> clause with traits \"\n-\t\t\t\t      \"and memory spaces\");\n+\t    {\n+\t      if (!n->sym->attr.referenced)\n+\t\tcontinue;\n+\t      tree node = build_omp_clause (input_location,\n+\t\t\t\t\t    OMP_CLAUSE_USES_ALLOCATORS);\n+\t      tree t;\n+\t      if (n->sym->attr.flavor == FL_VARIABLE)\n+\t\tt = gfc_get_symbol_decl (n->sym);\n+\t      else\n+\t\t{\n+\t\t  t = gfc_conv_mpz_to_tree (n->sym->value->value.integer,\n+\t\t\t\t\t    n->sym->ts.kind);\n+\t\t  t = fold_convert (ptr_type_node, t);\n+\t\t}\n+\t      OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(node) = t;\n+\t      if (n->u.memspace_sym)\n+\t\t{\n+\t\t  n->u.memspace_sym->attr.referenced = true;\n+\t\t  OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (node)\n+\t\t    = gfc_get_symbol_decl (n->u.memspace_sym);\n+\t\t}\n+\t      if (n->u2.traits_sym)\n+\t\t{\n+\t\t  n->u2.traits_sym->attr.referenced = true;\n+\t\t  OMP_CLAUSE_USES_ALLOCATORS_TRAITS (node)\n+\t\t    = gfc_get_symbol_decl (n->u2.traits_sym);\n+\t\t}\n+\t      omp_clauses = gfc_trans_add_clause (node, omp_clauses);\n+\t    }\n \t  break;\n \tdefault:\n \t  break;\ndiff --git a/gcc/gimplify.cc b/gcc/gimplify.cc\nindex 8bfc71315f9..585032d0960 100644\n--- a/gcc/gimplify.cc\n+++ b/gcc/gimplify.cc\n@@ -184,6 +184,9 @@ enum gimplify_omp_var_data\n   /* Flag for GOVD_FIRSTPRIVATE: OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT.  */\n   GOVD_FIRSTPRIVATE_IMPLICIT = 0x4000000,\n \n+  /* Flag to indicate this an allocator for the uses_allocators clause.  */\n+  GOVD_USES_ALLOCATORS_ALLOCATOR = 0x8000000,\n+\n   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE\n \t\t\t   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR\n \t\t\t   | GOVD_LOCAL)\n@@ -1450,18 +1453,46 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)\n \t\t dynamic_allocators clause is present in the same compilation\n \t\t unit.  */\n \t      bool missing_dyn_alloc = false;\n-\t      if (alloc == NULL_TREE\n-\t\t  && ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS)\n-\t\t      == 0))\n+\t      if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)\n \t\t{\n \t\t  /* This comes too early for omp_discover_declare_target...,\n \t\t     but should at least catch the most common cases.  */\n \t\t  missing_dyn_alloc\n-\t\t    = cgraph_node::get (current_function_decl)->offloadable;\n+\t\t    = (alloc == NULL_TREE\n+\t\t       && cgraph_node::get (current_function_decl)->offloadable);\n \t\t  for (struct gimplify_omp_ctx *ctx2 = ctx;\n \t\t       ctx2 && !missing_dyn_alloc; ctx2 = ctx2->outer_context)\n \t\t    if (ctx2->code == OMP_TARGET)\n-\t\t      missing_dyn_alloc = true;\n+\t\t      {\n+\t\t\tif (alloc == NULL_TREE)\n+\t\t\t  missing_dyn_alloc = true;\n+\t\t\telse if (TREE_CODE (alloc) != INTEGER_CST)\n+\t\t\t  {\n+\t\t\t    tree alloc2 = alloc;\n+\t\t\t    if (TREE_CODE (alloc2) == MEM_REF\n+\t\t\t\t|| TREE_CODE (alloc2) == INDIRECT_REF)\n+\t\t\t      alloc2 = TREE_OPERAND (alloc2, 0);\n+\t\t\t    tree c2;\n+\t\t\t    for (c2 = ctx2->clauses; c2;\n+\t\t\t\t c2 = OMP_CLAUSE_CHAIN (c2))\n+\t\t\t      if (OMP_CLAUSE_CODE (c2)\n+\t\t\t\t  == OMP_CLAUSE_USES_ALLOCATORS)\n+\t\t\t\t{\n+\t\t\t\t  tree t2\n+\t\t\t\t    = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c2);\n+\t\t\t\t  if (operand_equal_p (alloc2, t2))\n+\t\t\t\t    break;\n+\t\t\t\t}\n+\t\t\t    if (c2 == NULL_TREE)\n+\t\t\t      error_at (EXPR_LOC_OR_LOC (\n+\t\t\t\t\t  alloc, DECL_SOURCE_LOCATION (t)),\n+\t\t\t\t\t\"%qE in %<allocator%> clause inside a \"\n+\t\t\t\t\t\"target region must be specified in an \"\n+\t\t\t\t\t\"%<uses_allocators%> clause on the \"\n+\t\t\t\t\t\"%<target%> directive\", alloc2);\n+\t\t\t  }\n+\t\t\tbreak;\n+\t\t      }\n \t\t}\n \t      if (missing_dyn_alloc)\n \t\terror_at (DECL_SOURCE_LOCATION (t),\n@@ -10982,6 +11013,7 @@ omp_get_attachment (omp_mapping_group *grp)\n \t  case GOMP_MAP_FIRSTPRIVATE_POINTER:\n \t  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:\n \t  case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:\n+\t  case GOMP_MAP_USES_ALLOCATORS:\n \t    return NULL_TREE;\n \n \t  case GOMP_MAP_ATTACH_DETACH:\n@@ -11026,6 +11058,7 @@ omp_get_attachment (omp_mapping_group *grp)\n     case GOMP_MAP_FIRSTPRIVATE_INT:\n     case GOMP_MAP_USE_DEVICE_PTR:\n     case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:\n+    case GOMP_MAP_USES_ALLOCATORS:\n       return NULL_TREE;\n \n     default:\n@@ -11314,6 +11347,7 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained,\n     case GOMP_MAP_FIRSTPRIVATE_INT:\n     case GOMP_MAP_USE_DEVICE_PTR:\n     case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:\n+    case GOMP_MAP_USES_ALLOCATORS:\n       return NULL_TREE;\n \n     case GOMP_MAP_FIRSTPRIVATE_POINTER:\n@@ -14931,8 +14965,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,\n \t  break;\n \n \tcase OMP_CLAUSE_USES_ALLOCATORS:\n-\t  sorry_at (OMP_CLAUSE_LOCATION (c), \"%<uses_allocators%> clause\");\n-\t  remove = 1;\n+\t  if (TREE_CODE (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c))\n+\t      != INTEGER_CST)\n+\t    {\n+\t      decl = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);\n+\t      omp_add_variable (ctx, decl,\n+\t\t\t\tGOVD_SEEN | GOVD_USES_ALLOCATORS_ALLOCATOR);\n+\t    }\n+\t  else\n+\t    remove = true;\n \t  break;\n \n \tcase OMP_CLAUSE_ORDERED:\n@@ -15091,6 +15132,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,\n \t      remove = true;\n \t      break;\n \t    }\n+\t  if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0\n+\t      && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)\n+\t      && TREE_CODE (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) != INTEGER_CST)\n+\t    {\n+\t      tree allocator = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c);\n+\t      tree clauses = NULL_TREE;\n+\n+\t      /* Get clause list of the nearest enclosing target construct.  */\n+\t      if (ctx->code == OMP_TARGET)\n+\t\tclauses = *orig_list_p;\n+\t      else\n+\t\t{\n+\t\t  struct gimplify_omp_ctx *tctx = ctx->outer_context;\n+\t\t  while (tctx && tctx->code != OMP_TARGET)\n+\t\t    tctx = tctx->outer_context;\n+\t\t  if (tctx)\n+\t\t    clauses = tctx->clauses;\n+\t\t}\n+\n+\t      if (clauses)\n+\t\t{\n+\t\t  tree uc;\n+\t\t  if (TREE_CODE (allocator) == MEM_REF\n+\t\t      || TREE_CODE (allocator) == INDIRECT_REF)\n+\t\t    allocator = TREE_OPERAND (allocator, 0);\n+\t\t  for (uc = clauses; uc; uc = OMP_CLAUSE_CHAIN (uc))\n+\t\t    if (OMP_CLAUSE_CODE (uc) == OMP_CLAUSE_USES_ALLOCATORS)\n+\t\t      {\n+\t\t\ttree uc_allocator\n+\t\t\t  = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (uc);\n+\t\t\tif (operand_equal_p (allocator, uc_allocator))\n+\t\t\t  break;\n+\t\t      }\n+\t\t  if (uc == NULL_TREE)\n+\t\t    {\n+\t\t      error_at (OMP_CLAUSE_LOCATION (c), \"allocator %qE \"\n+\t\t\t\t\"requires %<uses_allocators(%E)%> clause in \"\n+\t\t\t\t\"target region\", allocator, allocator);\n+\t\t      remove = true;\n+\t\t      break;\n+\t\t    }\n+\t\t}\n+\t    }\n \t  if (gimplify_expr (&OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), pre_p, NULL,\n \t\t\t     is_gimple_val, fb_rvalue) == GS_ERROR)\n \t    {\n@@ -15409,6 +15493,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)\n       code = OMP_CLAUSE__CONDTEMP_;\n       gimple_add_tmp_var (decl);\n     }\n+  else if (flags & GOVD_USES_ALLOCATORS_ALLOCATOR)\n+    return 0;\n   else\n     gcc_unreachable ();\n \ndiff --git a/gcc/omp-low.cc b/gcc/omp-low.cc\nindex b93012107f1..990c46ecfff 100644\n--- a/gcc/omp-low.cc\n+++ b/gcc/omp-low.cc\n@@ -1177,6 +1177,36 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)\n \t    && omp_maybe_offloaded_ctx (ctx))\n \t  error_at (OMP_CLAUSE_LOCATION (c), \"%<allocate%> clause must\"\n \t\t    \" specify an allocator here\");\n+\tif ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0\n+\t    && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) != NULL_TREE\n+\t    && DECL_P (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c))\n+\t    && !DECL_ARTIFICIAL (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)))\n+\t  {\n+\t    tree alloc2 = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c);\n+\t    if (TREE_CODE (alloc2) == MEM_REF\n+\t\t|| TREE_CODE (alloc2) == INDIRECT_REF)\n+\t      alloc2 = TREE_OPERAND (alloc2, 0);\n+\t    omp_context *ctx2 = ctx;\n+\t    for (; ctx2; ctx2 = ctx2->outer)\n+\t      if (is_gimple_omp_offloaded (ctx2->stmt))\n+\t\tbreak;\n+\t    if (ctx2 != NULL)\n+\t      {\n+\t\ttree c2 = gimple_omp_target_clauses (ctx2->stmt);\n+\t\tfor (; c2; c2 = OMP_CLAUSE_CHAIN (c2))\n+\t\t  if (OMP_CLAUSE_CODE (c2) == OMP_CLAUSE_USES_ALLOCATORS\n+\t\t      && operand_equal_p (\n+\t\t\t   alloc2, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c2)))\n+\t\t    break;\n+\t\tif (c2 == NULL_TREE)\n+\t\t  error_at (EXPR_LOC_OR_LOC (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c),\n+\t\t\t\t\t     OMP_CLAUSE_LOCATION (c)),\n+\t\t\t    \"allocator %qE in %<allocate%> clause inside a \"\n+\t\t\t    \"target region must be specified in an \"\n+\t\t\t    \"%<uses_allocators%> clause on the %<target%> \"\n+\t\t\t    \"directive\", alloc2);\n+\t      }\n+\t  }\n \tif (ctx->allocate_map == NULL)\n \t  ctx->allocate_map = new hash_map<tree, tree>;\n \ttree val = integer_zero_node;\n@@ -1775,6 +1805,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)\n \tcase OMP_CLAUSE_DEVICE_TYPE:\n \t  break;\n \n+\tcase OMP_CLAUSE_USES_ALLOCATORS:\n+\t  decl = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);\n+\t  gcc_assert (DECL_P (decl));\n+\t  gcc_assert (is_gimple_omp_offloaded (ctx->stmt));\n+\t  install_var_field (decl, false, 3, ctx);\n+\t  install_var_local (decl, ctx);\n+\t  break;\n+\n \tcase OMP_CLAUSE_ALIGNED:\n \t  decl = OMP_CLAUSE_DECL (c);\n \t  if (is_global_var (decl)\n@@ -2000,6 +2038,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)\n \tcase OMP_CLAUSE_USE:\n \tcase OMP_CLAUSE_DESTROY:\n \tcase OMP_CLAUSE_DEVICE_TYPE:\n+\tcase OMP_CLAUSE_USES_ALLOCATORS:\n \t  break;\n \n \tcase OMP_CLAUSE__CACHE_:\n@@ -13173,6 +13212,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)\n \t    sorry_at (OMP_CLAUSE_LOCATION (c),\n \t\t      \"only the %<device_type(any)%> is supported\");\n \t  break;\n+      case OMP_CLAUSE_USES_ALLOCATORS:\n+\tallocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);\n+\ttree new_allocator = lookup_decl (allocator, ctx);\n+\tx = build_receiver_ref (allocator, false, ctx);\n+\tSET_DECL_VALUE_EXPR (new_allocator, x);\n+\tDECL_HAS_VALUE_EXPR_P (new_allocator) = 1;\n+\tmap_cnt++;\n+\tbreak;\n       }\n \n   if (offloaded)\n@@ -13905,6 +13952,80 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)\n \t    CONSTRUCTOR_APPEND_ELT (vkind, purpose,\n \t\t\t\t    build_int_cstu (tkind_type, tkind));\n \t    break;\n+\n+\t  case OMP_CLAUSE_USES_ALLOCATORS:\n+\t    tree allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);\n+\t    tree memspace = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);\n+\t    tree traits = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);\n+\n+\t    tree ntraits, traits_var;\n+\t    if (traits == NULL_TREE)\n+\t      {\n+\t\tntraits = integer_zero_node;\n+\t\ttraits_var = null_pointer_node;\n+\t      }\n+\t    else if (DECL_INITIAL (traits))\n+\t      {\n+\t\tlocation_t loc = OMP_CLAUSE_LOCATION (c);\n+\t\tntraits = array_type_nelts_top (TREE_TYPE (traits));\n+\t\ttree t = DECL_INITIAL (traits);\n+\t\tt = get_initialized_tmp_var (t, &ilist, NULL);\n+\t\ttraits_var = build_fold_addr_expr_loc (loc, t);\n+\t      }\n+\t    else\n+\t      {\n+\t\t/* This happens for VLAs, which probably aren't useful\n+\t\t   because they can't be const initialized in the same\n+\t\t   scope....  is there something else?  */\n+\t\tlocation_t loc = OMP_CLAUSE_LOCATION (c);\n+\t\tgcc_assert (TREE_CODE (TREE_TYPE (traits)) == ARRAY_TYPE);\n+\t\tntraits = array_type_nelts_top (TREE_TYPE (traits));\n+\t\ttraits_var = build_fold_addr_expr_loc (loc, traits);\n+\t      }\n+\n+\t    if (memspace == NULL_TREE)\n+\t      memspace = build_int_cst (pointer_sized_int_node, 0);\n+\t    else\n+\t      memspace = fold_convert (pointer_sized_int_node, memspace);\n+\n+\t    tree arr_type\n+\t      = build_array_type_nelts (pointer_sized_int_node, 3);\n+\t    tree uses_allocators_descr\n+\t      = create_tmp_var (arr_type, \"uses_allocator_descr\");\n+\t    tree ua_descr[3];\n+\t    for (int i = 0; i < 3; i++)\n+\t      ua_descr[i] = build4 (ARRAY_REF, pointer_sized_int_node,\n+\t\t\t\t    uses_allocators_descr,\n+\t\t\t\t    build_int_cst (size_type_node, i),\n+\t\t\t\t    NULL_TREE, NULL_TREE);\n+\t    gimplify_assign (ua_descr[0],\n+\t\t\t     fold_convert (pointer_sized_int_node, memspace),\n+\t\t\t     &ilist);\n+\t    gimplify_assign (ua_descr[1],\n+\t\t\t     fold_convert (pointer_sized_int_node, ntraits),\n+\t\t\t     &ilist);\n+\t    gimplify_assign (ua_descr[2],\n+\t\t\t     fold_convert (pointer_sized_int_node, traits_var),\n+\t\t\t     &ilist);\n+\n+\t    x = build_sender_ref (allocator, ctx);\n+\t    tree ptr = build_fold_addr_expr (uses_allocators_descr);\n+\t    gimplify_assign (x, fold_convert (TREE_TYPE (x), ptr), &ilist);\n+\n+\t    s = size_int (0);\n+\t    purpose = size_int (map_idx++);\n+\t    CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);\n+\t    tkind = GOMP_MAP_USES_ALLOCATORS;\n+\t    gcc_checking_assert (tkind\n+\t\t\t\t < (HOST_WIDE_INT_C (1U) << talign_shift));\n+\t    talign = TYPE_ALIGN_UNIT (pointer_sized_int_node);\n+\t    talign = ceil_log2 (talign);\n+\t    tkind |= talign << talign_shift;\n+\t    gcc_checking_assert (tkind\n+\t\t\t\t <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));\n+\t    CONSTRUCTOR_APPEND_ELT (vkind, purpose,\n+\t\t\t\t    build_int_cstu (tkind_type, tkind));\n+\t    break;\n \t  }\n \n       gcc_assert (map_idx == map_cnt);\ndiff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c\nindex df82cbbcba9..ec2f56e21d0 100644\n--- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c\n+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c\n@@ -1,3 +1,4 @@\n+/*\n typedef enum omp_allocator_handle_t\n #if __cplusplus >= 201103L\n : __UINTPTR_TYPE__\n@@ -15,17 +16,19 @@ typedef struct omp_alloctrait_t\n } omp_alloctrait_t;\n \n extern void *omp_alloc (__SIZE_TYPE__, omp_allocator_handle_t);\n+*/\n+#include <omp.h>\n \n void\n f (omp_allocator_handle_t my_alloc)\n {\n   #pragma omp target\n   {\n-    int a; /* { dg-error \"'my_alloc' in 'allocator' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive\" \"not yet implemented\" { xfail *-*-* } } */\n+    int a; /* { dg-error \"'my_alloc' in 'allocator' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive\" \"\" { target c } } */\n     #pragma omp allocate(a) allocator(my_alloc) /* { dg-message \"sorry, unimplemented: '#pragma omp allocate' not yet supported\" \"\" { target c++ } }  */\n     a  = 5;\n     void *prt = omp_alloc(32, my_alloc);\n-    #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a) /* { dg-error \"allocator 'my_alloc' in 'allocate' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive\" \"not yet implemented\" { xfail *-*-* } } */\n+    #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a) /* { dg-error \"allocator 'my_alloc' in 'allocate' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive\" } */\n       a = 7;\n   }\n }\n@@ -44,5 +47,3 @@ g (omp_allocator_handle_t my_alloc)\n       a = 7;\n   }\n }\n-\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 37 }\ndiff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c\nindex c766ea3559d..4dd1f13100a 100644\n--- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c\n+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c\n@@ -31,5 +31,3 @@ g ()\n    #pragma omp target uses_allocators(traits(trait) : my_alloc)\n      ;\n }\n-\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 31 }\ndiff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c\nnew file mode 100644\nindex 00000000000..1e4699d99c2\n--- /dev/null\n+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c\n@@ -0,0 +1,53 @@\n+/* { dg-do compile } */\n+/* { dg-additional-options \"-fdump-tree-original -fdump-tree-gimple\" } */\n+\n+#include <omp.h>\n+\n+omp_alloctrait_key_t k;\n+omp_alloctrait_value_t v;\n+\n+int n = 2;\n+\n+int main (void)\n+{\n+  omp_allocator_handle_t foo, bar;\n+  const omp_alloctrait_t foo_traits[] = { { omp_atk_pinned,    omp_atv_true },\n+  \t\t\t\t\t  { omp_atk_partition, omp_atv_nearest } };\n+\n+  #pragma omp target\n+    ;\n+  #pragma omp target uses_allocators (bar)\n+    ;\n+  #pragma omp target uses_allocators (foo (foo_traits)) /* { dg-warning \"the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\\\(traits\\\\)' is deprecated since OpenMP 5.2 \\\\\\[-Wdeprecated-openmp\\\\\\]\" } */\n+    ;\n+  #pragma omp target uses_allocators (foo (foo_traits), bar (foo_traits)) /* { dg-warning \"the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\\\(traits\\\\)' is deprecated since OpenMP 5.2 \\\\\\[-Wdeprecated-openmp\\\\\\]\" } */\n+    ;\n+  #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo)\n+    ;\n+  #pragma omp target uses_allocators (traits(foo_traits) : bar)\n+    ;\n+  #pragma omp target parallel uses_allocators (memspace(omp_high_bw_mem_space), traits(foo_traits) : bar)\n+    ;\n+  #pragma omp target parallel uses_allocators (traits(foo_traits), memspace(omp_high_bw_mem_space) : bar) uses_allocators(foo)\n+  {\n+    void *p = omp_alloc ((unsigned long) 32, bar);\n+    omp_free (p, bar);\n+  }\n+  return 0;\n+}\n+\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : bar\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(foo_traits\\\\) : foo\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(foo_traits\\\\) : bar\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(foo_traits\\\\) : foo\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(.+\\\\), traits\\\\(\\\\) : foo\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(foo_traits\\\\) : bar\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(.+\\\\), traits\\\\(foo_traits\\\\) : bar\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(.+\\\\), traits\\\\(foo_traits\\\\) : bar\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : foo\\\\)\" \"original\" } } */\n+\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : bar\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(foo_traits\\\\) : foo\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(foo_traits\\\\) : bar\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(foo_traits\\\\) : foo\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(.+\\\\), traits\\\\(\\\\) : foo\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(foo_traits\\\\) : bar\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(.+\\\\), traits\\\\(foo_traits\\\\) : bar\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(.+\\\\), traits\\\\(foo_traits\\\\) : bar\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : foo\\\\)\" \"gimple\" } } */\ndiff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c\nindex 5f3650157fc..a27e7208473 100644\n--- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c\n+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-4.c\n@@ -135,10 +135,3 @@ int f (const omp_alloctrait_t arg_traits[], int n)\n                                                                                          /* { dg-error \"legacy 'foo\\\\\\(foo_traits\\\\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers\" \"\" { target *-*-* } .-2 } */\n   return 0;\n }\n-\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 103 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target c } 111 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target c } 113 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target c } 117 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target c } 119 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 131 }\ndiff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-5.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-5.c\nnew file mode 100644\nindex 00000000000..eeac5f097cb\n--- /dev/null\n+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-5.c\n@@ -0,0 +1,34 @@\n+/* { dg-do compile } */\n+/* { dg-additional-options \"-fdump-tree-original -fdump-tree-gimple\" } */\n+\n+#include <omp.h>\n+\n+int main (void)\n+{\n+  omp_allocator_handle_t memspace, traits;\n+  const omp_alloctrait_t mytraits[] = { { omp_atk_pinned,    omp_atv_true },\n+\t\t\t\t\t{ omp_atk_partition, omp_atv_nearest } };\n+  #pragma omp target uses_allocators (memspace)\n+    ;\n+  #pragma omp target uses_allocators (traits)\n+    ;\n+  #pragma omp target uses_allocators (traits, memspace)\n+    ;\n+  #pragma omp target uses_allocators (traits (mytraits)) /* { dg-warning \"the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\\\(traits\\\\)' is deprecated since OpenMP 5.2\" } */\n+    ;\n+  #pragma omp target uses_allocators (memspace (mytraits), omp_default_mem_alloc) /* { dg-warning \"the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\\\(traits\\\\)' is deprecated since OpenMP 5.2\" } */\n+    ;\n+  return 0;\n+}\n+\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : memspace\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : traits\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : memspace\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : traits\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(mytraits\\\\) : traits\\\\)\" \"original\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(mytraits\\\\) : memspace\\\\)\" \"original\" } } */\n+\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : memspace\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : traits\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : memspace\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : traits\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(mytraits\\\\) : traits\\\\)\" \"gimple\" } } */\n+/* { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(mytraits\\\\) : memspace\\\\)\" \"gimple\" } } */\ndiff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-6.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-6.c\nnew file mode 100644\nindex 00000000000..23a831e8ba9\n--- /dev/null\n+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-6.c\n@@ -0,0 +1,53 @@\n+/* { dg-additional-options \"-fdump-tree-gimple\" } */\n+\n+#include <stdint.h>\n+#include <omp.h>\n+\n+int\n+main ()\n+{\n+  int x, *xbuf[10];\n+  omp_allocator_handle_t my_alloc;\n+  const omp_alloctrait_t trait[1]= {{omp_atk_alignment,128}};\n+\n+  #pragma omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none)\n+    #pragma omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf)\n+      {\n+\tif ((uintptr_t) &x % 128 != 0)\n+\t  __builtin_abort ();\n+\tif ((uintptr_t) xbuf % 128 != 0)\n+\t  __builtin_abort ();\n+      }\n+\n+  my_alloc = (omp_allocator_handle_t) 0xABCD;\n+\n+  #pragma omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf)\n+    #pragma omp parallel allocate(allocator(my_alloc): x, xbuf) if(0) firstprivate(x, xbuf)\n+      {\n+\tif ((uintptr_t) &x % 128 != 0)\n+\t  __builtin_abort ();\n+\tif ((uintptr_t) xbuf % 128 != 0)\n+\t  __builtin_abort ();\n+      }\n+\n+  if (my_alloc != (omp_allocator_handle_t) 0xABCD)\n+    __builtin_abort ();\n+\n+  /* The following creates an allocator with empty traits + default mem space. */\n+  #pragma omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none)\n+    #pragma omp parallel allocate(allocator(my_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf)\n+      {\n+\tif ((uintptr_t) &x % 128 != 0)\n+\t  __builtin_abort ();\n+\tif ((uintptr_t) xbuf % 128 != 0)\n+\t  __builtin_abort ();\n+      }\n+\n+  if (my_alloc != (omp_allocator_handle_t) 0xABCD)\n+    __builtin_abort ();\n+\n+  return 0;\n+}\n+\n+/* { dg-final { scan-tree-dump-times \"#pragma omp target .* uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(trait\\\\) : my_alloc\\\\)\" 1 \"gimple\" } } */\n+/* { dg-final { scan-tree-dump-times \"#pragma omp target .* uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : my_alloc\\\\)\" 1 \"gimple\" } } */\ndiff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-7.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-7.c\nindex ce43d70ecf2..66b257c7fbb 100644\n--- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-7.c\n+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-7.c\n@@ -76,8 +76,3 @@ const omp_alloctrait_t t[] = {};\n  #pragma omp target uses_allocators(memspace(omp_default_mem_space) : my, my(t)) // { dg-error \"legacy 'my\\\\(t\\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers\" }\n    ;\n }\n-\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 57 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 61 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 63 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 76 }\ndiff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-8.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-8.c\nindex 3ab87576c41..f03a52ef4ce 100644\n--- a/gcc/testsuite/c-c++-common/gomp/uses_allocators-8.c\n+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-8.c\n@@ -56,6 +56,3 @@ const omp_alloctrait_t t2[] = {};\n \n // { dg-final { scan-tree-dump \"#pragma omp target uses_allocators\\\\(memspace\\\\(1\\\\), traits\\\\(\\\\) : my4\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(t2\\\\) : my3\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : my2\\\\) uses_allocators\\\\(memspace\\\\(3\\\\), traits\\\\(t\\\\) : my\\\\)\" \"original\" { target c } } }\n // { dg-final { scan-tree-dump \"#pragma omp target uses_allocators\\\\(memspace\\\\(omp_large_cap_mem_space\\\\), traits\\\\(\\\\) : my4\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(t2\\\\) : my3\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : my2\\\\) uses_allocators\\\\(memspace\\\\(omp_high_bw_mem_space\\\\), traits\\\\(t\\\\) : my\\\\)\" \"original\" { target c++ } } }\n-\n-\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 53 }\ndiff --git a/gcc/testsuite/g++.dg/gomp/uses_allocators-1.C b/gcc/testsuite/g++.dg/gomp/uses_allocators-1.C\nindex 79e51638534..b4f9fe772f8 100644\n--- a/gcc/testsuite/g++.dg/gomp/uses_allocators-1.C\n+++ b/gcc/testsuite/g++.dg/gomp/uses_allocators-1.C\n@@ -99,7 +99,3 @@ void use2()\n \n   g2<omp_allocator_handle_t, const omp_alloctrait_t[]>(my); // OK\n }\n-\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 51 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 58 }\n-// { dg-message \"sorry, unimplemented: 'uses_allocators' clause\" \"\" { target *-*-* } 89 }\ndiff --git a/gcc/testsuite/gcc.dg/gomp/deprecate-2.c b/gcc/testsuite/gcc.dg/gomp/deprecate-2.c\nindex 46d286141ef..b7689dd1554 100644\n--- a/gcc/testsuite/gcc.dg/gomp/deprecate-2.c\n+++ b/gcc/testsuite/gcc.dg/gomp/deprecate-2.c\n@@ -42,6 +42,3 @@ void f()\n \n // { dg-warning \"42: the specification of arguments to 'uses_allocators' where each item is of the form 'allocator\\\\(traits\\\\)' is deprecated since OpenMP 5.2 \\\\\\[-Wdeprecated-openmp\\\\\\]\" \"\" { target *-*-* }  34 }\n }\n-\n-// { dg-excess-errors \"sorry, unimplemented: 'uses_allocators' clause\" }\n-// { dg-excess-errors \"sorry, unimplemented: 'uses_allocators' clause\" }\ndiff --git a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90\nindex 8bc6b768778..0463f0e0af9 100644\n--- a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90\n+++ b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90\n@@ -24,6 +24,10 @@ module omp_lib_kinds\n      parameter :: omp_pteam_mem_alloc = 7\n   integer (kind=omp_allocator_handle_kind), &\n      parameter :: omp_thread_mem_alloc = 8\n+\n+  integer, parameter :: omp_memspace_handle_kind = c_intptr_t\n+  integer (omp_memspace_handle_kind), &\n+     parameter :: omp_default_mem_space = 0\n end module\n \n subroutine bar (a, b, c)\n@@ -80,7 +84,8 @@ subroutine foo(x, y)\n   \n   !$omp target teams distribute parallel do private (x) firstprivate (y) &\n   !$omp allocate ((omp_default_mem_alloc + 0):z) allocate &\n-  !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r)\n+  !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) &\n+  !$omp uses_allocators(memspace(omp_default_mem_space) : h)\n   do i = 1, 10\n     call bar (0, x, z);\n     call bar2 (1, y, r);\ndiff --git a/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 b/gcc/testsuite/gfortran.dg/gomp/scope-6.f90\nindex 4c4f5e034f7..39a65904c33 100644\n--- a/gcc/testsuite/gfortran.dg/gomp/scope-6.f90\n+++ b/gcc/testsuite/gfortran.dg/gomp/scope-6.f90\n@@ -20,4 +20,4 @@ contains\n   end\n end\n \n-! { dg-final { scan-tree-dump \"omp scope private\\\\(a\\\\) firstprivate\\\\(b\\\\) reduction\\\\(\\\\+:c\\\\) allocate\\\\(allocator\\\\(D\\\\.\\[0-9\\]+\\\\):a\\\\) allocate\\\\(allocator\\\\(D\\\\.\\[0-9\\]+\\\\):b\\\\) allocate\\\\(allocator\\\\(D\\\\.\\[0-9\\]+\\\\):c\\\\)\" \"original\" } }\n+! { dg-final { scan-tree-dump \"omp scope private\\\\(a\\\\) firstprivate\\\\(b\\\\) reduction\\\\(\\\\+:c\\\\) allocate\\\\(allocator\\\\(h\\\\):a\\\\) allocate\\\\(allocator\\\\(h\\\\):b\\\\) allocate\\\\(allocator\\\\(h\\\\):c\\\\)\" \"original\" } }\ndiff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90\nnew file mode 100644\nindex 00000000000..0c6c5495729\n--- /dev/null\n+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90\n@@ -0,0 +1,17 @@\n+use iso_c_binding\n+use omp_lib\n+implicit none\n+contains\n+subroutine x\n+integer :: mem\n+type(omp_alloctrait), parameter:: mem2(1) = [omp_alloctrait(1,1)]\n+integer(omp_allocator_handle_kind) :: var\n+!$omp target uses_allocators(memspace(omp_default_mem_space), traits(mem2) : var) defaultmap(none)\n+block;\n+type(c_ptr) ::c\n+c = omp_alloc(omp_default_mem_space, 20_8)\n+end block\n+!$omp target uses_allocators(omp_default_mem_alloc, var(mem2))  ! { dg-warning \"The specification of arguments to 'uses_allocators' at \\\\(1\\\\) where each item is of the form 'allocator\\\\(traits\\\\)' is deprecated since OpenMP 5.2; instead use 'uses_allocators\\\\(traits\\\\(mem2\\\\): var\\\\)' \\\\\\[-Wdeprecated-openmp\\\\\\]\" }\n+block; end block\n+end\n+end\ndiff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_1.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90\nsimilarity index 100%\nrename from libgomp/testsuite/libgomp.fortran/uses_allocators_1.f90\nrename to gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90\ndiff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90\nsimilarity index 70%\nrename from libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90\nrename to gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90\nindex 0ab09975f49..005ead56ee4 100644\n--- a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90\n+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90\n@@ -3,8 +3,6 @@\n ! Minimal test for valid code:\n ! - predefined allocators do not need any special treatment in uses_allocators\n !   (as 'requires dynamic_allocators' is the default).\n-!\n-! - Non-predefined allocators are currently rejected ('sorry)'\n \n subroutine test\n   use omp_lib\n@@ -35,22 +33,22 @@ subroutine non_predef\n \n   integer(kind=omp_allocator_handle_kind) :: a1, a2, a3\n \n-  !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))\n   block; end block\n \n-  !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))\n   block; end block\n \n \n   !$omp target uses_allocators(traits(trait):a1) &\n-  !$omp&        uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp&        uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)\n   block; end block\n \n   !$omp target parallel uses_allocators(traits(trait):a1) &\n-  !$omp&        uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp&        uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)\n   block; end block\n \n-  !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3)  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3)\n   block; end block\n end subroutine\n \n@@ -62,7 +60,7 @@ subroutine trait_present\n   integer(kind=omp_allocator_handle_kind) :: a1\n \n   ! Invalid in OpenMP 5.0 / 5.1, but valid since 5.2 the same as omp_default_mem_space + emptry traits array\n-  !$omp target uses_allocators ( a1 )  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp target uses_allocators ( a1 )\n   block; end block\n end\n \n@@ -76,13 +74,13 @@ subroutine odd_names\n   integer(kind=omp_allocator_handle_kind) :: traits\n   integer(kind=omp_allocator_handle_kind) :: memspace\n \n-  !$omp target uses_allocators ( traits(trait1), memspace(trait1) )  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp target uses_allocators ( traits(trait1), memspace(trait1) )\n   block; end block\n \n-  !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space)  : traits)  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space)  : traits)\n   block; end block\n \n-  !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace)  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace)\n   block; end block\n end\n \n@@ -94,6 +92,6 @@ subroutine more_checks\n   integer(kind=omp_allocator_handle_kind) :: a1, a2(4)\n   integer(kind=1) :: a3\n \n-  !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 )  ! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" }\n+  !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 )\n   block; end block\n end\ndiff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-4.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-4.f90\nnew file mode 100644\nindex 00000000000..0efdc9ef54b\n--- /dev/null\n+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-4.f90\n@@ -0,0 +1,59 @@\n+! { dg-do compile }\n+! { dg-additional-options \"-fdump-tree-original -fdump-tree-gimple\" }\n+\n+program main\n+  use omp_lib\n+  implicit none\n+  integer, allocatable :: arr(:)\n+  integer (omp_allocator_handle_kind) :: bar, foo\n+\n+  type (omp_alloctrait), parameter :: traits_array(*) = &\n+       [omp_alloctrait(omp_atk_pinned,omp_atv_true),&\n+       omp_alloctrait(omp_atk_partition,omp_atv_nearest)]\n+\n+  !$omp target allocate(bar : arr) uses_allocators(bar)\n+  block\n+    allocate(arr(100))\n+  end block\n+\n+  !$omp target uses_allocators(omp_default_mem_alloc)\n+  block\n+  end block\n+\n+  !$omp target uses_allocators(bar(traits_array), foo (traits_array)) ! { dg-warning \"The specification of arguments to 'uses_allocators' at \\\\(1\\\\) where each item is of the form 'allocator\\\\(traits\\\\)' is deprecated since OpenMP 5.2; instead use 'uses_allocators\\\\(traits\\\\(traits_array\\\\): foo\\\\)' \\\\\\[-Wdeprecated-openmp\\\\\\]\" }\n+  block\n+    if (foo == 0) stop 1\n+  end block\n+\n+  !$omp target uses_allocators(traits(traits_array) : bar)\n+  block\n+  end block\n+\n+  !$omp target parallel uses_allocators(memspace (omp_low_lat_mem_space) : bar)\n+  block\n+  end block\n+\n+  !$omp target parallel uses_allocators(memspace (omp_high_bw_mem_space), traits(traits_array) : bar)\n+  block\n+    use iso_c_binding\n+    type(c_ptr) :: ptr\n+    integer(c_size_t) :: sz = 32\n+    ptr = omp_alloc (sz, bar)\n+    call omp_free (ptr, bar)\n+  end block\n+\n+end program main\n+\n+! { dg-final { scan-tree-dump \"pragma omp target allocate\\\\(allocator\\\\(bar\\\\):arr\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : bar\\\\)\" \"original\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target\" \"original\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(traits_array\\\\) : bar\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(traits_array\\\\) : foo\\\\)\" \"original\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(traits_array\\\\) : bar\\\\)\" \"original\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(omp_low_lat_mem_space\\\\), traits\\\\(\\\\) : bar\\\\)\" \"original\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target uses_allocators\\\\(memspace\\\\(omp_high_bw_mem_space\\\\), traits\\\\(traits_array\\\\) : bar\\\\)\" \"original\" } }\n+\n+! { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) allocate\\\\(allocator\\\\(bar\\\\):arr\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : bar\\\\)\" \"gimple\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target\" \"gimple\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(traits_array\\\\) : bar\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(traits_array\\\\) : foo\\\\)\" \"gimple\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(traits_array\\\\) : bar\\\\)\" \"gimple\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(omp_low_lat_mem_space\\\\), traits\\\\(\\\\) : bar\\\\)\" \"gimple\" } }\n+! { dg-final { scan-tree-dump \"pragma omp target num_teams\\\\(-2\\\\) thread_limit\\\\(0\\\\) uses_allocators\\\\(memspace\\\\(omp_high_bw_mem_space\\\\), traits\\\\(traits_array\\\\) : bar\\\\)\" \"gimple\" } }\ndiff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-5.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-5.f90\nnew file mode 100644\nindex 00000000000..86a86a08bf0\n--- /dev/null\n+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-5.f90\n@@ -0,0 +1,54 @@\n+! { dg-do compile }\n+\n+program main\n+  use omp_lib\n+  implicit none\n+  integer (omp_allocator_handle_kind) :: bar, foo\n+\n+  type (omp_alloctrait), parameter :: traits_array(*) = &\n+       [omp_alloctrait(omp_atk_pinned,omp_atv_true),&\n+       omp_alloctrait(omp_atk_partition,omp_atv_nearest)]\n+\n+  !$omp target uses_allocators(omp_non_existant_alloc) ! { dg-error \"Allocator 'omp_non_existant_alloc' at .1. in USES_ALLOCATORS must be a scalar integer of kind 'omp_allocator_handle_kind'\" }\n+  block  ! { dg-error \"Symbol 'omp_non_existant_alloc' at .1. has no IMPLICIT type; did you mean 'omp_const_mem_alloc'\\?\" \"\" { target *-*-* } .-1 }\n+  end block\n+\n+  !$omp target uses_allocators(bar(traits_array), foo (traits_array), ) ! { dg-error \"Invalid character in name\" }\n+  block\n+  end block\n+\n+  !$omp target uses_allocators(traits(xyz) : bar) ! { dg-error \"Symbol 'xyz' at .1. has no IMPLICIT type\" }\n+  block  ! { dg-error \"Traits array 'xyz' in USES_ALLOCATORS .1. must be a one-dimensional named constant array of type 'omp_alloctrait'\" \"\" { target *-*-* } .-1 }\n+  end block\n+\n+  !$omp target uses_allocators(memspace(omp_non_existant_mem_space) : foo) ! { dg-error \"Symbol 'omp_non_existant_mem_space' at .1. has no IMPLICIT type; did you mean 'omp_const_mem_space'\\?\" }\n+  ! { dg-error \"Memspace 'omp_non_existant_mem_space' at .1. in USES_ALLOCATORS must be a predefined memory space\" \"\" { target *-*-* } .-1 }\n+\n+  block\n+  end block\n+\n+  !$omp target uses_allocators(traits(traits_array), traits(traits_array) : bar) ! { dg-error \"Duplicate TRAITS modifier at .1. in USES_ALLOCATORS clause\" }\n+  block\n+  end block\n+\n+  !$omp target uses_allocators(memspace(omp_default_mem_space), memspace(omp_default_mem_space) : foo) ! { dg-error \"Duplicate MEMSPACE modifier at .1. in USES_ALLOCATORS clause\" }\n+  block\n+  end block\n+\n+  !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array), traits(traits_array) : foo) ! { dg-error \"Duplicate TRAITS modifier at .1. in USES_ALLOCATORS clause\" }\n+  block\n+  end block\n+\n+  !$omp target uses_allocators (omp_null_allocator)\n+  block\n+  end block\n+\n+  !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar)\n+  block\n+  end block\n+\n+  !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) ! { dg-error \"70:Unexpected '\\\\(' at .1.\" }\n+  block\n+  end block\n+\n+end program main\ndiff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-6.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-6.f90\nnew file mode 100644\nindex 00000000000..00f87109d2c\n--- /dev/null\n+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-6.f90\n@@ -0,0 +1,14 @@\n+! { dg-do compile }\n+\n+program main\n+  use omp_lib\n+  implicit none\n+  integer, allocatable :: arr(:)\n+  integer (omp_allocator_handle_kind) :: bar\n+\n+  !$omp target allocate(bar : arr) ! { dg-error \"allocator 'bar' requires 'uses_allocators.bar.' clause in target region\" }\n+  block\n+    allocate(arr(100))\n+  end block\n+\n+end program main\ndiff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-7.f90\nsimilarity index 65%\nrename from libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90\nrename to gcc/testsuite/gfortran.dg/gomp/uses_allocators-7.f90\nindex 65e2fca082e..5ca636e7c79 100644\n--- a/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90\n+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-7.f90\n@@ -1,5 +1,4 @@\n-! { dg-do compile }\n-! { dg-additional-options \"-fdump-tree-gimple\" }\n+! { dg-additional-options \"-fdump-tree-original -fdump-tree-gimple\" }\n \n program main\n   use iso_c_binding\n@@ -12,7 +11,7 @@ program main\n   type(omp_alloctrait), parameter :: t(*) = [omp_alloctrait:: ]\n   type(omp_alloctrait), parameter :: t2(*) = [omp_alloctrait:: ]\n \n-  ! FIXME - improve check that that ';' is handled\n+  !! FIXME - improve check that that ';' is handled\n   !$omp target uses_allocators(traits(t), memspace(omp_high_bw_mem_space) : my; omp_default_mem_alloc, omp_null_allocator; my2; traits(t2) : my3; memspace(omp_large_cap_mem_space) : my4)\n   block\n   end block\n@@ -54,16 +53,6 @@ program main\n     stop 8\n end\n \n-\n-! FIXME ENABLE: 'dg FIXME final' -> 'dg-final'\n-! { dg  FIXME  final { scan-tree-dump-times \"#pragma omp target .*private\\\\(my_alloc\\\\).*uses_allocators\\\\(my_alloc: memspace\\\\(\\\\), traits\\\\(trait\\\\)\\\\)\" 1 \"gimple\" } }\n-! { dg  FIXME  final { scan-tree-dump-times \"#pragma omp target .*private\\\\(my_alloc\\\\).*uses_allocators\\\\(my_alloc: memspace\\\\(\\\\), traits\\\\(\\\\)\\\\)\" 1 \"gimple\" } }\n-! { dg  FIXME  final { scan-tree-dump \"#pragma omp target uses_allocators\\\\(memspace\\\\(1\\\\), traits\\\\(\\\\) : my4\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(t2\\\\) : my3\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : my2\\\\) uses_allocators\\\\(memspace\\\\(3\\\\), traits\\\\(t\\\\) : my\\\\)\" 1 \"original\" } }\n-\n-\n-! FIXME ENABLE code above for \"gimple\" once it has been implemented:\n-! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" \"\" { target *-*-* } 16 }\n-! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" \"\" { target *-*-* } 31 }\n-! { dg-message \"sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces\" \"\" { target *-*-* } 44 }\n-! { dg-bogus \"'my_alloc' not specified in enclosing 'target'\" \"bogus issue because clause is ignored\" { xfail *-*-* } 32 }\n-! { dg-bogus \"'my_alloc' not specified in enclosing 'target'\" \"bogus issue because clause is ignored\" { xfail *-*-* } 45 }\n+! { dg-final { scan-tree-dump-times \"#pragma omp target .*uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(trait\\\\) : my_alloc\\\\)\" 1 \"gimple\" } }\n+! { dg-final { scan-tree-dump-times \"#pragma omp target .*uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : my_alloc\\\\)\" 1 \"gimple\" } }\n+! { dg-final { scan-tree-dump-times \"#pragma omp target .*uses_allocators\\\\(memspace\\\\(1\\\\), traits\\\\(\\\\) : my4\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(t2\\\\) : my3\\\\) uses_allocators\\\\(memspace\\\\(\\\\), traits\\\\(\\\\) : my2\\\\) uses_allocators\\\\(memspace\\\\(3\\\\), traits\\\\(t\\\\) : my\\\\)\" 1 \"original\" } }\ndiff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc\nindex 79e0a734dbb..f678f44f243 100644\n--- a/gcc/tree-pretty-print.cc\n+++ b/gcc/tree-pretty-print.cc\n@@ -1154,6 +1154,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)\n \tcase GOMP_MAP_POP_MAPPER_NAME:\n \t  pp_string (pp, \"pop_mapper\");\n \t  break;\n+\tcase GOMP_MAP_USES_ALLOCATORS:\n+\t  pp_string (pp, \"uses_allocators\");\n+\t  break;\n \tdefault:\n \t  gcc_unreachable ();\n \t}\ndiff --git a/include/gomp-constants.h b/include/gomp-constants.h\nindex 0a0761043f9..072e7a571bc 100644\n--- a/include/gomp-constants.h\n+++ b/include/gomp-constants.h\n@@ -195,6 +195,11 @@ enum gomp_map_kind\n     GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION\n       =\t\t\t\t\t(GOMP_MAP_DEEP_COPY | 2),\n \n+    /* Used to pass a descriptor for an OpenMP uses_allocators clause,\n+       also serves as an effective 'firstprivate' clause for the allocator\n+       variable.  */\n+    GOMP_MAP_USES_ALLOCATORS =\t\t(GOMP_MAP_DEEP_COPY | 3),\n+\n     /* Internal to GCC, not used in libgomp.  */\n     /* Do not map, but pointer assign a pointer instead.  */\n     GOMP_MAP_FIRSTPRIVATE_POINTER =\t(GOMP_MAP_LAST | 1),\ndiff --git a/libgomp/allocator.c b/libgomp/allocator.c\nindex 8fdaf9bd198..cbfbd6d0333 100644\n--- a/libgomp/allocator.c\n+++ b/libgomp/allocator.c\n@@ -243,14 +243,24 @@ struct omp_allocator_data\n   unsigned int fallback : 8;\n   unsigned int pinned : 1;\n   unsigned int partition : 7;\n-#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA)\n+  /* To unify the format of this type across host/accelerator, enable\n+     this field unconditionally when offload is enabled.  */\n+  #if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA) ||\t\\\n+    defined(LIBGOMP_OFFLOAD_HOST) || defined(LIBGOMP_OFFLOADED_ONLY)\n   unsigned int memkind : 8;\n #endif\n+  /* Note: we now require __sync builtins for offload host/accelerator,\n+     checked during configuration. This lock should never be enabled\n+     for offload configs.  */\n #ifndef HAVE_SYNC_BUILTINS\n   gomp_mutex_t lock;\n #endif\n };\n \n+/* Size of allocator data, exported within libgomp.  */\n+const size_t gomp_omp_allocator_data_size attribute_hidden\n+  = sizeof (struct omp_allocator_data);\n+\n struct omp_mem_header\n {\n   void *ptr;\n@@ -400,38 +410,44 @@ gomp_get_memkind (void)\n }\n #endif\n \n-omp_allocator_handle_t\n-omp_init_allocator (omp_memspace_handle_t memspace, int ntraits,\n-\t\t    const omp_alloctrait_t traits[])\n+static omp_allocator_handle_t\n+gomp_init_allocator_data (struct omp_allocator_data *data,\n+\t\t\t  omp_memspace_handle_t memspace, int ntraits,\n+\t\t\t  const omp_alloctrait_t traits[])\n {\n-  struct omp_allocator_data data\n-    = { memspace, 1, ~(uintptr_t) 0, 0, 0, omp_atv_contended, omp_atv_all,\n-\tomp_atv_default_mem_fb, omp_atv_false, omp_atv_environment,\n-#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA)\n-\tGOMP_MEMKIND_NONE\n+  data->memspace = memspace;\n+  data->alignment = 1;\n+  data->pool_size = ~(uintptr_t) 0;\n+  data->used_pool_size = 0;\n+  data->fb_data = 0;\n+  data->sync_hint = omp_atv_contended;\n+  data->access = omp_atv_all;\n+  data->fallback = omp_atv_default_mem_fb;\n+  data->pinned = omp_atv_false;\n+  data->partition = omp_atv_environment;\n+#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA) ||\t\\\n+  defined(LIBGOMP_OFFLOAD_HOST) || defined(LIBGOMP_OFFLOADED_ONLY)\n+  data->memkind = GOMP_MEMKIND_NONE;\n #endif\n-      };\n-  struct omp_allocator_data *ret;\n-  int i;\n \n   if (memspace > omp_max_predefined_mem_space\n       && (memspace < ompx_gnu_min_predefined_mem_space\n \t  || memspace > ompx_gnu_max_predefined_mem_space))\n     return omp_null_allocator;\n-  for (i = 0; i < ntraits; i++)\n+  for (int i = 0; i < ntraits; i++)\n     switch (traits[i].key)\n       {\n       case omp_atk_sync_hint:\n \tswitch (traits[i].value)\n \t  {\n \t  case omp_atv_default:\n-\t    data.sync_hint = omp_atv_contended;\n+\t    data->sync_hint = omp_atv_contended;\n \t    break;\n \t  case omp_atv_contended:\n \t  case omp_atv_uncontended:\n \t  case omp_atv_serialized:\n \t  case omp_atv_private:\n-\t    data.sync_hint = traits[i].value;\n+\t    data->sync_hint = traits[i].value;\n \t    break;\n \t  default:\n \t    return omp_null_allocator;\n@@ -440,25 +456,25 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits,\n       case omp_atk_alignment:\n         if (traits[i].value == omp_atv_default)\n \t  {\n-\t    data.alignment = 1;\n+\t    data->alignment = 1;\n \t    break;\n \t  }\n \tif ((traits[i].value & (traits[i].value - 1)) != 0\n \t    || !traits[i].value)\n \t  return omp_null_allocator;\n-\tdata.alignment = traits[i].value;\n+\tdata->alignment = traits[i].value;\n \tbreak;\n       case omp_atk_access:\n \tswitch (traits[i].value)\n \t  {\n \t  case omp_atv_default:\n-\t    data.access = omp_atv_all;\n+\t    data->access = omp_atv_all;\n \t    break;\n \t  case omp_atv_all:\n \t  case omp_atv_cgroup:\n \t  case omp_atv_pteam:\n \t  case omp_atv_thread:\n-\t    data.access = traits[i].value;\n+\t    data->access = traits[i].value;\n \t    break;\n \t  default:\n \t    return omp_null_allocator;\n@@ -466,38 +482,38 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits,\n \tbreak;\n       case omp_atk_pool_size:\n \tif (traits[i].value == omp_atv_default)\n-\t  data.pool_size = ~(uintptr_t) 0;\n+\t  data->pool_size = ~(uintptr_t) 0;\n \telse\n-\t  data.pool_size = traits[i].value;\n+\t  data->pool_size = traits[i].value;\n \tbreak;\n       case omp_atk_fallback:\n \tswitch (traits[i].value)\n \t  {\n \t  case omp_atv_default:\n-\t    data.fallback = omp_atv_default_mem_fb;\n+\t    data->fallback = omp_atv_default_mem_fb;\n \t    break;\n \t  case omp_atv_default_mem_fb:\n \t  case omp_atv_null_fb:\n \t  case omp_atv_abort_fb:\n \t  case omp_atv_allocator_fb:\n-\t    data.fallback = traits[i].value;\n+\t    data->fallback = traits[i].value;\n \t    break;\n \t  default:\n \t    return omp_null_allocator;\n \t  }\n \tbreak;\n       case omp_atk_fb_data:\n-\tdata.fb_data = traits[i].value;\n+\tdata->fb_data = traits[i].value;\n \tbreak;\n       case omp_atk_pinned:\n \tswitch (traits[i].value)\n \t  {\n \t  case omp_atv_default:\n \t  case omp_atv_false:\n-\t    data.pinned = omp_atv_false;\n+\t    data->pinned = omp_atv_false;\n \t    break;\n \t  case omp_atv_true:\n-\t    data.pinned = omp_atv_true;\n+\t    data->pinned = omp_atv_true;\n \t    break;\n \t  default:\n \t    return omp_null_allocator;\n@@ -507,13 +523,13 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits,\n \tswitch (traits[i].value)\n \t  {\n \t  case omp_atv_default:\n-\t    data.partition = omp_atv_environment;\n+\t    data->partition = omp_atv_environment;\n \t    break;\n \t  case omp_atv_environment:\n \t  case omp_atv_nearest:\n \t  case omp_atv_blocked:\n \t  case omp_atv_interleaved:\n-\t    data.partition = traits[i].value;\n+\t    data->partition = traits[i].value;\n \t    break;\n \t  default:\n \t    return omp_null_allocator;\n@@ -523,8 +539,8 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits,\n \treturn omp_null_allocator;\n       }\n \n-  if (data.alignment < sizeof (void *))\n-    data.alignment = sizeof (void *);\n+  if (data->alignment < sizeof (void *))\n+    data->alignment = sizeof (void *);\n \n   switch (memspace)\n     {\n@@ -532,57 +548,68 @@ omp_init_allocator (omp_memspace_handle_t memspace, int ntraits,\n     case omp_high_bw_mem_space:\n       struct gomp_memkind_data *memkind_data;\n       memkind_data = gomp_get_memkind ();\n-      if (data.partition == omp_atv_interleaved\n+      if (data->partition == omp_atv_interleaved\n \t  && memkind_data->kinds[GOMP_MEMKIND_HBW_INTERLEAVE])\n \t{\n-\t  data.memkind = GOMP_MEMKIND_HBW_INTERLEAVE;\n+\t  data->memkind = GOMP_MEMKIND_HBW_INTERLEAVE;\n \t  break;\n \t}\n       else if (memkind_data->kinds[GOMP_MEMKIND_HBW_PREFERRED])\n \t{\n-\t  data.memkind = GOMP_MEMKIND_HBW_PREFERRED;\n+\t  data->memkind = GOMP_MEMKIND_HBW_PREFERRED;\n \t  break;\n \t}\n       break;\n     case omp_large_cap_mem_space:\n       memkind_data = gomp_get_memkind ();\n       if (memkind_data->kinds[GOMP_MEMKIND_DAX_KMEM_ALL])\n-\tdata.memkind = GOMP_MEMKIND_DAX_KMEM_ALL;\n+\tdata->memkind = GOMP_MEMKIND_DAX_KMEM_ALL;\n       else if (memkind_data->kinds[GOMP_MEMKIND_DAX_KMEM])\n-\tdata.memkind = GOMP_MEMKIND_DAX_KMEM;\n+\tdata->memkind = GOMP_MEMKIND_DAX_KMEM;\n       break;\n #endif\n     default:\n #ifdef LIBGOMP_USE_MEMKIND\n-      if (data.partition == omp_atv_interleaved)\n+      if (data->partition == omp_atv_interleaved)\n \t{\n \t  memkind_data = gomp_get_memkind ();\n \t  if (memkind_data->kinds[GOMP_MEMKIND_INTERLEAVE])\n-\t    data.memkind = GOMP_MEMKIND_INTERLEAVE;\n+\t    data->memkind = GOMP_MEMKIND_INTERLEAVE;\n \t}\n #endif\n       break;\n     }\n \n #ifdef LIBGOMP_USE_LIBNUMA\n-  if (data.memkind == GOMP_MEMKIND_NONE && data.partition == omp_atv_nearest)\n+  if (data->memkind == GOMP_MEMKIND_NONE && data->partition == omp_atv_nearest)\n     {\n       libnuma_data = gomp_get_libnuma ();\n       if (libnuma_data->numa_alloc_local != NULL)\n-\tdata.memkind = GOMP_MEMKIND_LIBNUMA;\n+\tdata->memkind = GOMP_MEMKIND_LIBNUMA;\n     }\n #endif\n \n   /* Reject unsupported memory spaces.  */\n-  if (!MEMSPACE_VALIDATE (data.memspace, data.access, data.pinned))\n+  if (!MEMSPACE_VALIDATE (data->memspace, data->access, data->pinned))\n     return omp_null_allocator;\n \n-  ret = gomp_malloc (sizeof (struct omp_allocator_data));\n-  *ret = data;\n-#ifndef HAVE_SYNC_BUILTINS\n-  gomp_mutex_init (&ret->lock);\n-#endif\n-  return (omp_allocator_handle_t) ret;\n+  return (omp_allocator_handle_t) data;\n+}\n+\n+omp_allocator_handle_t\n+omp_init_allocator (omp_memspace_handle_t memspace, int ntraits,\n+\t\t    const omp_alloctrait_t traits[])\n+{\n+  struct omp_allocator_data *data\n+    = gomp_malloc (sizeof (struct omp_allocator_data));\n+\n+  omp_allocator_handle_t ret\n+    = gomp_init_allocator_data (data, memspace, ntraits, traits);\n+\n+  if (ret != omp_null_allocator)\n+    free (data);\n+\n+  return ret;\n }\n \n void\n@@ -1496,3 +1523,51 @@ fail:;\n     }\n   return NULL;\n }\n+\n+#if !defined(LIBGOMP_OFFLOADED_ONLY)\n+\n+/* Called from gomp_maps_vars, used to implement the effective\n+   omp_init_allocator() call in an uses_allocators clause.\n+\n+   The memspace_validate hook is also used here to check things for\n+   the offload accelerator on the host side.\n+*/\n+attribute_hidden uintptr_t\n+gomp_map_omp_init_allocator (struct gomp_device_descr *devicep,\n+\t\t\t     struct goacc_asyncqueue *aq,\n+\t\t\t     struct gomp_coalesce_buf *cbuf,\n+\t\t\t     void *allocator_data_devaddr, void *descr_ptr)\n+{\n+  uintptr_t *descr = (uintptr_t *) descr_ptr;\n+\n+  omp_memspace_handle_t memspace = (omp_memspace_handle_t) descr[0];\n+  int ntraits = (int) descr[1];\n+  omp_alloctrait_t *traits = (omp_alloctrait_t *) descr[2];\n+\n+  struct omp_allocator_data data;\n+  omp_allocator_handle_t ret = gomp_init_allocator_data (&data, memspace,\n+\t\t\t\t\t\t\t ntraits, traits);\n+  if (ret != omp_null_allocator)\n+    {\n+      if (devicep)\n+\t{\n+\t  /* Do memspace validation for the offload target, using the\n+\t     offload plugin.  */\n+\t  if (devicep->memspace_validate_func\n+\t      && !devicep->memspace_validate_func (data.memspace, data.access))\n+\t    return (uintptr_t) omp_null_allocator;\n+\n+\t  /* Copy to device. This is now the device-side omp_allocator_data.  */\n+\t  gomp_copy_host2dev (devicep, aq, allocator_data_devaddr, &data,\n+\t\t\t      sizeof (struct omp_allocator_data), true, cbuf);\n+\t}\n+      else\n+\t/* Used for host fallback case.  */\n+\tmemcpy (allocator_data_devaddr, &data,\n+\t\tsizeof (struct omp_allocator_data));\n+\n+      ret = (omp_allocator_handle_t) allocator_data_devaddr;\n+    }\n+  return (uintptr_t) ret;\n+}\n+#endif\ndiff --git a/libgomp/config/nvptx/allocator.c b/libgomp/config/nvptx/allocator.c\nindex ebbe9816909..8bb18daa0d2 100644\n--- a/libgomp/config/nvptx/allocator.c\n+++ b/libgomp/config/nvptx/allocator.c\n@@ -54,6 +54,9 @@ asm (\".extern .shared .u8 __nvptx_lowlat_pool[];\\n\");\n static void *\n nvptx_memspace_alloc (omp_memspace_handle_t memspace, size_t size)\n {\n+#if __PTX_ISA_VERSION_MAJOR__ > 4\t\t\t\t\t\\\n+    || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1)\n+  /* Low-latency memory is not available before PTX 4.1.  */\n   if (memspace == omp_low_lat_mem_space)\n     {\n       char *shared_pool;\n@@ -65,12 +68,16 @@ nvptx_memspace_alloc (omp_memspace_handle_t memspace, size_t size)\n     /* No non-standard memspaces are implemented for device-side nvptx.  */\n     return NULL;\n   else\n+#endif\n     return malloc (size);\n }\n \n static void *\n nvptx_memspace_calloc (omp_memspace_handle_t memspace, size_t size)\n {\n+#if __PTX_ISA_VERSION_MAJOR__ > 4\t\t\t\t\t\\\n+    || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1)\n+  /* Low-latency memory is not available before PTX 4.1.  */\n   if (memspace == omp_low_lat_mem_space)\n     {\n       char *shared_pool;\n@@ -82,12 +89,16 @@ nvptx_memspace_calloc (omp_memspace_handle_t memspace, size_t size)\n     /* No non-standard memspaces are implemented for device-side nvptx.  */\n     return NULL;\n   else\n+#endif\n     return calloc (1, size);\n }\n \n static void\n nvptx_memspace_free (omp_memspace_handle_t memspace, void *addr, size_t size)\n {\n+#if __PTX_ISA_VERSION_MAJOR__ > 4\t\t\t\t\t\\\n+    || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1)\n+  /* Low-latency memory is not available before PTX 4.1.  */\n   if (memspace == omp_low_lat_mem_space)\n     {\n       char *shared_pool;\n@@ -96,6 +107,7 @@ nvptx_memspace_free (omp_memspace_handle_t memspace, void *addr, size_t size)\n       __nvptx_lowlat_free (shared_pool, addr, size);\n     }\n   else\n+#endif\n     free (addr);\n }\n \n@@ -103,6 +115,9 @@ static void *\n nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr,\n \t\t\tsize_t oldsize, size_t size)\n {\n+#if __PTX_ISA_VERSION_MAJOR__ > 4 \\\n+    || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MINOR >= 1)\n+  /* Low-latency memory is not available before PTX 4.1.  */\n   if (memspace == omp_low_lat_mem_space)\n     {\n       char *shared_pool;\n@@ -110,11 +125,15 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr,\n \n       return __nvptx_lowlat_realloc (shared_pool, addr, oldsize, size);\n     }\n-  else if (memspace > GOMP_OMP_PREDEF_MEMSPACE_MAX)\n-    /* No non-standard memspaces are implemented for device-side nvptx.  */\n-    return NULL;\n   else\n-    return realloc (addr, size);\n+#endif\n+    {\n+      if (memspace > GOMP_OMP_PREDEF_MEMSPACE_MAX)\n+\t/* No non-standard memspaces are implemented for device-side nvptx.  */\n+\treturn NULL;\n+      else\n+\treturn realloc (addr, size);\n+    }\n }\n \n static inline int\ndiff --git a/libgomp/configure b/libgomp/configure\nindex 56cffe79634..3a9d6a76eef 100755\n--- a/libgomp/configure\n+++ b/libgomp/configure\n@@ -15782,6 +15782,12 @@ fi\n \n \n \n+if test x\"$offload_plugins\" != x; then\n+\n+$as_echo \"#define LIBGOMP_OFFLOAD_HOST 1\" >>confdefs.h\n+\n+fi\n+\n # Check for functions needed.\n for ac_func in getloadavg clock_gettime strtoull\n do :\n@@ -17120,6 +17126,17 @@ $as_echo \"#define HAVE_SYNC_BUILTINS 1\" >>confdefs.h\n \n   fi\n \n+if test x$libgomp_cv_have_sync_builtins = xno; then\n+  # We require accelerator targets to support __sync_* builtins.\n+  if test x$libgomp_offloaded_only = xyes; then\n+    as_fn_error $? \"accelerator targets require __sync_val_compare_and_swap to build libgomp.\" \"$LINENO\" 5\n+  fi\n+  # Same for offload hosts.\n+  if test x\"$offload_plugins\" = x; then\n+    as_fn_error $? \"offload hosts require __sync_val_compare_and_swap to build libgomp.\" \"$LINENO\" 5\n+  fi\n+fi\n+\n XCFLAGS=\"$XCFLAGS$XPCFLAGS\"\n \n # Add CET specific flags if CET is enabled\ndiff --git a/libgomp/configure.ac b/libgomp/configure.ac\nindex 1730c62c74c..7f6df0bfe85 100644\n--- a/libgomp/configure.ac\n+++ b/libgomp/configure.ac\n@@ -233,6 +233,11 @@ AC_CHECK_SIZEOF([void *])\n \n m4_include([plugin/configfrag.ac])\n \n+if test x\"$offload_plugins\" != x; then\n+  AC_DEFINE(LIBGOMP_OFFLOAD_HOST, 1,\n+    [Define to 1 if building libgomp for a offload host])\n+fi\n+\n # Check for functions needed.\n AC_CHECK_FUNCS(getloadavg clock_gettime strtoull)\n AC_CHECK_FUNCS(aligned_alloc posix_memalign memalign _aligned_malloc)\n@@ -365,6 +370,17 @@ CFLAGS=\"$save_CFLAGS $XCFLAGS\"\n # had a chance to set XCFLAGS.\n LIBGOMP_CHECK_SYNC_BUILTINS\n \n+if test x$libgomp_cv_have_sync_builtins = xno; then\n+  # We require accelerator targets to support __sync_* builtins.\n+  if test x$libgomp_offloaded_only = xyes; then\n+    AC_MSG_ERROR([accelerator targets require __sync_val_compare_and_swap to build libgomp.])\n+  fi\n+  # Same for offload hosts.\n+  if test x\"$offload_plugins\" = x; then\n+    AC_MSG_ERROR([offload hosts require __sync_val_compare_and_swap to build libgomp.])\n+  fi\n+fi\n+\n XCFLAGS=\"$XCFLAGS$XPCFLAGS\"\n \n # Add CET specific flags if CET is enabled\ndiff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h\nindex bb4d577b66d..358c188516c 100644\n--- a/libgomp/libgomp-plugin.h\n+++ b/libgomp/libgomp-plugin.h\n@@ -187,6 +187,7 @@ extern int GOMP_OFFLOAD_memcpy3d (int, int, size_t, size_t, size_t, void *,\n \t\t\t\t  const void *, size_t, size_t, size_t, size_t,\n \t\t\t\t  size_t);\n extern bool GOMP_OFFLOAD_memset (int, void *, int, size_t);\n+extern int GOMP_OFFLOAD_memspace_validate (omp_memspace_handle_t, unsigned int);\n extern bool GOMP_OFFLOAD_can_run (void *);\n extern void GOMP_OFFLOAD_run (int, void *, void *, void **);\n extern void GOMP_OFFLOAD_async_run (int, void *, void *, void **, void *);\ndiff --git a/libgomp/libgomp.h b/libgomp/libgomp.h\nindex 42f32439295..e8a8fb3e727 100644\n--- a/libgomp/libgomp.h\n+++ b/libgomp/libgomp.h\n@@ -614,6 +614,7 @@ extern bool gomp_display_affinity_var;\n extern char *gomp_affinity_format_var;\n extern size_t gomp_affinity_format_len;\n extern uintptr_t gomp_def_allocator;\n+extern const size_t gomp_omp_allocator_data_size;\n extern const struct gomp_default_icv gomp_default_icv_values;\n extern struct gomp_icv_list *gomp_initial_icv_list;\n extern struct gomp_offload_icv_list *gomp_offload_icv_list;\n@@ -1434,6 +1435,7 @@ struct gomp_device_descr\n   __typeof (GOMP_OFFLOAD_memcpy2d) *memcpy2d_func;\n   __typeof (GOMP_OFFLOAD_memcpy3d) *memcpy3d_func;\n   __typeof (GOMP_OFFLOAD_memset) *memset_func;\n+  __typeof (GOMP_OFFLOAD_memspace_validate) *memspace_validate_func;\n   __typeof (GOMP_OFFLOAD_can_run) *can_run_func;\n   __typeof (GOMP_OFFLOAD_run) *run_func;\n   __typeof (GOMP_OFFLOAD_async_run) *async_run_func;\n@@ -1503,6 +1505,13 @@ extern bool gomp_remove_var (struct gomp_device_descr *, splay_tree_key);\n extern void gomp_remove_var_async (struct gomp_device_descr *, splay_tree_key,\n \t\t\t\t   struct goacc_asyncqueue *);\n \n+/* allocator.c */\n+\n+extern uintptr_t gomp_map_omp_init_allocator (struct gomp_device_descr *,\n+\t\t\t\t\t      struct goacc_asyncqueue *,\n+\t\t\t\t\t      struct gomp_coalesce_buf *,\n+\t\t\t\t\t      void *, void *);\n+\n /* work.c */\n \n extern void gomp_init_work_share (struct gomp_work_share *, size_t, unsigned);\ndiff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c\nindex 028a5c943b7..77dec471fee 100644\n--- a/libgomp/oacc-host.c\n+++ b/libgomp/oacc-host.c\n@@ -286,6 +286,7 @@ static struct gomp_device_descr host_dispatch =\n     .host2dev_func = host_host2dev,\n     .memcpy2d_func = NULL,\n     .memcpy3d_func = NULL,\n+    .memspace_validate_func = NULL,\n     .run_func = host_run,\n \n     .mem_map = { NULL },\ndiff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c\nindex 32f573f1b7f..f52e4ac8092 100644\n--- a/libgomp/plugin/plugin-gcn.c\n+++ b/libgomp/plugin/plugin-gcn.c\n@@ -4718,6 +4718,19 @@ unlock:\n   return retval;\n }\n \n+int\n+GOMP_OFFLOAD_memspace_validate (omp_memspace_handle_t memspace, unsigned access)\n+{\n+  /* Disallow use of low-latency memory when it must be accessible by\n+     all threads.  */\n+  return (memspace != omp_low_lat_mem_space\n+\t  || access != omp_atv_all);\n+\n+  /* Otherwise, standard memspaces are accepted, even when we don't have\n+     anything special to do with them, and non-standard memspaces are assumed\n+     to need explicit support.  */\n+  return (memspace <= GOMP_OMP_PREDEF_MEMSPACE_MAX);\n+}\n \n static bool\n init_hip_runtime_functions (void)\ndiff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c\nindex a540e9d4cce..914b33f35d4 100644\n--- a/libgomp/plugin/plugin-nvptx.c\n+++ b/libgomp/plugin/plugin-nvptx.c\n@@ -2468,6 +2468,21 @@ GOMP_OFFLOAD_memset (int ord, void *ptr, int val, size_t count)\n   return true;\n }\n \n+int\n+GOMP_OFFLOAD_memspace_validate (omp_memspace_handle_t memspace, unsigned access)\n+{\n+  /* Disallow use of low-latency memory when it must be accessible by\n+     all threads.  */\n+  if (memspace == omp_low_lat_mem_space\n+      && access == omp_atv_all)\n+    return false;\n+\n+  /* Otherwise, standard memspaces are accepted, even when we don't have\n+     anything special to do with them, and non-standard memspaces are assumed\n+     to need explicit support.  */\n+  return (memspace <= GOMP_OMP_PREDEF_MEMSPACE_MAX);\n+}\n+\n bool\n GOMP_OFFLOAD_openacc_async_host2dev (int ord, void *dst, const void *src,\n \t\t\t\t     size_t n, struct goacc_asyncqueue *aq)\ndiff --git a/libgomp/target.c b/libgomp/target.c\nindex d562b0493ea..4fb2aaa222e 100644\n--- a/libgomp/target.c\n+++ b/libgomp/target.c\n@@ -1249,6 +1249,23 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,\n \t{\n \t  tgt->list[i].key = NULL;\n \t  tgt->list[i].offset = OFFSET_INLINED;\n+\t  has_firstprivate = true;\n+\t  continue;\n+\t}\n+      else if ((kind & typemask) == GOMP_MAP_USES_ALLOCATORS)\n+\t{\n+\t  tgt->list[i].key = NULL;\n+\t  tgt->list[i].offset = OFFSET_INLINED;\n+\n+\t  size_t align = (size_t) 1 << (kind >> rshift);\n+\t  if (tgt_align < align)\n+\t    tgt_align = align;\n+\t  tgt_size = (tgt_size + align - 1) & ~(align - 1);\n+\n+\t  /* Allocate space for omp_allocator_data.  */\n+\t  tgt_size += gomp_omp_allocator_data_size;\n+\n+\t  has_firstprivate = true;\n \t  continue;\n \t}\n       else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR\n@@ -1680,6 +1697,21 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,\n \t\t    tgt->list[i].offset = OFFSET_INLINED;\n \t\t  }\n \t\tcontinue;\n+\n+\t      case GOMP_MAP_USES_ALLOCATORS:\n+\t\talign = (size_t) 1 << (kind >> rshift);\n+\t\ttgt_size = (tgt_size + align - 1) & ~(align - 1);\n+\t\ttgt->list[i].offset = tgt_size;\n+\n+\t\tvoid *descr_ptr = hostaddrs[i];\n+\t\tvoid *devaddr = (void *) (tgt->tgt_start + tgt_size);\n+\t\tuintptr_t dev_allocator\n+\t\t  = gomp_map_omp_init_allocator (devicep, aq, cbufp, devaddr,\n+\t\t\t\t\t\t descr_ptr);\n+\t\thostaddrs[i] = (void *) dev_allocator;\n+\t\ttgt_size += gomp_omp_allocator_data_size;\n+\t\tcontinue;\n+\n \t      case GOMP_MAP_STRUCT_UNORD:\n \t\tif (sizes[i] > 1)\n \t\t  {\n@@ -3116,6 +3148,14 @@ calculate_firstprivate_requirements (size_t mapnum, size_t *sizes,\n \t*tgt_size = (*tgt_size + align - 1) & ~(align - 1);\n \t*tgt_size += sizes[i];\n       }\n+    else if ((kinds[i] & 0xff) == GOMP_MAP_USES_ALLOCATORS)\n+      {\n+\tsize_t align = (size_t) 1 << (kinds[i] >> 8);\n+\tif (*tgt_align < align)\n+\t  *tgt_align = align;\n+\t*tgt_size = (*tgt_size + align - 1) & ~(align - 1);\n+\t*tgt_size += gomp_omp_allocator_data_size;\n+      }\n }\n \n /* Copy data shared as GOMP_MAP_FIRSTPRIVATE to DST.  */\n@@ -3138,12 +3178,23 @@ copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs,\n \tmemcpy (tgt + tgt_size, hostaddrs[i], sizes[i]);\n \thostaddrs[i] = tgt + tgt_size;\n \ttgt_size = tgt_size + sizes[i];\n-\tif (i + 1 < mapnum && (kinds[i+1] & 0xff) == GOMP_MAP_ATTACH)\n+\tif (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH)\n \t  {\n-\t    *(*(uintptr_t**) hostaddrs[i+1] + sizes[i+1]) = (uintptr_t) hostaddrs[i];\n+\t    *(*(uintptr_t**) hostaddrs[i + 1] + sizes[i + 1])\n+\t      = (uintptr_t) hostaddrs[i];\n \t    ++i;\n \t  }\n       }\n+    else if ((kinds[i] & 0xff) == GOMP_MAP_USES_ALLOCATORS)\n+      {\n+\tassert (hostaddrs[i] != NULL);\n+\tsize_t align = (size_t) 1 << (kinds[i] >> 8);\n+\ttgt_size = (tgt_size + align - 1) & ~(align - 1);\n+\tuintptr_t allocator = gomp_map_omp_init_allocator (NULL, NULL, NULL,\n+\t\t\t\t\t\t\t   tgt + tgt_size,\n+\t\t\t\t\t\t\t   hostaddrs[i]);\n+\thostaddrs[i] = (void *) allocator;\n+      }\n }\n \n /* Helper function of GOMP_target{,_ext} routines.  */\n@@ -6056,6 +6107,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,\n   DLSYM (host2dev);\n   DLSYM_OPT (memcpy2d, memcpy2d);\n   DLSYM_OPT (memcpy3d, memcpy3d);\n+  DLSYM_OPT (memspace_validate, memspace_validate);\n   if (DLSYM_OPT (interop, interop))\n     {\n       DLSYM (get_interop_int);\n",
    "prefixes": [
        "OpenMP",
        "v2"
    ]
}