[og7] Properly handle alloca'd arrays in OpenACC data mappings

Message ID 8b5aa0ce-448a-0702-9c25-bd2a0a6c156b@codesourcery.com
State New
Headers show
Series
  • [og7] Properly handle alloca'd arrays in OpenACC data mappings
Related show

Commit Message

Cesar Philippidis Jan. 31, 2018, 11:07 p.m.
This patch teaches install_parm_decl (part of the PTX .param
optimization during omp-lowering) not to extract the identifier string
of artificial decls. Apparently, GCC's OpenACC testsuite did not have
any tests that contained local arrays which used alloca for storage
allocation, or else this problem would have been detected sooner. I'm
considering moving the PTX .param pass later, possible during
oaccdevlow. But that will have to wait for some other time.

I've applied this patch to openacc-gcc-7-branch.

Cesar

Patch

Properly handle alloca'd OpenACC data mappings

2018-01-31  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* omp-low.c (install_parm_decl): Don't extract identifiers from
	artifical decls.

	gcc/testsuite/
	* c-c++-common/goacc/large_array.c: New test.


diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index a2869e49ebd..aa04ed3151f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -669,7 +669,7 @@  install_parm_decl (tree var, tree type, omp_context *ctx)
   tree decl_name = NULL_TREE, t;
   location_t loc = UNKNOWN_LOCATION;
 
-  if (DECL_P (var))
+  if (DECL_P (var) && !DECL_ARTIFICIAL (var))
     {
       decl_name = get_identifier (get_name (var));
       loc = DECL_SOURCE_LOCATION (var);
diff --git a/gcc/testsuite/c-c++-common/goacc/large_array.c b/gcc/testsuite/c-c++-common/goacc/large_array.c
new file mode 100644
index 00000000000..ce0f4c12a74
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/large_array.c
@@ -0,0 +1,18 @@ 
+/* Ensure that alloca'ed arrays can be transferred to the
+   accelerator. */
+
+/* { dg-require-effective-target alloca } */
+
+int
+main ()
+{
+  int n = 100, m = 10, i, j;
+  float a[n][m];
+
+  #pragma acc parallel loop
+  for (i = 0; i < n; i++)
+    for (j = 0; j < m; j++)
+      a[i][j] = 0;
+
+  return 0;
+}