From patchwork Mon Feb 15 18:44:11 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Alexander Monakov X-Patchwork-Id: 583064 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 B11EF1401CA for ; Tue, 16 Feb 2016 05:45:42 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=wOu0ghHp; dkim-atps=neutral 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:date:message-id:in-reply-to:references; q=dns; s= default; b=BqOymd58DHAWT6uexvQ3QZiRhuHzKQmZ8ySFAXhzNV4uh2RJYVLwJ EVE7bLviQ4xPEwFuolSw/BbTIDYgWIT6Ab5Ri3Zd02Akb02Ci8eDMIGHsmC75+9g coBZK+eZ1yan1OE5BkvXLaIodgFYAssEe3smsnOCuDepzgn6nehIx0= 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:date:message-id:in-reply-to:references; s=default; bh=vU+4EyyA4IE7qEW/h8S0nMv5WCM=; b=wOu0ghHpBpSTfsS5iZXoXuzSCZ8z KvaXroxPueDmyTVnvXLsLUUhjmqgiYo9XN5ZWuQ6zKHPiHePk2uLt34h+JOLqbKq SqXwmcO7uqJxdkangOog50844WabmyWg6sjv2KSrktF7Yk03VU5NM1/C8AW+MCjU u5EpQ9KVZz48XPg= Received: (qmail 31051 invoked by alias); 15 Feb 2016 18:44:24 -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 30984 invoked by uid 89); 15 Feb 2016 18:44:23 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.5 required=5.0 tests=AWL, BAYES_00, KAM_LAZY_DOMAIN_SECURITY autolearn=no version=3.3.2 spammy=349, Hx-languages-length:1997, HX-HELO:eggs.gnu.org, Hx-spam-relays-external:208.118.235.92 X-HELO: eggs.gnu.org Received: from eggs.gnu.org (HELO eggs.gnu.org) (208.118.235.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-SHA encrypted) ESMTPS; Mon, 15 Feb 2016 18:44:21 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1aVO88-0003Gy-4h for gcc-patches@gcc.gnu.org; Mon, 15 Feb 2016 13:44:19 -0500 Received: from smtp.ispras.ru ([83.149.199.79]:35299) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1aVO87-0003GH-Tc for gcc-patches@gcc.gnu.org; Mon, 15 Feb 2016 13:44:16 -0500 Received: from condor.intra.ispras.ru (unknown [83.149.199.91]) by smtp.ispras.ru (Postfix) with ESMTP id 218B7226B9 for ; Mon, 15 Feb 2016 21:44:13 +0300 (MSK) Received: by condor.intra.ispras.ru (Postfix, from userid 23246) id 095AC1225D9F; Mon, 15 Feb 2016 21:44:13 +0300 (MSK) From: Alexander Monakov To: gcc-patches@gcc.gnu.org Subject: [gomp-nvptx 4/5] libgomp: remove __nvptx_stacks setup code Date: Mon, 15 Feb 2016 21:44:11 +0300 Message-Id: <1455561852-9237-5-git-send-email-amonakov@ispras.ru> In-Reply-To: <1455561852-9237-1-git-send-email-amonakov@ispras.ru> References: <1455561852-9237-1-git-send-email-amonakov@ispras.ru> X-detected-operating-system: by eggs.gnu.org: GNU/Linux 3.x X-Received-From: 83.149.199.79 X-IsSubscribed: yes This patch implements the NVPTX libgomp part of the transition to host-allocated soft stacks. The wrapper around gomp_nvptx_main previously responsible for that is no longer needed. This mostly reverts commit b408f1293e29a009ba70a3fda7b800277e1f310a. * config/nvptx/team.c: (gomp_nvptx_main_1): Rename back to... (gomp_nvptx_main): ...this; delete the wrapper. diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c index bc8c4e6..b9f9f9f 100644 --- a/libgomp/config/nvptx/team.c +++ b/libgomp/config/nvptx/team.c @@ -34,9 +34,12 @@ struct gomp_thread *nvptx_thrs __attribute__((shared)); static void gomp_thread_start (struct gomp_thread_pool *); -static void __attribute__((noinline)) -gomp_nvptx_main_1 (void (*fn) (void *), void *fn_data, int ntids, int tid) +void +gomp_nvptx_main (void (*fn) (void *), void *fn_data) { + int tid, ntids; + asm ("mov.u32 %0, %%tid.y;" : "=r" (tid)); + asm ("mov.u32 %0, %%ntid.y;" : "=r" (ntids)); if (tid == 0) { gomp_global_icv.nthreads_var = ntids; @@ -69,30 +72,6 @@ gomp_nvptx_main_1 (void (*fn) (void *), void *fn_data, int ntids, int tid) } } -void -gomp_nvptx_main (void (*fn) (void *), void *fn_data) -{ - int tid, ntids; - asm ("mov.u32 %0, %%tid.y;" : "=r" (tid)); - asm ("mov.u32 %0, %%ntid.y;" : "=r" (ntids)); - char *stacks = 0; - int *__nvptx_uni; - asm ("cvta.shared.u64 %0, __nvptx_uni;" : "=r" (__nvptx_uni)); - __nvptx_uni[tid] = 0; - if (tid == 0) - { - size_t stacksize = 131072; - stacks = gomp_malloc (stacksize * ntids); - char **__nvptx_stacks = 0; - asm ("cvta.shared.u64 %0, __nvptx_stacks;" : "=r" (__nvptx_stacks)); - for (int i = 0; i < ntids; i++) - __nvptx_stacks[i] = stacks + stacksize * (i + 1); - } - asm ("bar.sync 0;"); - gomp_nvptx_main_1 (fn, fn_data, ntids, tid); - free (stacks); -} - /* This function is a pthread_create entry point. This contains the idle loop in which a thread waits to be called up to become part of a team. */