Ping.

> I've noticed while re-reading that this patch incorrectly handled SIMT case
> in lower_lastprivate_clauses.  The code was changed to look for variables
> with "omp simt private" attribute, and was left under
> 'simduid && DECL_HAS_VALUE_EXPR_P (new_var)' condition.  This effectively
> constrained processing to privatized (i.e. addressable) variables, as
> non-addressable variables now have neither the value-expr nor the attribute.
> 
> This wasn't caught in testing, as apparently all testcases that have target
> simd loops with linear/lastprivate clauses also have the corresponding 
> variables
> mentioned in target map clause, which makes them addressable (is that 
> necessary?),
> and I didn't think to check something like that manually.
> 
> The following patch fixes it by adjusting the if's in 
> lower_lastprivate_clauses;
> alternatively it may be possible to keep that code as-is, and instead set up
> value-expr and "omp simt private" attributes for all formally private 
> variables,
> not just addressable ones, but to me that sounds less preferable.  OK for 
> trunk?
> 
> I think it's possible to improve test coverage for NVPTX by running all OpenMP
> testcases via nvptx-none-run (it'll need some changes, but shouldn't be hard).
> 
> gcc/
>       * omp-low.c (lower_lastprivate_clauses): Correct handling of linear and
>       lastprivate clauses in SIMT case.
> 
> libgomp/
>       * testsuite/libgomp.c/target-36.c: New testcase.
> 
> Thanks.
> Alexander
>       
> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> index 253dc85..02b681c 100644
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -4768,11 +4768,10 @@ lower_lastprivate_clauses (tree clauses, tree 
> predicate, gimple_seq *stmt_list,
>               TREE_NO_WARNING (new_var) = 1;
>           }
>  
> -       if (simduid && DECL_HAS_VALUE_EXPR_P (new_var))
> +       if (!maybe_simt && simduid && DECL_HAS_VALUE_EXPR_P (new_var))
>           {
>             tree val = DECL_VALUE_EXPR (new_var);
> -           if (!maybe_simt
> -               && TREE_CODE (val) == ARRAY_REF
> +           if (TREE_CODE (val) == ARRAY_REF
>                 && VAR_P (TREE_OPERAND (val, 0))
>                 && lookup_attribute ("omp simd array",
>                                      DECL_ATTRIBUTES (TREE_OPERAND (val,
> @@ -4792,26 +4791,26 @@ lower_lastprivate_clauses (tree clauses, tree 
> predicate, gimple_seq *stmt_list,
>                                   TREE_OPERAND (val, 0), lastlane,
>                                   NULL_TREE, NULL_TREE);
>               }
> -           else if (maybe_simt
> -                    && VAR_P (val)
> -                    && lookup_attribute ("omp simt private",
> -                                         DECL_ATTRIBUTES (val)))
> +         }
> +       else if (maybe_simt)
> +         {
> +           tree val = (DECL_HAS_VALUE_EXPR_P (new_var)
> +                       ? DECL_VALUE_EXPR (new_var)
> +                       : new_var);
> +           if (simtlast == NULL)
>               {
> -               if (simtlast == NULL)
> -                 {
> -                   simtlast = create_tmp_var (unsigned_type_node);
> -                   gcall *g = gimple_build_call_internal
> -                     (IFN_GOMP_SIMT_LAST_LANE, 1, simtcond);
> -                   gimple_call_set_lhs (g, simtlast);
> -                   gimple_seq_add_stmt (stmt_list, g);
> -                 }
> -               x = build_call_expr_internal_loc
> -                 (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_IDX,
> -                  TREE_TYPE (val), 2, val, simtlast);
> -               new_var = unshare_expr (new_var);
> -               gimplify_assign (new_var, x, stmt_list);
> -               new_var = unshare_expr (new_var);
> +               simtlast = create_tmp_var (unsigned_type_node);
> +               gcall *g = gimple_build_call_internal
> +                 (IFN_GOMP_SIMT_LAST_LANE, 1, simtcond);
> +               gimple_call_set_lhs (g, simtlast);
> +               gimple_seq_add_stmt (stmt_list, g);
>               }
> +           x = build_call_expr_internal_loc
> +             (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_IDX,
> +              TREE_TYPE (val), 2, val, simtlast);
> +           new_var = unshare_expr (new_var);
> +           gimplify_assign (new_var, x, stmt_list);
> +           new_var = unshare_expr (new_var);
>           }
>  
>         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
> diff --git a/libgomp/testsuite/libgomp.c/target-36.c 
> b/libgomp/testsuite/libgomp.c/target-36.c
> new file mode 100644
> index 0000000..6925a2a
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/target-36.c
> @@ -0,0 +1,18 @@
> +int
> +main ()
> +{
> +  int ah, bh, n = 1024;
> +#pragma omp target map(from: ah, bh)
> +  {
> +    int a, b;
> +#pragma omp simd lastprivate(b)
> +    for (a = 0; a < n; a++)
> +      {
> +     b = a + n + 1;
> +     asm volatile ("" : "+r"(b));
> +      }
> +    ah = a, bh = b;
> +  }
> +  if (ah != n || bh != 2 * n)
> +    __builtin_abort ();
> +}
> 

Reply via email to