From patchwork Thu Dec 12 12:07:41 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1208397 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=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-515797-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="HBjKsjQK"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 47YXcg36pDz9sP6 for ; Thu, 12 Dec 2019 23:08:01 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:subject:message-id:mime-version:content-type; q=dns; s= default; b=uyEku/hyO6XTquw+hFrQOTBkgc52MZHniMb94liYDsA79a8C8G8c3 0vmbe5SJgQAKIgiwIZu+pf0RGNwQh7RshDpkn/6jtHPT3BwMcvHDwfF+cHIWpzvh x+7i+3A80XzhwRnaDiLpofUAKTQqk/PTh4jYFUPX6mYHaK8vPrPonI= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:subject:message-id:mime-version:content-type; s= default; bh=TBVAIhsuNs9xwoOc3McHCQXSArY=; b=HBjKsjQKUkOnT/vF5BRV cmdu6lwIYYNBsmcRyhzzi3Ftv2ooG8wTo+O1ENR9K5O/bQGtdAv3b02Xb1zip7aV cT9Q2CqGG6X5C+hrEb0uaxEbYjSkj/Y+6sldgpjyWRpuzacZcRquzsEK6NdZdFcM 3UvGRH4FGXjQuC6tgfVSAj4= Received: (qmail 88088 invoked by alias); 12 Dec 2019 12:07:54 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 88070 invoked by uid 89); 12 Dec 2019 12:07:54 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-16.4 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy=17598, lookup_dev X-HELO: esa4.mentor.iphmx.com Received: from esa4.mentor.iphmx.com (HELO esa4.mentor.iphmx.com) (68.232.137.252) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 12 Dec 2019 12:07:51 +0000 IronPort-SDR: jFFEjL+GR3+NL01lQiZDbBjrCRAfHeNdTrw9eTVlQJYwnbjBCN3D02UCZHMxi8SU0JLgiSJnSn TwweBkici4aMIEQGRGDugTxSSkVLvgH41vO+xZkryeEcY4V/lfvPeF1ab3D+YDAvdEwnpOaS9m 3NR1ICZmmSYcojFez2g6yEEXdvhyR9UBQSL6EZyMMNDggTdo/Gm54eKg5pvN1S8dtHQKyfBdH/ uISRKrM0KRZce9ixvLGUb/Mphz5+ThBUSJcsVO0dFOTwSlCBS1dxOaLh7YeuQ7+oGJ8fg1WuH0 DZs= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 12 Dec 2019 04:07:49 -0800 IronPort-SDR: Rg4IxnYEO/DtEaf528cZsGWrtf5gNUfl5Tt0Drh8NkylbYaQgUKD3iFyR0B4TzYN9rl1GYISP9 EZoTlC6CdDPF66iD56qtqKulloBVP39gvYSzPS9VIPTCyIXBPXuXZAHkKddS/pS7qHQjVkfpPN L0jbRXLUrfSrihFkBvBBy09h62vnoc5BzrtgqM8p7Spxd2LBe9qIqjdJpCFbUcdsh4YpEKzP8m btRtKgen+rGNoESGlXiMR+3H3F6YqH1qIn/XZtWubpSRKyl7S9i+POTFFccJG/OO7PHNyorb1N AS8= Date: Thu, 12 Dec 2019 12:07:41 +0000 From: Julian Brown To: , , Jakub Jelinek Subject: [PATCH] OpenACC device-pointer lookup with globally-mapped variables (PR92888) Message-ID: <20191212120741.08be278b@squid.athome> MIME-Version: 1.0 X-IsSubscribed: yes Hi, This patch provides a fix for PR92888, wherein global variables mapped using an OpenACC 'declare' directive would not be visible to device-pointer lookups. Tested with offloading to nvptx. OK? Thanks, Julian ChangeLog 2019-12-12 Julian Brown PR libgomp/92888 libgomp/ * oacc-parallel.c (GOACC_parallel_keyed): Add tgt_start in target function address calculation. * target.c (gomp_load_image_to_device): Record address range for target_mem_desc for mapped functions and global variables, and adjust tgt_offsets to be within that range. (gomp_get_target_fn_addr): Add tgt_start in target function address calculation. * testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c: Remove XFAIL. commit 16e774d2ce86af90ff282b9126cf615e66e7efae Author: Julian Brown Date: Mon Dec 9 11:04:58 2019 -0800 Find address range for offloaded functions and global variables (PR92888) PR libgomp/92888 libgomp/ * oacc-parallel.c (GOACC_parallel_keyed): Add tgt_start in target function address calculation. * target.c (gomp_load_image_to_device): Record address range for target_mem_desc for mapped functions and global variables, and adjust offsets to be within that range. (gomp_get_target_fn_addr): Add tgt_start in target function address calculation. * testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c: Remove XFAIL. diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index f5ef5050bbd..5a5697cf6e6 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -377,7 +377,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), if (tgt_fn_key == NULL) gomp_fatal ("target function wasn't mapped"); - tgt_fn = (void (*)) tgt_fn_key->tgt_offset; + tgt_fn = (void (*)) (tgt_fn_key->tgt->tgt_start + tgt_fn_key->tgt_offset); } else tgt_fn = (void (*)) fn; diff --git a/libgomp/target.c b/libgomp/target.c index bb392dd1c8f..b023e3daf1a 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1759,6 +1759,8 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, tgt->device_descr = devicep; splay_tree_node array = tgt->array; + uintptr_t max_addr = 0, min_addr = ~(uintptr_t) 0; + for (i = 0; i < num_funcs; i++) { splay_tree_key k = &array->key; @@ -1766,6 +1768,10 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, k->host_end = k->host_start + 1; k->tgt = tgt; k->tgt_offset = target_table[i].start; + if (target_table[i].start < min_addr) + min_addr = target_table[i].start; + if (target_table[i].end > max_addr) + max_addr = target_table[i].end; k->refcount = REFCOUNT_INFINITY; k->virtual_refcount = 0; k->aux = NULL; @@ -1799,6 +1805,10 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, = k->host_start + (size_mask & (uintptr_t) host_var_table[i * 2 + 1]); k->tgt = tgt; k->tgt_offset = target_var->start; + if (target_var->start < min_addr) + min_addr = target_var->start; + if (target_var->end > max_addr) + max_addr = target_var->end; k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY; k->virtual_refcount = 0; k->aux = NULL; @@ -1808,6 +1818,17 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, array++; } + /* Make the tgt_mem_desc cover all of the functions and variables so that + oacc-mem.c:lookup_dev can find mapped global variables properly. */ + tgt->tgt_start = min_addr; + tgt->tgt_end = max_addr; + + for (array = tgt->array, i = 0; i < num_vars + num_funcs; i++, array++) + { + splay_tree_key k = &array->key; + k->tgt_offset -= min_addr; + } + free (target_table); } @@ -2170,7 +2191,7 @@ gomp_get_target_fn_addr (struct gomp_device_descr *devicep, if (tgt_fn == NULL) return NULL; - return (void *) tgt_fn->tgt_offset; + return (void *) (tgt_fn->tgt->tgt_start + tgt_fn->tgt_offset); } } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c index 7cd2936219a..0807bc9d694 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c @@ -24,5 +24,5 @@ main () /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "device address \\\[\[0-9a-fA-FxX\]+, \\\+8\\\] is already mapped" { xfail *-*-* } } TODO */ -/* { dg-shouldfail "TODO" { INV-AL-ID } } */ +/* { dg-output "device address \\\[\[0-9a-fA-FxX\]+, \\\+8\\\] is already mapped" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92888-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92888-1.c new file mode 100644 index 00000000000..0cd7f13656c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92888-1.c @@ -0,0 +1,19 @@ +/* Make sure that we can resolve back via 'acc_hostptr' an 'acc_deviceptr' + retrieved for a '#pragma acc declare'd variable. */ + +#include +#include + +double global_var; +#pragma acc declare create (global_var) + +int +main () +{ + void *global_var_p_d = acc_deviceptr (&global_var); + assert (acc_hostptr (global_var_p_d) == &global_var); + + return 0; +} + +/* { dg-xfail-run-if "PR92888" { ! openacc_host_selected } } */