From patchwork Wed Aug 2 17:00:35 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Andrew Stubbs X-Patchwork-Id: 129989 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:9f41:0:b0:3e4:2afc:c1 with SMTP id v1csp597495vqx; Wed, 2 Aug 2023 10:02:30 -0700 (PDT) X-Google-Smtp-Source: APBJJlGXYTjGdWv+CR3X04Q04tdhWmh/JQR0UBE3Kpf9mys6nqBFrsAYnlVEGKnR5MTWiZ6QnmbU X-Received: by 2002:a17:906:db:b0:98e:16b7:e038 with SMTP id 27-20020a17090600db00b0098e16b7e038mr5767792eji.23.1690995749846; Wed, 02 Aug 2023 10:02:29 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1690995749; cv=none; d=google.com; s=arc-20160816; b=tydXihDgdeN04iYsEPnadj8K9yHEAoxCha5BWrbZFXSQ+gCIANcUA3LxhF+xGt32JM IMy7MtZnQ+vRsMKnZQx39tAJwN8NsylbhqzXs8laXZWuWJLWIaodH2OYH+r7pPj7BP7m 4ZIKY1c5AvKDP2v5INmkTyu9XlbUAjHu9rwVP8jnWY7ipbUQGOaRtFbivFPQoIDWoH6A IpUd98GXD4yBWwwlZ7bNzKxiyqHlLD6mfkCiulKsrEdxhR3bR0ZKQrQo+nZLDVPhjPK7 tnhS8BsSCEtqyiEck35y9yvTT5lmxWKpJiJGqGs6M1vUuxo79dMfuku4HoXzFO6WJZ6b rrrQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:content-transfer-encoding :mime-version:references:in-reply-to:message-id:date:subject:to:from :ironport-sdr:dmarc-filter:delivered-to; bh=2fYk3NdhvhfUfhgKH9z9wWhFd++FX0XA1K8BrqqY2tk=; fh=9xD62naDDQ44CEuxrRV0ts5PrjZies6B3X/DCYdxTRw=; b=hiJidArtUDnksaSMfQgix1qqdZ2HHkHB2X0OVF8c8ce5nh779wKj/RwCATjrmaCOp4 ayKrkaSZpQJTZW2geClytIwnqYCEkeAywdCvD7R3ipAxIuHJ7+Z/YocNEaQrRDLlHqT4 kCKHt1R7V4HXjOFeuUV3OZTgiRE9Fvqp8uF96dhz8Sp8IdOHNUzet7TRhXoJEU1bUktg TErazNnY8sxOGzUsDaewAh03jYTQiOmAUeO6CK6FOwWFPkKg955MQEvHdZmiLwtjJiHv yiKqG6Z6Hsd4yqILa6I7S7JLKfKY9M0pZ5FEvGEoro1V4TNTrMmpG0JMB86SGMusooGQ kNaQ== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (ip-8-43-85-97.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id rk13-20020a170907214d00b0099bd67190c8si9977040ejb.291.2023.08.02.10.02.29 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 02 Aug 2023 10:02:29 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 1BE5B385B516 for ; Wed, 2 Aug 2023 17:01:36 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 972833858D33 for ; Wed, 2 Aug 2023 17:01:04 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 972833858D33 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,249,1684828800"; d="scan'208";a="13385825" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 02 Aug 2023 09:01:00 -0800 IronPort-SDR: d7vor+ZAdzKq3QSjJAGLMbEWnD0j+T7Azwcse36swAUXfVi/pyHJ6tVCqhSO3w5Clc4T5h20P6 3NRK6s2wD+I/8HbdRxxRTsQS52F5xTA37XTkUNDt/+4QJjWhy931YAlX7kQ1PsubmtMz/5bSVd Ca/XRZn1KbsF9w/S8aOX4pN7fzlvfZe0dl1J8xN/Swq0OjMsNaLz8P2D71oHgVkNgMqgwttwuL bWpsVFgtlF+ewG1a8/u8JVE2DsPSh/4u5v14pQwwIdZCd9/16VxuFB+U3q+xTsl2VtZ9kTcJFV i9I= From: Andrew Stubbs To: Subject: [PATCH v2 2/3] openmp, nvptx: low-lat memory access traits Date: Wed, 2 Aug 2023 18:00:35 +0100 Message-ID: <161001070f7573c98d2b72223933dbba49405fea.1690994309.git.ams@codesourcery.com> X-Mailer: git-send-email 2.41.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-12.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1773137559011227754 X-GMAIL-MSGID: 1773137559011227754 The NVPTX low latency memory is not accessible outside the team that allocates it, and therefore should be unavailable for allocators with the access trait "all". This change means that the omp_low_lat_mem_alloc predefined allocator now implicitly implies the "pteam" trait. libgomp/ChangeLog: * allocator.c (MEMSPACE_VALIDATE): New macro. (omp_aligned_alloc): Use MEMSPACE_VALIDATE. (omp_aligned_calloc): Likewise. (omp_realloc): Likewise. * config/nvptx/allocator.c (nvptx_memspace_validate): New function. (MEMSPACE_VALIDATE): New macro. * testsuite/libgomp.c/omp_alloc-4.c (main): Add access trait. * testsuite/libgomp.c/omp_alloc-6.c (main): Add access trait. * testsuite/libgomp.c/omp_alloc-traits.c: New test. --- libgomp/allocator.c | 16 +++++ libgomp/config/nvptx/allocator.c | 11 +++ libgomp/testsuite/libgomp.c/omp_alloc-4.c | 7 +- libgomp/testsuite/libgomp.c/omp_alloc-6.c | 7 +- .../testsuite/libgomp.c/omp_alloc-traits.c | 68 +++++++++++++++++++ 5 files changed, 103 insertions(+), 6 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-traits.c diff --git a/libgomp/allocator.c b/libgomp/allocator.c index fbf7b1ab061..35b8ec71480 100644 --- a/libgomp/allocator.c +++ b/libgomp/allocator.c @@ -56,6 +56,10 @@ #define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE) \ free (((void)(MEMSPACE), (void)(SIZE), (ADDR))) #endif +#ifndef MEMSPACE_VALIDATE +#define MEMSPACE_VALIDATE(MEMSPACE, ACCESS) \ + (((void)(MEMSPACE), (void)(ACCESS), 1)) +#endif /* Map the predefined allocators to the correct memory space. The index to this table is the omp_allocator_handle_t enum value. @@ -507,6 +511,10 @@ retry: if (__builtin_add_overflow (size, new_size, &new_size)) goto fail; + if (allocator_data + && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access)) + goto fail; + if (__builtin_expect (allocator_data && allocator_data->pool_size < ~(uintptr_t) 0, 0)) { @@ -817,6 +825,10 @@ retry: if (__builtin_add_overflow (size_temp, new_size, &new_size)) goto fail; + if (allocator_data + && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access)) + goto fail; + if (__builtin_expect (allocator_data && allocator_data->pool_size < ~(uintptr_t) 0, 0)) { @@ -1063,6 +1075,10 @@ retry: goto fail; old_size = data->size; + if (allocator_data + && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access)) + goto fail; + if (__builtin_expect (allocator_data && allocator_data->pool_size < ~(uintptr_t) 0, 0)) { diff --git a/libgomp/config/nvptx/allocator.c b/libgomp/config/nvptx/allocator.c index 6014fba177f..f19ac28d32a 100644 --- a/libgomp/config/nvptx/allocator.c +++ b/libgomp/config/nvptx/allocator.c @@ -108,6 +108,15 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr, return realloc (addr, size); } +static inline int +nvptx_memspace_validate (omp_memspace_handle_t memspace, unsigned access) +{ + /* Disallow use of low-latency memory when it must be accessible by + all threads. */ + return (memspace != omp_low_lat_mem_space + || access != omp_atv_all); +} + #define MEMSPACE_ALLOC(MEMSPACE, SIZE) \ nvptx_memspace_alloc (MEMSPACE, SIZE) #define MEMSPACE_CALLOC(MEMSPACE, SIZE) \ @@ -116,5 +125,7 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr, nvptx_memspace_realloc (MEMSPACE, ADDR, OLDSIZE, SIZE) #define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE) \ nvptx_memspace_free (MEMSPACE, ADDR, SIZE) +#define MEMSPACE_VALIDATE(MEMSPACE, ACCESS) \ + nvptx_memspace_validate (MEMSPACE, ACCESS) #include "../../allocator.c" diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-4.c b/libgomp/testsuite/libgomp.c/omp_alloc-4.c index 66e13c09234..9d169858151 100644 --- a/libgomp/testsuite/libgomp.c/omp_alloc-4.c +++ b/libgomp/testsuite/libgomp.c/omp_alloc-4.c @@ -23,10 +23,11 @@ main () #pragma omp target { /* Ensure that the memory we get *is* low-latency with a null-fallback. */ - omp_alloctrait_t traits[1] - = { { omp_atk_fallback, omp_atv_null_fb } }; + omp_alloctrait_t traits[2] + = { { omp_atk_fallback, omp_atv_null_fb }, + { omp_atk_access, omp_atv_pteam } }; omp_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space, - 1, traits); + 2, traits); int size = 4; diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-6.c b/libgomp/testsuite/libgomp.c/omp_alloc-6.c index 66bf69b0455..b5f0a296998 100644 --- a/libgomp/testsuite/libgomp.c/omp_alloc-6.c +++ b/libgomp/testsuite/libgomp.c/omp_alloc-6.c @@ -23,10 +23,11 @@ main () #pragma omp target { /* Ensure that the memory we get *is* low-latency with a null-fallback. */ - omp_alloctrait_t traits[1] - = { { omp_atk_fallback, omp_atv_null_fb } }; + omp_alloctrait_t traits[2] + = { { omp_atk_fallback, omp_atv_null_fb }, + { omp_atk_access, omp_atv_pteam } }; omp_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space, - 1, traits); + 2, traits); int size = 16; diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-traits.c b/libgomp/testsuite/libgomp.c/omp_alloc-traits.c new file mode 100644 index 00000000000..6294ba19c16 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/omp_alloc-traits.c @@ -0,0 +1,68 @@ +/* { dg-do run } */ + +/* { dg-require-effective-target offload_device } */ +/* { dg-xfail-if "not implemented" { ! offload_target_nvptx } } */ + +/* Test that GPU low-latency allocation is limited to team access. */ + +#include +#include + +#pragma omp requires dynamic_allocators + +int +main () +{ + #pragma omp target + { + /* Ensure that the memory we get *is* low-latency with a null-fallback. */ + omp_alloctrait_t traits[2] + = { { omp_atk_fallback, omp_atv_null_fb }, + { omp_atk_access, omp_atv_pteam } }; + omp_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space, + 2, traits); + + omp_alloctrait_t traits_all[2] + = { { omp_atk_fallback, omp_atv_null_fb }, + { omp_atk_access, omp_atv_all } }; + omp_allocator_handle_t lowlat_all + = omp_init_allocator (omp_low_lat_mem_space, 2, traits_all); + + omp_alloctrait_t traits_default[1] + = { { omp_atk_fallback, omp_atv_null_fb } }; + omp_allocator_handle_t lowlat_default + = omp_init_allocator (omp_low_lat_mem_space, 1, traits_default); + + void *a = omp_alloc (1, lowlat); // good + void *b = omp_alloc (1, lowlat_all); // bad + void *c = omp_alloc (1, lowlat_default); // bad + + if (!a || b || c) + __builtin_abort (); + + omp_free (a, lowlat); + + + a = omp_calloc (1, 1, lowlat); // good + b = omp_calloc (1, 1, lowlat_all); // bad + c = omp_calloc (1, 1, lowlat_default); // bad + + if (!a || b || c) + __builtin_abort (); + + omp_free (a, lowlat); + + + a = omp_realloc (NULL, 1, lowlat, lowlat); // good + b = omp_realloc (NULL, 1, lowlat_all, lowlat_all); // bad + c = omp_realloc (NULL, 1, lowlat_default, lowlat_default); // bad + + if (!a || b || c) + __builtin_abort (); + + omp_free (a, lowlat); + } + + return 0; +} +