Skip to content
Snippets Groups Projects
  • Tom de Vries's avatar
    3f2e15c2
    [openacc] Fix acc declare for VLAs · 3f2e15c2
    Tom de Vries authored
    Consider test-case test.c, with VLA A:
    ...
    int main (void) {
      int N = 1000;
      int A[N];
      #pragma acc declare copy(A)
      return 0;
    }
    ...
    compiled using:
    ...
    $ gcc test.c -fopenacc -S -fdump-tree-all
    ...
    
    At original, we have:
    ...
      #pragma acc declare map(tofrom:A);
    ...
    but at gimple, we have a map (to:A.1), but not a map (from:A.1):
    ...
      int[0:D.2074] * A.1;
    
      {
        int A[0:D.2074] [value-expr: *A.1];
    
        saved_stack.2 = __builtin_stack_save ();
        try
          {
            A.1 = __builtin_alloca_with_align (D.2078, 32);
            #pragma omp target oacc_declare map(to:(*A.1) [len: D.2076])
          }
        finally
          {
            __builtin_stack_restore (saved_stack.2);
          }
      }
    ...
    
    This is caused by the following incompatibility.  When storing the desired
    from clause in oacc_declare_returns, we use 'A.1' as the key:
    ...
    10898                 oacc_declare_returns->put (decl, c);
    (gdb) call debug_generic_expr (decl)
    A.1
    (gdb) call debug_generic_expr (c)
    map(from:(*A.1))
    ...
    but when looking it up, we use 'A' as the key:
    ...
    (gdb)
    1471                  tree *c = oacc_declare_returns->get (t);
    (gdb) call debug_generic_expr (t)
    A
    ...
    
    Fix this by extracing the 'A.1' lookup key from 'A' using the decl-expr.
    
    In addition, unshare the looked up value, to fix avoid running into
    an "incorrect sharing of tree nodes" error.
    
    Using these two fixes, we get our desired:
    ...
         finally
           {
    +        #pragma omp target oacc_declare map(from:(*A.1))
             __builtin_stack_restore (saved_stack.2);
           }
    ...
    
    Build on x86_64-linux with nvptx accelerator, tested libgomp.
    
    gcc/ChangeLog:
    
    2020-10-06  Tom de Vries  <tdevries@suse.de>
    
    	PR middle-end/90861
    	* gimplify.c (gimplify_bind_expr): Handle lookup in
    	oacc_declare_returns using key with decl-expr.
    
    libgomp/ChangeLog:
    
    2020-10-06  Tom de Vries  <tdevries@suse.de>
    
    	PR middle-end/90861
    	* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Remove xfail.
    3f2e15c2
    History
    [openacc] Fix acc declare for VLAs
    Tom de Vries authored
    Consider test-case test.c, with VLA A:
    ...
    int main (void) {
      int N = 1000;
      int A[N];
      #pragma acc declare copy(A)
      return 0;
    }
    ...
    compiled using:
    ...
    $ gcc test.c -fopenacc -S -fdump-tree-all
    ...
    
    At original, we have:
    ...
      #pragma acc declare map(tofrom:A);
    ...
    but at gimple, we have a map (to:A.1), but not a map (from:A.1):
    ...
      int[0:D.2074] * A.1;
    
      {
        int A[0:D.2074] [value-expr: *A.1];
    
        saved_stack.2 = __builtin_stack_save ();
        try
          {
            A.1 = __builtin_alloca_with_align (D.2078, 32);
            #pragma omp target oacc_declare map(to:(*A.1) [len: D.2076])
          }
        finally
          {
            __builtin_stack_restore (saved_stack.2);
          }
      }
    ...
    
    This is caused by the following incompatibility.  When storing the desired
    from clause in oacc_declare_returns, we use 'A.1' as the key:
    ...
    10898                 oacc_declare_returns->put (decl, c);
    (gdb) call debug_generic_expr (decl)
    A.1
    (gdb) call debug_generic_expr (c)
    map(from:(*A.1))
    ...
    but when looking it up, we use 'A' as the key:
    ...
    (gdb)
    1471                  tree *c = oacc_declare_returns->get (t);
    (gdb) call debug_generic_expr (t)
    A
    ...
    
    Fix this by extracing the 'A.1' lookup key from 'A' using the decl-expr.
    
    In addition, unshare the looked up value, to fix avoid running into
    an "incorrect sharing of tree nodes" error.
    
    Using these two fixes, we get our desired:
    ...
         finally
           {
    +        #pragma omp target oacc_declare map(from:(*A.1))
             __builtin_stack_restore (saved_stack.2);
           }
    ...
    
    Build on x86_64-linux with nvptx accelerator, tested libgomp.
    
    gcc/ChangeLog:
    
    2020-10-06  Tom de Vries  <tdevries@suse.de>
    
    	PR middle-end/90861
    	* gimplify.c (gimplify_bind_expr): Handle lookup in
    	oacc_declare_returns using key with decl-expr.
    
    libgomp/ChangeLog:
    
    2020-10-06  Tom de Vries  <tdevries@suse.de>
    
    	PR middle-end/90861
    	* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Remove xfail.