Hi Cesar! On Thu, 20 Mar 2014 15:42:48 +0100, I wrote: > On Tue, 18 Mar 2014 14:50:44 +0100, I wrote: > > On Tue, 18 Mar 2014 16:37:24 +0400, Ilmir Usmanov <i.usma...@samsung.com> > > wrote: > > > This patch introduces support of OpenACC loop directive (and combined > > > directives) in C front-end up to GENERIC. Currently no clause is allowed. > > > > Thanks! I had worked on a simpler patch, not yet dealing with combined > > clauses. Also, I have some work for the GIMPLE level, namely building on > > GIMPLE_OMP_FOR, adding a new GF_OMP_FOR_KIND_OACC_LOOP. I'll post this > > soon. > > Here are the patches, committed in r208702..4 to gomp-4_0-branch.
Cesar, I hope I'm not confusing things here, but I remember that you once
pointed out that in Fortran OpenACC, we have to explicitly specify the
loop iteration variable in a private clause, whereas the OpenACC
specification says this needs to happen automatically (predetermined data
attribute). I see gcc/fortran/openmp.c:gfc_resolve_do_iterator add this
private clause, which I assume we're using also for OpenACC loops.
Looking at -fdump-tree-all output for the following two test cases
complied in OpenMP and OpenACC mode, I see that indeed an explicit
private clause is added for Fortran code, but not for C. Why is that
required? (I have not yet spent any time on figuring this out myself.)
int
main(void)
{
int i;
#pragma acc parallel
#pragma acc loop
#pragma omp parallel for
for (i = 0; i < 10; ++i)
;
return 0;
}
program test
implicit none
integer :: i
!$acc parallel
!$acc loop
!$omp parallel do
DO i = 1, 10
ENDDO
!$omp end parallel do
!$acc end parallel
end
In light of this, please also review whether the following
gimplify_omp_for changes of mine (when I originally added the OACC_LOOP
support) are correct. Evidently, this code still has TODO markers in it;
this was well before you added preliminary support for the private
clause. That is, should we use GOVD_PRIVATE also for OACC_LOOP?
Why does this nevertheless currently work for C for loops without any
private clause for the loop iteration variable? Maybe because in C, the
loop iteration variable ends up in a register and that is not shared
between threads? I do see a private clause being added during
gimplification for the OpenMP C test case, but not for OpenACC -- which I
assume is precisely due to the code I'm quoting below.
Once this has been resolved, please also remove the explicit private
clauses from the existing test cases (Fortran and also C, as applicable)
to make sure that the predermined data attributes are tested there.
> --- gcc/gimplify.c
> +++ gcc/gimplify.c
> @@ -6683,14 +6683,36 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
> gimple_seq for_body, for_pre_body;
> int i;
> bool simd;
> + enum gimplify_omp_var_data govd_private;
> + enum omp_region_type ort;
> bitmap has_decl_expr = NULL;
>
> orig_for_stmt = for_stmt = *expr_p;
>
> - simd = TREE_CODE (for_stmt) == OMP_SIMD
> - || TREE_CODE (for_stmt) == CILK_SIMD;
> - gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p,
> - simd ? ORT_SIMD : ORT_WORKSHARE);
> + switch (TREE_CODE (for_stmt))
> + {
> + case OMP_FOR:
> + case OMP_DISTRIBUTE:
> + simd = false;
> + govd_private = GOVD_PRIVATE;
> + ort = ORT_WORKSHARE;
> + break;
> + case OACC_LOOP:
> + simd = false;
> + govd_private = /* TODO */ GOVD_LOCAL;
> + ort = /* TODO */ ORT_WORKSHARE;
> + break;
> + case OMP_SIMD:
> + case CILK_SIMD:
> + simd = true;
> + govd_private = GOVD_PRIVATE;
> + ort = ORT_SIMD;
> + break;
> + default:
> + gcc_unreachable ();
> + }
> +
> + gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort);
>
> /* Handle OMP_FOR_INIT. */
> for_pre_body = NULL;
> @@ -6722,6 +6744,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>
> if (OMP_FOR_INIT (for_stmt) == NULL_TREE)
> {
> + gcc_assert (TREE_CODE (for_stmt) != OACC_LOOP);
> for_stmt = walk_tree (&OMP_FOR_BODY (for_stmt), find_combined_omp_for,
> NULL, NULL);
> gcc_assert (for_stmt != NULL_TREE);
> @@ -6742,7 +6765,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
> gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
> || POINTER_TYPE_P (TREE_TYPE (decl)));
>
> - /* Make sure the iteration variable is private. */
> + /* Make sure the iteration variable is some kind of private. */
> tree c = NULL_TREE;
> if (orig_for_stmt != for_stmt)
> /* Do this only on innermost construct for combined ones. */;
> @@ -6768,6 +6791,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
> }
> else
> {
> + gcc_assert (govd_private == GOVD_PRIVATE);
> bool lastprivate
> = (!has_decl_expr
> || !bitmap_bit_p (has_decl_expr, DECL_UID (decl)));
> @@ -6785,7 +6809,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
> else if (omp_is_private (gimplify_omp_ctxp, decl, simd))
> omp_notice_variable (gimplify_omp_ctxp, decl, true);
> else
> - omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
> + omp_add_variable (gimplify_omp_ctxp, decl, govd_private | GOVD_SEEN);
>
> /* If DECL is not a gimple register, create a temporary variable to act
> as an iteration counter. This is valid, since DECL cannot be
> @@ -6799,7 +6823,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>
> gimplify_seq_add_stmt (&for_body, gimple_build_assign (decl, var));
>
> - omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
> + omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
> }
> else
> var = decl;
> @@ -6936,7 +6960,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
> t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
> decl = TREE_OPERAND (t, 0);
> var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
> - omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
> + omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
> TREE_OPERAND (t, 0) = var;
> t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
> TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1));
> @@ -6952,6 +6976,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
> case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break;
> case CILK_SIMD: kind = GF_OMP_FOR_KIND_CILKSIMD; break;
> case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break;
> + case OACC_LOOP: kind = GF_OMP_FOR_KIND_OACC_LOOP; break;
> default:
> gcc_unreachable ();
> }
Grüße,
Thomas
pgpDm91yCvz7n.pgp
Description: PGP signature
