[gomp-nvptx 0/7] Various fixes

[gomp-nvptx 0/7] Various fixes

[Alexander Monakov]

Hello,

I have committed the following patches to amonakov/gomp-nvptx branch to fix a
few bugs uncovered in recent testing (including testing on a 32-bit ARM
platform).  Patch 1 fixes an odd mispaste in bar.c, patches 2,5,6,7 address
32-bit portability issues, patch 3 works around a deadlock on error reporting
(this is a regression that is also visible on trunk with OpenACC offloading),
and patch 4 is a slightly more comprehensive fix to nvptx debuginfo generation.

Alexander Monakov (7):
  libgomp: remove paste error in gomp_team_barrier_wait_end
  nvptx libgcc: use attribute shared
  libgomp plugin: make cuMemFreeHost error non-fatal
  nvptx backend: re-enable line info generation
  nvptx backend: use POINTER_SIZE instead of BITS_PER_WORD
  nvptx backend: change mul.u32 to mul.lo.u32
  nvptx backend: define STACK_SIZE_MODE

 gcc/ChangeLog.gomp-nvptx      | 23 +++++++++++++++++++++++
 gcc/config/nvptx/nvptx.c      | 21 ++++++---------------
 gcc/config/nvptx/nvptx.h      |  1 +
 libgcc/ChangeLog.gomp-nvptx   |  8 ++++++++
 libgcc/config/nvptx/crt0.c    | 10 ++++------
 libgcc/config/nvptx/stacks.c  |  9 ++-------
 libgomp/ChangeLog.gomp-nvptx  |  9 +++++++++
 libgomp/config/nvptx/bar.c    |  2 --
 libgomp/plugin/plugin-nvptx.c |  2 +-
 9 files changed, 54 insertions(+), 31 deletions(-)

[gomp-nvptx 1/7] libgomp: remove paste error in gomp_team_barrier_wait_end

[Alexander Monakov]

	* config/nvptx/bar.c: Remove wrong invocation of
	gomp_barrier_wait_end from gomp_team_barrier_wait_end.
---
 libgomp/ChangeLog.gomp-nvptx | 5 +++++
 libgomp/config/nvptx/bar.c   | 2 --
 2 files changed, 5 insertions(+), 2 deletions(-)

diff --git a/libgomp/config/nvptx/bar.c b/libgomp/config/nvptx/bar.c
index e6e8daa..a0d8a44 100644
--- a/libgomp/config/nvptx/bar.c
+++ b/libgomp/config/nvptx/bar.c
@@ -80,8 +80,6 @@ gomp_team_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
 {
   unsigned int generation, gen;
 
-  gomp_barrier_wait_end (bar, state);
-
   if (__builtin_expect (state & BAR_WAS_LAST, 0))
     {
       /* Next time we'll be awaiting TOTAL threads again.  */

[gomp-nvptx 4/7] nvptx backend: re-enable line info generation

[Alexander Monakov]

	* config/nvptx/nvptx.c (nvptx_option_override): Remove custom handling
	of debug info options.
---
 gcc/ChangeLog.gomp-nvptx | 5 +++++
 gcc/config/nvptx/nvptx.c | 9 ---------
 2 files changed, 5 insertions(+), 9 deletions(-)

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 81dd9a2..e69e0be 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -156,15 +156,6 @@ nvptx_option_override (void)
   /* Assumes that it will see only hard registers.  */
   flag_var_tracking = 0;
 
-  if (write_symbols == DBX_DEBUG)
-    /* The stabs testcases want to know stabs isn't supported.  */
-    sorry ("stabs debug format not supported");
-
-  /* Actually we don't have any debug format, but don't be
-     unneccesarily noisy.  */
-  write_symbols = NO_DEBUG;
-  debug_info_level = DINFO_LEVEL_NONE;
-
   if (nvptx_optimize < 0)
     nvptx_optimize = optimize > 0;
 

[gomp-nvptx 2/7] nvptx libgcc: use attribute shared

[Alexander Monakov]

	* config/nvptx/crt0.c (__nvptx_stacks): Define in C.  Use it...
	(__nvptx_uni): Ditto.
	(__main): ...here instead of inline asm.
	* config/nvptx/stacks.c (__nvptx_stacks): Define in C.
	(__nvptx_uni): Ditto.
---
 libgcc/ChangeLog.gomp-nvptx  |  8 ++++++++
 libgcc/config/nvptx/crt0.c   | 10 ++++------
 libgcc/config/nvptx/stacks.c |  9 ++-------
 3 files changed, 14 insertions(+), 13 deletions(-)

diff --git a/libgcc/config/nvptx/crt0.c b/libgcc/config/nvptx/crt0.c
index 5e04b0f..9e9a25e 100644
--- a/libgcc/config/nvptx/crt0.c
+++ b/libgcc/config/nvptx/crt0.c
@@ -41,10 +41,8 @@ abort (void)
   exit (255);
 }
 
-asm ("// BEGIN GLOBAL VAR DECL: __nvptx_stacks");
-asm (".extern .shared .u64 __nvptx_stacks[32];");
-asm ("// BEGIN GLOBAL VAR DECL: __nvptx_uni");
-asm (".extern .shared .u32 __nvptx_uni[32];");
+extern char *__nvptx_stacks[32] __attribute__((shared));
+extern unsigned __nvptx_uni[32] __attribute__((shared));
 
 extern int main (int argc, char *argv[]);
 
@@ -54,8 +52,8 @@ __main (int *__retval, int __argc, char *__argv[])
   __exitval = __retval;
 
   static char gstack[131072] __attribute__((aligned(8)));
-  asm ("st.shared.u64 [__nvptx_stacks], %0;" : : "r" (gstack + sizeof gstack));
-  asm ("st.shared.u32 [__nvptx_uni], %0;" : : "r" (0));
+  __nvptx_stacks[0] = gstack + sizeof gstack;
+  __nvptx_uni[0] = 0;
 
   exit (main (__argc, __argv));
 }
diff --git a/libgcc/config/nvptx/stacks.c b/libgcc/config/nvptx/stacks.c
index a7e640a..4640fc9 100644
--- a/libgcc/config/nvptx/stacks.c
+++ b/libgcc/config/nvptx/stacks.c
@@ -21,10 +21,5 @@
    see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
    <http://www.gnu.org/licenses/>.  */
 
-/* __shared__ char *__nvptx_stacks[32];  */
-asm ("// BEGIN GLOBAL VAR DEF: __nvptx_stacks");
-asm (".visible .shared .u64 __nvptx_stacks[32];");
-
-/* __shared__ unsigned __nvptx_uni[32];  */
-asm ("// BEGIN GLOBAL VAR DEF: __nvptx_uni");
-asm (".visible .shared .u32 __nvptx_uni[32];");
+char *__nvptx_stacks[32] __attribute__((shared)) = { 0 };
+unsigned __nvptx_uni[32] __attribute__((shared)) = { 0 };

[gomp-nvptx 3/7] libgomp plugin: make cuMemFreeHost error non-fatal

[Alexander Monakov]

Unlike cuMemFree and other resource-releasing functions called on exit,
cuMemFreeHost appears to re-report errors encountered in kernel launch.
This leads to a deadlock after GOMP_PLUGIN_fatal is reentered.

While the behavior on libgomp side is suboptimal (there's no need to
call resource-releasing functions if we're about to destroy the CUDA
context anyway), this behavior on cuMemFreeHost part is not useful
and just makes error "recovery" harder.  This was reported to NVIDIA
(bug ref. 1737876), but we can work around it by simply reporting the
error without making it fatal.

	* plugin/plugin-nvptx.c (map_fini): Make cuMemFreeHost error non-fatal.
---
 libgomp/ChangeLog.gomp-nvptx  | 4 ++++
 libgomp/plugin/plugin-nvptx.c | 2 +-
 2 files changed, 5 insertions(+), 1 deletion(-)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index adf57b1..4e44242 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -135,7 +135,7 @@ map_fini (struct ptx_stream *s)
 
   r = cuMemFreeHost (s->h);
   if (r != CUDA_SUCCESS)
-    GOMP_PLUGIN_fatal ("cuMemFreeHost error: %s", cuda_error (r));
+    GOMP_PLUGIN_error ("cuMemFreeHost error: %s", cuda_error (r));
 }
 
 static void

[gomp-nvptx 7/7] nvptx backend: define STACK_SIZE_MODE

[Alexander Monakov]

Default definition of STACK_SIZE_MODE is word_mode, which is DImode on NVPTX.
However, stack pointer mode matches pointer mode, so needs to be SImode on
32-bit NVPTX ABI.  Define it to Pmode to fix 32-bit code generation.

	* config/nvptx/nvptx.h (STACK_SIZE_MODE): Define.
---
 gcc/ChangeLog.gomp-nvptx | 4 ++++
 gcc/config/nvptx/nvptx.h | 1 +
 2 files changed, 5 insertions(+)

diff --git a/gcc/config/nvptx/nvptx.h b/gcc/config/nvptx/nvptx.h
index 7810cca..6da4d06 100644
--- a/gcc/config/nvptx/nvptx.h
+++ b/gcc/config/nvptx/nvptx.h
@@ -83,6 +83,7 @@
 
 #define POINTER_SIZE (TARGET_ABI64 ? 64 : 32)
 #define Pmode (TARGET_ABI64 ? DImode : SImode)
+#define STACK_SIZE_MODE Pmode
 
 /* Registers.  Since ptx is a virtual target, we just define a few
    hard registers for special purposes and leave pseudos unallocated.

[gomp-nvptx 5/7] nvptx backend: use POINTER_SIZE instead of BITS_PER_WORD

[Alexander Monakov]

POINTER_SIZE is the proper macro to retrieve pointer size in bits for the
target ABI, but new code incorrectly used BITS_PER_WORD, breaking 32-bit
code generation.

	* config/nvptx/nvptx.c (nvptx_init_unisimt_predicate): Use
	POINTER_SIZE instead of BITS_PER_WORD.
	(nvptx_declare_function_name): Ditto.
	(nvptx_output_return): Ditto.
	(nvptx_file_end): Ditto.
---
 gcc/ChangeLog.gomp-nvptx | 8 ++++++++
 gcc/config/nvptx/nvptx.c | 8 ++++----
 2 files changed, 12 insertions(+), 4 deletions(-)

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index e69e0be..93bf781 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -946,7 +946,7 @@ nvptx_init_axis_predicate (FILE *file, int regno, const char *name)
 static void
 nvptx_init_unisimt_predicate (FILE *file)
 {
-  int bits = BITS_PER_WORD;
+  int bits = POINTER_SIZE;
   int master = REGNO (cfun->machine->unisimt_master);
   int pred = REGNO (cfun->machine->unisimt_predicate);
   fprintf (file, "\t{\n");
@@ -1108,7 +1108,7 @@ nvptx_declare_function_name (FILE *file, const char *name, const_tree decl)
       /* Maintain 64-bit stack alignment.  */
       int keep_align = BIGGEST_ALIGNMENT / BITS_PER_UNIT;
       sz = ROUND_UP (sz, keep_align);
-      int bits = BITS_PER_WORD;
+      int bits = POINTER_SIZE;
       fprintf (file, "\t.reg.u%d %%frame;\n", bits);
       fprintf (file, "\t.reg.u32 %%fstmp0;\n");
       fprintf (file, "\t.reg.u%d %%fstmp1;\n", bits);
@@ -1177,7 +1177,7 @@ nvptx_output_return (void)
 
   if (cfun->machine->using_softstack)
     fprintf (asm_out_file, "\tst.shared.u%d [%%fstmp2], %%fstmp1;\n",
-	     BITS_PER_WORD);
+	     POINTER_SIZE);
 
   if (mode != VOIDmode)
     fprintf (asm_out_file, "\tst.param%s\t[%s_out], %s;\n",
@@ -4191,7 +4191,7 @@ nvptx_file_end (void)
     {
       write_var_marker (asm_out_file, false, true, "__nvptx_stacks");
       fprintf (asm_out_file, ".extern .shared .u%d __nvptx_stacks[32];\n",
-	       BITS_PER_WORD);
+	       POINTER_SIZE);
     }
   if (need_unisimt_decl)
     {

[gomp-nvptx 6/7] nvptx backend: change mul.u32 to mul.lo.u32

[Alexander Monakov]

Recent testing uncovered that PTX JIT may reject attempts to use 'mul.u32' as
a non-widening 32-bit multiply instruction.  Use 'mul.lo.u32' to fix 32-bit
code generation and conform to the PTX spec better.

	* config/nvptx/nvptx.c (nvptx_init_unisimt_predicate): Emit
	'mul.lo.u32' instead of 'mul.u32' for 32-bit ABI target.
	(nvptx_declare_function_name): Ditto.
---
 gcc/ChangeLog.gomp-nvptx | 6 ++++++
 gcc/config/nvptx/nvptx.c | 4 ++--
 2 files changed, 8 insertions(+), 2 deletions(-)

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 93bf781..bc187ea 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -955,7 +955,7 @@ nvptx_init_unisimt_predicate (FILE *file)
   fprintf (file, "\t\t.reg.u%d %%ustmp2;\n", bits);
   fprintf (file, "\t\tmov.u32 %%ustmp0, %%tid.y;\n");
   fprintf (file, "\t\tmul%s.u32 %%ustmp1, %%ustmp0, 4;\n",
-	   bits == 64 ? ".wide" : "");
+	   bits == 64 ? ".wide" : ".lo");
   fprintf (file, "\t\tmov.u%d %%ustmp2, __nvptx_uni;\n", bits);
   fprintf (file, "\t\tadd.u%d %%ustmp2, %%ustmp2, %%ustmp1;\n", bits);
   fprintf (file, "\t\tld.shared.u32 %%r%d, [%%ustmp2];\n", master);
@@ -1115,7 +1115,7 @@ nvptx_declare_function_name (FILE *file, const char *name, const_tree decl)
       fprintf (file, "\t.reg.u%d %%fstmp2;\n", bits);
       fprintf (file, "\tmov.u32 %%fstmp0, %%tid.y;\n");
       fprintf (file, "\tmul%s.u32 %%fstmp1, %%fstmp0, %d;\n",
-	       bits == 64 ? ".wide" : "", bits / 8);
+	       bits == 64 ? ".wide" : ".lo", bits / 8);
       fprintf (file, "\tmov.u%d %%fstmp2, __nvptx_stacks;\n", bits);
       /* fstmp2 = &__nvptx_stacks[tid.y];  */
       fprintf (file, "\tadd.u%d %%fstmp2, %%fstmp2, %%fstmp1;\n", bits);