Patch Detail
get:
Show a patch.
patch:
Update a patch.
put:
Update a patch.
GET /api/patches/2216549/?format=api
{ "id": 2216549, "url": "http://patchwork.ozlabs.org/api/patches/2216549/?format=api", "web_url": "http://patchwork.ozlabs.org/project/gcc/patch/2cdb2430-557a-447f-b6da-27ce2373194f@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": "<2cdb2430-557a-447f-b6da-27ce2373194f@baylibre.com>", "list_archive_url": null, "date": "2026-03-26T17:43:34", "name": "libgomp: fix omp_target_is_present and omp_get_mapped_ptr", "commit_ref": null, "pull_url": null, "state": "new", "archived": false, "hash": "a5d0a11e7b10c41458ba638a27cd6486622464f7", "submitter": { "id": 87873, "url": "http://patchwork.ozlabs.org/api/people/87873/?format=api", "name": "Tobias Burnus", "email": "tburnus@baylibre.com" }, "delegate": null, "mbox": "http://patchwork.ozlabs.org/project/gcc/patch/2cdb2430-557a-447f-b6da-27ce2373194f@baylibre.com/mbox/", "series": [ { "id": 497635, "url": "http://patchwork.ozlabs.org/api/series/497635/?format=api", "web_url": "http://patchwork.ozlabs.org/project/gcc/list/?series=497635", "date": "2026-03-26T17:43:34", "name": "libgomp: fix omp_target_is_present and omp_get_mapped_ptr", "version": 1, "mbox": "http://patchwork.ozlabs.org/series/497635/mbox/" } ], "comments": "http://patchwork.ozlabs.org/api/patches/2216549/comments/", "check": "pending", "checks": "http://patchwork.ozlabs.org/api/patches/2216549/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.20230601.gappssmtp.com\n header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256\n header.s=20230601 header.b=Kf+ovSsr;\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.20230601.gappssmtp.com\n header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256\n header.s=20230601 header.b=Kf+ovSsr", "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.49" ], "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 4fhWNt0Wh8z1y1G\n\tfor <incoming@patchwork.ozlabs.org>; Fri, 27 Mar 2026 04:44:14 +1100 (AEDT)", "from vm01.sourceware.org (localhost [127.0.0.1])\n\tby sourceware.org (Postfix) with ESMTP id 032FE4BA23E6\n\tfor <incoming@patchwork.ozlabs.org>; Thu, 26 Mar 2026 17:44:12 +0000 (GMT)", "from mail-wm1-f49.google.com (mail-wm1-f49.google.com\n [209.85.128.49])\n by sourceware.org (Postfix) with ESMTPS id 7069E4BA23F9\n for <gcc-patches@gcc.gnu.org>; Thu, 26 Mar 2026 17:43:38 +0000 (GMT)", "by mail-wm1-f49.google.com with SMTP id\n 5b1f17b1804b1-4853e1ce427so14455005e9.3\n for <gcc-patches@gcc.gnu.org>; Thu, 26 Mar 2026 10:43:38 -0700 (PDT)", "from [10.1.5.249] (88-127-129-70.subs.proxad.net. [88.127.129.70])\n by smtp.gmail.com with ESMTPSA id\n 5b1f17b1804b1-48722be608bsm56551345e9.0.2026.03.26.10.43.35\n (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128);\n Thu, 26 Mar 2026 10:43:36 -0700 (PDT)" ], "DKIM-Filter": [ "OpenDKIM Filter v2.11.0 sourceware.org 032FE4BA23E6", "OpenDKIM Filter v2.11.0 sourceware.org 7069E4BA23F9" ], "DMARC-Filter": "OpenDMARC Filter v1.4.2 sourceware.org 7069E4BA23F9", "ARC-Filter": "OpenARC Filter v1.0.0 sourceware.org 7069E4BA23F9", "ARC-Seal": "i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1774547018; cv=none;\n b=GFiVRB7nv/BQpBO4/vq4H63Rm9VNpqL2+xBf98nRWl7kHznyL/e/Rgn7AOj4+qjMvhGCxGv9hY4gRmBEjDWvY5Zl7rmaOKmihMcZwkh+UR+mMRrGp3VCv4+M9sfbQ6wHLMrEKXYrbdjLPfkFjQAD/xCc/RvxFFL8+Blmg1OYH7w=", "ARC-Message-Signature": "i=1; a=rsa-sha256; d=sourceware.org; s=key;\n t=1774547018; c=relaxed/simple;\n bh=8mPun/95wOG3SbqGD+VX9XCpbot2WxFSbhJ7htmEKRk=;\n h=DKIM-Signature:Message-ID:Date:MIME-Version:To:From:Subject;\n b=GeeUSPSQ0M0/rTmlr7YdDS+NVSY1sy6vSe5CnJUYG65vfwsNmJ83GP0SKMf88b2F4+JTLyNZmaoxv7oBv/1WpAjgEpQ11cnm7WnDeCmkw0OrVRhEJ9RSt1X+8juVz8tzglqY2dbnlyWCHpnaucLcJ8J6j/XSqPOnEzH1BORkGsU=", "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=1774547017;\n x=1775151817;\n darn=gcc.gnu.org;\n h=subject:cc:from:to:content-language:user-agent:mime-version:date\n :message-id:from:to:cc:subject:date:message-id:reply-to;\n bh=fKTT1uvqAZGmGfqH6uYen76MBigH/RBlfFSA3ENyX2M=;\n b=Kf+ovSsrfluSDCL0EiZ0oLmPaUkSezWnVKfdyGeFK6Rr/5QNqxBOUgKPfX1jsfAQUR\n yFu+rF6XjX35wtctcXS+KZ9ultfH902YMZlKYF3UhD6QXR81g1bdcfUobkvkkn1gX3gq\n KpG0xW3zJWML95GWm7tYHNh0j+G+rCMP933kNn9saLkhA7L8xU4lN0Fc783InkP3ODkI\n 0xhOdfBiOxvnuBsGLzYdOciQsbNBdTGLYsfLi2Q3TIE65bAQ8IIZKATouLB/hyWEQeTV\n 4dLomwBCuqqF9uAKbAE5a6slRl4ssehnsi2dVaM2oy//o8GbHNx7LNz5M/dZtYgosWRP\n 6LbQ==", "X-Google-DKIM-Signature": "v=1; a=rsa-sha256; c=relaxed/relaxed;\n d=1e100.net; s=20251104; t=1774547017; x=1775151817;\n h=subject:cc: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=fKTT1uvqAZGmGfqH6uYen76MBigH/RBlfFSA3ENyX2M=;\n b=TosJxJM+OMeylgdPAz+fYFfdaODIHCptpto25A7HDHxNvIUsDYOQr5mweeGZVVoWT5\n pXdJc2gYzO70uL+LMG7maomIMGSe1iytYiDG86t81e8n7HPj+GmlXV8nqmvyJrCeMbES\n /p+9S9Hjq9o7ZjUA4d9UZrHai/uYn+66hYo9iai2FluSLY8N72xhVgdHE1OJXe0TgJKO\n MnDxCdxdM3M13oD3p5ThFVh5qNf9efYwFTBj7NoHbywY2nrdS4UKqCKET0ygz+T7KJIH\n cx4odMEpW1+/Sz7CxwsqpeduLlhtgwS+SS7m84vck8+xhB+dVrt+D+e7wCSDl0NkaoHY\n dytg==", "X-Gm-Message-State": "AOJu0YwDf+43nusNwU6cd1heeb6sKXJv+T4+ZomMffJnL+Ml/2JwhYe9\n PAeR3viG9+dUoITHT2gMkV2Cml+LAqBQSr5B0pr4AM10sfAETv2jPeBoHkew4FdvtniXv5ctbyT\n ItCa+p4U=", "X-Gm-Gg": "ATEYQzz6XiUIIEI1U/Ta+XGCOSgz5eXaaMZIdZeBv9laayR/lKmpVSd6X0RcDg7whSW\n 7HWnXo8c+R7t/iZyNkfS2AtJnU5IyyORF48eQkxIpwJoWwZU5Tv71a/b7H2resI5JymyMUXqKOG\n tK0N+J4YcytXu1j5R3XTgNf/kXPEGZEwwvC5XYoctJxuoLWIMtyIIok4Z00+E2TbN/OS50sJpHX\n l5jMbdy8WlnY+CeIxDLifw6lgD7phIj8ucoY3llrEH+jabiOK3Sb9N4/3SFP+t1J0XG2e6N96qu\n 58hb1CvJ/rCDMBDOf2zRoid4vJtccYNOgN5PFca1l9erYV8juajsnomTKoOcgI/LRrVC8nH42Tu\n vsEhILtVl1pzq1aFM4y4yTueHbfI977gvXnjFe9EF0YXfAAimbdoTh5dyd1I8yqMUhxOErfwR1j\n tg3T/5gfaiK+QtEGTWWqmoXyJ6ngH/G6rVIhOsJnFdPAUOyqbSlfLrQVQ=", "X-Received": "by 2002:a05:600c:a12:b0:485:481c:e7bb with SMTP id\n 5b1f17b1804b1-4871605645cmr128230715e9.20.1774547017310;\n Thu, 26 Mar 2026 10:43:37 -0700 (PDT)", "Content-Type": "multipart/mixed; boundary=\"------------pW84lvU43dv5FK17OHGtRptw\"", "Message-ID": "<2cdb2430-557a-447f-b6da-27ce2373194f@baylibre.com>", "Date": "Thu, 26 Mar 2026 18:43:34 +0100", "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>", "Cc": "Sandra Loosemore <sloosemore@baylibre.com>", "Subject": "[Patch] libgomp: fix omp_target_is_present and omp_get_mapped_ptr", "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 patch fixes some corner cases issues with the two routines; it\nis rather simple and only affects those (mainly: NULL and USM use).\nPlus tweaks the documentation a bit.\n\nSeemingly, there was some confusion whether being host accessible also\ncounts as being 'present', but the routines should only work with\nactually mapped (corresponding storage). Additionally, some corner cases\nlike whether a NULL pointer is present were clarified (it is not).\n(For the get_mapped function, it has only been moved to avoid going\nto locking etc. when the result is known: NULL is always NULL, whether\nbeing the original pointer or being a way to denote an error.)\n\nAnd for unified-shared memory, some tweaking was required, especially as\nsome memory still has corresponding storage: all memory with 'declare target',\nexcept for the 'link' clause for which we link to the host memory.\n\nThe 'link' clause handling has also been clarified in the manual; it\nis not completely clear to me whether the pointer or pointee should\nbe regarded in this case, either kind of makes sense. Documented is\nnow what GCC does, which seems to make sense.\n\n\nIn particular, the commit fixes a wrong-code and potential crash issue\nfor code like the following; the first one has been taken from such a\ncode and failed with 'requires unified_shared_memory':\n\n if (omp_target_is_present(D, default_device_omp_target_call)) {\n int rc = omp_target_disassociate_ptr(D, default_device_omp_target_call);\n\nalso affected would be code such as:\n\n if (omp_target_is_present (...)\n omp_target_memcpy (omp_get_mapped_ptr (ptr), ptr, ...);\n\nComment, questions, concerns before it gets applied?\n\nTobias", "diff": "libgomp: fix omp_target_is_present and omp_get_mapped_ptr\n\nThere were a few minor issues with the two routines, partially because of\nnot handling corner cases and partially some clarifications are only in\nnewer versions of the spec.\n\nIn particular, for omp_target_is_present\n* NULL pointer aren't regarded as present\n* For (unified-)shared memory, claiming that something has always corresponding\n storage is wrong - it mostly never has. (but it is omp_target_is_accessible).\n* Even with shared memory, 'declare target' usually has device memory. For\n 'link' it is made to point to the host, i.e. it is not mapped, all others\n are still mapped. (With 'requires self_mapping', 'enter' should also not be\n mapped (and turned internally to 'link'), only 'local' needs to be mapped.)\n\nFor omp_get_mapped_ptr\n* For NULL we can return NULL early also for devices.\n* For shared memory, we shouldn't touch link (it is not counting as mapped);\n hence return NULL for it.\n\nThe documentation was updated add some missing cross references as the more\nuseful ones were missing. Additionally, the description for the two modified\nroutines has been updated.\n\nlibgomp/ChangeLog:\n\n\t* target.c (omp_target_is_present, omp_get_mapped_ptr): Update handling\n\tfor nullptr and shared-memory devices.\n\t* libgomp.texi (omp_target_is_present, omp_get_mapped_ptr): Update\n\tdescription, add see-also @refs.\n\t(omp_target_is_accessible, omp_target_associate_ptr): Add see-also\n\t@refs.\n\t* testsuite/libgomp.c/omp_target_is_present.c: New test.\n\t* testsuite/libgomp.c/omp_target_is_present-2.c: New test.\n\n libgomp/libgomp.texi | 35 ++--\n libgomp/target.c | 37 ++--\n .../testsuite/libgomp.c/omp_target_is_present-2.c | 7 +\n .../testsuite/libgomp.c/omp_target_is_present.c | 206 +++++++++++++++++++++\n 4 files changed, 254 insertions(+), 31 deletions(-)\n\ndiff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi\nindex 258ea8a7619..18877b23c73 100644\n--- a/libgomp/libgomp.texi\n+++ b/libgomp/libgomp.texi\n@@ -2122,18 +2122,9 @@ is not supported.\n @item @emph{Description}:\n This routine tests whether storage, identified by the host pointer @var{ptr}\n is mapped to the device specified by @var{device_num}. If so, it returns\n-a nonzero value and otherwise zero.\n-\n-In GCC, this includes self mapping such that @code{omp_target_is_present}\n-returns @emph{true} when @var{device_num} specifies the host or when the host\n-and the device share memory. If @var{ptr} is a null pointer, @var{true} is\n-returned and if @var{device_num} is an invalid device number, @var{false} is\n-returned.\n-\n-If those conditions do not apply, @emph{true} is returned if the association has\n-been established by an explicit or implicit @code{map} clause, the\n-@code{declare target} directive or a call to the @code{omp_target_associate_ptr}\n-routine.\n+a nonzero value and otherwise zero. In particular, it always returns zero\n+for the null pointer and for invalid device numbers; for the host device,\n+a nonzero value is returned for all non-null pointers.\n \n Running this routine in a @code{target} region except on the initial device\n is not supported.\n@@ -2154,7 +2145,7 @@ is not supported.\n @end multitable\n \n @item @emph{See also}:\n-@ref{omp_target_associate_ptr}\n+@ref{omp_get_mapped_ptr}, @ref{omp_target_associate_ptr}, @ref{omp_target_is_accessible}\n \n @item @emph{Reference}:\n @uref{https://www.openmp.org, OpenMP specification v5.1}, Section 3.8.3\n@@ -2204,7 +2195,7 @@ is not supported.\n @end multitable\n \n @item @emph{See also}:\n-@ref{omp_target_associate_ptr}\n+@ref{omp_target_is_present}\n \n @item @emph{Reference}:\n @uref{https://www.openmp.org, OpenMP specification v5.1}, Section 3.8.4\n@@ -2599,7 +2590,8 @@ is not supported.\n \n @item @emph{See also}:\n @ref{omp_target_disassociate_ptr}, @ref{omp_target_is_present},\n-@ref{omp_target_alloc}\n+@ref{omp_get_mapped_ptr}, @ref{omp_target_alloc}\n+\n \n @item @emph{Reference}:\n @uref{https://www.openmp.org, OpenMP specification v5.1}, Section 3.8.9\n@@ -2657,13 +2649,20 @@ is not supported.\n @subsection @code{omp_get_mapped_ptr} -- Return device pointer to a host pointer\n @table @asis\n @item @emph{Description}:\n-If the device number is refers to the initial device or to a device with\n-memory accessible from the host (shared memory), the @code{omp_get_mapped_ptr}\n+If the device number is refers to the initial device, the @code{omp_get_mapped_ptr}\n routines returns the value of the passed @var{ptr}. Otherwise, if associated\n storage to the passed host pointer @var{ptr} exists on device associated with\n @var{device_num}, it returns that pointer. In all other cases and in cases of\n an error, a null pointer is returned.\n \n+If the device number is not the initial device and the pointer points to a\n+variable that is specified in a @code{declare target} directive: When\n+requiring @code{unified_shared_memory} or @code{self_maps}, a null pointer is\n+returned if the variable appeared in a @code{link} or @code{enter} clause.\n+Otherwise, the corresponding device memory is returned; with the @code{link}\n+clause, GCC returns the address of the pointer-typed link variable on the device\n+not to the data that is mapped to that variable.\n+\n The association of storage location is established either via an explicit or\n implicit @code{map} clause, the @code{declare target} directive or the\n @code{omp_target_associate_ptr} routine.\n@@ -2685,7 +2684,7 @@ is not supported.\n @end multitable\n \n @item @emph{See also}:\n-@ref{omp_target_associate_ptr}\n+@ref{omp_target_is_present}, @ref{omp_target_associate_ptr}\n \n @item @emph{Reference}:\n @uref{https://www.openmp.org, OpenMP specification v5.1}, Section 3.8.11\ndiff --git a/libgomp/target.c b/libgomp/target.c\nindex 29e9a2c6367..d3e2480b40e 100644\n--- a/libgomp/target.c\n+++ b/libgomp/target.c\n@@ -4862,9 +4862,24 @@ gomp_page_locked_host_free (void *ptr)\n \t\tdevice->name);\n }\n \n+/* Check whether corresponding storage exists on the device.\n+ - NULL pointer or invalid device: return 0\n+ - host device: return 1\n+ - Has corresponding storage: return 1\n+ - Otherwise: return 0\n+\n+ Note that for GOMP_OFFLOAD_CAP_SHARED_MEM self mapping is used and\n+ omp_target_associate_ptr is disabled; the only corresponding storage\n+ exists then for declare_target with other clauses than an explicit or\n+ implicit 'link' clause.\n+ However, the link cause with shared memory does not count as mapped. */\n+\n int\n omp_target_is_present (const void *ptr, int device_num)\n {\n+ if (ptr == NULL)\n+ return 0;\n+\n if (device_num == omp_default_device)\n device_num = gomp_get_default_device ();\n \n@@ -4876,13 +4891,8 @@ omp_target_is_present (const void *ptr, int device_num)\n if (devicep == NULL)\n return 0;\n \n- if (ptr == NULL)\n- return 1;\n-\n- if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)\n- || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)\n- return 1;\n-\n+ bool is_shared = (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)\n+\t\t || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM);\n gomp_mutex_lock (&devicep->lock);\n struct splay_tree_s *mem_map = &devicep->mem_map;\n struct splay_tree_key_s cur_node;\n@@ -4890,7 +4900,7 @@ omp_target_is_present (const void *ptr, int device_num)\n cur_node.host_start = (uintptr_t) ptr;\n cur_node.host_end = cur_node.host_start;\n splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);\n- int ret = n != NULL;\n+ int ret = n != NULL && (!is_shared || n->refcount != REFCOUNT_LINK);\n gomp_mutex_unlock (&devicep->lock);\n return ret;\n }\n@@ -5558,7 +5568,8 @@ omp_get_mapped_ptr (const void *ptr, int device_num)\n if (device_num == omp_default_device)\n device_num = gomp_get_default_device ();\n \n- if (device_num == omp_initial_device\n+ if (ptr == NULL\n+ || device_num == omp_initial_device\n || device_num == omp_get_initial_device ())\n return (void *) ptr;\n \n@@ -5566,10 +5577,8 @@ omp_get_mapped_ptr (const void *ptr, int device_num)\n if (devicep == NULL)\n return NULL;\n \n- if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)\n- || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)\n- return (void *) ptr;\n-\n+ bool is_shared = (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)\n+\t\t || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM);\n gomp_mutex_lock (&devicep->lock);\n \n struct splay_tree_s *mem_map = &devicep->mem_map;\n@@ -5584,6 +5593,8 @@ omp_get_mapped_ptr (const void *ptr, int device_num)\n {\n uintptr_t offset = cur_node.host_start - n->host_start;\n ret = (void *) (n->tgt->tgt_start + n->tgt_offset + offset);\n+ if (is_shared && n->refcount == REFCOUNT_LINK)\n+ ret = NULL;\n }\n \n gomp_mutex_unlock (&devicep->lock);\ndiff --git a/libgomp/testsuite/libgomp.c/omp_target_is_present-2.c b/libgomp/testsuite/libgomp.c/omp_target_is_present-2.c\nnew file mode 100644\nindex 00000000000..ee4d8215f33\n--- /dev/null\n+++ b/libgomp/testsuite/libgomp.c/omp_target_is_present-2.c\n@@ -0,0 +1,7 @@\n+// { dg-do run }\n+\n+#define REQ_SELF_MAPS 1\n+\n+#pragma omp requires self_maps\n+\n+#include \"omp_target_is_present.c\"\ndiff --git a/libgomp/testsuite/libgomp.c/omp_target_is_present.c b/libgomp/testsuite/libgomp.c/omp_target_is_present.c\nnew file mode 100644\nindex 00000000000..87468dda5c9\n--- /dev/null\n+++ b/libgomp/testsuite/libgomp.c/omp_target_is_present.c\n@@ -0,0 +1,206 @@\n+// { dg-do run }\n+\n+// Check mainly omp_target_is_present - but also some related functions\n+\n+/* omp_target_is_present is only 1 if device == host or when there is corresponding\n+ storage on the device, which implies ptr != omp_get_mapped_ptr (ptr, dev).\n+\n+ Note that a NULL ptr is regarded as not being present. */\n+\n+#include <omp.h>\n+\n+#ifndef REQ_SELF_MAPS\n+ #define REQ_SELF_MAPS 0\n+#endif\n+\n+// FIXME: change enter to link clause for gLink, once implemented\n+\n+int gEnter = 3, gLink = 4, gLocal = 5;\n+#pragma omp declare target enter(gEnter) link(gLink) enter(gLocal)\n+\n+void check_routines (int dev)\n+{\n+ int A = 1, B = 2;\n+\n+ int dev2 = dev;\n+ if (dev2 == omp_default_device)\n+ dev2 = omp_get_default_device ();\n+\n+ bool initial_dev = dev2 == omp_initial_device || dev2 == omp_get_num_devices();\n+ bool self_mapping = false;\n+ bool invalid_dev = dev == omp_invalid_device;\n+ if (!invalid_dev && !initial_dev)\n+ {\n+ #pragma omp target map(to: self_mapping) device(dev)\n+ self_mapping = true;\n+ if (REQ_SELF_MAPS && !self_mapping)\n+ __builtin_abort ();\n+ }\n+\n+ if (omp_target_is_present (nullptr, dev) != 0)\n+ __builtin_abort ();\n+\n+ if (omp_target_is_accessible (nullptr, 0, dev) != 0)\n+ __builtin_abort ();\n+\n+\n+ if (invalid_dev)\n+ return; // Will otherwise fail with: libgomp: omp_invalid_device encountered\n+\n+\n+ if (omp_target_is_present (&A, dev) != initial_dev)\n+ __builtin_abort ();\n+\n+ // For link, it points to the pointer var - FIXME: update for self_maps implying 'link'\n+ if (omp_target_is_present (&gEnter, dev) != !invalid_dev)\n+ __builtin_abort ();\n+\n+ if (omp_target_is_present (&gLink, dev) != (!invalid_dev && (initial_dev || !REQ_SELF_MAPS)))\n+ __builtin_abort ();\n+\n+ if (omp_target_is_present (&gLocal, dev) != !invalid_dev)\n+ __builtin_abort ();\n+\n+ int *ptr = (int*) 0xDEEDBEEF;\n+ if (!invalid_dev)\n+ {\n+ #pragma omp target enter data map(to: A) device(dev)\n+ #pragma omp target enter data map(to: gEnter) device(dev)\n+ #pragma omp target enter data map(to: gLink) device(dev)\n+ #pragma omp target enter data map(to: gLocal) device(dev)\n+\n+ ptr = omp_target_alloc (sizeof (int), dev);\n+ if (ptr == nullptr || !omp_target_is_accessible (ptr, sizeof (int), dev))\n+\t__builtin_abort ();\n+ }\n+\n+ // Invalid\n+ if ((initial_dev || invalid_dev) && omp_target_associate_ptr (ptr, ptr, sizeof (int), 0, dev) == 0)\n+ __builtin_abort ();\n+ if ((initial_dev || invalid_dev) && omp_target_associate_ptr (((char*)ptr) + 2, ptr, sizeof (int)-2, 2, dev) == 0)\n+ __builtin_abort ();\n+\n+ // Should yield 0/success except for self mapping, host or invalid device\n+ // use !! to convert the result to 0 or 1, as errors can also be, e.g. EINVAL\n+ if (!!omp_target_associate_ptr (&B, ptr, sizeof (int), 0, dev)\n+ != (self_mapping || initial_dev || invalid_dev))\n+ __builtin_abort ();\n+\n+ // Try again, should still work as it is the same pointer\n+ if (!!omp_target_associate_ptr (&B, ptr, sizeof (int), 0, dev)\n+ != (self_mapping || initial_dev || invalid_dev))\n+ __builtin_abort ();\n+\n+ if (!!omp_target_is_present (&A, dev)\n+ != (initial_dev || (!self_mapping && !invalid_dev)))\n+ __builtin_abort ();\n+\n+ if (!!omp_target_is_present (&B, dev)\n+ != (initial_dev || (!self_mapping && !invalid_dev)))\n+ __builtin_abort ();\n+\n+ if (!!omp_target_is_present (&gEnter, dev)\n+ != (initial_dev || (/* !self_mapping && */ !invalid_dev)))\n+ __builtin_abort ();\n+\n+ if (!!omp_target_is_present (&gLink, dev)\n+ != (initial_dev || (!self_mapping && !invalid_dev)))\n+ __builtin_abort ();\n+\n+ if (!!omp_target_is_present (&gLocal, dev) != !invalid_dev)\n+ __builtin_abort ();\n+\n+ int *ptr2 = omp_get_mapped_ptr (&A, dev);\n+ if (initial_dev)\n+ {\n+ if (ptr2 != &A)\n+\t__builtin_abort ();\n+ }\n+ else if (invalid_dev || self_mapping)\n+ {\n+ if (ptr2 != nullptr)\n+\t__builtin_abort ();\n+ }\n+ else if (ptr2 == &A || ptr2 == nullptr)\n+ __builtin_abort ();\n+\n+ ptr2 = omp_get_mapped_ptr (&B, dev);\n+ if (initial_dev)\n+ {\n+ if (ptr2 != &B)\n+\t__builtin_abort ();\n+ }\n+ else if (invalid_dev || self_mapping)\n+ {\n+ if (ptr2 != nullptr)\n+\t__builtin_abort ();\n+ }\n+ else if (ptr2 != ptr)\n+ __builtin_abort ();\n+\n+ ptr2 = omp_get_mapped_ptr (&gEnter, dev);\n+ if (initial_dev)\n+ {\n+ if (ptr2 != &gEnter)\n+\t__builtin_abort ();\n+ }\n+ else if (invalid_dev /* FIXME: || self_mapping */)\n+ {\n+ if (ptr2 != nullptr)\n+\t__builtin_abort ();\n+ }\n+ else if (ptr2 == &gEnter || ptr2 == nullptr)\n+ __builtin_abort ();\n+\n+ ptr2 = omp_get_mapped_ptr (&gLink, dev);\n+ if (initial_dev)\n+ {\n+ if (ptr2 != &gLink)\n+\t__builtin_abort ();\n+ }\n+ else if (invalid_dev || self_mapping)\n+ {\n+ if (ptr2 != nullptr)\n+\t__builtin_abort ();\n+ }\n+ else if (ptr2 == ptr || ptr2 == nullptr)\n+ __builtin_abort ();\n+\n+ ptr2 = omp_get_mapped_ptr (&gLocal, dev);\n+ if (initial_dev)\n+ {\n+ if (ptr2 != &gLocal)\n+\t__builtin_abort ();\n+ }\n+ else if (invalid_dev)\n+ {\n+ if (ptr2 != nullptr)\n+\t__builtin_abort ();\n+ }\n+ else if (ptr2 == &gLocal || ptr2 == nullptr)\n+ __builtin_abort ();\n+\n+ if (!invalid_dev)\n+ {\n+ omp_target_free (ptr, dev);\n+ #pragma omp target exit data map(release: A) device(dev)\n+ #pragma omp target exit data map(release: gLink) device(dev)\n+ #pragma omp target exit data map(release: gEnter) device(dev)\n+ #pragma omp target exit data map(release: gLocal) device(dev)\n+ }\n+}\n+\n+int main()\n+{\n+ for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)\n+ check_routines (dev);\n+\n+ check_routines (omp_invalid_device);\n+ check_routines (omp_default_device);\n+\n+ for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++)\n+ {\n+ omp_set_default_device (dev);\n+ check_routines (omp_default_device);\n+ }\n+}\n", "prefixes": [] }