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
