[PATCH, 0/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS,JIT} in libgomp nvptx plugin

[PATCH, 0/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS,JIT} in libgomp nvptx plugin

[Tom de Vries]

Hi,

I've written a patch series to facilitate debugging libgomp openacc 
testcase failures on the nvptx accelerator.


When running an openacc test-case on an nvptx accelerator, the following 
happens:
- the plugin obtains the ptx assembly for the acceleration kernels
- it calls the cuda jit to compile and link the ptx into a module
- it loads the module
- it starts an acceleration kernel

The patch series adds these environment variables:
- GOMP_OPENACC_NVPTX_SAVE_TEMPS: a means to save the resulting module
   such that it can be investigated using nvdisasm and cuobjdump.
- GOMP_OPENACC_NVPTX_DISASM: a means to see the resulting module in
   the debug output,  by writing it into a file and calling nvdisasm on
   it
- GOMP_OPENACC_NVPTX_JIT: a means to set parameters of the
   compilation/linking process, currently supporting:
   * -O[0-4], mapping onto CU_JIT_OPTIMIZATION_LEVEL
   * -ori, mapping onto CU_JIT_NEW_SM3X_OPT


The patch series consists of these patches:

1. Show value of GOMP_OPENACC_DIM in libgomp nvptx plugin
2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
3. Handle GOMP_OPENACC_NVPTX_JIT=-O[0-4] in libgomp nvptx plugin
4. Handle GOMP_OPENACC_NVPTX_JIT=-ori in libgomp nvptx plugin


I've tested the patch series on top of gomp-4_0-branch, by running an 
openacc testcase from the command line and defining the various 
environment variables.

[ A relevant difference between gomp-4_0-branch and master is that:
- master defines and includes ./libgomp/plugin/cuda/cuda.h, so I had to
   add the CU_JIT constants there, while
- gomp-4_0-branch doesn't define that local minimal cuda.h file but
   includes cuda's cuda.h. My setup linked against cuda 6.5 which defines
   CU_JIT_OPTIMIZATION_LEVEL but not yet CU_JIT_NEW_SM3X_OPT (that seems
   to have been introduced at cuda 8.0), so I had to hardcode the latter.
]


OK for trunk if bootstrap and reg-test on x86_64 with nvidia accelerator 
succeeds?

Thanks,
- Tom

[PATCH, 1/4] Show value of GOMP_OPENACC_DIM in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1286 bytes --]

On 06/26/2017 01:24 PM, Tom de Vries wrote:
> Hi,
> 
> I've written a patch series to facilitate debugging libgomp openacc 
> testcase failures on the nvptx accelerator.
> 
> 
> When running an openacc test-case on an nvptx accelerator, the following 
> happens:
> - the plugin obtains the ptx assembly for the acceleration kernels
> - it calls the cuda jit to compile and link the ptx into a module
> - it loads the module
> - it starts an acceleration kernel
> 
> The patch series adds these environment variables:
> - GOMP_OPENACC_NVPTX_SAVE_TEMPS: a means to save the resulting module
>    such that it can be investigated using nvdisasm and cuobjdump.
> - GOMP_OPENACC_NVPTX_DISASM: a means to see the resulting module in
>    the debug output,  by writing it into a file and calling nvdisasm on
>    it
> - GOMP_OPENACC_NVPTX_JIT: a means to set parameters of the
>    compilation/linking process, currently supporting:
>    * -O[0-4], mapping onto CU_JIT_OPTIMIZATION_LEVEL
>    * -ori, mapping onto CU_JIT_NEW_SM3X_OPT
> 
> 
> The patch series consists of these patches:
> 
> 1. Show value of GOMP_OPENACC_DIM in libgomp nvptx plugin

This patch adds a debug message (for GOMP_DEBUG=1) about the value of 
the GOMP_OPENACC_DIM variable read from the environment.

Thanks,
- Tom

[-- Attachment #2: 0001-Show-value-of-GOMP_OPENACC_DIM-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 1475 bytes --]

Show value of GOMP_OPENACC_DIM in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (notify_var): New function.
	(nvptx_exec): Use notify_var for GOMP_OPENACC_DIM.

---
 libgomp/plugin/plugin-nvptx.c | 12 +++++++++++-
 1 file changed, 11 insertions(+), 1 deletion(-)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 0e1b3e2..71630b5 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -867,6 +867,14 @@ nvptx_get_num_devices (void)
   return n;
 }
 
+static void
+notify_var (const char *var_name, const char *env_var)
+{
+  if (env_var == NULL)
+    GOMP_PLUGIN_debug (0, "%s: <Not defined>\n", var_name);
+  else
+    GOMP_PLUGIN_debug (0, "%s: '%s'\n", var_name, env_var);
+}
 
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
@@ -1089,10 +1097,12 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
       pthread_mutex_lock (&ptx_dev_lock);
       if (!default_dims[0])
 	{
+	  const char *var_name = "GOMP_OPENACC_DIM";
 	  /* We only read the environment variable once.  You can't
 	     change it in the middle of execution.  The syntax  is
 	     the same as for the -fopenacc-dim compilation option.  */
-	  const char *env_var = getenv ("GOMP_OPENACC_DIM");
+	  const char *env_var = getenv (var_name);
+	  notify_var (var_name, env_var);
 	  if (env_var)
 	    {
 	      const char *pos = env_var;

[PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1348 bytes --]

On 06/26/2017 01:24 PM, Tom de Vries wrote:
> Hi,
> 
> I've written a patch series to facilitate debugging libgomp openacc 
> testcase failures on the nvptx accelerator.
> 
> 
> When running an openacc test-case on an nvptx accelerator, the following 
> happens:
> - the plugin obtains the ptx assembly for the acceleration kernels
> - it calls the cuda jit to compile and link the ptx into a module
> - it loads the module
> - it starts an acceleration kernel
> 
> The patch series adds these environment variables:
> - GOMP_OPENACC_NVPTX_SAVE_TEMPS: a means to save the resulting module
>    such that it can be investigated using nvdisasm and cuobjdump.
> - GOMP_OPENACC_NVPTX_DISASM: a means to see the resulting module in
>    the debug output,  by writing it into a file and calling nvdisasm on
>    it
> - GOMP_OPENACC_NVPTX_JIT: a means to set parameters of the
>    compilation/linking process, currently supporting:
>    * -O[0-4], mapping onto CU_JIT_OPTIMIZATION_LEVEL
>    * -ori, mapping onto CU_JIT_NEW_SM3X_OPT
> 
> 
> The patch series consists of these patches:
> 
> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

This patch adds handling of:
- GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
- GOMP_OPENACC_NVPTX_DISASM=[01]

The filename used for dumping the module is plugin-nvptx.<pid>.cubin.

Thanks,
- Tom

[-- Attachment #2: 0002-Handle-GOMP_OPENACC_NVPTX_-DISASM-SAVE_TEMPS-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 3864 bytes --]

Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (do_prog, debug_linkout): New function.
	(link_ptx): Use debug_linkout.

---
 libgomp/plugin/plugin-nvptx.c | 103 ++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 103 insertions(+)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 71630b5..df1bfdd 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -47,6 +47,9 @@
 #include <unistd.h>
 #include <assert.h>
 #include <errno.h>
+#include <stdlib.h>
+#include <sys/types.h>
+#include <sys/wait.h>
 
 #if PLUGIN_NVPTX_DYNAMIC
 # include <dlfcn.h>
@@ -876,6 +879,104 @@ notify_var (const char *var_name, const char *env_var)
     GOMP_PLUGIN_debug (0, "%s: '%s'\n", var_name, env_var);
 }
 
+static void
+do_prog (const char *prog, const char *arg)
+{
+  pid_t pid = fork ();
+
+  if (pid == -1)
+    {
+      GOMP_PLUGIN_error ("Fork failed");
+      return;
+    }
+  else if (pid > 0)
+    {
+      int status;
+      waitpid (pid, &status, 0);
+      if (!WIFEXITED (status))
+	GOMP_PLUGIN_error ("Running %s %s failed", prog, arg);
+    }
+  else
+    {
+      execlp (prog, prog /* argv[0] */, arg, NULL);
+      abort ();
+    }
+}
+
+static void
+debug_linkout (void *linkout, size_t linkoutsize)
+{
+  static int gomp_openacc_nvptx_disasm = -1;
+  if (gomp_openacc_nvptx_disasm == -1)
+    {
+      const char *var_name = "GOMP_OPENACC_NVPTX_DISASM";
+      const char *env_var = getenv (var_name);
+      notify_var (var_name, env_var);
+      gomp_openacc_nvptx_disasm
+	= ((env_var != NULL && env_var[0] == '1' && env_var[1] == '\0')
+	   ? 1 : 0);
+    }
+
+  static int gomp_openacc_nvptx_save_temps = -1;
+  if (gomp_openacc_nvptx_save_temps == -1)
+    {
+      const char *var_name = "GOMP_OPENACC_NVPTX_SAVE_TEMPS";
+      const char *env_var = getenv (var_name);
+      notify_var (var_name, env_var);
+      gomp_openacc_nvptx_save_temps
+	= ((env_var != NULL && env_var[0] == '1' && env_var[1] == '\0')
+	   ? 1 : 0);
+    }
+
+  if (gomp_openacc_nvptx_disasm == 0
+      && gomp_openacc_nvptx_save_temps == 0)
+    return;
+
+  const char *prefix = "plugin-nvptx.";
+  const char *postfix = ".cubin";
+  const int len =	(strlen (prefix)
+			 + 20 /* %lld.  */
+			 + strlen (postfix)
+			 + 1  /* '\0'.  */);
+  char file_name[len];
+  int res = snprintf (file_name, len, "%s%lld%s", prefix,
+		      (long long)getpid (), postfix);
+  assert (res < len); /* Assert there's no truncation.  */
+
+  GOMP_PLUGIN_debug (0, "Generating %s with size %zu\n",
+		     file_name, linkoutsize);
+  FILE *cubin_file = fopen (file_name, "wb");
+  if (cubin_file == NULL)
+    {
+      GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+      return;
+    }
+
+  fwrite (linkout, linkoutsize, 1, cubin_file);
+  unsigned int write_succeeded = ferror (cubin_file) == 0;
+  if (!write_succeeded)
+    GOMP_PLUGIN_debug (0, "Writing %s failed\n", file_name);
+
+  res = fclose (cubin_file);
+  if (res != 0)
+    GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+
+  if (!write_succeeded)
+    return;
+
+  if (gomp_openacc_nvptx_disasm == 1)
+    {
+      GOMP_PLUGIN_debug (0, "Disassembling %s\n", file_name);
+      do_prog ("nvdisasm", file_name);
+    }
+
+  if (gomp_openacc_nvptx_save_temps == 0)
+    {
+      GOMP_PLUGIN_debug (0, "Removing %s\n", file_name);
+      remove (file_name);
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
@@ -939,6 +1040,8 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
       return false;
     }
 
+  debug_linkout (linkout, linkoutsize);
+
   CUDA_CALL (cuModuleLoadData, module, linkout);
   CUDA_CALL (cuLinkDestroy, linkstate);
   return true;

Re: [PATCH, 0/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS,JIT} in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1232 bytes --]

On 06/26/2017 01:24 PM, Tom de Vries wrote:
> Hi,
> 
> I've written a patch series to facilitate debugging libgomp openacc 
> testcase failures on the nvptx accelerator.
> 
> 
> When running an openacc test-case on an nvptx accelerator, the following 
> happens:
> - the plugin obtains the ptx assembly for the acceleration kernels
> - it calls the cuda jit to compile and link the ptx into a module
> - it loads the module
> - it starts an acceleration kernel
> 
> The patch series adds these environment variables:
> - GOMP_OPENACC_NVPTX_SAVE_TEMPS: a means to save the resulting module
>    such that it can be investigated using nvdisasm and cuobjdump.
> - GOMP_OPENACC_NVPTX_DISASM: a means to see the resulting module in
>    the debug output,  by writing it into a file and calling nvdisasm on
>    it
> - GOMP_OPENACC_NVPTX_JIT: a means to set parameters of the
>    compilation/linking process, currently supporting:
>    * -O[0-4], mapping onto CU_JIT_OPTIMIZATION_LEVEL
>    * -ori, mapping onto CU_JIT_NEW_SM3X_OPT
> 
> 
> The patch series consists of these patches:
> 
> 3. Handle GOMP_OPENACC_NVPTX_JIT=-O[0-4] in libgomp nvptx plugin

This patch adds handling of Handle GOMP_OPENACC_NVPTX_JIT=-O[0-4].

Thanks,
- Tom

[-- Attachment #2: 0003-Handle-GOMP_OPENACC_NVPTX_JIT-O-0-4-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 2593 bytes --]

Handle GOMP_OPENACC_NVPTX_JIT=-O[0-4] in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/cuda/cuda.h (enum CUjit_option): Add CU_JIT_OPTIMIZATION_LEVEL.
	* plugin/plugin-nvptx.c (process_GOMP_OPENACC_NVPTX_JIT): New function.
	(link_ptx): Add CU_JIT_OPTIMIZATION_LEVEL to opts.

---
 libgomp/plugin/cuda/cuda.h    |  1 +
 libgomp/plugin/plugin-nvptx.c | 44 ++++++++++++++++++++++++++++++++++++++++---
 2 files changed, 42 insertions(+), 3 deletions(-)

diff --git a/libgomp/plugin/cuda/cuda.h b/libgomp/plugin/cuda/cuda.h
index 25d5d19..75dfe3d 100644
--- a/libgomp/plugin/cuda/cuda.h
+++ b/libgomp/plugin/cuda/cuda.h
@@ -88,6 +88,7 @@ typedef enum {
   CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES = 4,
   CU_JIT_ERROR_LOG_BUFFER = 5,
   CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES = 6,
+  CU_JIT_OPTIMIZATION_LEVEL = 7,
   CU_JIT_LOG_VERBOSE = 12
 } CUjit_option;
 
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index df1bfdd..3cd5557 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -977,12 +977,43 @@ debug_linkout (void *linkout, size_t linkoutsize)
     }
 }
 
+static void
+process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o)
+{
+  const char *var_name = "GOMP_OPENACC_NVPTX_JIT";
+  const char *env_var = getenv (var_name);
+  notify_var (var_name, env_var);
+
+  *gomp_openacc_nvptx_o = 4;
+  if (env_var == NULL)
+    return;
+
+  const char *c = env_var;
+  while (*c != '\0')
+    {
+      while (*c == ' ')
+	c++;
+
+      if (c[0] == '-' && c[1] == 'O'
+	  && '0' <= c[2] && c[2] <= '4'
+	  && (c[3] == '\0' || c[3] == ' '))
+	{
+	  *gomp_openacc_nvptx_o = c[2] - '0';
+	  c += 3;
+	  continue;
+	}
+
+      GOMP_PLUGIN_error ("Error parsing %s", var_name);
+      break;
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
 {
-  CUjit_option opts[6];
-  void *optvals[6];
+  CUjit_option opts[7];
+  void *optvals[7];
   float elapsed = 0.0;
   char elog[1024];
   char ilog[16384];
@@ -1009,7 +1040,14 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
   opts[5] = CU_JIT_LOG_VERBOSE;
   optvals[5] = (void *) 1;
 
-  CUDA_CALL (cuLinkCreate, 6, opts, optvals, &linkstate);
+  static intptr_t gomp_openacc_nvptx_o = -1;
+  if (gomp_openacc_nvptx_o == -1)
+    process_GOMP_OPENACC_NVPTX_JIT (&gomp_openacc_nvptx_o);
+
+  opts[6] = CU_JIT_OPTIMIZATION_LEVEL;
+  optvals[6] = (void *) gomp_openacc_nvptx_o;
+
+  CUDA_CALL (cuLinkCreate, 7, opts, optvals, &linkstate);
 
   for (; num_objs--; ptx_objs++)
     {

[PATCH, 4/4] Handle GOMP_OPENACC_NVPTX_JIT=-ori in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1219 bytes --]

On 06/26/2017 01:24 PM, Tom de Vries wrote:
> Hi,
> 
> I've written a patch series to facilitate debugging libgomp openacc 
> testcase failures on the nvptx accelerator.
> 
> 
> When running an openacc test-case on an nvptx accelerator, the following 
> happens:
> - the plugin obtains the ptx assembly for the acceleration kernels
> - it calls the cuda jit to compile and link the ptx into a module
> - it loads the module
> - it starts an acceleration kernel
> 
> The patch series adds these environment variables:
> - GOMP_OPENACC_NVPTX_SAVE_TEMPS: a means to save the resulting module
>    such that it can be investigated using nvdisasm and cuobjdump.
> - GOMP_OPENACC_NVPTX_DISASM: a means to see the resulting module in
>    the debug output,  by writing it into a file and calling nvdisasm on
>    it
> - GOMP_OPENACC_NVPTX_JIT: a means to set parameters of the
>    compilation/linking process, currently supporting:
>    * -O[0-4], mapping onto CU_JIT_OPTIMIZATION_LEVEL
>    * -ori, mapping onto CU_JIT_NEW_SM3X_OPT
> 
> 
> The patch series consists of these patches:
> 
> 4. Handle GOMP_OPENACC_NVPTX_JIT=-ori in libgomp nvptx plugin

This patch adds handling of GOMP_OPENACC_NVPTX_JIT=-ori.

Thanks,
- Tom

[-- Attachment #2: 0004-Handle-GOMP_OPENACC_NVPTX_JIT-ori-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 3081 bytes --]

Handle GOMP_OPENACC_NVPTX_JIT=-ori in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/cuda/cuda.h (enum CUjit_option): Add CU_JIT_NEW_SM3X_OPT.
	* plugin/plugin-nvptx.c (process_GOMP_OPENACC_NVPTX_JIT): Add
	gomp_openacc_nvptx_ori parameter.  Handle -ori.
	(link_ptx): Add CU_JIT_NEW_SM3X_OPT to opts.

---
 libgomp/plugin/cuda/cuda.h    |  3 ++-
 libgomp/plugin/plugin-nvptx.c | 30 +++++++++++++++++++++++++-----
 2 files changed, 27 insertions(+), 6 deletions(-)

diff --git a/libgomp/plugin/cuda/cuda.h b/libgomp/plugin/cuda/cuda.h
index 75dfe3d..4644870 100644
--- a/libgomp/plugin/cuda/cuda.h
+++ b/libgomp/plugin/cuda/cuda.h
@@ -89,7 +89,8 @@ typedef enum {
   CU_JIT_ERROR_LOG_BUFFER = 5,
   CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES = 6,
   CU_JIT_OPTIMIZATION_LEVEL = 7,
-  CU_JIT_LOG_VERBOSE = 12
+  CU_JIT_LOG_VERBOSE = 12,
+  CU_JIT_NEW_SM3X_OPT = 15
 } CUjit_option;
 
 typedef enum {
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 3cd5557..a8548fb 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -978,13 +978,15 @@ debug_linkout (void *linkout, size_t linkoutsize)
 }
 
 static void
-process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o)
+process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o,
+				intptr_t *gomp_openacc_nvptx_ori)
 {
   const char *var_name = "GOMP_OPENACC_NVPTX_JIT";
   const char *env_var = getenv (var_name);
   notify_var (var_name, env_var);
 
   *gomp_openacc_nvptx_o = 4;
+  *gomp_openacc_nvptx_ori = 0;
   if (env_var == NULL)
     return;
 
@@ -1003,6 +1005,14 @@ process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o)
 	  continue;
 	}
 
+      if (c[0] == '-' && c[1] == 'o' && c[2] == 'r' && c[3] == 'i'
+	  && (c[4] == '\0' || c[4] == ' '))
+	{
+	  *gomp_openacc_nvptx_ori = 1;
+	  c += 4;
+	  continue;
+	}
+
       GOMP_PLUGIN_error ("Error parsing %s", var_name);
       break;
     }
@@ -1012,8 +1022,8 @@ static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
 {
-  CUjit_option opts[7];
-  void *optvals[7];
+  CUjit_option opts[8];
+  void *optvals[8];
   float elapsed = 0.0;
   char elog[1024];
   char ilog[16384];
@@ -1041,13 +1051,23 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
   optvals[5] = (void *) 1;
 
   static intptr_t gomp_openacc_nvptx_o = -1;
+  static intptr_t gomp_openacc_nvptx_ori = -1;
   if (gomp_openacc_nvptx_o == -1)
-    process_GOMP_OPENACC_NVPTX_JIT (&gomp_openacc_nvptx_o);
+    process_GOMP_OPENACC_NVPTX_JIT (&gomp_openacc_nvptx_o,
+				    &gomp_openacc_nvptx_ori);
 
   opts[6] = CU_JIT_OPTIMIZATION_LEVEL;
   optvals[6] = (void *) gomp_openacc_nvptx_o;
 
-  CUDA_CALL (cuLinkCreate, 7, opts, optvals, &linkstate);
+  int nopts = 7;
+  if (gomp_openacc_nvptx_ori)
+    {
+      opts[nopts] = CU_JIT_NEW_SM3X_OPT;
+      optvals[nopts] = (void *) gomp_openacc_nvptx_ori;
+      nopts++;
+    }
+
+  CUDA_CALL (cuLinkCreate, nopts, opts, optvals, &linkstate);
 
   for (; num_objs--; ptx_objs++)
     {

[PATCH, 3/4] Handle GOMP_OPENACC_NVPTX_JIT=-O[0-4] in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1338 bytes --]

[ reposting with proper subject ]

On 06/26/2017 01:42 PM, Tom de Vries wrote:
> On 06/26/2017 01:24 PM, Tom de Vries wrote:
>> Hi,
>>
>> I've written a patch series to facilitate debugging libgomp openacc 
>> testcase failures on the nvptx accelerator.
>>
>>
>> When running an openacc test-case on an nvptx accelerator, the 
>> following happens:
>> - the plugin obtains the ptx assembly for the acceleration kernels
>> - it calls the cuda jit to compile and link the ptx into a module
>> - it loads the module
>> - it starts an acceleration kernel
>>
>> The patch series adds these environment variables:
>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS: a means to save the resulting module
>>    such that it can be investigated using nvdisasm and cuobjdump.
>> - GOMP_OPENACC_NVPTX_DISASM: a means to see the resulting module in
>>    the debug output,  by writing it into a file and calling nvdisasm on
>>    it
>> - GOMP_OPENACC_NVPTX_JIT: a means to set parameters of the
>>    compilation/linking process, currently supporting:
>>    * -O[0-4], mapping onto CU_JIT_OPTIMIZATION_LEVEL
>>    * -ori, mapping onto CU_JIT_NEW_SM3X_OPT
>>
>>
>> The patch series consists of these patches:
>>
>> 3. Handle GOMP_OPENACC_NVPTX_JIT=-O[0-4] in libgomp nvptx plugin
> 
> This patch adds handling of Handle GOMP_OPENACC_NVPTX_JIT=-O[0-4].

Thanks,
- Tom

[-- Attachment #2: 0003-Handle-GOMP_OPENACC_NVPTX_JIT-O-0-4-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 2593 bytes --]

Handle GOMP_OPENACC_NVPTX_JIT=-O[0-4] in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/cuda/cuda.h (enum CUjit_option): Add CU_JIT_OPTIMIZATION_LEVEL.
	* plugin/plugin-nvptx.c (process_GOMP_OPENACC_NVPTX_JIT): New function.
	(link_ptx): Add CU_JIT_OPTIMIZATION_LEVEL to opts.

---
 libgomp/plugin/cuda/cuda.h    |  1 +
 libgomp/plugin/plugin-nvptx.c | 44 ++++++++++++++++++++++++++++++++++++++++---
 2 files changed, 42 insertions(+), 3 deletions(-)

diff --git a/libgomp/plugin/cuda/cuda.h b/libgomp/plugin/cuda/cuda.h
index 25d5d19..75dfe3d 100644
--- a/libgomp/plugin/cuda/cuda.h
+++ b/libgomp/plugin/cuda/cuda.h
@@ -88,6 +88,7 @@ typedef enum {
   CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES = 4,
   CU_JIT_ERROR_LOG_BUFFER = 5,
   CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES = 6,
+  CU_JIT_OPTIMIZATION_LEVEL = 7,
   CU_JIT_LOG_VERBOSE = 12
 } CUjit_option;
 
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index df1bfdd..3cd5557 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -977,12 +977,43 @@ debug_linkout (void *linkout, size_t linkoutsize)
     }
 }
 
+static void
+process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o)
+{
+  const char *var_name = "GOMP_OPENACC_NVPTX_JIT";
+  const char *env_var = getenv (var_name);
+  notify_var (var_name, env_var);
+
+  *gomp_openacc_nvptx_o = 4;
+  if (env_var == NULL)
+    return;
+
+  const char *c = env_var;
+  while (*c != '\0')
+    {
+      while (*c == ' ')
+	c++;
+
+      if (c[0] == '-' && c[1] == 'O'
+	  && '0' <= c[2] && c[2] <= '4'
+	  && (c[3] == '\0' || c[3] == ' '))
+	{
+	  *gomp_openacc_nvptx_o = c[2] - '0';
+	  c += 3;
+	  continue;
+	}
+
+      GOMP_PLUGIN_error ("Error parsing %s", var_name);
+      break;
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
 {
-  CUjit_option opts[6];
-  void *optvals[6];
+  CUjit_option opts[7];
+  void *optvals[7];
   float elapsed = 0.0;
   char elog[1024];
   char ilog[16384];
@@ -1009,7 +1040,14 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
   opts[5] = CU_JIT_LOG_VERBOSE;
   optvals[5] = (void *) 1;
 
-  CUDA_CALL (cuLinkCreate, 6, opts, optvals, &linkstate);
+  static intptr_t gomp_openacc_nvptx_o = -1;
+  if (gomp_openacc_nvptx_o == -1)
+    process_GOMP_OPENACC_NVPTX_JIT (&gomp_openacc_nvptx_o);
+
+  opts[6] = CU_JIT_OPTIMIZATION_LEVEL;
+  optvals[6] = (void *) gomp_openacc_nvptx_o;
+
+  CUDA_CALL (cuLinkCreate, 7, opts, optvals, &linkstate);
 
   for (; num_objs--; ptx_objs++)
     {

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Joseph Myers]

On Mon, 26 Jun 2017, Tom de Vries wrote:

> > 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
> 
> This patch adds handling of:
> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
> - GOMP_OPENACC_NVPTX_DISASM=[01]
> 
> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.

Are you sure this use of getenv and writing to that file is safe for 
setuid/setgid programs?  I'd expect you to need to use secure_getenv as in 
plugin-hsa.c; certainly for anything that could results in writes to a 
file like that.

-- 
Joseph S. Myers
joseph@codesourcery.com

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Jakub Jelinek]

On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
> On Mon, 26 Jun 2017, Tom de Vries wrote:
> 
> > > 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
> > 
> > This patch adds handling of:
> > - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
> > - GOMP_OPENACC_NVPTX_DISASM=[01]
> > 
> > The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
> 
> Are you sure this use of getenv and writing to that file is safe for 
> setuid/setgid programs?  I'd expect you to need to use secure_getenv as in 
> plugin-hsa.c; certainly for anything that could results in writes to a 
> file like that.

Yeah, definitely it should be using secure_getenv/__secure_getenv.
And IMNSHO GOMP_DEBUG too.

	Jakub

[PATCH] Use secure_getenv for GOMP_DEBUG

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1410 bytes --]

[ was: Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in 
libgomp nvptx plugin ]

On 06/26/2017 05:29 PM, Jakub Jelinek wrote:
> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>
>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
>>>
>>> This patch adds handling of:
>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
>>>
>>> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
>>
>> Are you sure this use of getenv and writing to that file is safe for
>> setuid/setgid programs?  I'd expect you to need to use secure_getenv as in
>> plugin-hsa.c; certainly for anything that could results in writes to a
>> file like that.
> 
> Yeah, definitely it should be using secure_getenv/__secure_getenv.
> And IMNSHO GOMP_DEBUG too.
> 

This patch uses secure_getenv for GOMP_DEBUG.

It factors out the secure_getenv code from plugin-hsa.c into 
libgomp/secure_getenv.h, and reuses it in env.c.

I've added _GNU_SOURCE before the libgomp.h include in env.c to make 
sure that secure_getenv (imported from stdlib.h) is available.

I've also added a test-case that sets GOMP_DEBUG to 1 and verifies that 
some output is generated.

Build for c-only on x86_64 without accelerator, tested libgomp -m64/-m32.

OK if x86_64 bootstrap and reg-test succeeds?

Thanks,
- Tom

[-- Attachment #2: 0001-Use-secure_getenv-for-GOMP_DEBUG.patch --]
[-- Type: text/x-patch, Size: 7441 bytes --]

Use secure_getenv for GOMP_DEBUG

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* env.c (parse_unsigned_long_1): Factor out of ...
	(parse_unsigned_long): ... here.
	(parse_int_1): Factor out of ...
	(parse_int): ... here.
	(parse_int_secure): New function.
	(initialize_env): Use parse_int_secure for GOMP_DEBUG.
	* secure_getenv.h: Factor out of ...
	* plugin/plugin-hsa.c: ... here.
	* testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c: New test.

---
 libgomp/env.c                                      | 44 +++++++++++++++---
 libgomp/plugin/plugin-hsa.c                        | 27 +----------
 libgomp/secure_getenv.h                            | 53 ++++++++++++++++++++++
 .../libgomp.oacc-c-c++-common/gomp-debug-env.c     | 13 ++++++
 4 files changed, 104 insertions(+), 33 deletions(-)

diff --git a/libgomp/env.c b/libgomp/env.c
index ced752d..802c73b 100644
--- a/libgomp/env.c
+++ b/libgomp/env.c
@@ -26,6 +26,7 @@
 /* This file defines the OpenMP internal control variables and arranges
    for them to be initialized from environment variables at startup.  */
 
+#define _GNU_SOURCE
 #include "libgomp.h"
 #include "gomp-constants.h"
 #include <limits.h>
@@ -58,6 +59,8 @@
 #endif
 #endif /* LIBGOMP_OFFLOADED_ONLY */
 
+#include "secure_getenv.h"
+
 struct gomp_task_icv gomp_global_icv = {
   .nthreads_var = 1,
   .thread_limit_var = UINT_MAX,
@@ -171,15 +174,17 @@ parse_schedule (void)
 }
 
 /* Parse an unsigned long environment variable.  Return true if one was
-   present and it was successfully parsed.  */
+   present and it was successfully parsed.  If SECURE, use secure_getenv to the
+   environment variable.  */
 
 static bool
-parse_unsigned_long (const char *name, unsigned long *pvalue, bool allow_zero)
+parse_unsigned_long_1 (const char *name, unsigned long *pvalue, bool allow_zero,
+		       bool secure)
 {
   char *env, *end;
   unsigned long value;
 
-  env = getenv (name);
+  env = (secure ? secure_getenv (name) : getenv (name));
   if (env == NULL)
     return false;
 
@@ -206,14 +211,23 @@ parse_unsigned_long (const char *name, unsigned long *pvalue, bool allow_zero)
   return false;
 }
 
+/* As parse_unsigned_long_1, but always use getenv.  */
+
+static bool
+parse_unsigned_long (const char *name, unsigned long *pvalue, bool allow_zero)
+{
+  return parse_unsigned_long_1 (name, pvalue, allow_zero, false);
+}
+
 /* Parse a positive int environment variable.  Return true if one was
-   present and it was successfully parsed.  */
+   present and it was successfully parsed.  If SECURE, use secure_getenv to the
+   environment variable.  */
 
 static bool
-parse_int (const char *name, int *pvalue, bool allow_zero)
+parse_int_1 (const char *name, int *pvalue, bool allow_zero, bool secure)
 {
   unsigned long value;
-  if (!parse_unsigned_long (name, &value, allow_zero))
+  if (!parse_unsigned_long_1 (name, &value, allow_zero, secure))
     return false;
   if (value > INT_MAX)
     {
@@ -224,6 +238,22 @@ parse_int (const char *name, int *pvalue, bool allow_zero)
   return true;
 }
 
+/* As parse_int_1, but use getenv.  */
+
+static bool
+parse_int (const char *name, int *pvalue, bool allow_zero)
+{
+  return parse_int_1 (name, pvalue, allow_zero, false);
+}
+
+/* As parse_int_1, but use getenv_secure.  */
+
+static bool
+parse_int_secure (const char *name, int *pvalue, bool allow_zero)
+{
+  return parse_int_1 (name, pvalue, allow_zero, true);
+}
+
 /* Parse an unsigned long list environment variable.  Return true if one was
    present and it was successfully parsed.  */
 
@@ -1207,7 +1237,7 @@ initialize_env (void)
       gomp_global_icv.thread_limit_var
 	= thread_limit_var > INT_MAX ? UINT_MAX : thread_limit_var;
     }
-  parse_int ("GOMP_DEBUG", &gomp_debug_var, true);
+  parse_int_secure ("GOMP_DEBUG", &gomp_debug_var, true);
 #ifndef HAVE_SYNC_BUILTINS
   gomp_mutex_init (&gomp_managed_threads_lock);
 #endif
diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index 90ca247..adb07ac 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -39,32 +39,7 @@
 #include <dlfcn.h>
 #include "libgomp-plugin.h"
 #include "gomp-constants.h"
-
-/* Secure getenv() which returns NULL if running as SUID/SGID.  */
-#ifndef HAVE_SECURE_GETENV
-#ifdef HAVE___SECURE_GETENV
-#define secure_getenv __secure_getenv
-#elif defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
-  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
-
-#include <unistd.h>
-
-/* Implementation of secure_getenv() for targets where it is not provided but
-   we have at least means to test real and effective IDs. */
-
-static char *
-secure_getenv (const char *name)
-{
-  if ((getuid () == geteuid ()) && (getgid () == getegid ()))
-    return getenv (name);
-  else
-    return NULL;
-}
-
-#else
-#define secure_getenv getenv
-#endif
-#endif
+#include "secure-getenv.h"
 
 /* As an HSA runtime is dlopened, following structure defines function
    pointers utilized by the HSA plug-in.  */
diff --git a/libgomp/secure_getenv.h b/libgomp/secure_getenv.h
new file mode 100644
index 0000000..6804e61
--- /dev/null
+++ b/libgomp/secure_getenv.h
@@ -0,0 +1,53 @@
+/* Copyright (C) 2017 Free Software Foundation, Inc.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify
+it under the terms of the GNU General Public License as published by
+the Free Software Foundation; either version 3, or (at your option)
+any later version.
+
+GCC is distributed in the hope that it will be useful,
+but WITHOUT ANY WARRANTY; without even the implied warranty of
+MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+GNU General Public License for more details.
+
+Under Section 7 of GPL version 3, you are granted additional
+permissions described in the GCC Runtime Library Exception, version
+3.1, as published by the Free Software Foundation.
+
+You should have received a copy of the GNU General Public License and
+a copy of the GCC Runtime Library Exception along with this program;
+see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+<http://www.gnu.org/licenses/>.  */
+
+#ifndef _SECURE_GETENV_H
+#define _SECURE_GETENV_H 1
+
+/* Secure getenv() which returns NULL if running as SUID/SGID.  */
+#ifndef HAVE_SECURE_GETENV
+#ifdef HAVE___SECURE_GETENV
+#define secure_getenv __secure_getenv
+#elif defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
+  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
+
+#include <unistd.h>
+
+/* Implementation of secure_getenv() for targets where it is not provided but
+   we have at least means to test real and effective IDs.  */
+
+static char *
+secure_getenv (const char *name)
+{
+  if ((getuid () == geteuid ()) && (getgid () == getegid ()))
+    return getenv (name);
+  else
+    return NULL;
+}
+
+#else
+#define secure_getenv getenv
+#endif
+#endif
+
+#endif /* _SECURE_GETENV_H.  */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c
new file mode 100644
index 0000000..3fc3503
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c
@@ -0,0 +1,13 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var GOMP_DEBUG "1" } */
+
+/* Check that GOMP_DEBUG=1 triggers some output.  */
+
+int
+main (void)
+{
+#pragma acc parallel
+  ;
+}
+
+/* { dg-output "GOACC_parallel_keyed" } */

Re: [PATCH] Use secure_getenv for GOMP_DEBUG

[Jakub Jelinek]

On Tue, Jun 27, 2017 at 09:17:57AM +0200, Tom de Vries wrote:
> This patch uses secure_getenv for GOMP_DEBUG.
> 
> It factors out the secure_getenv code from plugin-hsa.c into
> libgomp/secure_getenv.h, and reuses it in env.c.
> 
> I've added _GNU_SOURCE before the libgomp.h include in env.c to make sure
> that secure_getenv (imported from stdlib.h) is available.
> 
> I've also added a test-case that sets GOMP_DEBUG to 1 and verifies that some
> output is generated.
> 
> Build for c-only on x86_64 without accelerator, tested libgomp -m64/-m32.
> 
> OK if x86_64 bootstrap and reg-test succeeds?
> 
> Thanks,
> - Tom

> Use secure_getenv for GOMP_DEBUG
> 

> --- /dev/null
> +++ b/libgomp/secure_getenv.h
> @@ -0,0 +1,53 @@
> +/* Copyright (C) 2017 Free Software Foundation, Inc.
> +
> +This file is part of GCC.
> +
> +GCC is free software; you can redistribute it and/or modify
> +it under the terms of the GNU General Public License as published by
> +the Free Software Foundation; either version 3, or (at your option)
> +any later version.
> +
> +GCC is distributed in the hope that it will be useful,
> +but WITHOUT ANY WARRANTY; without even the implied warranty of
> +MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
> +GNU General Public License for more details.
> +
> +Under Section 7 of GPL version 3, you are granted additional
> +permissions described in the GCC Runtime Library Exception, version
> +3.1, as published by the Free Software Foundation.
> +
> +You should have received a copy of the GNU General Public License and
> +a copy of the GCC Runtime Library Exception along with this program;
> +see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
> +<http://www.gnu.org/licenses/>.  */
> +
> +#ifndef _SECURE_GETENV_H
> +#define _SECURE_GETENV_H 1
> +
> +/* Secure getenv() which returns NULL if running as SUID/SGID.  */
> +#ifndef HAVE_SECURE_GETENV
> +#ifdef HAVE___SECURE_GETENV
> +#define secure_getenv __secure_getenv
> +#elif defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
> +  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
> +
> +#include <unistd.h>
> +
> +/* Implementation of secure_getenv() for targets where it is not provided but
> +   we have at least means to test real and effective IDs.  */
> +
> +static char *
> +secure_getenv (const char *name)

Shouldn't this be static inline char * ?  I mean, even at -O0 we don't want
it to be emitted into every TU.

Another thing is that we probably want to follow what libgfortran does for
the case where secure_getenv isn't available, but __secure_getenv is -
in particular emit this function (if geteuid and getegid are present), but
emit a weakref call to __secure_getenv first if non-NULL.

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c
> @@ -0,0 +1,13 @@
> +/* { dg-do run } */
> +/* { dg-set-target-env-var GOMP_DEBUG "1" } */
> +
> +/* Check that GOMP_DEBUG=1 triggers some output.  */
> +
> +int
> +main (void)
> +{
> +#pragma acc parallel
> +  ;
> +}
> +
> +/* { dg-output "GOACC_parallel_keyed" } */

Does dg-set-target-env-var work for remote testing?

	Jakub

[PATCH, 5/4] Handle GOMP_OPENACC_NVPTX_PTXRW in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1697 bytes --]

[ was: Re: [PATCH, 0/4] Handle 
GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS,JIT} in libgomp nvptx plugin ]

On 06/26/2017 01:24 PM, Tom de Vries wrote:
> Hi,
> 
> I've written a patch series to facilitate debugging libgomp openacc 
> testcase failures on the nvptx accelerator.
> 
> 
> When running an openacc test-case on an nvptx accelerator, the following 
> happens:
> - the plugin obtains the ptx assembly for the acceleration kernels
> - it calls the cuda jit to compile and link the ptx into a module
> - it loads the module
> - it starts an acceleration kernel

A typical scenario when developing the compiler is:
- run gcc test.c -save-temps
- run a.out
- edit test.s to fix bug or make code faster or smaller
- run gcc test.s
- run a.out
- edit compiler sources to make the compiler do the same as the .s edit

With openacc test-cases, this scenario is currently not available for 
ptx assembly. Using -save-temps -foffload=-save-temps we can get a .s 
containing ptx. But to insert the edited .s back into the compilation 
flow is difficult.

This patch facilitates such a scenario in the nvptx plugin.
- we define GOMP_OPENACC_NVPTX_PTXRW == 'w', and the plugin writes the
   ptx assembly into a series of files
- we edit one of those files
- we define GOMP_OPENACC_NVPTX_PTXRW == 'r', and the plugin reads the
   ptx assembly back from those files, and uses that instead of the ptx
   in the executable.

I've tested this patch series on top of gomp-4_0-branch, by running an
openacc testcase from the command line and going through the 
write-edit-readscenario with an observable ptx edit.

OK for trunk if bootstrap and reg-test on x86_64 with nvidia accelerator
succeeds?

Thanks,
- Tom

[-- Attachment #2: 0006-Handle-GOMP_OPENACC_NVPTX_PTXRW-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 4629 bytes --]

Handle GOMP_OPENACC_NVPTX_PTXRW in libgomp nvptx plugin

2017-06-27  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (post_process_ptx): New function.
	(link_ptx): Call post_process_ptx.

---
 libgomp/plugin/plugin-nvptx.c | 129 +++++++++++++++++++++++++++++++++++++++++-
 1 file changed, 127 insertions(+), 2 deletions(-)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index db42292..26e453f 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1020,6 +1020,128 @@ process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o,
     }
 }
 
+/* If environment variable GOMP_OPENACC_NVPTX_PTXRW=[Ww], write *RES_CODE to
+   file plugin-nvptx.<NUM>.ptx.  If it is [Rr], read *RES_CODE from file
+   instead.  */
+
+static void
+post_process_ptx (unsigned num, const char **res_code, size_t *res_size)
+{
+  static int gomp_openacc_nvptx_ptxrw = -1;
+
+  if (gomp_openacc_nvptx_ptxrw == -1)
+    {
+      const char *var_name = "GOMP_OPENACC_NVPTX_PTXRW";
+      const char *env_var = secure_getenv (var_name);
+      notify_var (var_name, env_var);
+
+      gomp_openacc_nvptx_ptxrw = 0;
+      if (env_var == NULL)
+	;
+      else if ((env_var[0] == 'w' || env_var[0] == 'W')
+	       && env_var[1] == '\0')
+	gomp_openacc_nvptx_ptxrw = 1;
+      else if ((env_var[0] == 'r' || env_var[0] == 'R')
+	       && env_var[1] == '\0')
+	gomp_openacc_nvptx_ptxrw = 2;
+      else
+	GOMP_PLUGIN_error ("Error parsing %s", var_name);
+    }
+
+  if (gomp_openacc_nvptx_ptxrw == 0)
+    return;
+
+  const char *prefix = "plugin-nvptx.";
+  const char *postfix = ".ptx";
+  const int len =	(strlen (prefix)
+			 + 10 /* %u.  */
+			 + strlen (postfix)
+			 + 1  /* '\0'.  */);
+  char file_name[len];
+  int res = snprintf (file_name, len, "%s%u%s", prefix,
+		      num, postfix);
+  assert (res < len); /* Assert there's no truncation.  */
+
+  GOMP_PLUGIN_debug (0, "%s %s \n",
+		     (gomp_openacc_nvptx_ptxrw == 1 ? "Writing" : "Reading"),
+		     file_name);
+
+  if (gomp_openacc_nvptx_ptxrw == 1)
+    {
+      FILE *ptx_file = fopen (file_name, "w");
+      if (ptx_file == NULL)
+	{
+	  GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+	  return;
+	}
+
+      int res = fprintf (ptx_file, "%s", code);
+      unsigned int write_succeeded = res == size - 1;
+      if (!write_succeeded)
+	GOMP_PLUGIN_debug (0,
+			   "Writing %s failed: written %d but expected %zu\n",
+			   file_name, res, size - 1);
+
+      res = fclose (ptx_file);
+      if (res != 0)
+	GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+
+      return;
+    }
+
+  if (gomp_openacc_nvptx_ptxrw == 2)
+    {
+      FILE *ptx_file = fopen (file_name, "r");
+      if (ptx_file == NULL)
+	{
+	  GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+	  return;
+	}
+
+      if (fseek (ptx_file, 0L, SEEK_END) != 0)
+	{
+	  GOMP_PLUGIN_debug (0, "Seeking end of %s failed\n", file_name);
+	  return;
+	}
+
+      long bufsize = ftell (ptx_file);
+      if (bufsize == -1)
+	{
+	  GOMP_PLUGIN_debug (0, "ftell of %s failed\n", file_name);
+	  return;
+	}
+
+      if (fseek (ptx_file, 0L, SEEK_SET) != 0)
+	{
+	  GOMP_PLUGIN_debug (0, "Seeking start of %s failed\n", file_name);
+	  return;
+	}
+
+	char *new_code = GOMP_PLUGIN_malloc (sizeof (char) * (bufsize + 1));
+
+	size_t new_size = fread (new_code, sizeof (char), bufsize, ptx_file);
+	if (ferror (ptx_file) != 0)
+	  {
+	    GOMP_PLUGIN_debug (0, "Reading %s failed\n", file_name);
+	    return;
+	  }
+
+	assert (new_size < bufsize + 1);
+	new_code[new_size++] = '\0';
+
+	int res = fclose (ptx_file);
+	if (res != 0)
+	  {
+	    GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+	    return;
+	  }
+
+	*res_code = new_code;
+	*res_size = new_size;
+	return;
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
@@ -1073,11 +1195,14 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 
   for (; num_objs--; ptx_objs++)
     {
+      const char *ptx_code = ptx_objs->code;
+      size_t ptx_size = ptx_objs->size;
+      post_process_ptx (num_objs, &ptx_code, &ptx_size);
+      GOMP_PLUGIN_debug (0, "Loading:\n---\n%s\n---\n", ptx_code);
       /* cuLinkAddData's 'data' argument erroneously omits the const
 	 qualifier.  */
-      GOMP_PLUGIN_debug (0, "Loading:\n---\n%s\n---\n", ptx_objs->code);
       r = CUDA_CALL_NOCHECK (cuLinkAddData, linkstate, CU_JIT_INPUT_PTX,
-			     (char *) ptx_objs->code, ptx_objs->size,
+			     (char *) ptx_code, ptx_size,
 			     0, 0, 0, 0);
       if (r != CUDA_SUCCESS)
 	{

Re: [PATCH] Use secure_getenv for GOMP_DEBUG

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 3642 bytes --]

On 06/27/2017 09:38 AM, Jakub Jelinek wrote:
> On Tue, Jun 27, 2017 at 09:17:57AM +0200, Tom de Vries wrote:
>> This patch uses secure_getenv for GOMP_DEBUG.
>>
>> It factors out the secure_getenv code from plugin-hsa.c into
>> libgomp/secure_getenv.h, and reuses it in env.c.
>>
>> I've added _GNU_SOURCE before the libgomp.h include in env.c to make sure
>> that secure_getenv (imported from stdlib.h) is available.
>>
>> I've also added a test-case that sets GOMP_DEBUG to 1 and verifies that some
>> output is generated.
>>
>> Build for c-only on x86_64 without accelerator, tested libgomp -m64/-m32.
>>
>> OK if x86_64 bootstrap and reg-test succeeds?
>>
>> Thanks,
>> - Tom
> 
>> Use secure_getenv for GOMP_DEBUG
>>
> 
>> --- /dev/null
>> +++ b/libgomp/secure_getenv.h
>> @@ -0,0 +1,53 @@
>> +/* Copyright (C) 2017 Free Software Foundation, Inc.
>> +
>> +This file is part of GCC.
>> +
>> +GCC is free software; you can redistribute it and/or modify
>> +it under the terms of the GNU General Public License as published by
>> +the Free Software Foundation; either version 3, or (at your option)
>> +any later version.
>> +
>> +GCC is distributed in the hope that it will be useful,
>> +but WITHOUT ANY WARRANTY; without even the implied warranty of
>> +MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
>> +GNU General Public License for more details.
>> +
>> +Under Section 7 of GPL version 3, you are granted additional
>> +permissions described in the GCC Runtime Library Exception, version
>> +3.1, as published by the Free Software Foundation.
>> +
>> +You should have received a copy of the GNU General Public License and
>> +a copy of the GCC Runtime Library Exception along with this program;
>> +see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
>> +<http://www.gnu.org/licenses/>.  */
>> +
>> +#ifndef _SECURE_GETENV_H
>> +#define _SECURE_GETENV_H 1
>> +
>> +/* Secure getenv() which returns NULL if running as SUID/SGID.  */
>> +#ifndef HAVE_SECURE_GETENV
>> +#ifdef HAVE___SECURE_GETENV
>> +#define secure_getenv __secure_getenv
>> +#elif defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
>> +  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
>> +
>> +#include <unistd.h>
>> +
>> +/* Implementation of secure_getenv() for targets where it is not provided but
>> +   we have at least means to test real and effective IDs.  */
>> +
>> +static char *
>> +secure_getenv (const char *name)
> 
> Shouldn't this be static inline char * ?  I mean, even at -O0 we don't want
> it to be emitted into every TU.
> 

Done.

> Another thing is that we probably want to follow what libgfortran does for
> the case where secure_getenv isn't available, but __secure_getenv is -
> in particular emit this function (if geteuid and getegid are present), but
> emit a weakref call to __secure_getenv first if non-NULL.
> 

I've copied the approach from libgfortran.

Build and tested:
- once as is, and
- once with '#ifndef HAVE_SECURE_GETENV' replaced with '#if 1' to
   trigger the fallback.

>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c
>> @@ -0,0 +1,13 @@
>> +/* { dg-do run } */
>> +/* { dg-set-target-env-var GOMP_DEBUG "1" } */
>> +
>> +/* Check that GOMP_DEBUG=1 triggers some output.  */
>> +
>> +int
>> +main (void)
>> +{
>> +#pragma acc parallel
>> +  ;
>> +}
>> +
>> +/* { dg-output "GOACC_parallel_keyed" } */
> 
> Does dg-set-target-env-var work for remote testing?

No, it'll make the testcase unsupported for remote testing.
[ Discussed before at f.i. 
https://gcc.gnu.org/ml/gcc-patches/2015-07/msg01808.html ]

Thanks,
- Tom



[-- Attachment #2: 0001-Use-secure_getenv-for-GOMP_DEBUG.patch --]
[-- Type: text/x-patch, Size: 7667 bytes --]

Use secure_getenv for GOMP_DEBUG

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* env.c (parse_unsigned_long_1): Factor out of ...
	(parse_unsigned_long): ... here.
	(parse_int_1): Factor out of ...
	(parse_int): ... here.
	(parse_int_secure): New function.
	(initialize_env): Use parse_int_secure for GOMP_DEBUG.
	* secure_getenv.h: Factor out of ...
	* plugin/plugin-hsa.c: ... here.
	* testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c: New test.

---
 libgomp/env.c                                      | 44 +++++++++++++---
 libgomp/plugin/plugin-hsa.c                        | 27 +---------
 libgomp/secure_getenv.h                            | 60 ++++++++++++++++++++++
 .../libgomp.oacc-c-c++-common/gomp-debug-env.c     | 13 +++++
 4 files changed, 111 insertions(+), 33 deletions(-)

diff --git a/libgomp/env.c b/libgomp/env.c
index ced752d..802c73b 100644
--- a/libgomp/env.c
+++ b/libgomp/env.c
@@ -26,6 +26,7 @@
 /* This file defines the OpenMP internal control variables and arranges
    for them to be initialized from environment variables at startup.  */
 
+#define _GNU_SOURCE
 #include "libgomp.h"
 #include "gomp-constants.h"
 #include <limits.h>
@@ -58,6 +59,8 @@
 #endif
 #endif /* LIBGOMP_OFFLOADED_ONLY */
 
+#include "secure_getenv.h"
+
 struct gomp_task_icv gomp_global_icv = {
   .nthreads_var = 1,
   .thread_limit_var = UINT_MAX,
@@ -171,15 +174,17 @@ parse_schedule (void)
 }
 
 /* Parse an unsigned long environment variable.  Return true if one was
-   present and it was successfully parsed.  */
+   present and it was successfully parsed.  If SECURE, use secure_getenv to the
+   environment variable.  */
 
 static bool
-parse_unsigned_long (const char *name, unsigned long *pvalue, bool allow_zero)
+parse_unsigned_long_1 (const char *name, unsigned long *pvalue, bool allow_zero,
+		       bool secure)
 {
   char *env, *end;
   unsigned long value;
 
-  env = getenv (name);
+  env = (secure ? secure_getenv (name) : getenv (name));
   if (env == NULL)
     return false;
 
@@ -206,14 +211,23 @@ parse_unsigned_long (const char *name, unsigned long *pvalue, bool allow_zero)
   return false;
 }
 
+/* As parse_unsigned_long_1, but always use getenv.  */
+
+static bool
+parse_unsigned_long (const char *name, unsigned long *pvalue, bool allow_zero)
+{
+  return parse_unsigned_long_1 (name, pvalue, allow_zero, false);
+}
+
 /* Parse a positive int environment variable.  Return true if one was
-   present and it was successfully parsed.  */
+   present and it was successfully parsed.  If SECURE, use secure_getenv to the
+   environment variable.  */
 
 static bool
-parse_int (const char *name, int *pvalue, bool allow_zero)
+parse_int_1 (const char *name, int *pvalue, bool allow_zero, bool secure)
 {
   unsigned long value;
-  if (!parse_unsigned_long (name, &value, allow_zero))
+  if (!parse_unsigned_long_1 (name, &value, allow_zero, secure))
     return false;
   if (value > INT_MAX)
     {
@@ -224,6 +238,22 @@ parse_int (const char *name, int *pvalue, bool allow_zero)
   return true;
 }
 
+/* As parse_int_1, but use getenv.  */
+
+static bool
+parse_int (const char *name, int *pvalue, bool allow_zero)
+{
+  return parse_int_1 (name, pvalue, allow_zero, false);
+}
+
+/* As parse_int_1, but use getenv_secure.  */
+
+static bool
+parse_int_secure (const char *name, int *pvalue, bool allow_zero)
+{
+  return parse_int_1 (name, pvalue, allow_zero, true);
+}
+
 /* Parse an unsigned long list environment variable.  Return true if one was
    present and it was successfully parsed.  */
 
@@ -1207,7 +1237,7 @@ initialize_env (void)
       gomp_global_icv.thread_limit_var
 	= thread_limit_var > INT_MAX ? UINT_MAX : thread_limit_var;
     }
-  parse_int ("GOMP_DEBUG", &gomp_debug_var, true);
+  parse_int_secure ("GOMP_DEBUG", &gomp_debug_var, true);
 #ifndef HAVE_SYNC_BUILTINS
   gomp_mutex_init (&gomp_managed_threads_lock);
 #endif
diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index 90ca247..adb07ac 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -39,32 +39,7 @@
 #include <dlfcn.h>
 #include "libgomp-plugin.h"
 #include "gomp-constants.h"
-
-/* Secure getenv() which returns NULL if running as SUID/SGID.  */
-#ifndef HAVE_SECURE_GETENV
-#ifdef HAVE___SECURE_GETENV
-#define secure_getenv __secure_getenv
-#elif defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
-  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
-
-#include <unistd.h>
-
-/* Implementation of secure_getenv() for targets where it is not provided but
-   we have at least means to test real and effective IDs. */
-
-static char *
-secure_getenv (const char *name)
-{
-  if ((getuid () == geteuid ()) && (getgid () == getegid ()))
-    return getenv (name);
-  else
-    return NULL;
-}
-
-#else
-#define secure_getenv getenv
-#endif
-#endif
+#include "secure-getenv.h"
 
 /* As an HSA runtime is dlopened, following structure defines function
    pointers utilized by the HSA plug-in.  */
diff --git a/libgomp/secure_getenv.h b/libgomp/secure_getenv.h
new file mode 100644
index 0000000..1d03eb7
--- /dev/null
+++ b/libgomp/secure_getenv.h
@@ -0,0 +1,60 @@
+/* Copyright (C) 2017 Free Software Foundation, Inc.
+
+This file is part of GCC.
+
+GCC is free software; you can redistribute it and/or modify
+it under the terms of the GNU General Public License as published by
+the Free Software Foundation; either version 3, or (at your option)
+any later version.
+
+GCC is distributed in the hope that it will be useful,
+but WITHOUT ANY WARRANTY; without even the implied warranty of
+MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+GNU General Public License for more details.
+
+Under Section 7 of GPL version 3, you are granted additional
+permissions described in the GCC Runtime Library Exception, version
+3.1, as published by the Free Software Foundation.
+
+You should have received a copy of the GNU General Public License and
+a copy of the GCC Runtime Library Exception along with this program;
+see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+<http://www.gnu.org/licenses/>.  */
+
+#ifndef _SECURE_GETENV_H
+#define _SECURE_GETENV_H 1
+
+/* Secure getenv() which returns NULL if running as SUID/SGID.  */
+#ifndef HAVE_SECURE_GETENV
+#if defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
+  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
+
+#include <unistd.h>
+
+#if SUPPORTS_WEAKREF && defined(HAVE___SECURE_GETENV)
+static char* weak_secure_getenv (const char*)
+  __attribute__((__weakref__("__secure_getenv")));
+#endif
+
+/* Implementation of secure_getenv() for targets where it is not provided but
+   we have at least means to test real and effective IDs.  */
+
+static inline char *
+secure_getenv (const char *name)
+{
+#if SUPPORTS_WEAKREF && defined(HAVE___SECURE_GETENV)
+  if (weak_secure_getenv)
+    return weak_secure_getenv (name);
+#endif
+
+  if ((getuid () == geteuid ()) && (getgid () == getegid ()))
+    return getenv (name);
+  else
+    return NULL;
+}
+#else
+#define secure_getenv getenv
+#endif
+#endif
+
+#endif /* _SECURE_GETENV_H.  */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c
new file mode 100644
index 0000000..3fc3503
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/gomp-debug-env.c
@@ -0,0 +1,13 @@
+/* { dg-do run } */
+/* { dg-set-target-env-var GOMP_DEBUG "1" } */
+
+/* Check that GOMP_DEBUG=1 triggers some output.  */
+
+int
+main (void)
+{
+#pragma acc parallel
+  ;
+}
+
+/* { dg-output "GOACC_parallel_keyed" } */

Re: [PATCH] Use secure_getenv for GOMP_DEBUG

[Jakub Jelinek]

On Tue, Jun 27, 2017 at 01:10:19PM +0200, Tom de Vries wrote:
> +/* Secure getenv() which returns NULL if running as SUID/SGID.  */
> +#ifndef HAVE_SECURE_GETENV
> +#if defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
> +  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)

Can you please us consistent formatting (i.e. space between defined and (
everywhere in the patch)?  The above line will be too long, so you'll need
to split it into 3 lines.

Ok with that change, thanks.

	Jakub

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 835 bytes --]

On 06/26/2017 05:29 PM, Jakub Jelinek wrote:
> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>
>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
>>>
>>> This patch adds handling of:
>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
>>>
>>> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
>>
>> Are you sure this use of getenv and writing to that file is safe for
>> setuid/setgid programs?  I'd expect you to need to use secure_getenv as in
>> plugin-hsa.c; certainly for anything that could results in writes to a
>> file like that.
> 
> Yeah, definitely it should be using secure_getenv/__secure_getenv.
> And IMNSHO GOMP_DEBUG too.
> 

Updated patch using secure_getenv.h.

Thanks,
- Tom

[-- Attachment #2: 0003-Handle-GOMP_OPENACC_NVPTX_-DISASM-SAVE_TEMPS-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 4186 bytes --]

Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (do_prog, debug_linkout): New function.
	(link_ptx): Use debug_linkout.

---
 libgomp/plugin/plugin-nvptx.c | 105 ++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 105 insertions(+)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 71630b5..7aa2b3b 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -47,6 +47,9 @@
 #include <unistd.h>
 #include <assert.h>
 #include <errno.h>
+#include <stdlib.h>
+#include <sys/types.h>
+#include <sys/wait.h>
 
 #if PLUGIN_NVPTX_DYNAMIC
 # include <dlfcn.h>
@@ -138,6 +141,8 @@ init_cuda_lib (void)
 # define init_cuda_lib() true
 #endif
 
+#include "secure_getenv.h"
+
 /* Convenience macros for the frequently used CUDA library call and
    error handling sequence as well as CUDA library calls that
    do the error checking themselves or don't do it at all.  */
@@ -876,6 +881,104 @@ notify_var (const char *var_name, const char *env_var)
     GOMP_PLUGIN_debug (0, "%s: '%s'\n", var_name, env_var);
 }
 
+static void
+do_prog (const char *prog, const char *arg)
+{
+  pid_t pid = fork ();
+
+  if (pid == -1)
+    {
+      GOMP_PLUGIN_error ("Fork failed");
+      return;
+    }
+  else if (pid > 0)
+    {
+      int status;
+      waitpid (pid, &status, 0);
+      if (!WIFEXITED (status))
+	GOMP_PLUGIN_error ("Running %s %s failed", prog, arg);
+    }
+  else
+    {
+      execlp (prog, prog /* argv[0] */, arg, NULL);
+      abort ();
+    }
+}
+
+static void
+debug_linkout (void *linkout, size_t linkoutsize)
+{
+  static int gomp_openacc_nvptx_disasm = -1;
+  if (gomp_openacc_nvptx_disasm == -1)
+    {
+      const char *var_name = "GOMP_OPENACC_NVPTX_DISASM";
+      const char *env_var = secure_getenv (var_name);
+      notify_var (var_name, env_var);
+      gomp_openacc_nvptx_disasm
+	= ((env_var != NULL && env_var[0] == '1' && env_var[1] == '\0')
+	   ? 1 : 0);
+    }
+
+  static int gomp_openacc_nvptx_save_temps = -1;
+  if (gomp_openacc_nvptx_save_temps == -1)
+    {
+      const char *var_name = "GOMP_OPENACC_NVPTX_SAVE_TEMPS";
+      const char *env_var = secure_getenv (var_name);
+      notify_var (var_name, env_var);
+      gomp_openacc_nvptx_save_temps
+	= ((env_var != NULL && env_var[0] == '1' && env_var[1] == '\0')
+	   ? 1 : 0);
+    }
+
+  if (gomp_openacc_nvptx_disasm == 0
+      && gomp_openacc_nvptx_save_temps == 0)
+    return;
+
+  const char *prefix = "plugin-nvptx.";
+  const char *postfix = ".cubin";
+  const int len =	(strlen (prefix)
+			 + 20 /* %lld.  */
+			 + strlen (postfix)
+			 + 1  /* '\0'.  */);
+  char file_name[len];
+  int res = snprintf (file_name, len, "%s%lld%s", prefix,
+		      (long long)getpid (), postfix);
+  assert (res < len); /* Assert there's no truncation.  */
+
+  GOMP_PLUGIN_debug (0, "Generating %s with size %zu\n",
+		     file_name, linkoutsize);
+  FILE *cubin_file = fopen (file_name, "wb");
+  if (cubin_file == NULL)
+    {
+      GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+      return;
+    }
+
+  fwrite (linkout, linkoutsize, 1, cubin_file);
+  unsigned int write_succeeded = ferror (cubin_file) == 0;
+  if (!write_succeeded)
+    GOMP_PLUGIN_debug (0, "Writing %s failed\n", file_name);
+
+  res = fclose (cubin_file);
+  if (res != 0)
+    GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+
+  if (!write_succeeded)
+    return;
+
+  if (gomp_openacc_nvptx_disasm == 1)
+    {
+      GOMP_PLUGIN_debug (0, "Disassembling %s\n", file_name);
+      do_prog ("nvdisasm", file_name);
+    }
+
+  if (gomp_openacc_nvptx_save_temps == 0)
+    {
+      GOMP_PLUGIN_debug (0, "Removing %s\n", file_name);
+      remove (file_name);
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
@@ -939,6 +1042,8 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
       return false;
     }
 
+  debug_linkout (linkout, linkoutsize);
+
   CUDA_CALL (cuModuleLoadData, module, linkout);
   CUDA_CALL (cuLinkDestroy, linkstate);
   return true;

Re: [PATCH, 1/4] Show value of GOMP_OPENACC_DIM in libgomp nvptx plugin

[Tom de Vries]

On 06/26/2017 01:31 PM, Tom de Vries wrote:
> On 06/26/2017 01:24 PM, Tom de Vries wrote:
>> Hi,
>>
>> I've written a patch series to facilitate debugging libgomp openacc 
>> testcase failures on the nvptx accelerator.
>>
>>
>> When running an openacc test-case on an nvptx accelerator, the 
>> following happens:
>> - the plugin obtains the ptx assembly for the acceleration kernels
>> - it calls the cuda jit to compile and link the ptx into a module
>> - it loads the module
>> - it starts an acceleration kernel
>>
>> The patch series adds these environment variables:
>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS: a means to save the resulting module
>>    such that it can be investigated using nvdisasm and cuobjdump.
>> - GOMP_OPENACC_NVPTX_DISASM: a means to see the resulting module in
>>    the debug output,  by writing it into a file and calling nvdisasm on
>>    it
>> - GOMP_OPENACC_NVPTX_JIT: a means to set parameters of the
>>    compilation/linking process, currently supporting:
>>    * -O[0-4], mapping onto CU_JIT_OPTIMIZATION_LEVEL
>>    * -ori, mapping onto CU_JIT_NEW_SM3X_OPT
>>
>>
>> The patch series consists of these patches:
>>
>> 1. Show value of GOMP_OPENACC_DIM in libgomp nvptx plugin
> 
> This patch adds a debug message (for GOMP_DEBUG=1) about the value of 
> the GOMP_OPENACC_DIM variable read from the environment.
> 

Committed as trivial.

Thanks,
- Tom

> Thanks,
> - Tom
> 
> 0001-Show-value-of-GOMP_OPENACC_DIM-in-libgomp-nvptx-plugin.patch
> 
> 
> Show value of GOMP_OPENACC_DIM in libgomp nvptx plugin
> 
> 2017-06-26  Tom de Vries  <tom@codesourcery.com>
> 
> 	* plugin/plugin-nvptx.c (notify_var): New function.
> 	(nvptx_exec): Use notify_var for GOMP_OPENACC_DIM.
> 
> ---
>   libgomp/plugin/plugin-nvptx.c | 12 +++++++++++-
>   1 file changed, 11 insertions(+), 1 deletion(-)
> 
> diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
> index 0e1b3e2..71630b5 100644
> --- a/libgomp/plugin/plugin-nvptx.c
> +++ b/libgomp/plugin/plugin-nvptx.c
> @@ -867,6 +867,14 @@ nvptx_get_num_devices (void)
>     return n;
>   }
>   
> +static void
> +notify_var (const char *var_name, const char *env_var)
> +{
> +  if (env_var == NULL)
> +    GOMP_PLUGIN_debug (0, "%s: <Not defined>\n", var_name);
> +  else
> +    GOMP_PLUGIN_debug (0, "%s: '%s'\n", var_name, env_var);
> +}
>   
>   static bool
>   link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
> @@ -1089,10 +1097,12 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
>         pthread_mutex_lock (&ptx_dev_lock);
>         if (!default_dims[0])
>   	{
> +	  const char *var_name = "GOMP_OPENACC_DIM";
>   	  /* We only read the environment variable once.  You can't
>   	     change it in the middle of execution.  The syntax  is
>   	     the same as for the -fopenacc-dim compilation option.  */
> -	  const char *env_var = getenv ("GOMP_OPENACC_DIM");
> +	  const char *env_var = getenv (var_name);
> +	  notify_var (var_name, env_var);
>   	  if (env_var)
>   	    {
>   	      const char *pos = env_var;
> 

Re: [PATCH, 4/4] Handle GOMP_OPENACC_NVPTX_JIT=-ori in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1812 bytes --]

On 06/26/2017 01:44 PM, Tom de Vries wrote:
> On 06/26/2017 01:24 PM, Tom de Vries wrote:
>> Hi,
>>
>> I've written a patch series to facilitate debugging libgomp openacc 
>> testcase failures on the nvptx accelerator.
>>
>>
>> When running an openacc test-case on an nvptx accelerator, the 
>> following happens:
>> - the plugin obtains the ptx assembly for the acceleration kernels
>> - it calls the cuda jit to compile and link the ptx into a module
>> - it loads the module
>> - it starts an acceleration kernel
>>
>> The patch series adds these environment variables:
>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS: a means to save the resulting module
>>    such that it can be investigated using nvdisasm and cuobjdump.
>> - GOMP_OPENACC_NVPTX_DISASM: a means to see the resulting module in
>>    the debug output,  by writing it into a file and calling nvdisasm on
>>    it
>> - GOMP_OPENACC_NVPTX_JIT: a means to set parameters of the
>>    compilation/linking process, currently supporting:
>>    * -O[0-4], mapping onto CU_JIT_OPTIMIZATION_LEVEL
>>    * -ori, mapping onto CU_JIT_NEW_SM3X_OPT
>>
>>
>> The patch series consists of these patches:
>>
>> 4. Handle GOMP_OPENACC_NVPTX_JIT=-ori in libgomp nvptx plugin
> 
> This patch adds handling of GOMP_OPENACC_NVPTX_JIT=-ori.
> 
> Thanks,
> - Tom
> 
> 0004-Handle-GOMP_OPENACC_NVPTX_JIT-ori-in-libgomp-nvptx-plugin.patch
> 

> -  CU_JIT_LOG_VERBOSE = 12
> +  CU_JIT_LOG_VERBOSE = 12,
> +  CU_JIT_NEW_SM3X_OPT = 15
>   } CUjit_option;

Adding the constant to plugin/cuda/cuda.h makes sure the constant is 
available when not linking the plugin against cuda.

But when linking against cuda 7.5 and earlier, this still fails because 
the constant is not available yet in cuda.h. Fixed by hardcoding the 
constant if not available in the cuda version.

Thanks,
- Tom

[-- Attachment #2: 0003-Handle-GOMP_OPENACC_NVPTX_JIT-ori-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 3422 bytes --]

Handle GOMP_OPENACC_NVPTX_JIT=-ori in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/cuda/cuda.h (enum CUjit_option): Add CU_JIT_NEW_SM3X_OPT.
	* plugin/plugin-nvptx.c (process_GOMP_OPENACC_NVPTX_JIT): Add
	gomp_openacc_nvptx_ori parameter.  Handle -ori.
	(link_ptx): Add CU_JIT_NEW_SM3X_OPT to opts.

---
 libgomp/plugin/cuda/cuda.h    |  3 ++-
 libgomp/plugin/plugin-nvptx.c | 34 +++++++++++++++++++++++++++++-----
 2 files changed, 31 insertions(+), 6 deletions(-)

diff --git a/libgomp/plugin/cuda/cuda.h b/libgomp/plugin/cuda/cuda.h
index 75dfe3d..4644870 100644
--- a/libgomp/plugin/cuda/cuda.h
+++ b/libgomp/plugin/cuda/cuda.h
@@ -89,7 +89,8 @@ typedef enum {
   CU_JIT_ERROR_LOG_BUFFER = 5,
   CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES = 6,
   CU_JIT_OPTIMIZATION_LEVEL = 7,
-  CU_JIT_LOG_VERBOSE = 12
+  CU_JIT_LOG_VERBOSE = 12,
+  CU_JIT_NEW_SM3X_OPT = 15
 } CUjit_option;
 
 typedef enum {
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 594ca39..41ecfec 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -143,6 +143,10 @@ init_cuda_lib (void)
 
 #include "secure_getenv.h"
 
+#if CUDA_VERSION < 8000
+#define CU_JIT_NEW_SM3X_OPT 15
+#endif
+
 /* Convenience macros for the frequently used CUDA library call and
    error handling sequence as well as CUDA library calls that
    do the error checking themselves or don't do it at all.  */
@@ -980,13 +984,15 @@ debug_linkout (void *linkout, size_t linkoutsize)
 }
 
 static void
-process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o)
+process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o,
+				intptr_t *gomp_openacc_nvptx_ori)
 {
   const char *var_name = "GOMP_OPENACC_NVPTX_JIT";
   const char *env_var = getenv (var_name);
   notify_var (var_name, env_var);
 
   *gomp_openacc_nvptx_o = 4;
+  *gomp_openacc_nvptx_ori = 0;
   if (env_var == NULL)
     return;
 
@@ -1005,6 +1011,14 @@ process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o)
 	  continue;
 	}
 
+      if (c[0] == '-' && c[1] == 'o' && c[2] == 'r' && c[3] == 'i'
+	  && (c[4] == '\0' || c[4] == ' '))
+	{
+	  *gomp_openacc_nvptx_ori = 1;
+	  c += 4;
+	  continue;
+	}
+
       GOMP_PLUGIN_error ("Error parsing %s", var_name);
       break;
     }
@@ -1014,8 +1028,8 @@ static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
 {
-  CUjit_option opts[7];
-  void *optvals[7];
+  CUjit_option opts[8];
+  void *optvals[8];
   float elapsed = 0.0;
   char elog[1024];
   char ilog[16384];
@@ -1043,13 +1057,23 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
   optvals[5] = (void *) 1;
 
   static intptr_t gomp_openacc_nvptx_o = -1;
+  static intptr_t gomp_openacc_nvptx_ori = -1;
   if (gomp_openacc_nvptx_o == -1)
-    process_GOMP_OPENACC_NVPTX_JIT (&gomp_openacc_nvptx_o);
+    process_GOMP_OPENACC_NVPTX_JIT (&gomp_openacc_nvptx_o,
+				    &gomp_openacc_nvptx_ori);
 
   opts[6] = CU_JIT_OPTIMIZATION_LEVEL;
   optvals[6] = (void *) gomp_openacc_nvptx_o;
 
-  CUDA_CALL (cuLinkCreate, 7, opts, optvals, &linkstate);
+  int nopts = 7;
+  if (gomp_openacc_nvptx_ori)
+    {
+      opts[nopts] = CU_JIT_NEW_SM3X_OPT;
+      optvals[nopts] = (void *) gomp_openacc_nvptx_ori;
+      nopts++;
+    }
+
+  CUDA_CALL (cuLinkCreate, nopts, opts, optvals, &linkstate);
 
   for (; num_objs--; ptx_objs++)
     {

Re: [PATCH, 5/4] Handle GOMP_OPENACC_NVPTX_PTXRW in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1867 bytes --]

On 06/27/2017 11:16 AM, Tom de Vries wrote:
> [ was: Re: [PATCH, 0/4] Handle 
> GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS,JIT} in libgomp nvptx plugin ]
> 
> On 06/26/2017 01:24 PM, Tom de Vries wrote:
>> Hi,
>>
>> I've written a patch series to facilitate debugging libgomp openacc 
>> testcase failures on the nvptx accelerator.
>>
>>
>> When running an openacc test-case on an nvptx accelerator, the 
>> following happens:
>> - the plugin obtains the ptx assembly for the acceleration kernels
>> - it calls the cuda jit to compile and link the ptx into a module
>> - it loads the module
>> - it starts an acceleration kernel
> 
> A typical scenario when developing the compiler is:
> - run gcc test.c -save-temps
> - run a.out
> - edit test.s to fix bug or make code faster or smaller
> - run gcc test.s
> - run a.out
> - edit compiler sources to make the compiler do the same as the .s edit
> 
> With openacc test-cases, this scenario is currently not available for 
> ptx assembly. Using -save-temps -foffload=-save-temps we can get a .s 
> containing ptx. But to insert the edited .s back into the compilation 
> flow is difficult.
> 
> This patch facilitates such a scenario in the nvptx plugin.
> - we define GOMP_OPENACC_NVPTX_PTXRW == 'w', and the plugin writes the
>    ptx assembly into a series of files
> - we edit one of those files
> - we define GOMP_OPENACC_NVPTX_PTXRW == 'r', and the plugin reads the
>    ptx assembly back from those files, and uses that instead of the ptx
>    in the executable.
> 
> I've tested this patch series on top of gomp-4_0-branch, by running an
> openacc testcase from the command line and going through the 
> write-edit-readscenario with an observable ptx edit.
> 
> OK for trunk if bootstrap and reg-test on x86_64 with nvidia accelerator
> succeeds?
> 

This updated patch fixes a trivial build error.

Thanks,
- Tom


[-- Attachment #2: 0004-Handle-GOMP_OPENACC_NVPTX_PTXRW-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 4692 bytes --]

Handle GOMP_OPENACC_NVPTX_PTXRW in libgomp nvptx plugin

2017-06-27  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (post_process_ptx): New function.
	(link_ptx): Call post_process_ptx.

---
 libgomp/plugin/plugin-nvptx.c | 132 +++++++++++++++++++++++++++++++++++++++++-
 1 file changed, 130 insertions(+), 2 deletions(-)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 41ecfec..365c787 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1024,6 +1024,131 @@ process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o,
     }
 }
 
+/* If environment variable GOMP_OPENACC_NVPTX_PTXRW=[Ww], write *RES_CODE to
+   file plugin-nvptx.<NUM>.ptx.  If it is [Rr], read *RES_CODE from file
+   instead.  */
+
+static void
+post_process_ptx (unsigned num, const char **res_code, size_t *res_size)
+{
+  static int gomp_openacc_nvptx_ptxrw = -1;
+
+  if (gomp_openacc_nvptx_ptxrw == -1)
+    {
+      const char *var_name = "GOMP_OPENACC_NVPTX_PTXRW";
+      const char *env_var = secure_getenv (var_name);
+      notify_var (var_name, env_var);
+
+      gomp_openacc_nvptx_ptxrw = 0;
+      if (env_var == NULL)
+	;
+      else if ((env_var[0] == 'w' || env_var[0] == 'W')
+	       && env_var[1] == '\0')
+	gomp_openacc_nvptx_ptxrw = 1;
+      else if ((env_var[0] == 'r' || env_var[0] == 'R')
+	       && env_var[1] == '\0')
+	gomp_openacc_nvptx_ptxrw = 2;
+      else
+	GOMP_PLUGIN_error ("Error parsing %s", var_name);
+    }
+
+  if (gomp_openacc_nvptx_ptxrw == 0)
+    return;
+
+  const char *code = *res_code;
+  size_t size = *res_size;
+
+  const char *prefix = "plugin-nvptx.";
+  const char *postfix = ".ptx";
+  const int len =	(strlen (prefix)
+			 + 10 /* %u.  */
+			 + strlen (postfix)
+			 + 1  /* '\0'.  */);
+  char file_name[len];
+  int res = snprintf (file_name, len, "%s%u%s", prefix,
+		      num, postfix);
+  assert (res < len); /* Assert there's no truncation.  */
+
+  GOMP_PLUGIN_debug (0, "%s %s \n",
+		     (gomp_openacc_nvptx_ptxrw == 1 ? "Writing" : "Reading"),
+		     file_name);
+
+  if (gomp_openacc_nvptx_ptxrw == 1)
+    {
+      FILE *ptx_file = fopen (file_name, "w");
+      if (ptx_file == NULL)
+	{
+	  GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+	  return;
+	}
+
+      int res = fprintf (ptx_file, "%s", code);
+      unsigned int write_succeeded = res == size - 1;
+      if (!write_succeeded)
+	GOMP_PLUGIN_debug (0,
+			   "Writing %s failed: written %d but expected %zu\n",
+			   file_name, res, size - 1);
+
+      res = fclose (ptx_file);
+      if (res != 0)
+	GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+
+      return;
+    }
+
+  if (gomp_openacc_nvptx_ptxrw == 2)
+    {
+      FILE *ptx_file = fopen (file_name, "r");
+      if (ptx_file == NULL)
+	{
+	  GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+	  return;
+	}
+
+      if (fseek (ptx_file, 0L, SEEK_END) != 0)
+	{
+	  GOMP_PLUGIN_debug (0, "Seeking end of %s failed\n", file_name);
+	  return;
+	}
+
+      long bufsize = ftell (ptx_file);
+      if (bufsize == -1)
+	{
+	  GOMP_PLUGIN_debug (0, "ftell of %s failed\n", file_name);
+	  return;
+	}
+
+      if (fseek (ptx_file, 0L, SEEK_SET) != 0)
+	{
+	  GOMP_PLUGIN_debug (0, "Seeking start of %s failed\n", file_name);
+	  return;
+	}
+
+	char *new_code = GOMP_PLUGIN_malloc (sizeof (char) * (bufsize + 1));
+
+	size_t new_size = fread (new_code, sizeof (char), bufsize, ptx_file);
+	if (ferror (ptx_file) != 0)
+	  {
+	    GOMP_PLUGIN_debug (0, "Reading %s failed\n", file_name);
+	    return;
+	  }
+
+	assert (new_size < bufsize + 1);
+	new_code[new_size++] = '\0';
+
+	int res = fclose (ptx_file);
+	if (res != 0)
+	  {
+	    GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+	    return;
+	  }
+
+	*res_code = new_code;
+	*res_size = new_size;
+	return;
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
@@ -1077,11 +1202,14 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 
   for (; num_objs--; ptx_objs++)
     {
+      const char *ptx_code = ptx_objs->code;
+      size_t ptx_size = ptx_objs->size;
+      post_process_ptx (num_objs, &ptx_code, &ptx_size);
+      GOMP_PLUGIN_debug (0, "Loading:\n---\n%s\n---\n", ptx_code);
       /* cuLinkAddData's 'data' argument erroneously omits the const
 	 qualifier.  */
-      GOMP_PLUGIN_debug (0, "Loading:\n---\n%s\n---\n", ptx_objs->code);
       r = CUDA_CALL_NOCHECK (cuLinkAddData, linkstate, CU_JIT_INPUT_PTX,
-			     (char *) ptx_objs->code, ptx_objs->size,
+			     (char *) ptx_code, ptx_size,
 			     0, 0, 0, 0);
       if (r != CUDA_SUCCESS)
 	{

[PATCH, 6/4] Handle GOMP_OPENACC_NVPTX_JIT=-arch=<n> in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 717 bytes --]

[ was: Re: [PATCH, 0/4] Handle 
GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS,JIT} in libgomp nvptx plugin ]
On 06/26/2017 01:24 PM, Tom de Vries wrote:
> Hi,
> 
> I've written a patch series to facilitate debugging libgomp openacc 
> testcase failures on the nvptx accelerator.
> 
> 
> When running an openacc test-case on an nvptx accelerator, the following 
> happens:
> - the plugin obtains the ptx assembly for the acceleration kernels
> - it calls the cuda jit to compile and link the ptx into a module
> - it loads the module
> - it starts an acceleration kernel
> 

This patch adds handling of GOMP_OPENACC_NVPTX_JIT=-arch=<n> in libgomp 
nvptx plugin.

F.i. GOMP_OPENACC_NVPTX_JIT=-arch=60 for sm_60.

Thanks,
- Tom

[-- Attachment #2: 0005-Handle-GOMP_OPENACC_NVPTX_JIT-arch-n-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 2542 bytes --]

libgomp/ChangeLog:

2017-06-30  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (parse_number):
	(process_GOMP_OPENACC_NVPTX_JIT):
	(link_ptx):


Handle GOMP_OPENACC_NVPTX_JIT=-arch=<n> in libgomp nvptx plugin

---
 libgomp/plugin/plugin-nvptx.c | 40 ++++++++++++++++++++++++++++++++++++++--
 1 file changed, 38 insertions(+), 2 deletions(-)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 365c787..4cca0c7 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -983,9 +983,24 @@ debug_linkout (void *linkout, size_t linkoutsize)
     }
 }
 
+static bool
+parse_number (const char *c, unsigned long* resp, char **end)
+{
+  unsigned long res;
+
+  errno = 0;
+  res = strtoul (c, end, 10);
+  if (errno)
+    return false;
+
+  *resp = res;
+  return true;
+}
+
 static void
 process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o,
-				intptr_t *gomp_openacc_nvptx_ori)
+				intptr_t *gomp_openacc_nvptx_ori,
+				uintptr_t *gomp_openacc_nvptx_target)
 {
   const char *var_name = "GOMP_OPENACC_NVPTX_JIT";
   const char *env_var = getenv (var_name);
@@ -1019,6 +1034,19 @@ process_GOMP_OPENACC_NVPTX_JIT (intptr_t *gomp_openacc_nvptx_o,
 	  continue;
 	}
 
+      if (c[0] == '-' && c[1] == 'a' && c[2] == 'r' && c[3] == 'c'
+	  && c[4] == 'h' && c[5] == '=')
+	{
+	  const char *end;
+	  unsigned long val;
+	  if (parse_number (&c[6], &val, (char**)&end))
+	    {
+	      *gomp_openacc_nvptx_target = val;
+	      c = end;
+	      continue;
+	    }
+	}
+
       GOMP_PLUGIN_error ("Error parsing %s", var_name);
       break;
     }
@@ -1183,9 +1211,11 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 
   static intptr_t gomp_openacc_nvptx_o = -1;
   static intptr_t gomp_openacc_nvptx_ori = -1;
+  static uintptr_t gomp_openacc_nvptx_target = 0;
   if (gomp_openacc_nvptx_o == -1)
     process_GOMP_OPENACC_NVPTX_JIT (&gomp_openacc_nvptx_o,
-				    &gomp_openacc_nvptx_ori);
+				    &gomp_openacc_nvptx_ori,
+				    &gomp_openacc_nvptx_target);
 
   opts[6] = CU_JIT_OPTIMIZATION_LEVEL;
   optvals[6] = (void *) gomp_openacc_nvptx_o;
@@ -1197,6 +1227,12 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
       optvals[nopts] = (void *) gomp_openacc_nvptx_ori;
       nopts++;
     }
+  if (gomp_openacc_nvptx_target)
+    {
+      opts[nopts] = CU_JIT_TARGET;
+      optvals[nopts] = (void *) gomp_openacc_nvptx_target;
+      nopts++;
+    }
 
   CUDA_CALL (cuLinkCreate, nopts, opts, optvals, &linkstate);
 

Re: [PATCH] Use secure_getenv for GOMP_DEBUG

[Franz Sirl]

Am 27.06.17 um 13:10 schrieb Tom de Vries:
> --- a/libgomp/plugin/plugin-hsa.c
> +++ b/libgomp/plugin/plugin-hsa.c
> @@ -39,32 +39,7 @@
>   #include <dlfcn.h>
>   #include "libgomp-plugin.h"
>   #include "gomp-constants.h"
> -
> -/* Secure getenv() which returns NULL if running as SUID/SGID.  */
> -#ifndef HAVE_SECURE_GETENV
> -#ifdef HAVE___SECURE_GETENV
> -#define secure_getenv __secure_getenv
> -#elif defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && defined(HAVE_GETEUID) \
> -  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
> -
> -#include <unistd.h>
> -
> -/* Implementation of secure_getenv() for targets where it is not provided but
> -   we have at least means to test real and effective IDs. */
> -
> -static char *
> -secure_getenv (const char *name)
> -{
> -  if ((getuid () == geteuid ()) && (getgid () == getegid ()))
> -    return getenv (name);
> -  else
> -    return NULL;
> -}
> -
> -#else
> -#define secure_getenv getenv
> -#endif
> -#endif
> +#include "secure-getenv.h"

Hi,

that should be secure_getenv.h (underscore instead of dash).

Franz

Re: [PATCH] Use secure_getenv for GOMP_DEBUG

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 1249 bytes --]

On 07/03/2017 02:26 PM, Franz Sirl wrote:
> Am 27.06.17 um 13:10 schrieb Tom de Vries:
>> --- a/libgomp/plugin/plugin-hsa.c
>> +++ b/libgomp/plugin/plugin-hsa.c
>> @@ -39,32 +39,7 @@
>>   #include <dlfcn.h>
>>   #include "libgomp-plugin.h"
>>   #include "gomp-constants.h"
>> -
>> -/* Secure getenv() which returns NULL if running as SUID/SGID.  */
>> -#ifndef HAVE_SECURE_GETENV
>> -#ifdef HAVE___SECURE_GETENV
>> -#define secure_getenv __secure_getenv
>> -#elif defined (HAVE_UNISTD_H) && defined(HAVE_GETUID) && 
>> defined(HAVE_GETEUID) \
>> -  && defined(HAVE_GETGID) && defined(HAVE_GETEGID)
>> -
>> -#include <unistd.h>
>> -
>> -/* Implementation of secure_getenv() for targets where it is not 
>> provided but
>> -   we have at least means to test real and effective IDs. */
>> -
>> -static char *
>> -secure_getenv (const char *name)
>> -{
>> -  if ((getuid () == geteuid ()) && (getgid () == getegid ()))
>> -    return getenv (name);
>> -  else
>> -    return NULL;
>> -}
>> -
>> -#else
>> -#define secure_getenv getenv
>> -#endif
>> -#endif
>> +#include "secure-getenv.h"
> 
> Hi,
> 
> that should be secure_getenv.h (underscore instead of dash).

Hi Franz,

sorry for the breakage.

Fixed in attached patch.

Committed.

Thanks,
- Tom


[-- Attachment #2: 0001-Fix-secure_getenv.h-include-in-plugin-hsa.c.patch --]
[-- Type: text/x-patch, Size: 678 bytes --]

Fix secure_getenv.h include in plugin-hsa.c

2017-07-03  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-hsa.c: Fix secure_getenv.h include.

---
 libgomp/plugin/plugin-hsa.c | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index adb07ac..fc08f5d 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -39,7 +39,7 @@
 #include <dlfcn.h>
 #include "libgomp-plugin.h"
 #include "gomp-constants.h"
-#include "secure-getenv.h"
+#include "secure_getenv.h"
 
 /* As an HSA runtime is dlopened, following structure defines function
    pointers utilized by the HSA plug-in.  */

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Thomas Schwinge]

Hi!

On Mon, 26 Jun 2017 17:29:11 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
> > On Mon, 26 Jun 2017, Tom de Vries wrote:
> > 
> > > > 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
> > > 
> > > This patch adds handling of:
> > > - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
> > > - GOMP_OPENACC_NVPTX_DISASM=[01]

Why the "OPENACC" in these names?  Doesn't this debugging aid apply to
any variant of offloading?

> > > The filename used for dumping the module is plugin-nvptx.<pid>.cubin.

Also, I suggest to make these names similar to their controlling options,
that is: "gomp-nvptx*", for example.

> > Are you sure this use of getenv and writing to that file is safe for 
> > setuid/setgid programs?  I'd expect you to need to use secure_getenv as in 
> > plugin-hsa.c; certainly for anything that could results in writes to a 
> > file like that.
> 
> Yeah, definitely it should be using secure_getenv/__secure_getenv.

ACK.

> And IMNSHO GOMP_DEBUG too.

But why that?  Isn't GOMP_DEBUG just controlling terminal debugging
output (that you'd also like to see in setuid/setgid programs)?


Grüße
 Thomas

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Jakub Jelinek]

On Mon, Jul 03, 2017 at 04:08:10PM +0200, Thomas Schwinge wrote:
> > And IMNSHO GOMP_DEBUG too.
> 
> But why that?  Isn't GOMP_DEBUG just controlling terminal debugging
> output (that you'd also like to see in setuid/setgid programs)?

The output could go into stderr, which could very well be redirected into
some file and some other program could be expecting specific content in
there.  So allowing an attacker to add there other stuff is really
dangerous.  If you want to use GOMP_DEBUG on suid/sgid processes, just
run them under root.

	Jakub

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Tom de Vries]

On 07/03/2017 04:08 PM, Thomas Schwinge wrote:
> Hi!
> 
> On Mon, 26 Jun 2017 17:29:11 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
>> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>>
>>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
>>>>
>>>> This patch adds handling of:
>>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
> 
> Why the "OPENACC" in these names?

I took the format from 'GOMP_OPENACC_DIM'.

> Doesn't this debugging aid apply to
> any variant of offloading?

I guess you're right. These environment variables would also be 
applicable for f.i. offloading via openmp on nvptx. I'll strip the 
'OPENACC_' bit from the variables.

>>>> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
> 
> Also, I suggest to make these names similar to their controlling options,
> that is: "gomp-nvptx*", for example.
> 

Makes sense, will do.

Thanks,
- Tom

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Tom de Vries]

On 07/03/2017 04:24 PM, Tom de Vries wrote:
> On 07/03/2017 04:08 PM, Thomas Schwinge wrote:
>> Hi!
>>
>> On Mon, 26 Jun 2017 17:29:11 +0200, Jakub Jelinek <jakub@redhat.com> 
>> wrote:
>>> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>>>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>>>
>>>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx 
>>>>>> plugin
>>>>>
>>>>> This patch adds handling of:
>>>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
>>
>> Why the "OPENACC" in these names?
> 
> I took the format from 'GOMP_OPENACC_DIM'.
> 
>> Doesn't this debugging aid apply to
>> any variant of offloading?
> 
> I guess you're right. These environment variables would also be 
> applicable for f.i. offloading via openmp on nvptx. I'll strip the 
> 'OPENACC_' bit from the variables.
> 
>>>>> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
>>
>> Also, I suggest to make these names similar to their controlling options,
>> that is: "gomp-nvptx*", for example.
>>
> 
> Makes sense, will do.

Changes in the patch series:
- removed OPENACC_ from environment variable names
- made temp files use gomp-nvptx prefix.
- fixed build error due to missing _GNU_SOURCE in libgomp-nvptx.c.
- merged the three GOMP_NVPTX_JIT patches into one
- rewrote GOMP_NVPTX_JIT to add no extra flags to the JIT compiler
   invocation if GOMP_NVPTX_JIT if not defined, removing the need for
   hardcoding default values
- added CU_JIT_TARGET to plugin/cuda/cuda.h

Build on x86_64 with nvptx offloading enabled (using plugin/cuda/cuda.h).

The patch series now looks like:
1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
    plugin

I'll repost the patch series in reply to this email.

Thanks,
- Tom
	

[PATCH, 1/3] Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 2186 bytes --]

On 07/04/2017 12:05 PM, Tom de Vries wrote:
> On 07/03/2017 04:24 PM, Tom de Vries wrote:
>> On 07/03/2017 04:08 PM, Thomas Schwinge wrote:
>>> Hi!
>>>
>>> On Mon, 26 Jun 2017 17:29:11 +0200, Jakub Jelinek <jakub@redhat.com> 
>>> wrote:
>>>> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>>>>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>>>>
>>>>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx 
>>>>>>> plugin
>>>>>>
>>>>>> This patch adds handling of:
>>>>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>>>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
>>>
>>> Why the "OPENACC" in these names?
>>
>> I took the format from 'GOMP_OPENACC_DIM'.
>>
>>> Doesn't this debugging aid apply to
>>> any variant of offloading?
>>
>> I guess you're right. These environment variables would also be 
>> applicable for f.i. offloading via openmp on nvptx. I'll strip the 
>> 'OPENACC_' bit from the variables.
>>
>>>>>> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
>>>
>>> Also, I suggest to make these names similar to their controlling 
>>> options,
>>> that is: "gomp-nvptx*", for example.
>>>
>>
>> Makes sense, will do.
> 
> Changes in the patch series:
> - removed OPENACC_ from environment variable names
> - made temp files use gomp-nvptx prefix.
> - fixed build error due to missing _GNU_SOURCE in libgomp-nvptx.c.
> - merged the three GOMP_NVPTX_JIT patches into one
> - rewrote GOMP_NVPTX_JIT to add no extra flags to the JIT compiler
>    invocation if GOMP_NVPTX_JIT if not defined, removing the need for
>    hardcoding default values
> - added CU_JIT_TARGET to plugin/cuda/cuda.h
> 
> Build on x86_64 with nvptx offloading enabled (using plugin/cuda/cuda.h).
> 
> The patch series now looks like:
> 1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
> 2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
> 3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
>     plugin
> 
> I'll repost the patch series in reply to this email.

1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
    ( original submission at
      https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01918.html )

Thanks,
- Tom

[-- Attachment #2: 0001-Handle-GOMP_NVPTX_-DISASM-SAVE_TEMPS-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 4309 bytes --]

Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (do_prog, debug_linkout): New function.
	(link_ptx): Use debug_linkout.

---
 libgomp/plugin/plugin-nvptx.c | 106 ++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 106 insertions(+)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 71630b5..3e33c5b 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -31,6 +31,7 @@
    is not clear as to what that state might be.  Or how one might
    propagate it from one thread to another.  */
 
+#define _GNU_SOURCE
 #include "openacc.h"
 #include "config.h"
 #include "libgomp-plugin.h"
@@ -47,6 +48,9 @@
 #include <unistd.h>
 #include <assert.h>
 #include <errno.h>
+#include <stdlib.h>
+#include <sys/types.h>
+#include <sys/wait.h>
 
 #if PLUGIN_NVPTX_DYNAMIC
 # include <dlfcn.h>
@@ -138,6 +142,8 @@ init_cuda_lib (void)
 # define init_cuda_lib() true
 #endif
 
+#include "secure_getenv.h"
+
 /* Convenience macros for the frequently used CUDA library call and
    error handling sequence as well as CUDA library calls that
    do the error checking themselves or don't do it at all.  */
@@ -876,6 +882,104 @@ notify_var (const char *var_name, const char *env_var)
     GOMP_PLUGIN_debug (0, "%s: '%s'\n", var_name, env_var);
 }
 
+static void
+do_prog (const char *prog, const char *arg)
+{
+  pid_t pid = fork ();
+
+  if (pid == -1)
+    {
+      GOMP_PLUGIN_error ("Fork failed");
+      return;
+    }
+  else if (pid > 0)
+    {
+      int status;
+      waitpid (pid, &status, 0);
+      if (!WIFEXITED (status))
+	GOMP_PLUGIN_error ("Running %s %s failed", prog, arg);
+    }
+  else
+    {
+      execlp (prog, prog /* argv[0] */, arg, NULL);
+      abort ();
+    }
+}
+
+static void
+debug_linkout (void *linkout, size_t linkoutsize)
+{
+  static int gomp_nvptx_disasm = -1;
+  if (gomp_nvptx_disasm == -1)
+    {
+      const char *var_name = "GOMP_NVPTX_DISASM";
+      const char *env_var = secure_getenv (var_name);
+      notify_var (var_name, env_var);
+      gomp_nvptx_disasm
+	= ((env_var != NULL && env_var[0] == '1' && env_var[1] == '\0')
+	   ? 1 : 0);
+    }
+
+  static int gomp_nvptx_save_temps = -1;
+  if (gomp_nvptx_save_temps == -1)
+    {
+      const char *var_name = "GOMP_NVPTX_SAVE_TEMPS";
+      const char *env_var = secure_getenv (var_name);
+      notify_var (var_name, env_var);
+      gomp_nvptx_save_temps
+	= ((env_var != NULL && env_var[0] == '1' && env_var[1] == '\0')
+	   ? 1 : 0);
+    }
+
+  if (gomp_nvptx_disasm == 0
+      && gomp_nvptx_save_temps == 0)
+    return;
+
+  const char *prefix = "gomp-nvptx.";
+  const char *postfix = ".cubin";
+  const int len =	(strlen (prefix)
+			 + 20 /* %lld.  */
+			 + strlen (postfix)
+			 + 1  /* '\0'.  */);
+  char file_name[len];
+  int res = snprintf (file_name, len, "%s%lld%s", prefix,
+		      (long long)getpid (), postfix);
+  assert (res < len); /* Assert there's no truncation.  */
+
+  GOMP_PLUGIN_debug (0, "Generating %s with size %zu\n",
+		     file_name, linkoutsize);
+  FILE *cubin_file = fopen (file_name, "wb");
+  if (cubin_file == NULL)
+    {
+      GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+      return;
+    }
+
+  fwrite (linkout, linkoutsize, 1, cubin_file);
+  unsigned int write_succeeded = ferror (cubin_file) == 0;
+  if (!write_succeeded)
+    GOMP_PLUGIN_debug (0, "Writing %s failed\n", file_name);
+
+  res = fclose (cubin_file);
+  if (res != 0)
+    GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+
+  if (!write_succeeded)
+    return;
+
+  if (gomp_nvptx_disasm == 1)
+    {
+      GOMP_PLUGIN_debug (0, "Disassembling %s\n", file_name);
+      do_prog ("nvdisasm", file_name);
+    }
+
+  if (gomp_nvptx_save_temps == 0)
+    {
+      GOMP_PLUGIN_debug (0, "Removing %s\n", file_name);
+      remove (file_name);
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
@@ -939,6 +1043,8 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
       return false;
     }
 
+  debug_linkout (linkout, linkoutsize);
+
   CUDA_CALL (cuModuleLoadData, module, linkout);
   CUDA_CALL (cuLinkDestroy, linkstate);
   return true;

[PATCH, 2/3] Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 2172 bytes --]

On 07/04/2017 12:05 PM, Tom de Vries wrote:
> On 07/03/2017 04:24 PM, Tom de Vries wrote:
>> On 07/03/2017 04:08 PM, Thomas Schwinge wrote:
>>> Hi!
>>>
>>> On Mon, 26 Jun 2017 17:29:11 +0200, Jakub Jelinek <jakub@redhat.com> 
>>> wrote:
>>>> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>>>>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>>>>
>>>>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx 
>>>>>>> plugin
>>>>>>
>>>>>> This patch adds handling of:
>>>>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>>>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
>>>
>>> Why the "OPENACC" in these names?
>>
>> I took the format from 'GOMP_OPENACC_DIM'.
>>
>>> Doesn't this debugging aid apply to
>>> any variant of offloading?
>>
>> I guess you're right. These environment variables would also be 
>> applicable for f.i. offloading via openmp on nvptx. I'll strip the 
>> 'OPENACC_' bit from the variables.
>>
>>>>>> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
>>>
>>> Also, I suggest to make these names similar to their controlling 
>>> options,
>>> that is: "gomp-nvptx*", for example.
>>>
>>
>> Makes sense, will do.
> 
> Changes in the patch series:
> - removed OPENACC_ from environment variable names
> - made temp files use gomp-nvptx prefix.
> - fixed build error due to missing _GNU_SOURCE in libgomp-nvptx.c.
> - merged the three GOMP_NVPTX_JIT patches into one
> - rewrote GOMP_NVPTX_JIT to add no extra flags to the JIT compiler
>    invocation if GOMP_NVPTX_JIT if not defined, removing the need for
>    hardcoding default values
> - added CU_JIT_TARGET to plugin/cuda/cuda.h
> 
> Build on x86_64 with nvptx offloading enabled (using plugin/cuda/cuda.h).
> 
> The patch series now looks like:
> 1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
> 2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
> 3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
>     plugin
> 
> I'll repost the patch series in reply to this email.

2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
    ( original submission at
      https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01992.html )

Thanks,
- Tom

[-- Attachment #2: 0002-Handle-GOMP_NVPTX_PTXRW-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 4576 bytes --]

Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin

2017-06-27  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (post_process_ptx): New function.
	(link_ptx): Call post_process_ptx.

---
 libgomp/plugin/plugin-nvptx.c | 132 +++++++++++++++++++++++++++++++++++++++++-
 1 file changed, 130 insertions(+), 2 deletions(-)

diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 3e33c5b..cc2ee5e 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -980,6 +980,131 @@ debug_linkout (void *linkout, size_t linkoutsize)
     }
 }
 
+/* If environment variable GOMP_NVPTX_PTXRW=[Ww], write *RES_CODE to file
+   gomp-nvptx.<NUM>.ptx.  If it is [Rr], read *RES_CODE from file
+   instead.  */
+
+static void
+post_process_ptx (unsigned num, const char **res_code, size_t *res_size)
+{
+  static int gomp_nvptx_ptxrw = -1;
+
+  if (gomp_nvptx_ptxrw == -1)
+    {
+      const char *var_name = "GOMP_NVPTX_PTXRW";
+      const char *env_var = secure_getenv (var_name);
+      notify_var (var_name, env_var);
+
+      gomp_nvptx_ptxrw = 0;
+      if (env_var == NULL)
+	;
+      else if ((env_var[0] == 'w' || env_var[0] == 'W')
+	       && env_var[1] == '\0')
+	gomp_nvptx_ptxrw = 1;
+      else if ((env_var[0] == 'r' || env_var[0] == 'R')
+	       && env_var[1] == '\0')
+	gomp_nvptx_ptxrw = 2;
+      else
+	GOMP_PLUGIN_error ("Error parsing %s", var_name);
+    }
+
+  if (gomp_nvptx_ptxrw == 0)
+    return;
+
+  const char *code = *res_code;
+  size_t size = *res_size;
+
+  const char *prefix = "gomp-nvptx.";
+  const char *postfix = ".ptx";
+  const int len =	(strlen (prefix)
+			 + 10 /* %u.  */
+			 + strlen (postfix)
+			 + 1  /* '\0'.  */);
+  char file_name[len];
+  int res = snprintf (file_name, len, "%s%u%s", prefix,
+		      num, postfix);
+  assert (res < len); /* Assert there's no truncation.  */
+
+  GOMP_PLUGIN_debug (0, "%s %s \n",
+		     (gomp_nvptx_ptxrw == 1 ? "Writing" : "Reading"),
+		     file_name);
+
+  if (gomp_nvptx_ptxrw == 1)
+    {
+      FILE *ptx_file = fopen (file_name, "w");
+      if (ptx_file == NULL)
+	{
+	  GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+	  return;
+	}
+
+      int res = fprintf (ptx_file, "%s", code);
+      unsigned int write_succeeded = res == size - 1;
+      if (!write_succeeded)
+	GOMP_PLUGIN_debug (0,
+			   "Writing %s failed: written %d but expected %zu\n",
+			   file_name, res, size - 1);
+
+      res = fclose (ptx_file);
+      if (res != 0)
+	GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+
+      return;
+    }
+
+  if (gomp_nvptx_ptxrw == 2)
+    {
+      FILE *ptx_file = fopen (file_name, "r");
+      if (ptx_file == NULL)
+	{
+	  GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+	  return;
+	}
+
+      if (fseek (ptx_file, 0L, SEEK_END) != 0)
+	{
+	  GOMP_PLUGIN_debug (0, "Seeking end of %s failed\n", file_name);
+	  return;
+	}
+
+      long bufsize = ftell (ptx_file);
+      if (bufsize == -1)
+	{
+	  GOMP_PLUGIN_debug (0, "ftell of %s failed\n", file_name);
+	  return;
+	}
+
+      if (fseek (ptx_file, 0L, SEEK_SET) != 0)
+	{
+	  GOMP_PLUGIN_debug (0, "Seeking start of %s failed\n", file_name);
+	  return;
+	}
+
+	char *new_code = GOMP_PLUGIN_malloc (sizeof (char) * (bufsize + 1));
+
+	size_t new_size = fread (new_code, sizeof (char), bufsize, ptx_file);
+	if (ferror (ptx_file) != 0)
+	  {
+	    GOMP_PLUGIN_debug (0, "Reading %s failed\n", file_name);
+	    return;
+	  }
+
+	assert (new_size < bufsize + 1);
+	new_code[new_size++] = '\0';
+
+	int res = fclose (ptx_file);
+	if (res != 0)
+	  {
+	    GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+	    return;
+	  }
+
+	*res_code = new_code;
+	*res_size = new_size;
+	return;
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
@@ -1016,11 +1141,14 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 
   for (; num_objs--; ptx_objs++)
     {
+      const char *ptx_code = ptx_objs->code;
+      size_t ptx_size = ptx_objs->size;
+      post_process_ptx (num_objs, &ptx_code, &ptx_size);
+      GOMP_PLUGIN_debug (0, "Loading:\n---\n%s\n---\n", ptx_code);
       /* cuLinkAddData's 'data' argument erroneously omits the const
 	 qualifier.  */
-      GOMP_PLUGIN_debug (0, "Loading:\n---\n%s\n---\n", ptx_objs->code);
       r = CUDA_CALL_NOCHECK (cuLinkAddData, linkstate, CU_JIT_INPUT_PTX,
-			     (char *) ptx_objs->code, ptx_objs->size,
+			     (char *) ptx_code, ptx_size,
 			     0, 0, 0, 0);
       if (r != CUDA_SUCCESS)
 	{

[PATCH, 3/3] Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 2370 bytes --]

On 07/04/2017 12:05 PM, Tom de Vries wrote:
> On 07/03/2017 04:24 PM, Tom de Vries wrote:
>> On 07/03/2017 04:08 PM, Thomas Schwinge wrote:
>>> Hi!
>>>
>>> On Mon, 26 Jun 2017 17:29:11 +0200, Jakub Jelinek <jakub@redhat.com> 
>>> wrote:
>>>> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>>>>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>>>>
>>>>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx 
>>>>>>> plugin
>>>>>>
>>>>>> This patch adds handling of:
>>>>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>>>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
>>>
>>> Why the "OPENACC" in these names?
>>
>> I took the format from 'GOMP_OPENACC_DIM'.
>>
>>> Doesn't this debugging aid apply to
>>> any variant of offloading?
>>
>> I guess you're right. These environment variables would also be 
>> applicable for f.i. offloading via openmp on nvptx. I'll strip the 
>> 'OPENACC_' bit from the variables.
>>
>>>>>> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
>>>
>>> Also, I suggest to make these names similar to their controlling 
>>> options,
>>> that is: "gomp-nvptx*", for example.
>>>
>>
>> Makes sense, will do.
> 
> Changes in the patch series:
> - removed OPENACC_ from environment variable names
> - made temp files use gomp-nvptx prefix.
> - fixed build error due to missing _GNU_SOURCE in libgomp-nvptx.c.
> - merged the three GOMP_NVPTX_JIT patches into one
> - rewrote GOMP_NVPTX_JIT to add no extra flags to the JIT compiler
>    invocation if GOMP_NVPTX_JIT if not defined, removing the need for
>    hardcoding default values
> - added CU_JIT_TARGET to plugin/cuda/cuda.h
> 
> Build on x86_64 with nvptx offloading enabled (using plugin/cuda/cuda.h).
> 
> The patch series now looks like:
> 1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
> 2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
> 3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
>     plugin
> 
> I'll repost the patch series in reply to this email.
> 

3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
    plugin
    ( combination of 3 GOMP_NVPTX_JIT patches originally submitted at:
      https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01921.html
      https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01920.html
      https://gcc.gnu.org/ml/gcc-patches/2017-06/msg02407.html )

Thanks,
- Tom

[-- Attachment #2: 0003-Handle-GOMP_NVPTX_JIT-O-0-4-ori-arch-n-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 4399 bytes --]

Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx plugin

2017-06-26  Tom de Vries  <tom@codesourcery.com>

	* plugin/cuda/cuda.h (enum CUjit_option): Add CU_JIT_OPTIMIZATION_LEVEL,
	CU_JIT_NEW_SM3X_OPT and CU_JIT_TARGET.
	* plugin/plugin-nvptx.c (parse_number): New function.
	(process_GOMP_NVPTX_JIT): New function.
	(link_ptx): Add CU_JIT_OPTIMIZATION_LEVEL, CU_JIT_NEW_SM3X_OPT and
	CU_JIT_TARGET to opts if specified.

---
 libgomp/plugin/cuda/cuda.h    |   5 +-
 libgomp/plugin/plugin-nvptx.c | 108 ++++++++++++++++++++++++++++++++++++++++--
 2 files changed, 109 insertions(+), 4 deletions(-)

diff --git a/libgomp/plugin/cuda/cuda.h b/libgomp/plugin/cuda/cuda.h
index 25d5d19..7d190f1 100644
--- a/libgomp/plugin/cuda/cuda.h
+++ b/libgomp/plugin/cuda/cuda.h
@@ -88,7 +88,10 @@ typedef enum {
   CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES = 4,
   CU_JIT_ERROR_LOG_BUFFER = 5,
   CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES = 6,
-  CU_JIT_LOG_VERBOSE = 12
+  CU_JIT_OPTIMIZATION_LEVEL = 7,
+  CU_JIT_TARGET = 9,
+  CU_JIT_LOG_VERBOSE = 12,
+  CU_JIT_NEW_SM3X_OPT = 15
 } CUjit_option;
 
 typedef enum {
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index cc2ee5e..f5b9502 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -144,6 +144,10 @@ init_cuda_lib (void)
 
 #include "secure_getenv.h"
 
+#if CUDA_VERSION < 8000
+#define CU_JIT_NEW_SM3X_OPT 15
+#endif
+
 /* Convenience macros for the frequently used CUDA library call and
    error handling sequence as well as CUDA library calls that
    do the error checking themselves or don't do it at all.  */
@@ -1106,11 +1110,77 @@ post_process_ptx (unsigned num, const char **res_code, size_t *res_size)
 }
 
 static bool
+parse_number (const char *c, unsigned long* resp, char **end)
+{
+  unsigned long res;
+
+  errno = 0;
+  res = strtoul (c, end, 10);
+  if (errno)
+    return false;
+
+  *resp = res;
+  return true;
+}
+
+static void
+process_GOMP_NVPTX_JIT (intptr_t *gomp_nvptx_o, intptr_t *gomp_nvptx_ori,
+			uintptr_t *gomp_nvptx_target)
+{
+  const char *var_name = "GOMP_NVPTX_JIT";
+  const char *env_var = getenv (var_name);
+  notify_var (var_name, env_var);
+
+  if (env_var == NULL)
+    return;
+
+  const char *c = env_var;
+  while (*c != '\0')
+    {
+      while (*c == ' ')
+	c++;
+
+      if (c[0] == '-' && c[1] == 'O'
+	  && '0' <= c[2] && c[2] <= '4'
+	  && (c[3] == '\0' || c[3] == ' '))
+	{
+	  *gomp_nvptx_o = c[2] - '0';
+	  c += 3;
+	  continue;
+	}
+
+      if (c[0] == '-' && c[1] == 'o' && c[2] == 'r' && c[3] == 'i'
+	  && (c[4] == '\0' || c[4] == ' '))
+	{
+	  *gomp_nvptx_ori = 1;
+	  c += 4;
+	  continue;
+	}
+
+      if (c[0] == '-' && c[1] == 'a' && c[2] == 'r' && c[3] == 'c'
+	  && c[4] == 'h' && c[5] == '=')
+	{
+	  const char *end;
+	  unsigned long val;
+	  if (parse_number (&c[6], &val, (char**)&end))
+	    {
+	      *gomp_nvptx_target = val;
+	      c = end;
+	      continue;
+	    }
+	}
+
+      GOMP_PLUGIN_error ("Error parsing %s", var_name);
+      break;
+    }
+}
+
+static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
 {
-  CUjit_option opts[6];
-  void *optvals[6];
+  CUjit_option opts[9];
+  void *optvals[9];
   float elapsed = 0.0;
   char elog[1024];
   char ilog[16384];
@@ -1137,7 +1207,39 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
   opts[5] = CU_JIT_LOG_VERBOSE;
   optvals[5] = (void *) 1;
 
-  CUDA_CALL (cuLinkCreate, 6, opts, optvals, &linkstate);
+  static intptr_t gomp_nvptx_o = -1;
+  static intptr_t gomp_nvptx_ori = -1;
+  static uintptr_t gomp_nvptx_target = 0;
+
+  static bool init_done = false;
+  if (!init_done)
+    {
+      process_GOMP_NVPTX_JIT (&gomp_nvptx_o, &gomp_nvptx_ori,
+			      &gomp_nvptx_target);
+      init_done = true;
+  }
+
+  int nopts = 6;
+  if (gomp_nvptx_o != -1)
+    {
+      opts[nopts] = CU_JIT_OPTIMIZATION_LEVEL;
+      optvals[nopts] = (void *) gomp_nvptx_o;
+      nopts++;
+    }
+  if (gomp_nvptx_ori != -1)
+    {
+      opts[nopts] = CU_JIT_NEW_SM3X_OPT;
+      optvals[nopts] = (void *) gomp_nvptx_ori;
+      nopts++;
+    }
+  if (gomp_nvptx_target != 0)
+    {
+      opts[nopts] = CU_JIT_TARGET;
+      optvals[nopts] = (void *) gomp_nvptx_target;
+      nopts++;
+    }
+
+  CUDA_CALL (cuLinkCreate, nopts, opts, optvals, &linkstate);
 
   for (; num_objs--; ptx_objs++)
     {

[PING] [PATCH, 3/3] Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx plugin

[Tom de Vries]

On 07/04/2017 12:23 PM, Tom de Vries wrote:
> On 07/04/2017 12:05 PM, Tom de Vries wrote:
>> On 07/03/2017 04:24 PM, Tom de Vries wrote:
>>> On 07/03/2017 04:08 PM, Thomas Schwinge wrote:
>>>> Hi!
>>>>
>>>> On Mon, 26 Jun 2017 17:29:11 +0200, Jakub Jelinek <jakub@redhat.com> 
>>>> wrote:
>>>>> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>>>>>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>>>>>
>>>>>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp 
>>>>>>>> nvptx plugin
>>>>>>>
>>>>>>> This patch adds handling of:
>>>>>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>>>>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
>>>>
>>>> Why the "OPENACC" in these names?
>>>
>>> I took the format from 'GOMP_OPENACC_DIM'.
>>>
>>>> Doesn't this debugging aid apply to
>>>> any variant of offloading?
>>>
>>> I guess you're right. These environment variables would also be 
>>> applicable for f.i. offloading via openmp on nvptx. I'll strip the 
>>> 'OPENACC_' bit from the variables.
>>>
>>>>>>> The filename used for dumping the module is 
>>>>>>> plugin-nvptx.<pid>.cubin.
>>>>
>>>> Also, I suggest to make these names similar to their controlling 
>>>> options,
>>>> that is: "gomp-nvptx*", for example.
>>>>
>>>
>>> Makes sense, will do.
>>
>> Changes in the patch series:
>> - removed OPENACC_ from environment variable names
>> - made temp files use gomp-nvptx prefix.
>> - fixed build error due to missing _GNU_SOURCE in libgomp-nvptx.c.
>> - merged the three GOMP_NVPTX_JIT patches into one
>> - rewrote GOMP_NVPTX_JIT to add no extra flags to the JIT compiler
>>    invocation if GOMP_NVPTX_JIT if not defined, removing the need for
>>    hardcoding default values
>> - added CU_JIT_TARGET to plugin/cuda/cuda.h
>>
>> Build on x86_64 with nvptx offloading enabled (using plugin/cuda/cuda.h).
>>
>> The patch series now looks like:
>> 1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
>> 2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
>> 3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
>>     plugin
>>
>> I'll repost the patch series in reply to this email.
>>
> 
> 3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
>     plugin
>     ( combination of 3 GOMP_NVPTX_JIT patches originally submitted at:
>       https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01921.html
>       https://gcc.gnu.org/ml/gcc-patches/2017-06/msg01920.html
>       https://gcc.gnu.org/ml/gcc-patches/2017-06/msg02407.html )
> 


Ping. I'd like to use GOMP_NVPTX_JIT in a workaround for a cuda JIT bug 
triggered in libgomp.c/for-5.c (see PR81805), like this:
...
  /* { dg-set-target-env-var GOMP_NVPTX_JIT "-O0" } */
...

Thanks,- Tom

> 
> 0003-Handle-GOMP_NVPTX_JIT-O-0-4-ori-arch-n-in-libgomp-nvptx-plugin.patch
> 
> 
> Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx plugin
> 
> 2017-06-26  Tom de Vries  <tom@codesourcery.com>
> 
> 	* plugin/cuda/cuda.h (enum CUjit_option): Add CU_JIT_OPTIMIZATION_LEVEL,
> 	CU_JIT_NEW_SM3X_OPT and CU_JIT_TARGET.
> 	* plugin/plugin-nvptx.c (parse_number): New function.
> 	(process_GOMP_NVPTX_JIT): New function.
> 	(link_ptx): Add CU_JIT_OPTIMIZATION_LEVEL, CU_JIT_NEW_SM3X_OPT and
> 	CU_JIT_TARGET to opts if specified.
> 
> ---
>   libgomp/plugin/cuda/cuda.h    |   5 +-
>   libgomp/plugin/plugin-nvptx.c | 108 ++++++++++++++++++++++++++++++++++++++++--
>   2 files changed, 109 insertions(+), 4 deletions(-)
> 
> diff --git a/libgomp/plugin/cuda/cuda.h b/libgomp/plugin/cuda/cuda.h
> index 25d5d19..7d190f1 100644
> --- a/libgomp/plugin/cuda/cuda.h
> +++ b/libgomp/plugin/cuda/cuda.h
> @@ -88,7 +88,10 @@ typedef enum {
>     CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES = 4,
>     CU_JIT_ERROR_LOG_BUFFER = 5,
>     CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES = 6,
> -  CU_JIT_LOG_VERBOSE = 12
> +  CU_JIT_OPTIMIZATION_LEVEL = 7,
> +  CU_JIT_TARGET = 9,
> +  CU_JIT_LOG_VERBOSE = 12,
> +  CU_JIT_NEW_SM3X_OPT = 15
>   } CUjit_option;
>   
>   typedef enum {
> diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
> index cc2ee5e..f5b9502 100644
> --- a/libgomp/plugin/plugin-nvptx.c
> +++ b/libgomp/plugin/plugin-nvptx.c
> @@ -144,6 +144,10 @@ init_cuda_lib (void)
>   
>   #include "secure_getenv.h"
>   
> +#if CUDA_VERSION < 8000
> +#define CU_JIT_NEW_SM3X_OPT 15
> +#endif
> +
>   /* Convenience macros for the frequently used CUDA library call and
>      error handling sequence as well as CUDA library calls that
>      do the error checking themselves or don't do it at all.  */
> @@ -1106,11 +1110,77 @@ post_process_ptx (unsigned num, const char **res_code, size_t *res_size)
>   }
>   
>   static bool
> +parse_number (const char *c, unsigned long* resp, char **end)
> +{
> +  unsigned long res;
> +
> +  errno = 0;
> +  res = strtoul (c, end, 10);
> +  if (errno)
> +    return false;
> +
> +  *resp = res;
> +  return true;
> +}
> +
> +static void
> +process_GOMP_NVPTX_JIT (intptr_t *gomp_nvptx_o, intptr_t *gomp_nvptx_ori,
> +			uintptr_t *gomp_nvptx_target)
> +{
> +  const char *var_name = "GOMP_NVPTX_JIT";
> +  const char *env_var = getenv (var_name);
> +  notify_var (var_name, env_var);
> +
> +  if (env_var == NULL)
> +    return;
> +
> +  const char *c = env_var;
> +  while (*c != '\0')
> +    {
> +      while (*c == ' ')
> +	c++;
> +
> +      if (c[0] == '-' && c[1] == 'O'
> +	  && '0' <= c[2] && c[2] <= '4'
> +	  && (c[3] == '\0' || c[3] == ' '))
> +	{
> +	  *gomp_nvptx_o = c[2] - '0';
> +	  c += 3;
> +	  continue;
> +	}
> +
> +      if (c[0] == '-' && c[1] == 'o' && c[2] == 'r' && c[3] == 'i'
> +	  && (c[4] == '\0' || c[4] == ' '))
> +	{
> +	  *gomp_nvptx_ori = 1;
> +	  c += 4;
> +	  continue;
> +	}
> +
> +      if (c[0] == '-' && c[1] == 'a' && c[2] == 'r' && c[3] == 'c'
> +	  && c[4] == 'h' && c[5] == '=')
> +	{
> +	  const char *end;
> +	  unsigned long val;
> +	  if (parse_number (&c[6], &val, (char**)&end))
> +	    {
> +	      *gomp_nvptx_target = val;
> +	      c = end;
> +	      continue;
> +	    }
> +	}
> +
> +      GOMP_PLUGIN_error ("Error parsing %s", var_name);
> +      break;
> +    }
> +}
> +
> +static bool
>   link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
>   	  unsigned num_objs)
>   {
> -  CUjit_option opts[6];
> -  void *optvals[6];
> +  CUjit_option opts[9];
> +  void *optvals[9];
>     float elapsed = 0.0;
>     char elog[1024];
>     char ilog[16384];
> @@ -1137,7 +1207,39 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
>     opts[5] = CU_JIT_LOG_VERBOSE;
>     optvals[5] = (void *) 1;
>   
> -  CUDA_CALL (cuLinkCreate, 6, opts, optvals, &linkstate);
> +  static intptr_t gomp_nvptx_o = -1;
> +  static intptr_t gomp_nvptx_ori = -1;
> +  static uintptr_t gomp_nvptx_target = 0;
> +
> +  static bool init_done = false;
> +  if (!init_done)
> +    {
> +      process_GOMP_NVPTX_JIT (&gomp_nvptx_o, &gomp_nvptx_ori,
> +			      &gomp_nvptx_target);
> +      init_done = true;
> +  }
> +
> +  int nopts = 6;
> +  if (gomp_nvptx_o != -1)
> +    {
> +      opts[nopts] = CU_JIT_OPTIMIZATION_LEVEL;
> +      optvals[nopts] = (void *) gomp_nvptx_o;
> +      nopts++;
> +    }
> +  if (gomp_nvptx_ori != -1)
> +    {
> +      opts[nopts] = CU_JIT_NEW_SM3X_OPT;
> +      optvals[nopts] = (void *) gomp_nvptx_ori;
> +      nopts++;
> +    }
> +  if (gomp_nvptx_target != 0)
> +    {
> +      opts[nopts] = CU_JIT_TARGET;
> +      optvals[nopts] = (void *) gomp_nvptx_target;
> +      nopts++;
> +    }
> +
> +  CUDA_CALL (cuLinkCreate, nopts, opts, optvals, &linkstate);
>   
>     for (; num_objs--; ptx_objs++)
>       {
> 

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Cesar Philippidis]

On 07/04/2017 03:05 AM, Tom de Vries wrote:
> On 07/03/2017 04:24 PM, Tom de Vries wrote:
>> On 07/03/2017 04:08 PM, Thomas Schwinge wrote:
>>> Hi!
>>>
>>> On Mon, 26 Jun 2017 17:29:11 +0200, Jakub Jelinek <jakub@redhat.com>
>>> wrote:
>>>> On Mon, Jun 26, 2017 at 03:26:57PM +0000, Joseph Myers wrote:
>>>>> On Mon, 26 Jun 2017, Tom de Vries wrote:
>>>>>
>>>>>>> 2. Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx
>>>>>>> plugin
>>>>>>
>>>>>> This patch adds handling of:
>>>>>> - GOMP_OPENACC_NVPTX_SAVE_TEMPS=[01], and
>>>>>> - GOMP_OPENACC_NVPTX_DISASM=[01]
>>>
>>> Why the "OPENACC" in these names?
>>
>> I took the format from 'GOMP_OPENACC_DIM'.
>>
>>> Doesn't this debugging aid apply to
>>> any variant of offloading?
>>
>> I guess you're right. These environment variables would also be
>> applicable for f.i. offloading via openmp on nvptx. I'll strip the
>> 'OPENACC_' bit from the variables.
>>
>>>>>> The filename used for dumping the module is plugin-nvptx.<pid>.cubin.
>>>
>>> Also, I suggest to make these names similar to their controlling
>>> options,
>>> that is: "gomp-nvptx*", for example.
>>>
>>
>> Makes sense, will do.
> 
> Changes in the patch series:
> - removed OPENACC_ from environment variable names
> - made temp files use gomp-nvptx prefix.
> - fixed build error due to missing _GNU_SOURCE in libgomp-nvptx.c.
> - merged the three GOMP_NVPTX_JIT patches into one
> - rewrote GOMP_NVPTX_JIT to add no extra flags to the JIT compiler
>   invocation if GOMP_NVPTX_JIT if not defined, removing the need for
>   hardcoding default values
> - added CU_JIT_TARGET to plugin/cuda/cuda.h
> 
> Build on x86_64 with nvptx offloading enabled (using plugin/cuda/cuda.h).
> 
> The patch series now looks like:
> 1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
> 2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
> 3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
>    plugin
> 
> I'll repost the patch series in reply to this email.

Ping.

Can we get this patch series into trunk and og7? The ability to easily
modify PTX code, via GOMP_NVPTX_PTXRW, is extremely helpful. It helped
me isolate one problem already.

Thanks,
Cesar

Re: [PATCH, 2/4] Handle GOMP_OPENACC_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin

[Jakub Jelinek]

On Tue, Nov 07, 2017 at 06:48:25AM -0800, Cesar Philippidis wrote:
> > Changes in the patch series:
> > - removed OPENACC_ from environment variable names
> > - made temp files use gomp-nvptx prefix.
> > - fixed build error due to missing _GNU_SOURCE in libgomp-nvptx.c.
> > - merged the three GOMP_NVPTX_JIT patches into one
> > - rewrote GOMP_NVPTX_JIT to add no extra flags to the JIT compiler
> >   invocation if GOMP_NVPTX_JIT if not defined, removing the need for
> >   hardcoding default values
> > - added CU_JIT_TARGET to plugin/cuda/cuda.h
> > 
> > Build on x86_64 with nvptx offloading enabled (using plugin/cuda/cuda.h).
> > 
> > The patch series now looks like:
> > 1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
> > 2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
> > 3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
> >    plugin
> > 
> > I'll repost the patch series in reply to this email.
> 
> Ping.
> 
> Can we get this patch series into trunk and og7? The ability to easily
> modify PTX code, via GOMP_NVPTX_PTXRW, is extremely helpful. It helped
> me isolate one problem already.

It can be helpful for debugging, but I'm afraid about having such code in
production, I think something like this would be very easy to exploit.
Sure, running a suid or sgid program with offloading is probably very
dangerous anyway, but it could be just some minor priviledge escalation
in the app (SELinux, ACLs, whatever else) and this stuff would allow anyone
to run anything else.
So, IMNSHO if it should be added, only enabled by non-default configure
option.

	Jakub

[PATCH] Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin

[Tom de Vries]

[-- Attachment #1: Type: text/plain, Size: 2277 bytes --]

On 11/07/2017 03:54 PM, Jakub Jelinek wrote:
> On Tue, Nov 07, 2017 at 06:48:25AM -0800, Cesar Philippidis wrote:
>>> Changes in the patch series:
>>> - removed OPENACC_ from environment variable names
>>> - made temp files use gomp-nvptx prefix.
>>> - fixed build error due to missing _GNU_SOURCE in libgomp-nvptx.c.
>>> - merged the three GOMP_NVPTX_JIT patches into one
>>> - rewrote GOMP_NVPTX_JIT to add no extra flags to the JIT compiler
>>>    invocation if GOMP_NVPTX_JIT if not defined, removing the need for
>>>    hardcoding default values
>>> - added CU_JIT_TARGET to plugin/cuda/cuda.h
>>>
>>> Build on x86_64 with nvptx offloading enabled (using plugin/cuda/cuda.h).
>>>
>>> The patch series now looks like:
>>> 1. Handle GOMP_NVPTX_{DISASM,SAVE_TEMPS} in libgomp nvptx plugin
>>> 2. Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin
>>> 3. Handle GOMP_NVPTX_JIT={-O[0-4],-ori,-arch=<n>} in libgomp nvptx
>>>     plugin
>>>
>>> I'll repost the patch series in reply to this email.
>>
>> Ping.
>>
>> Can we get this patch series into trunk and og7? The ability to easily
>> modify PTX code, via GOMP_NVPTX_PTXRW, is extremely helpful. It helped
>> me isolate one problem already.
> 
> It can be helpful for debugging, but I'm afraid about having such code in
> production, I think something like this would be very easy to exploit.
> Sure, running a suid or sgid program with offloading is probably very
> dangerous anyway, but it could be just some minor priviledge escalation
> in the app (SELinux, ACLs, whatever else) and this stuff would allow anyone
> to run anything else.
> So, IMNSHO if it should be added, only enabled by non-default configure
> option.

Hi,

I've made the GOMP_NVPTX_PTXRW patch stand-alone, and added an 
off-by-default libgomp configure option 
--enable-libgomp-plugin-developer-only-options, which sets a config.h 
macro LIBGOMP_PLUGIN_DEVELOPER_ONLY_OPTIONS, which is used to 
enable/disable the GOMP_NVPTX_PTXRW functionality.

I've build this on x86_64 for nvptx accelerator, both with and without 
the configure option, and confirmed that in one case using 
GOMP_NVPTX_PTXRW=w generates a gomp-nvptx.0.ptx file, and in the other 
case it doesn't.

OK for trunk if x86_64 bootstrap and reg-test succeeds?

Thanks,
- Tom

[-- Attachment #2: 0001-Handle-GOMP_NVPTX_PTXRW-in-libgomp-nvptx-plugin.patch --]
[-- Type: text/x-patch, Size: 10663 bytes --]

Handle GOMP_NVPTX_PTXRW in libgomp nvptx plugin

2017-11-21  Tom de Vries  <tom@codesourcery.com>

	* plugin/plugin-nvptx.c (_GNU_SOURCE): Define.
	(gomp_nvptx_ptxrw): New static variable.
	(parse_gomp_nvptx_ptxrw, post_process_ptx_write, post_process_ptx_read)
	(post_process_ptx): New function.
	(link_ptx): Call post_process_ptx.
	* configure.ac: Add configure option
	--enable-libgomp-plugin-developer-only-options.
	* config.h.in: Regenerate.
	* configure: Same.

---
 libgomp/config.h.in           |   3 +
 libgomp/configure             |  32 ++++++++-
 libgomp/configure.ac          |  11 +++
 libgomp/plugin/plugin-nvptx.c | 160 +++++++++++++++++++++++++++++++++++++++++-
 4 files changed, 202 insertions(+), 4 deletions(-)

diff --git a/libgomp/config.h.in b/libgomp/config.h.in
index e7bc4d97374..68cccea4186 100644
--- a/libgomp/config.h.in
+++ b/libgomp/config.h.in
@@ -118,6 +118,9 @@
 /* Define to 1 if building libgomp for an accelerator-only target. */
 #undef LIBGOMP_OFFLOADED_ONLY
 
+/* Define to 1 if libgomp plugin developer-only options are enabled. */
+#undef LIBGOMP_PLUGIN_DEVELOPER_ONLY_OPTIONS
+
 /* Define to 1 if libgomp should use POSIX threads. */
 #undef LIBGOMP_USE_PTHREADS
 
diff --git a/libgomp/configure b/libgomp/configure
index e7842b5519f..14e39d7fbec 100755
--- a/libgomp/configure
+++ b/libgomp/configure
@@ -780,6 +780,7 @@ ac_subst_files=''
 ac_user_opts='
 enable_option_checking
 enable_version_specific_runtime_libs
+enable_libgomp_plugin_developer_only_options
 enable_generated_files_in_srcdir
 enable_multilib
 enable_dependency_tracking
@@ -1434,6 +1435,9 @@ Optional Features:
   --enable-version-specific-runtime-libs
                           Specify that runtime libraries should be installed
                           in a compiler-specific directory [default=no]
+  --enable-libgomp-plugin-developer-only-options
+                          Specify that libgomp plugins should be build with
+                          additional developer-only options [default=no]
   --enable-generated-files-in-srcdir
                           put copies of generated files in source dir intended
                           for creating source tarballs for users without
@@ -2627,6 +2631,30 @@ fi
 { $as_echo "$as_me:${as_lineno-$LINENO}: result: $enable_version_specific_runtime_libs" >&5
 $as_echo "$enable_version_specific_runtime_libs" >&6; }
 
+{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for --enable-libgomp-plugin-developer-only-options" >&5
+$as_echo_n "checking for --enable-libgomp-plugin-developer-only-options... " >&6; }
+ # Check whether --enable-libgomp-plugin-developer-only-options was given.
+if test "${enable_libgomp_plugin_developer_only_options+set}" = set; then :
+  enableval=$enable_libgomp_plugin_developer_only_options;
+      case "$enableval" in
+       yes|no) ;;
+       *) as_fn_error "Unknown argument to enable/disable libgomp-plugin-developer-only-options" "$LINENO" 5 ;;
+                          esac
+
+else
+  enable_libgomp_plugin_developer_only_options=no
+fi
+
+
+{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $enable_libgomp_plugin_developer_only_options" >&5
+$as_echo "$enable_libgomp_plugin_developer_only_options" >&6; }
+if test x$enable_libgomp_plugin_developer_only_options != xno; then
+
+$as_echo "#define LIBGOMP_PLUGIN_DEVELOPER_ONLY_OPTIONS 1" >>confdefs.h
+
+fi
+
+
 # We would like our source tree to be readonly. However when releases or
 # pre-releases are generated, the flex/bison generated files as well as the
 # various formats of manuals need to be included along with the rest of the
@@ -11158,7 +11186,7 @@ else
   lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
   lt_status=$lt_dlunknown
   cat > conftest.$ac_ext <<_LT_EOF
-#line 11161 "configure"
+#line 11189 "configure"
 #include "confdefs.h"
 
 #if HAVE_DLFCN_H
@@ -11264,7 +11292,7 @@ else
   lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
   lt_status=$lt_dlunknown
   cat > conftest.$ac_ext <<_LT_EOF
-#line 11267 "configure"
+#line 11295 "configure"
 #include "confdefs.h"
 
 #if HAVE_DLFCN_H
diff --git a/libgomp/configure.ac b/libgomp/configure.ac
index 4e0bc8166a9..c6cfb1f0796 100644
--- a/libgomp/configure.ac
+++ b/libgomp/configure.ac
@@ -15,6 +15,17 @@ LIBGOMP_ENABLE(version-specific-runtime-libs, no, ,
    permit yes|no)
 AC_MSG_RESULT($enable_version_specific_runtime_libs)
 
+AC_MSG_CHECKING([for --enable-libgomp-plugin-developer-only-options])
+LIBGOMP_ENABLE(libgomp-plugin-developer-only-options, no, ,
+   [Specify that libgomp plugins should be build with additional developer-only options],
+   permit yes|no)
+AC_MSG_RESULT($enable_libgomp_plugin_developer_only_options)
+if test x$enable_libgomp_plugin_developer_only_options != xno; then
+  AC_DEFINE(LIBGOMP_PLUGIN_DEVELOPER_ONLY_OPTIONS, 1,
+            [Define to 1 if libgomp plugin developer-only options are enabled.])
+fi
+
+
 # We would like our source tree to be readonly. However when releases or
 # pre-releases are generated, the flex/bison generated files as well as the
 # various formats of manuals need to be included along with the rest of the
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 71630b57355..0cbc8fc197d 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -31,6 +31,7 @@
    is not clear as to what that state might be.  Or how one might
    propagate it from one thread to another.  */
 
+#define _GNU_SOURCE
 #include "openacc.h"
 #include "config.h"
 #include "libgomp-plugin.h"
@@ -138,6 +139,8 @@ init_cuda_lib (void)
 # define init_cuda_lib() true
 #endif
 
+#include "secure_getenv.h"
+
 /* Convenience macros for the frequently used CUDA library call and
    error handling sequence as well as CUDA library calls that
    do the error checking themselves or don't do it at all.  */
@@ -876,6 +879,156 @@ notify_var (const char *var_name, const char *env_var)
     GOMP_PLUGIN_debug (0, "%s: '%s'\n", var_name, env_var);
 }
 
+#ifdef LIBGOMP_PLUGIN_DEVELOPER_ONLY_OPTIONS
+static int gomp_nvptx_ptxrw = -1;
+#else
+static int gomp_nvptx_ptxrw = 0;
+#endif
+
+/* Parse environment variable GOMP_NVPTX_PTXRW=[WwRr].  */
+
+static void
+parse_gomp_nvptx_ptxrw (void)
+{
+  gomp_nvptx_ptxrw = 0;
+
+  const char *var_name = "GOMP_NVPTX_PTXRW";
+  const char *env_var = secure_getenv (var_name);
+  notify_var (var_name, env_var);
+
+  if (env_var == NULL)
+    ;
+  else if ((env_var[0] == 'w' || env_var[0] == 'W')
+	   && env_var[1] == '\0')
+    gomp_nvptx_ptxrw = 1;
+  else if ((env_var[0] == 'r' || env_var[0] == 'R')
+	   && env_var[1] == '\0')
+    gomp_nvptx_ptxrw = 2;
+  else
+    GOMP_PLUGIN_error ("Error parsing %s", var_name);
+}
+
+/* Write CODE with length SIZE to file FILE_NAME.  */
+
+static void
+post_process_ptx_write (char *file_name, const char *code, size_t size)
+{
+  FILE *ptx_file = fopen (file_name, "w");
+  if (ptx_file == NULL)
+    {
+      GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+      return;
+    }
+
+  int res = fprintf (ptx_file, "%s", code);
+  unsigned int write_succeeded = res == size - 1;
+  if (!write_succeeded)
+    GOMP_PLUGIN_debug (0,
+		       "Writing %s failed: written %d but expected %zu\n",
+		       file_name, res, size - 1);
+
+  res = fclose (ptx_file);
+  if (res != 0)
+    GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+}
+
+/* Read the contents of FILE_NAME into *RES_CODE and save the file size
+   in *RES_SIZE.  */
+
+static void
+post_process_ptx_read (char *file_name, const char **res_code, size_t *res_size)
+{
+  FILE *ptx_file = fopen (file_name, "r");
+  if (ptx_file == NULL)
+    {
+      GOMP_PLUGIN_debug (0, "Opening %s failed\n", file_name);
+      return;
+    }
+
+  if (fseek (ptx_file, 0L, SEEK_END) != 0)
+    {
+      GOMP_PLUGIN_debug (0, "Seeking end of %s failed\n", file_name);
+      return;
+    }
+
+  long bufsize = ftell (ptx_file);
+  if (bufsize == -1)
+    {
+      GOMP_PLUGIN_debug (0, "ftell of %s failed\n", file_name);
+      return;
+    }
+
+  if (fseek (ptx_file, 0L, SEEK_SET) != 0)
+    {
+      GOMP_PLUGIN_debug (0, "Seeking start of %s failed\n", file_name);
+      return;
+    }
+
+  char *new_code = GOMP_PLUGIN_malloc (sizeof (char) * (bufsize + 1));
+
+  size_t new_size = fread (new_code, sizeof (char), bufsize, ptx_file);
+  if (ferror (ptx_file) != 0)
+    {
+      GOMP_PLUGIN_debug (0, "Reading %s failed\n", file_name);
+      return;
+    }
+
+  assert (new_size < bufsize + 1);
+  new_code[new_size++] = '\0';
+
+  int res = fclose (ptx_file);
+  if (res != 0)
+    {
+      GOMP_PLUGIN_debug (0, "Closing %s failed\n", file_name);
+      return;
+    }
+
+  *res_code = new_code;
+  *res_size = new_size;
+}
+
+/* If environment variable GOMP_NVPTX_PTXRW=[Ww], write *RES_CODE to file
+   gomp-nvptx.<NUM>.ptx.  If it is [Rr], read *RES_CODE from file
+   instead.  */
+
+static void
+post_process_ptx (unsigned num, const char **res_code, size_t *res_size)
+{
+  if (gomp_nvptx_ptxrw == -1)
+    parse_gomp_nvptx_ptxrw ();
+
+  if (gomp_nvptx_ptxrw == 0)
+    return;
+
+  const char *code = *res_code;
+  size_t size = *res_size;
+
+  const char *prefix = "gomp-nvptx.";
+  const char *postfix = ".ptx";
+  const int len = (strlen (prefix)
+		   + 10 /* %u.  */
+		   + strlen (postfix)
+		   + 1  /* '\0'.  */);
+  char file_name[len];
+  int res = snprintf (file_name, len, "%s%u%s", prefix,
+		      num, postfix);
+  assert (res < len); /* Assert there's no truncation.  */
+
+  GOMP_PLUGIN_debug (0, "%s %s \n",
+		     (gomp_nvptx_ptxrw == 1 ? "Writing" : "Reading"),
+		     file_name);
+
+  switch (gomp_nvptx_ptxrw)
+    {
+    case 1:
+      post_process_ptx_write (file_name, code, size);
+      break;
+    case 2:
+      post_process_ptx_read (file_name, res_code, res_size);
+      break;
+    }
+}
+
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
@@ -912,11 +1065,14 @@ link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 
   for (; num_objs--; ptx_objs++)
     {
+      const char *ptx_code = ptx_objs->code;
+      size_t ptx_size = ptx_objs->size;
+      post_process_ptx (num_objs, &ptx_code, &ptx_size);
+      GOMP_PLUGIN_debug (0, "Loading:\n---\n%s\n---\n", ptx_code);
       /* cuLinkAddData's 'data' argument erroneously omits the const
 	 qualifier.  */
-      GOMP_PLUGIN_debug (0, "Loading:\n---\n%s\n---\n", ptx_objs->code);
       r = CUDA_CALL_NOCHECK (cuLinkAddData, linkstate, CU_JIT_INPUT_PTX,
-			     (char *) ptx_objs->code, ptx_objs->size,
+			     (char *) ptx_code, ptx_size,
 			     0, 0, 0, 0);
       if (r != CUDA_SUCCESS)
 	{