From patchwork Thu Jun 4 18:17:33 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1303683 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49dDXW5L9Hz9sSg for ; Fri, 5 Jun 2020 04:17:46 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B841C388F061; Thu, 4 Jun 2020 18:17:44 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 7A947385DC14 for ; Thu, 4 Jun 2020 18:17:42 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 7A947385DC14 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: kA3R4nYYQg++M658lGKyOlcX5COcCs1L4FUkDVjRBJezpL9HWNqBZWvpL+OTXh8Xft5Qq3m6XY y1b6tAeBnGk4YwSgWs+u1b/3ZijHOwNy+nESMe5D4BlPIbmGlzAg85Yx+M6T0ZKOgtLEWTpgxq 6Puta4saqM3EvvdtKvHV+RMJSLofH8pzhazo/vTnkXR9yYl0SpBamh9mnie6MI757xgteujokx C0C9Lsqyyy3yMbzrDq5DkqVR78J+PvvOXs60G61UjYses9eNmWIHb6VsxftYLEohoyK64gvgbq p5U= X-IronPort-AV: E=Sophos;i="5.73,472,1583222400"; d="scan'208,223";a="51583120" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 04 Jun 2020 10:17:41 -0800 IronPort-SDR: RdDopHUzerZFY9VRoLxNEbr3JZmQYJcPkXyDlMZVxnHgdpSTI7Sx8XLq2Z2Yf5B1JEsCMIZvoV q/UrTUx3UkiFFL9x5l3DNt0HGfCM6Zhh4dqxgaRI5/HASrIqZbRKXretcWkwEOFBaUsuhzNatE 2A3RueK1DV1IeQMrg3ccm6KZgX8MTqvToEWioI1gST44/SOxOVoihW2EYxsX96gTz36gP9fGHr w9r8/YpX0VKIcI9dFp8hi69U7BYKHoYT4xneWW3WddC9qQRKS/wznUJBhMTmDxxKkRYzJlOtu2 Ga0= From: Thomas Schwinge To: Subject: Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] (was: [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c') In-Reply-To: <87lfrl3ee5.fsf@euler.schwinge.homeip.net> References: <87lfrl3ee5.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Thu, 4 Jun 2020 20:17:33 +0200 Message-ID: <87k10msnk2.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Spam-Status: No, score=-11.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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@gcc.gnu.org Sender: "Gcc-patches" Hi! On 2019-12-09T12:52:02+0100, I wrote: > See attached "[PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c'", > committed to trunk in r279120, "to document the status quo", which does > match my understanding of the OpenACC 2.6 semantics. I've now pushed "Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854]" to master branch in commit af8fd1a99d9a21f8088ebb11250cd06a3f275052, and releases/gcc-10 branch in commit 364f46de9f02dc00e8ff51cc9e2662ae37520389, see attached. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter From 364f46de9f02dc00e8ff51cc9e2662ae37520389 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 29 May 2020 14:11:27 +0200 Subject: [PATCH] Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] libgomp/ PR libgomp/92854 * testsuite/libgomp.oacc-c-c++-common/pr92854-1.c: Extend some more. (cherry picked from commit af8fd1a99d9a21f8088ebb11250cd06a3f275052) --- .../libgomp.oacc-c-c++-common/pr92854-1.c | 64 ++++++++++++++----- 1 file changed, 47 insertions(+), 17 deletions(-) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c index 6ba96b6bf8f9..79cebf65c348 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c @@ -1,31 +1,61 @@ -/* Verify that 'acc_unmap_data' unmaps even in presence of dynamic reference - counts. */ +/* Verify that 'acc_unmap_data' unmaps even in presence of structured and + dynamic reference counts, but the device memory remains allocated. */ /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ #include #include +#include #include int main () { const int N = 180; - - char *h = (char *) malloc (N); - char *d = (char *) acc_malloc (N); - if (!d) - abort (); - acc_map_data (h, d, N); - - char *d_ = (char *) acc_create (h + 3, N - 77); - assert (d_ == d + 3); - - d_ = (char *) acc_create (h, N); - assert (d_ == d); - - acc_unmap_data (h); - assert (!acc_is_present (h, N)); + const int N_i = 537; + const int C = 37; + + unsigned char *h = (unsigned char *) malloc (N); + assert (h); + unsigned char *d = (unsigned char *) acc_malloc (N); + assert (d); + + for (int i = 0; i < N_i; ++i) + { + acc_map_data (h, d, N); + assert (acc_is_present (h, N)); +#pragma acc parallel present(h[0:N]) + { + if (i == 0) + memset (h, C, N); + } + + unsigned char *d_ = (unsigned char *) acc_create (h + 3, N - 77); + assert (d_ == d + 3); + +#pragma acc data create(h[6:N - 44]) + { + d_ = (unsigned char *) acc_create (h, N); + assert (d_ == d); + +#pragma acc enter data create(h[0:N]) + + assert (acc_is_present (h, N)); + acc_unmap_data (h); + assert (!acc_is_present (h, N)); + } + + /* We can however still access the device memory. */ +#pragma acc parallel loop deviceptr(d) + for (int j = 0; j < N; ++j) + d[j] += i * j; + } + + acc_memcpy_from_device(h, d, N); + for (int j = 0; j < N; ++j) + assert (h[j] == ((C + N_i * (N_i - 1) / 2 * j) % 256)); + + acc_free (d); return 0; } -- 2.26.2