This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp-nvptx 4/9] nvptx backend: add -mgomp option and multilib
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Alexander Monakov <amonakov at ispras dot ru>
- Cc: gcc-patches at gcc dot gnu dot org, Bernd Schmidt <bschmidt at redhat dot com>, Dmitry Melnik <dm at ispras dot ru>
- Date: Wed, 2 Dec 2015 11:56:17 +0100
- Subject: Re: [gomp-nvptx 4/9] nvptx backend: add -mgomp option and multilib
- Authentication-results: sourceware.org; auth=none
- References: <1448983707-18854-1-git-send-email-amonakov at ispras dot ru> <1448983707-18854-5-git-send-email-amonakov at ispras dot ru>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
On Tue, Dec 01, 2015 at 06:28:22PM +0300, Alexander Monakov wrote:
> Since OpenMP offloading requires both soft-stacks and "uniform SIMT", both
> non-traditional codegen variants, I'm building a multilib variant with those
> enabled. This patch adds option -mgomp which enables -msoft-stack plus
> -muniform-simt, and builds a multilib with it.
>
> * config/nvptx/nvptx.c (nvptx_option_override): Handle TARGET_GOMP.
> * config/nvptx/nvptx.opt (mgomp): New option.
> * config/nvptx/t-nvptx (MULTILIB_OPTIONS): New.
> * doc/invoke.texi (mgomp): Document.
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.
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.
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?
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)?
Jakub