From patchwork Fri Mar 24 16:35:15 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 66869 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id EB30C384D169 for ; Fri, 24 Mar 2023 16:35:51 +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 711EC3858D1E for ; Fri, 24 Mar 2023 16:35:22 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 711EC3858D1E 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="5.98,288,1673942400"; d="scan'208,223";a="281704" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 24 Mar 2023 08:35:21 -0800 IronPort-SDR: 6fP/iR1WOgIuS6ohHW7bxLF9j2ctEfQJHgxNMlg12GrQCD/GUClEqecmEjLfVPdCFTW6YY5Usv 6D/voEQukcXlraw0WiIQjFZKX4PT5UmzukFpS9ZIJQnSHE0IVhW/bYpa+JvpwFAplkwgdmgVUs qIeIBvjNRfUEhkanaxiTGO4JN58t407g2uODFpyiv13Mhnb4Xtm00nPnEQJg1XZi6eG0EY5R5M EfMosDoxWvFfIW3fN2iXxExI8LgLD8qyuqNe2sfBh/2dRos40hjB3a7qkBbWi+2NbcRuRscEI/ 78g= From: Thomas Schwinge To: CC: Andrew Stubbs Subject: [og12] Add 'libgomp.c/alloc-ompx_host_mem_alloc-1.c' User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/28.2 (x86_64-pc-linux-gnu) Date: Fri, 24 Mar 2023 17:35:15 +0100 Message-ID: <87zg8285vg.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-11.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi! This had fallen out of some earlier work of mine; I've now pushed to devel/omp/gcc-12 branch commit ae2dca26602678f8b70e22da1bce8302c0751b75 "Add 'libgomp.c/alloc-ompx_host_mem_alloc-1.c'", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 From ae2dca26602678f8b70e22da1bce8302c0751b75 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 15 Feb 2023 11:27:55 +0100 Subject: [PATCH] Add 'libgomp.c/alloc-ompx_host_mem_alloc-1.c' OpenMP 'ompx_host_mem_alloc' is available for host and nvptx offloading as of og12 commit 84914e197d91a67b3d27db0e4c69a433462983a5 "openmp, nvptx: ompx_unified_shared_mem_alloc", and for GCN offloading as of og12 commit c77c45a641fedc3fe770e909cc010fb1735bdbbd "amdgcn, libgomp: low-latency allocator". libgomp/ * testsuite/libgomp.c/alloc-ompx_host_mem_alloc-1.c: New. --- libgomp/ChangeLog.omp | 2 + .../libgomp.c/alloc-ompx_host_mem_alloc-1.c | 77 +++++++++++++++++++ 2 files changed, 79 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c/alloc-ompx_host_mem_alloc-1.c diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index ef957e3d2d8..6b816e46cd2 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,5 +1,7 @@ 2023-03-24 Thomas Schwinge + * testsuite/libgomp.c/alloc-ompx_host_mem_alloc-1.c: New. + * config/gcn/allocator.c (gcn_memspace_free): Explicitly handle 'memspace == ompx_host_mem_space'. diff --git a/libgomp/testsuite/libgomp.c/alloc-ompx_host_mem_alloc-1.c b/libgomp/testsuite/libgomp.c/alloc-ompx_host_mem_alloc-1.c new file mode 100644 index 00000000000..683b7aff3d4 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/alloc-ompx_host_mem_alloc-1.c @@ -0,0 +1,77 @@ +/* Verify that on the host we can but on a device we cannot allocate 'ompx_host_mem_alloc' memory. */ + +/* { dg-additional-options -DOFFLOAD_DEVICE { target offload_device } } */ + +#include + +#pragma omp requires dynamic_allocators + +int main() +{ +#pragma omp target + { + char *c, *c_; + + c = omp_alloc(1, ompx_host_mem_alloc); +#ifdef OFFLOAD_DEVICE + if (c) + __builtin_abort (); +#else + if (!c) + __builtin_abort (); +#endif + omp_free(c, ompx_host_mem_alloc); + + c = omp_aligned_alloc(128, 256, ompx_host_mem_alloc); +#ifdef OFFLOAD_DEVICE + if (c) + __builtin_abort (); +#else + if (!c) + __builtin_abort (); +#endif + omp_free(c, omp_null_allocator); + + c = omp_calloc(1, 1, ompx_host_mem_alloc); +#ifdef OFFLOAD_DEVICE + if (c) + __builtin_abort (); +#else + if (!c) + __builtin_abort (); +#endif + c_ = omp_realloc(c, 2, ompx_host_mem_alloc, ompx_host_mem_alloc); +#ifdef OFFLOAD_DEVICE + if (c_) + __builtin_abort (); +#else + if (!c_) + __builtin_abort (); +#endif + c = omp_realloc(c_, 0, ompx_host_mem_alloc, ompx_host_mem_alloc); + if (c) + __builtin_abort (); + + c = omp_aligned_calloc(64, 1, 512, ompx_host_mem_alloc); +#ifdef OFFLOAD_DEVICE + if (c) + __builtin_abort (); +#else + if (!c) + __builtin_abort (); +#endif + c_ = omp_realloc(c, 2, c ? omp_null_allocator : ompx_host_mem_alloc, omp_null_allocator); +#ifdef OFFLOAD_DEVICE + if (c_) + __builtin_abort (); +#else + if (!c_) + __builtin_abort (); +#endif + c = omp_realloc(c_, 0, omp_null_allocator, omp_null_allocator); + if (c) + __builtin_abort (); + } + + return 0; +} -- 2.25.1