From patchwork Sun May 20 19:46:33 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 917268 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-478024-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="U3z9UvgD"; 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 40psqq3c5Kz9s3X for ; Mon, 21 May 2018 05:47:03 +1000 (AEST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type:content-transfer-encoding; q=dns; s=default; b=ZlF AWdiTJ7djnJ0zoVoxWjoowqKEBFxcRsclJc46+krdXBilk4WwcedX4lD9sFbR+Ta sJEScKVDerZQ786xfoBjFMYNSAEE2Ehri2izb2jFwGNcqD/EwCbQyurIa+mOW1gE 0bmI+236O5U6Ll1PwLbDgGeuuTab81kGpN33xsIg= 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:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type:content-transfer-encoding; s=default; bh=HIgsfQnYk Xc+qD9UNLKVFTpMd8k=; b=U3z9UvgDlKofqE+Y7QnHiA2bxLUJvfzj4GLq7Hy72 wEzAKb9J2v9b6hTDdXRRWUePLPvl8HxAQqvPKBuwvwgV/I0dmn3bo6E+IxascbP6 4zB4O4a4471WqIoILJ0elq1LtIqZRlneWVHxQPNjNhG8WFJYg10GBGLxwpQMUBuC +8= Received: (qmail 42360 invoked by alias); 20 May 2018 19:46:53 -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 42328 invoked by uid 89); 20 May 2018 19:46:53 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.5 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=luck, Joseph, difficult, cesar X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sun, 20 May 2018 19:46:48 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=svr-ies-mbx-01.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1fKUI0-0003GC-RY from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Sun, 20 May 2018 12:46:45 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.87) by svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Sun, 20 May 2018 20:46:40 +0100 From: Thomas Schwinge To: Subject: [og7] Re: Forwarding -foffload=[...] from the driver (compile-time) to libgomp (run-time) In-Reply-To: References: <20141020111935.GA9362@msticlxl57.ims.intel.com> <20141024141601.GA62562@msticlxl57.ims.intel.com> <20141024142028.GD10376@tucnak.redhat.com> <20141028193047.GA17865@msticlxl57.ims.intel.com> <20141103092447.GO5026@tucnak.redhat.com> <20141105124655.GA42356@msticlxl57.ims.intel.com> <87egjopgh0.fsf@kepler.schwinge.homeip.net> <20150731142007.GA64740@msticlxl57.ims.intel.com> <20150805150904.GA3211@msticlxl57.ims.intel.com> <87bneatd5q.fsf@schwinge.name> <87lhddsfs4.fsf@schwinge.name> <87oai4r2ok.fsf@schwinge.name> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.5.1 (x86_64-pc-linux-gnu) Date: Sun, 20 May 2018 21:46:33 +0200 Message-ID: <87zi0u5cva.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) Hi! (This whole idea/patch still needs an overall re-work, as discussed, but here is a small incremental improvement/bug fix.) On Thu, 20 Aug 2015 22:52:58 +0000, Joseph Myers wrote: > On Tue, 18 Aug 2015, Thomas Schwinge wrote: > > [...] here is my current messy WIP patch [...] > +/* List of offload targets, separated by colon. Defaults to the list > + determined when configuring libgomp. */ > +static const char *gomp_offload_targets = OFFLOAD_TARGETS; > +static bool gomp_offload_targets_init = false; > + > +/* Override the list of offload targets. This must be called early, and only > + once. */ > + > +void > +GOMP_set_offload_targets (const char *offload_targets) > +{ > + gomp_debug (0, "%s (\"%s\")\n", __FUNCTION__, offload_targets); > + > + /* Make sure this gets called early. */ > + assert (gomp_is_initialized == PTHREAD_ONCE_INIT); > + /* Make sure this only gets called once. */ > + assert (!gomp_offload_targets_init); > + gomp_offload_targets_init = true; > + gomp_offload_targets = offload_targets; > +} This will obviously fail as soon as there are shared libraries involved, compiled for offloading, which contain additional GOMP_set_offload_targets constructor calls. Thus pushed to openacc-gcc-7-branch: commit 917e247055a37f912129ed545719182de0046adb Author: Thomas Schwinge Date: Sun May 20 21:31:01 2018 +0200 [PR81886] Avoid "GOMP_set_offload_targets: Assertion `!gomp_offload_targets_init' failed" PR libgomp/81886 * openacc.h (enum acc_device_t): Add _acc_device_intel_mic, _acc_device_hsa. * oacc-init.c (get_openacc_name): Handle these. (resolve_device): Debugging output. * target.c (resolve_device, gomp_init_device) (gomp_offload_target_available_p): Likewise. (GOMP_set_offload_targets): Rewrite. * testsuite/libgomp.oacc-c++/c++.exp: Provide offload target in "-DACC_DEVICE_TYPE_host", and "-DACC_DEVICE_TYPE_nvidia". * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. * testsuite/libgomp.oacc-c/offload-targets-1.c: New file. * testsuite/libgomp.oacc-c/offload-targets-2.c: Likewise. * testsuite/libgomp.oacc-c/offload-targets-3.c: Likewise. * testsuite/libgomp.oacc-c/offload-targets-4.c: Likewise. * testsuite/libgomp.oacc-c/offload-targets-5.c: Likewise. * testsuite/libgomp.oacc-c/offload-targets-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise. * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise. * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise. --- libgomp/ChangeLog.openacc | 34 ++++ libgomp/oacc-init.c | 7 + libgomp/openacc.h | 2 + libgomp/target.c | 178 +++++++++++++++++++-- libgomp/testsuite/libgomp.oacc-c++/c++.exp | 4 +- .../libgomp.oacc-c-c++-common/acc-on-device-2.c | 2 +- .../libgomp.oacc-c-c++-common/acc_on_device-1.c | 4 +- .../libgomp.oacc-c-c++-common/pr85381-2.c | 3 +- .../libgomp.oacc-c-c++-common/pr85381-3.c | 3 +- .../libgomp.oacc-c-c++-common/pr85381-4.c | 3 +- .../libgomp.oacc-c-c++-common/pr85381-5.c | 3 +- .../testsuite/libgomp.oacc-c-c++-common/pr85381.c | 3 +- .../libgomp.oacc-c-c++-common/pr85486-2.c | 3 +- .../libgomp.oacc-c-c++-common/pr85486-3.c | 3 +- .../testsuite/libgomp.oacc-c-c++-common/pr85486.c | 3 +- libgomp/testsuite/libgomp.oacc-c/c.exp | 4 +- .../testsuite/libgomp.oacc-c/offload-targets-1.c | 119 ++++++++++++++ .../testsuite/libgomp.oacc-c/offload-targets-2.c | 2 + .../testsuite/libgomp.oacc-c/offload-targets-3.c | 10 ++ .../testsuite/libgomp.oacc-c/offload-targets-4.c | 11 ++ .../testsuite/libgomp.oacc-c/offload-targets-5.c | 10 ++ .../testsuite/libgomp.oacc-c/offload-targets-6.c | 11 ++ .../libgomp.oacc-fortran/acc_on_device-1-1.f90 | 4 +- .../libgomp.oacc-fortran/acc_on_device-1-2.f | 4 +- .../libgomp.oacc-fortran/acc_on_device-1-3.f | 4 +- libgomp/testsuite/libgomp.oacc-fortran/fortran.exp | 4 +- 26 files changed, 400 insertions(+), 38 deletions(-) Grüße Thomas diff --git libgomp/ChangeLog.openacc libgomp/ChangeLog.openacc index d43b259..48b1f96 100644 --- libgomp/ChangeLog.openacc +++ libgomp/ChangeLog.openacc @@ -1,3 +1,37 @@ +2018-05-20 Thomas Schwinge + + PR libgomp/81886 + * openacc.h (enum acc_device_t): Add _acc_device_intel_mic, + _acc_device_hsa. + * oacc-init.c (get_openacc_name): Handle these. + (resolve_device): Debugging output. + * target.c (resolve_device, gomp_init_device) + (gomp_offload_target_available_p): Likewise. + (GOMP_set_offload_targets): Rewrite. + * testsuite/libgomp.oacc-c++/c++.exp: Provide offload target in + "-DACC_DEVICE_TYPE_host", and "-DACC_DEVICE_TYPE_nvidia". + * testsuite/libgomp.oacc-c/c.exp: Likewise. + * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. + * testsuite/libgomp.oacc-c/offload-targets-1.c: New file. + * testsuite/libgomp.oacc-c/offload-targets-2.c: Likewise. + * testsuite/libgomp.oacc-c/offload-targets-3.c: Likewise. + * testsuite/libgomp.oacc-c/offload-targets-4.c: Likewise. + * testsuite/libgomp.oacc-c/offload-targets-5.c: Likewise. + * testsuite/libgomp.oacc-c/offload-targets-6.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c: Adjust. + * testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/pr85381.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise. + * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise. + * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise. + * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise. + 2018-05-18 Cesar Philippidis Backport from mainline diff --git libgomp/oacc-init.c libgomp/oacc-init.c index d8348c0..19c2687 100644 --- libgomp/oacc-init.c +++ libgomp/oacc-init.c @@ -92,6 +92,8 @@ goacc_register (struct gomp_device_descr *disp) static const char * get_openacc_name (const char *name) { + /* not supported: _acc_device_intel_mic */ + /* not supported: _acc_device_hsa */ if (strcmp (name, "nvptx") == 0) return "nvidia"; else @@ -108,6 +110,8 @@ name_of_acc_device_t (enum acc_device_t type) case acc_device_host: return "host"; case acc_device_not_host: return "not_host"; case acc_device_nvidia: return "nvidia"; + case /* not supported */ _acc_device_intel_mic: + case /* not supported */ _acc_device_hsa: default: gomp_fatal ("unknown device type %u", (unsigned) type); } } @@ -119,6 +123,8 @@ name_of_acc_device_t (enum acc_device_t type) static struct gomp_device_descr * resolve_device (acc_device_t d, bool fail_is_error) { + gomp_debug (0, "%s (%d)\n", __FUNCTION__, (int) d); + acc_device_t d_arg = d; switch (d) @@ -203,6 +209,7 @@ resolve_device (acc_device_t d, bool fail_is_error) gomp_fatal ("device type %s not supported", name_of_acc_device_t (d)); } + gomp_debug (0, " %s: %d: %p\n", __FUNCTION__, (int) d, dispatchers[d]); return dispatchers[d]; } diff --git libgomp/openacc.h libgomp/openacc.h index 102723a..3d6d57e 100644 --- libgomp/openacc.h +++ libgomp/openacc.h @@ -55,6 +55,8 @@ typedef enum acc_device_t { /* acc_device_host_nonshm = 3 removed. */ acc_device_not_host = 4, acc_device_nvidia = 5, + /* not supported */ _acc_device_intel_mic = 6, + /* not supported */ _acc_device_hsa = 7, _ACC_device_hwm, /* Ensure enumeration is layout compatible with int. */ _ACC_highest = __INT_MAX__, diff --git libgomp/target.c libgomp/target.c index aa27dc8..b5f86c8 100644 --- libgomp/target.c +++ libgomp/target.c @@ -108,6 +108,8 @@ gomp_get_num_devices (void) static struct gomp_device_descr * resolve_device (int device) { + gomp_debug (0, "%s (%d)\n", __FUNCTION__, device); + int device_id; if (device == GOMP_DEVICE_ICV) { @@ -137,6 +139,7 @@ resolve_device (int device) && !gomp_offload_target_available_p (devices[device_id].type)) return NULL; + gomp_debug (0, " %s (%d): %d\n", __FUNCTION__, device, device_id); return &devices[device_id]; } @@ -1883,6 +1886,9 @@ GOMP_offload_unregister (const void *host_table, int target_type, attribute_hidden void gomp_init_device (struct gomp_device_descr *devicep) { + gomp_debug (0, "%s (%s; %d; %d)\n", __FUNCTION__, + devicep->name, (int) devicep->type, devicep->target_id); + int i; if (!devicep->init_device_func (devicep->target_id)) { @@ -1946,6 +1952,8 @@ gomp_unload_device (struct gomp_device_descr *devicep) attribute_hidden bool gomp_offload_target_available_p (int type) { + gomp_debug (0, "%s (%d)\n", __FUNCTION__, type); + bool available = false; /* Has the offload target already been initialized? */ @@ -1987,6 +1995,7 @@ gomp_offload_target_available_p (int type) gomp_mutex_unlock (®ister_lock); } + gomp_debug (0, " %s (%d): %d\n", __FUNCTION__, type, (int) available); return available; } @@ -3157,25 +3166,170 @@ offload_target_to_plugin_name (const char *offload_target) gomp_fatal ("Unknown offload target: %s", offload_target); } -/* List of offload targets, separated by colon. Defaults to the list +/* List of requested offload targets, separated by colon. Defaults to the list determined when configuring libgomp. */ static const char *gomp_offload_targets = OFFLOAD_TARGETS; -static bool gomp_offload_targets_init = false; +static bool gomp_offload_targets_set = false; +static bool gomp_offload_targets_malloced = false; -/* Override the list of offload targets. This must be called early, and only - once. */ +/* This function frees gomp_offload_targets. */ + +static void +free_gomp_offload_targets (void) +{ + free ((char *) gomp_offload_targets); +} + +/* Override the list of requested offload targets. This must be called + early, before gomp_target_init. */ void GOMP_set_offload_targets (const char *offload_targets) { - gomp_debug (0, "%s (\"%s\")\n", __FUNCTION__, offload_targets); - - /* Make sure this gets called early. */ - assert (gomp_is_initialized == PTHREAD_ONCE_INIT); - /* Make sure this only gets called once. */ - assert (!gomp_offload_targets_init); - gomp_offload_targets_init = true; - gomp_offload_targets = offload_targets; + gomp_debug (0, "%s (\"%s\"): %s\n", __FUNCTION__, + offload_targets, gomp_offload_targets); + + /* TODO: multithreading, locking. */ + /* TODO: this should not (sometimes) keep a copy of the offload_target + pointer, so that the caller knows what to expect. */ + /* TODO: What actually is supposed to happen if some parts of a program are + compiled with, for example, "-foffload=disable" (that is, when called with + the empty string for offload_targets), and others for other actual + (possibly different) offload targets? */ + if (gomp_is_initialized == PTHREAD_ONCE_INIT) + { + /* If we have not yet initialized, we capture all the offload targets + requested. We do not worry that the set of requested offload targets + vs. the set of available offload data will eventually match; any such + inconsistencies would be user error. (See also + gomp_offload_target_available_p.) */ + if (!gomp_offload_targets_set) + gomp_offload_targets = offload_targets; + else if (gomp_offload_targets == offload_targets + || strcmp (gomp_offload_targets, offload_targets) == 0) + /* Nothing to do if the same. */; + else + { + /* Merge offload_targets into gomp_offload_targets. */ + /* TODO: this could be simpler if we had the data available in a + different form. */ + size_t gomp_offload_targets_len = strlen (gomp_offload_targets); + /* Maximum length. */ + size_t len = (gomp_offload_targets_len + /* ":" */ 1 + + strlen (offload_targets) + /* '\0' */ 1); + char *gomp_offload_targets_new = gomp_malloc (len); + memcpy (gomp_offload_targets_new, + gomp_offload_targets, gomp_offload_targets_len); + char *gomp_offload_targets_new_next + = gomp_offload_targets_new + gomp_offload_targets_len; + *gomp_offload_targets_new_next = '\0'; + const char *cur = offload_targets; + while (*cur) + { + const char *cur_end = strchr (cur, ':'); + /* If no other offload target following... */ + if (cur_end == NULL) + /* ..., point to the terminating NUL character. */ + cur_end = cur + strlen (cur); + size_t cur_len = cur_end - cur; + + /* Do we already have this one listed? */ + const char *haystack = gomp_offload_targets_new; + while (haystack != NULL) + { + if (strncmp (haystack, cur, cur_len) == 0) + break; + else + { + haystack = strchr (haystack, ':'); + if (haystack != NULL) + haystack += /* ':' */ 1; + } + } + if (haystack == NULL) + { + /* Not yet listed; add it. */ + if (gomp_offload_targets_new_next != gomp_offload_targets_new) + *gomp_offload_targets_new_next++ = ':'; + assert (gomp_offload_targets_new_next + cur_len + /* '\0' */ 1 + <= gomp_offload_targets_new + len); + memcpy (gomp_offload_targets_new_next, cur, cur_len); + gomp_offload_targets_new_next += cur_len; + *gomp_offload_targets_new_next = '\0'; + } + + if (*cur_end == '\0') + break; + cur = cur_end + /* : */ 1; + } + + if (gomp_offload_targets_malloced) + free ((char *) gomp_offload_targets); + else + { + if (atexit (free_gomp_offload_targets) != 0) + gomp_fatal ("atexit failed"); + } + + gomp_offload_targets = gomp_offload_targets_new; + gomp_offload_targets_malloced = true; + } + } + else + { + /* If we have already initialized (which can happen only if a shared + library with another GOMP_set_offload_targets constructor call gets + loaded dynamically), and the user is now requesting offload targets + that were not requested previously, then we're out of luck: we can't + load new plugins now. Otherwise, we're all set. */ + if (gomp_offload_targets == offload_targets + || strcmp (gomp_offload_targets, offload_targets) == 0) + /* All fine if the same. */; + else + { + /* Check offload_targets against gomp_offload_targets. */ + /* TODO: this could be simpler if we had the data available in a + different form. */ + const char *cur = offload_targets; + while (*cur) + { + const char *cur_end = strchr (cur, ':'); + /* If no other offload target following... */ + if (cur_end == NULL) + /* ..., point to the terminating NUL character. */ + cur_end = cur + strlen (cur); + size_t cur_len = cur_end - cur; + + /* Do we have this one listed? */ + const char *haystack = gomp_offload_targets; + while (haystack != NULL) + { + if (strncmp (haystack, cur, cur_len) == 0) + break; + else + { + haystack = strchr (haystack, ':'); + if (haystack != NULL) + haystack += /* ':' */ 1; + } + } + if (haystack == NULL) + { + /* Not listed. */ + gomp_fatal ("Can't satisfy request for offload targets: %s; have loaded: %s", + offload_targets, gomp_offload_targets); + } + + if (*cur_end == '\0') + break; + cur = cur_end + /* : */ 1; + } + } + } + gomp_offload_targets_set = true; + + gomp_debug (0, " %s (\"%s\"): %s\n", __FUNCTION__, + offload_targets, gomp_offload_targets); } /* This function initializes the runtime needed for offloading. diff --git libgomp/testsuite/libgomp.oacc-c++/c++.exp libgomp/testsuite/libgomp.oacc-c++/c++.exp index 695b96d..2e17504 100644 --- libgomp/testsuite/libgomp.oacc-c++/c++.exp +++ libgomp/testsuite/libgomp.oacc-c++/c++.exp @@ -86,7 +86,7 @@ if { $lang_test_file_found } { switch -glob $offload_target_openacc { disable { set acc_mem_shared 1 - set tagopt "-DACC_DEVICE_TYPE_host=1" + set tagopt "-DACC_DEVICE_TYPE_host=\"\"" } nvptx* { if { ![check_effective_target_openacc_nvidia_accel_present] } { @@ -102,7 +102,7 @@ if { $lang_test_file_found } { lappend ALWAYS_CFLAGS "additional_flags=-I${srcdir}/libgomp.oacc-c-c++-common" set acc_mem_shared 0 - set tagopt "-DACC_DEVICE_TYPE_nvidia=1" + set tagopt "-DACC_DEVICE_TYPE_nvidia=\"$offload_target_openacc\"" } default { set acc_mem_shared 0 diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c index bfcb67d..758b1fc 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c @@ -14,7 +14,7 @@ int main () int expect = 1; -#if ACC_DEVICE_TYPE_host +#ifdef ACC_DEVICE_TYPE_host expect = 0; #endif diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c index 8112745..0270d06 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c @@ -37,7 +37,7 @@ main (int argc, char *argv[]) } -#if !ACC_DEVICE_TYPE_host +#ifndef ACC_DEVICE_TYPE_host /* Offloaded. */ @@ -49,7 +49,7 @@ main (int argc, char *argv[]) abort (); if (!acc_on_device (acc_device_not_host)) abort (); -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia if (!acc_on_device (acc_device_nvidia)) abort (); #else diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c index e5d02cf..6570c64 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c @@ -1,5 +1,6 @@ /* { dg-additional-options "-save-temps" } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } + { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ int main (void) diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c index 7d9ba1b..c5d1c5a 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-3.c @@ -1,5 +1,6 @@ /* { dg-additional-options "-save-temps -w" } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } + { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ int a; #pragma acc declare create(a) diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c index 477297d..d955d79 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c @@ -1,5 +1,6 @@ /* { dg-additional-options "-save-temps -w" } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } + { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ #define n 1024 diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c index 4653009..61e7e48 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c @@ -1,5 +1,6 @@ /* { dg-additional-options "-save-temps" } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } + { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ #define n 1024 diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c index f585ae5..2864dfc 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c @@ -1,5 +1,6 @@ /* { dg-additional-options "-save-temps" } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1 -O2" } } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } + { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ int main (void) diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c index a92b5dd..0f74921 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c @@ -1,5 +1,4 @@ -/* { dg-do run } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1" } } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-fopenacc-dim=-:-:128" } */ /* Minimized from ref-1.C. */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c index ae62206..b4ef878 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c @@ -1,5 +1,4 @@ -/* { dg-do run } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1" } } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-fopenacc-dim=-:-:-" } */ /* { dg-set-target-env-var "GOMP_OPENACC_DIM" "-:-:128" } */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c index f91dee0..99c0805 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c @@ -1,5 +1,4 @@ -/* { dg-do run } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_DEVICE_TYPE_nvidia=1" } } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ /* Minimized from ref-1.C. */ diff --git libgomp/testsuite/libgomp.oacc-c/c.exp libgomp/testsuite/libgomp.oacc-c/c.exp index 16f8295..73a7a5a 100644 --- libgomp/testsuite/libgomp.oacc-c/c.exp +++ libgomp/testsuite/libgomp.oacc-c/c.exp @@ -47,7 +47,7 @@ foreach offload_target_openacc $offload_targets_s_openacc { switch -glob $offload_target_openacc { disable { set acc_mem_shared 1 - set tagopt "-DACC_DEVICE_TYPE_host=1" + set tagopt "-DACC_DEVICE_TYPE_host=\"\"" } nvptx* { if { ![check_effective_target_openacc_nvidia_accel_present] } { @@ -63,7 +63,7 @@ foreach offload_target_openacc $offload_targets_s_openacc { lappend ALWAYS_CFLAGS "additional_flags=-I${srcdir}/libgomp.oacc-c-c++-common" set acc_mem_shared 0 - set tagopt "-DACC_DEVICE_TYPE_nvidia=1" + set tagopt "-DACC_DEVICE_TYPE_nvidia=\"$offload_target_openacc\"" } default { set acc_mem_shared 0 diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c new file mode 100644 index 0000000..b62a587 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/offload-targets-1.c @@ -0,0 +1,119 @@ +/* Test what happens for repeated GOMP_set_offload_targets calls, which happens + when shared libraries are involved, for example. As in the libgomp + testsuite infrastructure, it is difficult to build and link against shared + libraries, we simulate that by replicating some relevant + GOMP_set_offload_targets calls. */ + +#include +#include +#include +#include +#include "libgomp_g.h" + +int main () +{ + /* Before getting here, GOMP_set_offload_targets already got called via a + constructor. */ + + bool acc_device_types_requested[_ACC_device_hwm]; + for (int i = 0; i < _ACC_device_hwm; ++i) + acc_device_types_requested[i] = false; + + /* We're building for only one offload target ("-foffload=[...]") which is + the following. */ + const char *offload_target_requested; + acc_device_t acc_device_type_requested; +#if defined ACC_DEVICE_TYPE_nvidia + offload_target_requested = ACC_DEVICE_TYPE_nvidia; + acc_device_type_requested = acc_device_nvidia; +#elif defined ACC_DEVICE_TYPE_host + offload_target_requested = ACC_DEVICE_TYPE_host; + acc_device_type_requested = acc_device_host; +#else +# error Not ported to this ACC_DEVICE_TYPE +#endif + acc_device_types_requested[acc_device_type_requested] = true; + +#ifdef OFFLOAD_TARGETS_SAME_AGAIN + /* Call again; will have no noticeable difference. */ + GOMP_set_offload_targets (offload_target_requested); +#endif + +#ifdef OFFLOAD_TARGETS_ADD_EARLY + /* Request a (non-existing) offloading target (which will result in a + non-fatal diagnostic). */ + GOMP_set_offload_targets (OFFLOAD_TARGETS_ADD); +#endif + +#ifdef OFFLOAD_TARGETS_SAME_AGAIN + /* Call again; will have no noticeable difference. */ + GOMP_set_offload_targets (offload_target_requested); + char *s; + { + size_t len = 3 * (strlen (offload_target_requested) + 1); +# ifdef OFFLOAD_TARGETS_ADD_EARLY + len += 3 * (strlen (OFFLOAD_TARGETS_ADD) + 1); +# endif + s = malloc (len); + if (s == NULL) + __builtin_abort (); + size_t len_; +# ifndef OFFLOAD_TARGETS_ADD_EARLY + len_ = sprintf (s, "%s:%s:%s", + offload_target_requested, + offload_target_requested, + offload_target_requested); +# else + len_ = sprintf (s, "%s:%s:%s:%s:%s:%s", + offload_target_requested, + offload_target_requested, + OFFLOAD_TARGETS_ADD, + OFFLOAD_TARGETS_ADD, + offload_target_requested, + OFFLOAD_TARGETS_ADD); +# endif + if (len_ + 1 != len) + __builtin_abort (); + GOMP_set_offload_targets (s); + } +#endif + + /* Calling acc_get_num_devices will implicitly initialize offloading. */ +#if defined OFFLOAD_TARGETS_ADD_EARLY + fprintf (stderr, "CheCKpOInT1\n"); +#endif + /* acc_device_host is always available. */ + if ((acc_get_num_devices (acc_device_host) > 0) == false) + __builtin_abort (); +#if defined OFFLOAD_TARGETS_ADD_EARLY + fprintf (stderr, "WrONg WAy1\n"); +#endif + for (acc_device_t acc_device_type = acc_device_not_host + 1; + acc_device_type < _ACC_device_hwm; + ++acc_device_type) + { + /* The requested device type must be available. Any other device types + must not be available. */ + if ((acc_get_num_devices (acc_device_type) > 0) + != acc_device_types_requested[acc_device_type]) + __builtin_abort (); + } + +#ifdef OFFLOAD_TARGETS_SAME_AGAIN + /* Request the same again; will have no noticeable difference. */ + GOMP_set_offload_targets (offload_target_requested); +#endif +#if defined OFFLOAD_TARGETS_ADD_LATE + fprintf (stderr, "CheCKpOInT2\n"); + GOMP_set_offload_targets (OFFLOAD_TARGETS_ADD); + fprintf (stderr, "WrONg WAy2\n"); +#endif +#ifdef OFFLOAD_TARGETS_SAME_AGAIN + GOMP_set_offload_targets (s); + + /* Implementation defail: OK to "free (s)", in this case. */ + free (s); +#endif + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c new file mode 100644 index 0000000..977c559 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/offload-targets-2.c @@ -0,0 +1,2 @@ +#define OFFLOAD_TARGETS_SAME_AGAIN +#include "offload-targets-1.c" diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c new file mode 100644 index 0000000..1eb080b --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/offload-targets-3.c @@ -0,0 +1,10 @@ +#define OFFLOAD_TARGETS_ADD "XYZ" +#define OFFLOAD_TARGETS_ADD_EARLY +#include "offload-targets-1.c" + +/* + { dg-output "CheCKpOInT1(\n|\r\n|\r)+" } + { dg-output "libgomp: Unknown offload target: XYZ(\n|\r\n|\r)+" } + { dg-output "$" } + { dg-shouldfail "" } +*/ diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c new file mode 100644 index 0000000..2bb7204 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/offload-targets-4.c @@ -0,0 +1,11 @@ +#define OFFLOAD_TARGETS_SAME_AGAIN +#define OFFLOAD_TARGETS_ADD "XYZ" +#define OFFLOAD_TARGETS_ADD_EARLY +#include "offload-targets-1.c" + +/* + { dg-output "CheCKpOInT1(\n|\r\n|\r)+" } + { dg-output "libgomp: Unknown offload target: XYZ(\n|\r\n|\r)+" } + { dg-output "$" } + { dg-shouldfail "" } +*/ diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c new file mode 100644 index 0000000..8ba0792 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/offload-targets-5.c @@ -0,0 +1,10 @@ +#define OFFLOAD_TARGETS_ADD "XYZ" +#define OFFLOAD_TARGETS_ADD_LATE +#include "offload-targets-1.c" + +/* + { dg-output "CheCKpOInT2(\n|\r\n|\r)+" } + { dg-output "libgomp: Can't satisfy request for offload targets: XYZ; have loaded: \[a-z-\]*(\n|\r\n|\r)+" } + { dg-output "$" } + { dg-shouldfail "" } +*/ diff --git libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c new file mode 100644 index 0000000..4b15582 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/offload-targets-6.c @@ -0,0 +1,11 @@ +#define OFFLOAD_TARGETS_SAME_AGAIN +#define OFFLOAD_TARGETS_ADD "XYZ" +#define OFFLOAD_TARGETS_ADD_LATE +#include "offload-targets-1.c" + +/* + { dg-output "CheCKpOInT2(\n|\r\n|\r)+" } + { dg-output "libgomp: Can't satisfy request for offload targets: XYZ; have loaded: \[a-z-\]*(\n|\r\n|\r)+" } + { dg-output "$" } + { dg-shouldfail "" } +*/ diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 index 1a10f32..f57a2f2 100644 --- libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 +++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 @@ -25,7 +25,7 @@ if (acc_on_device (acc_device_nvidia)) call abort !$acc end parallel -#if !ACC_DEVICE_TYPE_host +#ifndef ACC_DEVICE_TYPE_host ! Offloaded. @@ -33,7 +33,7 @@ if (acc_on_device (acc_device_nvidia)) call abort if (acc_on_device (acc_device_none)) call abort if (acc_on_device (acc_device_host)) call abort if (.not. acc_on_device (acc_device_not_host)) call abort -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia if (.not. acc_on_device (acc_device_nvidia)) call abort #else if (acc_on_device (acc_device_nvidia)) call abort diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f index cbd1dd9..6209d12 100644 --- libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f +++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f @@ -26,7 +26,7 @@ !$ACC END PARALLEL -#if !ACC_DEVICE_TYPE_host +#ifndef ACC_DEVICE_TYPE_host ! Offloaded. @@ -34,7 +34,7 @@ IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT #else IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f index c391776..90d567f 100644 --- libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f +++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f @@ -25,7 +25,7 @@ !$ACC END PARALLEL -#if !ACC_DEVICE_TYPE_host +#ifndef ACC_DEVICE_TYPE_host ! Offloaded. @@ -33,7 +33,7 @@ IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT -#if ACC_DEVICE_TYPE_nvidia +#ifdef ACC_DEVICE_TYPE_nvidia IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT #else IF (ACC_ON_DEVICE (ACC_DEVICE_NVIDIA)) CALL ABORT diff --git libgomp/testsuite/libgomp.oacc-fortran/fortran.exp libgomp/testsuite/libgomp.oacc-fortran/fortran.exp index d78ce55..865c704 100644 --- libgomp/testsuite/libgomp.oacc-fortran/fortran.exp +++ libgomp/testsuite/libgomp.oacc-fortran/fortran.exp @@ -71,7 +71,7 @@ if { $lang_test_file_found } { switch -glob $offload_target_openacc { disable { set acc_mem_shared 1 - set tagopt "-DACC_DEVICE_TYPE_host=1" + set tagopt "-DACC_DEVICE_TYPE_host=\"\"" } nvptx* { if { ![check_effective_target_openacc_nvidia_accel_present] } { @@ -81,7 +81,7 @@ if { $lang_test_file_found } { } set acc_mem_shared 0 - set tagopt "-DACC_DEVICE_TYPE_nvidia=1" + set tagopt "-DACC_DEVICE_TYPE_nvidia=\"$offload_target_openacc\"" } default { set acc_mem_shared 0