From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 120538 invoked by alias); 2 Dec 2015 14:18:21 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 120518 invoked by uid 89); 2 Dec 2015 14:18:20 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.6 required=5.0 tests=AWL,BAYES_00,KAM_LAZY_DOMAIN_SECURITY,RCVD_IN_DNSWL_LOW,T_RP_MATCHES_RCVD autolearn=no version=3.3.2 X-HELO: smtp.ispras.ru Received: from smtp.ispras.ru (HELO smtp.ispras.ru) (83.149.199.79) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 02 Dec 2015 14:18:17 +0000 Received: from [10.10.3.121] (unknown [83.149.199.91]) by smtp.ispras.ru (Postfix) with ESMTP id B2349203FB; Wed, 2 Dec 2015 17:18:13 +0300 (MSK) Date: Wed, 02 Dec 2015 14:18:00 -0000 From: Alexander Monakov To: Jakub Jelinek cc: gcc-patches@gcc.gnu.org, Bernd Schmidt , Dmitry Melnik Subject: Re: [gomp-nvptx 4/9] nvptx backend: add -mgomp option and multilib In-Reply-To: <20151202105617.GH5675@tucnak.redhat.com> Message-ID: References: <1448983707-18854-1-git-send-email-amonakov@ispras.ru> <1448983707-18854-5-git-send-email-amonakov@ispras.ru> <20151202105617.GH5675@tucnak.redhat.com> User-Agent: Alpine 2.20 (LNX 67 2015-01-07) MIME-Version: 1.0 Content-Type: text/plain; charset=US-ASCII X-SW-Source: 2015-12/txt/msg00272.txt.bz2 On Wed, 2 Dec 2015, Jakub Jelinek wrote: > I thought the MULTILIB* vars allow you to multilib on none of > -msoft-stack/-muniform-simt and both -msoft-stack/-muniform-simt, without > building other variants, so you wouldn't need this. The nice effect of having -mgomp is better factorization: if I need to change what OpenMP needs, e.g. for going with your suggestion below and dropping -msoft-stack, I need to only change one line. Otherwise I'd have to change mkoffload too. > Furthermore, as I said, I believe for e.g. most of newlib libc / libm > I think it is enough if they are built as -muniform-simt -mno-soft-stack, > if those functions are leaf or don't call user routines that could have > #pragma omp parallel. -msoft-stack would unnecessarily slow the routines > down. Not obviously so. Outside of SIMD regions, running on hard stacks pointlessly amplifies cache/memory traffic for stack references, so there would have to be some evaluation before deciding. > So perhaps just multilib on -muniform-simt, and document that -muniform-simt > built code requires also that the soft-stack var is set up and thus > -msoft-stack can be used when needed? It's an interesting point, but I have doubts. Is that something you'd want me to address short-term? > Can you post sample code with assembly for -msoft-stack and -muniform-simt > showing how are short interesting cases expanded? > Is there really no way even in direct PTX assembly to have .local file scope > vars (rather than the global arrays indexed by %tid)? Allow me to post samples a bit later; as for .local, the PTX documentation explicitely states it must not be done: 5.1.5. Local State Space [...] When compiling to use the Application Binary Interface (ABI), .local state-space variables must be declared within function scope and are allocated on the stack. In implementations that do not support a stack, all local memory variables are stored at fixed addresses, recursive function calls are not supported, and .local variables may be declared at module scope. When compiling legacy PTX code (ISA versions prior to 3.0) containing module-scoped .local variables, the compiler silently disables use of the ABI. (while I'm unsure as to what exactly "compiling to use the ABI" is defined, I'm assuming that's what we want in GCC, and otherwise linking may not work) Thanks. Alexander