12-12-2013 Ilmir Usmanov <i.usmanov@samsung.com>

        Consider #pragma acc loop as construct in GENERIC and high
        GIMPLE

        * gcc/c/c-parser.c(c_parser_acc_clause_check_no_duplicate):
        Beautification
        (c_begin_acc_loop, c_finish_acc_loop): New functions
        (check_acc_loop_incr_expr, c_finish_acc_loop, 
        c_parser_acc_loop_1): Remove
        (c_parser_acc_loop, c_parser_acc_parallel, c_parser_acc_kernels): 
        Consider #pragma acc loop as construct
        * gcc/fortran/trans-openacc.c(trans_acc_loop, 
        trans_acc_parallel_loop, trans_acc_kernels_loop, 
        gfc_trans_acc_directive): Likewise
        * gcc/gimple-oacc.c(gimple_build_acc_loop): Likewise
        * gcc/gimplify.c(gimplify_acc_loop): Likewise
        * gcc/tree-pretty-print.c(dump_generic_node): Likewise
        * gcc/tree.def(ACC_LOOP): Likewise
        * gcc/tree.h(ACC_LOOP_INIT, ACC_LOOP_COND, ACC_LOOP_INCR,
        ACC_LOOP_PRE_BODY): Remove
          


git-svn-id: https://gcc.gnu.org/svn/gcc/branches/openacc-1_0-branch@205928 138bc75d-0d04-0410-961f-82ee72b054a4

9 files changed, 113 insertions, 962 deletions

diff --git a/ChangeLog.ACC b/ChangeLog.ACC
index 3fef54d5943..32945a4a9c4 100644
--- a/ChangeLog.ACC
+++ b/ChangeLog.ACC
@@ -1,4 +1,25 @@
 12-12-2013 Ilmir Usmanov <i.usmanov@samsung.com>
+	Consider #pragma acc loop as construct in GENERIC and high
+	GIMPLE
+	
+	* gcc/c/c-parser.c(c_parser_acc_clause_check_no_duplicate):
+	Beautification
+	(c_begin_acc_loop, c_finish_acc_loop): New functions
+	(check_acc_loop_incr_expr, c_finish_acc_loop, 
+	c_parser_acc_loop_1): Remove
+	(c_parser_acc_loop, c_parser_acc_parallel, c_parser_acc_kernels): 
+	Consider #pragma acc loop as construct 
+	* gcc/fortran/trans-openacc.c(trans_acc_loop, 
+	trans_acc_parallel_loop, trans_acc_kernels_loop, 
+	gfc_trans_acc_directive): Likewise
+	* gcc/gimple-oacc.c(gimple_build_acc_loop): Likewise
+	* gcc/gimplify.c(gimplify_acc_loop): Likewise
+	* gcc/tree-pretty-print.c(dump_generic_node): Likewise
+	* gcc/tree.def(ACC_LOOP): Likewise
+	* gcc/tree.h(ACC_LOOP_INIT, ACC_LOOP_COND, ACC_LOOP_INCR,
+	ACC_LOOP_PRE_BODY): Remove
+
+12-12-2013 Ilmir Usmanov <i.usmanov@samsung.com>
 	Fix parsing of scalar integer expressions and variable lists
 	in C FE
 
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 7e4ff2ffd14..4767b79d703 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11045,8 +11045,8 @@ c_parser_acc_clause_name (c_parser *parser)
 /* Validate that a clause of the given type does not already exist.  */
 static void
 c_parser_acc_clause_check_no_duplicate(tree clauses,
-                                           enum acc_clause_code code,
-                                           const char *name) {
+                                       enum acc_clause_code code,
+                                       const char *name) {
   tree c;
 
   for (c = clauses; c; c = ACC_CLAUSE_CHAIN(c)) {
@@ -11894,6 +11894,19 @@ c_begin_acc_host_data (void)
   return block;
 }
 
+/* Unlike OMP, in ACC we consider LOOP directive as construct.
+ */
+tree
+c_begin_acc_loop (void)
+{
+  tree block;
+
+  keep_next_level ();
+  block = c_begin_compound_stmt (true);
+
+  return block;
+}
+
 /* Generate ACC_PARALLEL, with CLAUSES and BLOCK as its compound
    statement. */
 tree
@@ -11966,563 +11979,22 @@ c_finish_acc_host_data (location_t loc, tree clauses, tree block)
   return add_stmt (stmt);
 }
 
-
-/* FIXME: split with check_omp_for_incr_expr */
-static tree
-check_acc_loop_incr_expr (location_t loc, tree exp, tree decl)
-{
-  tree t;
-
-  if (!INTEGRAL_TYPE_P (TREE_TYPE (exp))
-      || TYPE_PRECISION (TREE_TYPE (exp)) < TYPE_PRECISION (TREE_TYPE (decl)))
-    return error_mark_node;
-
-  if (exp == decl)
-    return build_int_cst (TREE_TYPE (exp), 0);
-
-  switch (TREE_CODE (exp))
-    {
-    CASE_CONVERT:
-      t = check_acc_loop_incr_expr (loc, TREE_OPERAND (exp, 0), decl);
-      if (t != error_mark_node)
-        return fold_convert_loc (loc, TREE_TYPE (exp), t);
-      break;
-    case MINUS_EXPR:
-      t = check_acc_loop_incr_expr (loc, TREE_OPERAND (exp, 0), decl);
-      if (t != error_mark_node)
-        return fold_build2_loc (loc, MINUS_EXPR,
-                            TREE_TYPE (exp), t, TREE_OPERAND (exp, 1));
-      break;
-    case PLUS_EXPR:
-      t = check_acc_loop_incr_expr (loc, TREE_OPERAND (exp, 0), decl);
-      if (t != error_mark_node)
-        return fold_build2_loc (loc, PLUS_EXPR,
-                            TREE_TYPE (exp), t, TREE_OPERAND (exp, 1));
-      t = check_acc_loop_incr_expr (loc, TREE_OPERAND (exp, 1), decl);
-      if (t != error_mark_node)
-        return fold_build2_loc (loc, PLUS_EXPR,
-                            TREE_TYPE (exp), TREE_OPERAND (exp, 0), t);
-      break;
-    case COMPOUND_EXPR:
-      {
-        /* cp_build_modify_expr forces preevaluation of the RHS to make
-           sure that it is evaluated before the lvalue-rvalue conversion
-           is applied to the LHS.  Reconstruct the original expression.  */
-        tree op0 = TREE_OPERAND (exp, 0);
-        if (TREE_CODE (op0) == TARGET_EXPR
-            && !VOID_TYPE_P (TREE_TYPE (op0)))
-          {
-            tree op1 = TREE_OPERAND (exp, 1);
-            tree temp = TARGET_EXPR_SLOT (op0);
-            if (TREE_CODE_CLASS (TREE_CODE (op1)) == tcc_binary
-                && TREE_OPERAND (op1, 1) == temp)
-              {
-                op1 = copy_node (op1);
-                TREE_OPERAND (op1, 1) = TARGET_EXPR_INITIAL (op0);
-                return check_acc_loop_incr_expr (loc, op1, decl);
-              }
-          }
-        break;
-      }
-    default:
-      break;
-    }
-
-  return error_mark_node;
-}
-
-/* FIXME: split with c_finish_omp_for */
-static tree
-c_finish_acc_loop (location_t locus, tree declv, tree initv, tree condv,
-                   tree incrv, tree body, tree pre_body)
-{
-  location_t elocus;
-  bool fail = false;
-  int i;
-
-  gcc_assert (TREE_VEC_LENGTH (declv) == TREE_VEC_LENGTH (initv));
-  gcc_assert (TREE_VEC_LENGTH (declv) == TREE_VEC_LENGTH (condv));
-  gcc_assert (TREE_VEC_LENGTH (declv) == TREE_VEC_LENGTH (incrv));
-  for (i = 0; i < TREE_VEC_LENGTH (declv); i++)
-    {
-      tree decl = TREE_VEC_ELT (declv, i);
-      tree init = TREE_VEC_ELT (initv, i);
-      tree cond = TREE_VEC_ELT (condv, i);
-      tree incr = TREE_VEC_ELT (incrv, i);
-
-      elocus = locus;
-      if (EXPR_HAS_LOCATION (init))
-        elocus = EXPR_LOCATION (init);
-
-      /* Validate the iteration variable.  */
-      if (!INTEGRAL_TYPE_P (TREE_TYPE (decl))
-          && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE)
-        {
-          error_at (elocus, "invalid type for iteration variable %qE", decl);
-          fail = true;
-        }
-
-      /* In the case of "for (int i = 0...)", init will be a decl.  It should
-         have a DECL_INITIAL that we can turn into an assignment.  */
-      if (init == decl)
-        {
-          elocus = DECL_SOURCE_LOCATION (decl);
-
-          init = DECL_INITIAL (decl);
-          if (init == NULL)
-            {
-              error_at (elocus, "%qE is not initialized", decl);
-              init = integer_zero_node;
-              fail = true;
-            }
-
-          init = build_modify_expr (elocus, decl, NULL_TREE, NOP_EXPR,
-                                    /* FIXME diagnostics: This should
-                                       be the location of the INIT.  */
-                                    elocus,
-                                    init,
-                                    NULL_TREE);
-        }
-      gcc_assert (TREE_CODE (init) == MODIFY_EXPR);
-      gcc_assert (TREE_OPERAND (init, 0) == decl);
-
-      if (cond == NULL_TREE)
-        {
-          error_at (elocus, "missing controlling predicate");
-          fail = true;
-        }
-      else
-        {
-          bool cond_ok = false;
-
-          if (EXPR_HAS_LOCATION (cond))
-            elocus = EXPR_LOCATION (cond);
-
-          if (TREE_CODE (cond) == LT_EXPR
-              || TREE_CODE (cond) == LE_EXPR
-              || TREE_CODE (cond) == GT_EXPR
-              || TREE_CODE (cond) == GE_EXPR
-              || TREE_CODE (cond) == NE_EXPR
-              || TREE_CODE (cond) == EQ_EXPR)
-            {
-              tree op0 = TREE_OPERAND (cond, 0);
-              tree op1 = TREE_OPERAND (cond, 1);
-
-              /* 2.5.1.  The comparison in the condition is computed in
-                 the type of DECL, otherwise the behavior is undefined.
-
-                 For example:
-                 long n; int i;
-                 i < n;
-
-                 according to ISO will be evaluated as:
-                 (long)i < n;
-
-                 We want to force:
-                 i < (int)n;  */
-              if (TREE_CODE (op0) == NOP_EXPR
-                  && decl == TREE_OPERAND (op0, 0))
-                {
-                  TREE_OPERAND (cond, 0) = TREE_OPERAND (op0, 0);
-                  TREE_OPERAND (cond, 1)
-                    = fold_build1_loc (elocus, NOP_EXPR, TREE_TYPE (decl),
-                                   TREE_OPERAND (cond, 1));
-                }
-              else if (TREE_CODE (op1) == NOP_EXPR
-                       && decl == TREE_OPERAND (op1, 0))
-                {
-                  TREE_OPERAND (cond, 1) = TREE_OPERAND (op1, 0);
-                  TREE_OPERAND (cond, 0)
-                    = fold_build1_loc (elocus, NOP_EXPR, TREE_TYPE (decl),
-                                   TREE_OPERAND (cond, 0));
-                }
-
-              if (decl == TREE_OPERAND (cond, 0))
-                cond_ok = true;
-              else if (decl == TREE_OPERAND (cond, 1))
-                {
-                  TREE_SET_CODE (cond,
-                                 swap_tree_comparison (TREE_CODE (cond)));
-                  TREE_OPERAND (cond, 1) = TREE_OPERAND (cond, 0);
-                  TREE_OPERAND (cond, 0) = decl;
-                  cond_ok = true;
-                }
-
-              if (TREE_CODE (cond) == NE_EXPR
-                  || TREE_CODE (cond) == EQ_EXPR)
-                {
-                  if (!INTEGRAL_TYPE_P (TREE_TYPE (decl)))
-                    cond_ok = false;
-                  else if (operand_equal_p (TREE_OPERAND (cond, 1),
-                                            TYPE_MIN_VALUE (TREE_TYPE (decl)),
-                                            0))
-                    TREE_SET_CODE (cond, TREE_CODE (cond) == NE_EXPR
-                                         ? GT_EXPR : LE_EXPR);
-                  else if (operand_equal_p (TREE_OPERAND (cond, 1),
-                                            TYPE_MAX_VALUE (TREE_TYPE (decl)),
-                                            0))
-                    TREE_SET_CODE (cond, TREE_CODE (cond) == NE_EXPR
-                                         ? LT_EXPR : GE_EXPR);
-                  else
-                    cond_ok = false;
-                }
-            }
-
-          if (!cond_ok)
-            {
-              error_at (elocus, "invalid controlling predicate");
-              fail = true;
-            }
-        }
-
-      if (incr == NULL_TREE)
-        {
-          error_at (elocus, "missing increment expression");
-          fail = true;
-        }
-      else
-        {
-          bool incr_ok = false;
-
-          if (EXPR_HAS_LOCATION (incr))
-            elocus = EXPR_LOCATION (incr);
-
-          /* Check all the valid increment expressions: v++, v--, ++v, --v,
-             v = v + incr, v = incr + v and v = v - incr.  */
-          switch (TREE_CODE (incr))
-            {
-            case POSTINCREMENT_EXPR:
-            case PREINCREMENT_EXPR:
-            case POSTDECREMENT_EXPR:
-            case PREDECREMENT_EXPR:
-              if (TREE_OPERAND (incr, 0) != decl)
-                break;
-
-              incr_ok = true;
-              if (POINTER_TYPE_P (TREE_TYPE (decl))
-                  && TREE_OPERAND (incr, 1))
-                {
-                  tree t = fold_convert_loc (elocus,
-                                             sizetype, TREE_OPERAND (incr, 1));
-
-                  if (TREE_CODE (incr) == POSTDECREMENT_EXPR
-                      || TREE_CODE (incr) == PREDECREMENT_EXPR)
-                    t = fold_build1_loc (elocus, NEGATE_EXPR, sizetype, t);
-                  t = fold_build_pointer_plus (decl, t);
-                  incr = build2 (MODIFY_EXPR, void_type_node, decl, t);
-                }
-              break;
-
-            case MODIFY_EXPR:
-              if (TREE_OPERAND (incr, 0) != decl)
-                break;
-              if (TREE_OPERAND (incr, 1) == decl)
-                break;
-              if (TREE_CODE (TREE_OPERAND (incr, 1)) == PLUS_EXPR
-                  && (TREE_OPERAND (TREE_OPERAND (incr, 1), 0) == decl
-                      || TREE_OPERAND (TREE_OPERAND (incr, 1), 1) == decl))
-                incr_ok = true;
-              else if ((TREE_CODE (TREE_OPERAND (incr, 1)) == MINUS_EXPR
-                        || (TREE_CODE (TREE_OPERAND (incr, 1))
-                            == POINTER_PLUS_EXPR))
-                       && TREE_OPERAND (TREE_OPERAND (incr, 1), 0) == decl)
-                incr_ok = true;
-              else
-                {
-                  tree t = check_acc_loop_incr_expr (elocus,
-                                                    TREE_OPERAND (incr, 1),
-                                                    decl);
-                  if (t != error_mark_node)
-                    {
-                      incr_ok = true;
-                      t = build2 (PLUS_EXPR, TREE_TYPE (decl), decl, t);
-                      incr = build2 (MODIFY_EXPR, void_type_node, decl, t);
-                    }
-                }
-              break;
-
-            default:
-              break;
-            }
-          if (!incr_ok)
-            {
-              error_at (elocus, "invalid increment expression");
-              fail = true;
-            }
-        }
-
-      TREE_VEC_ELT (initv, i) = init;
-      TREE_VEC_ELT (incrv, i) = incr;
-    }
-
-  if (fail)
-    return NULL;
-  else
-    {
-      tree t = make_node (ACC_LOOP);
-
-      TREE_TYPE (t) = void_type_node;
-      ACC_LOOP_INIT (t) = initv;
-      ACC_LOOP_COND (t) = condv;
-      ACC_LOOP_INCR (t) = incrv;
-      ACC_LOOP_BODY (t) = body;
-      ACC_LOOP_PRE_BODY (t) = pre_body;
-
-      SET_EXPR_LOCATION (t, locus);
-      return add_stmt (t);
-    }
-}
-
-static tree
-c_parser_acc_loop_1 (location_t loc,
-                     c_parser *parser,
-                     tree clauses,
-                     tree *par_clauses)
+/* Generate ACC_LOOP, with CLAUSES and BLOCK as its compound
+   statement. */
+tree
+c_finish_acc_loop (location_t loc, tree clauses, tree block)
 {
-  tree decl, cond, incr, save_break, save_cont, body, init, stmt, cl;
-  tree declv, condv, incrv, initv, ret = NULL;
-  bool fail = false, open_brace_parsed = false;
-  int i, collapse = 1, nbraces = 0;
-  location_t for_loc;
-  vec<tree, va_gc> *for_block = make_tree_vector ();
-
-  for (cl = clauses; cl; cl = ACC_CLAUSE_CHAIN (cl))
-  {
-    if (ACC_CLAUSE_CODE (cl) == ACC_CLAUSE_COLLAPSE)
-    {
-      collapse = tree_low_cst (ACC_CLAUSE_COLLAPSE_EXPR (cl), 0);
-    }
-  }
-
-  gcc_assert (collapse >= 1);
-
-  declv = make_tree_vec (collapse);
-  initv = make_tree_vec (collapse);
-  condv = make_tree_vec (collapse);
-  incrv = make_tree_vec (collapse);
-
-  if (!c_parser_next_token_is_keyword (parser, RID_FOR))
-  {
-    c_parser_error (parser, "for statement expected");
-    return NULL;
-  }
-
-  for_loc = c_parser_peek_token (parser)->location;
-  c_parser_consume_token (parser);
-
-  for (i = 0; i < collapse; i++)
-  {
-    int bracecount = 0;
-
-    if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
-      goto pop_scopes;
-
-    /* Parse the initialization declaration or expression.  */
-    if (c_parser_next_tokens_start_declaration (parser))
-      {
-        if (i > 0)
-          vec_safe_push (for_block, c_begin_compound_stmt (true));
-        c_parser_declaration_or_fndef (parser, true, true, true, true, true, NULL);
-        decl = check_for_loop_decls (for_loc, flag_isoc99);
-        if (decl == NULL)
-          goto error_init;
-        if (DECL_INITIAL (decl) == error_mark_node)
-          decl = error_mark_node;
-        init = decl;
-      }
-    else if (c_parser_next_token_is (parser, CPP_NAME)
-             && c_parser_peek_2nd_token (parser)->type == CPP_EQ)
-      {
-        struct c_expr decl_exp;
-        struct c_expr init_exp;
-        location_t init_loc;
-
-        decl_exp = c_parser_postfix_expression (parser);
-        decl = decl_exp.value;
-
-        c_parser_require (parser, CPP_EQ, "expected %<=%>");
-
-        init_loc = c_parser_peek_token (parser)->location;
-        init_exp = c_parser_expr_no_commas (parser, NULL);
-        init_exp = default_function_array_read_conversion (init_loc,
-                                                           init_exp);
-        init = build_modify_expr (init_loc, decl, decl_exp.original_type,
-                                  NOP_EXPR, init_loc, init_exp.value,
-                                  init_exp.original_type);
-        init = c_process_expr_stmt (init_loc, init);
-
-        c_parser_skip_until_found (parser, CPP_SEMICOLON, "expected %<;%>");
-      }
-    else
-      {
-      error_init:
-        c_parser_error (parser,
-                        "expected iteration declaration or initialization");
-        c_parser_skip_until_found (parser, CPP_CLOSE_PAREN,
-                                   "expected %<)%>");
-        fail = true;
-        goto parse_next;
-      }
-
-    /* Parse the loop condition.  */
-    cond = NULL_TREE;
-    if (c_parser_next_token_is_not (parser, CPP_SEMICOLON))
-      {
-        location_t cond_loc = c_parser_peek_token (parser)->location;
-        struct c_expr cond_expr = c_parser_binary_expression (parser, NULL,
-                                                              PREC_NONE);
-
-        cond = cond_expr.value;
-        cond = c_objc_common_truthvalue_conversion (cond_loc, cond);
-        cond = c_fully_fold (cond, false, NULL);
-        switch (cond_expr.original_code)
-          {
-          case GT_EXPR:
-          case GE_EXPR:
-          case LT_EXPR:
-          case LE_EXPR:
-            break;
-          default:
-            /* Can't be cond = error_mark_node, because we want to preserve
-               the location until c_finish_omp_for.  */
-            cond = build1 (NOP_EXPR, boolean_type_node, error_mark_node);
-            break;
-          }
-        protected_set_expr_location (cond, cond_loc);
-      }
-    c_parser_skip_until_found (parser, CPP_SEMICOLON, "expected %<;%>");
-
-    /* Parse the increment expression.  */
-    incr = NULL_TREE;
-    if (c_parser_next_token_is_not (parser, CPP_CLOSE_PAREN))
-      {
-        location_t incr_loc = c_parser_peek_token (parser)->location;
-
-        incr = c_process_expr_stmt (incr_loc,
-                                    c_parser_expression (parser).value);
-      }
-    c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>");
-
-    if (decl == NULL || decl == error_mark_node || init == error_mark_node)
-      fail = true;
-    else
-      {
-        TREE_VEC_ELT (declv, i) = decl;
-        TREE_VEC_ELT (initv, i) = init;
-        TREE_VEC_ELT (condv, i) = cond;
-        TREE_VEC_ELT (incrv, i) = incr;
-      }
-
-  parse_next:
-    if (i == collapse - 1)
-      break;
-
-    do
-      {
-        if (c_parser_next_token_is_keyword (parser, RID_FOR))
-          {
-            c_parser_consume_token (parser);
-            break;
-          }
-        else if (c_parser_next_token_is (parser, CPP_OPEN_BRACE))
-          {
-            c_parser_consume_token (parser);
-            bracecount++;
-          }
-        else if (bracecount
-                 && c_parser_next_token_is (parser, CPP_SEMICOLON))
-          c_parser_consume_token (parser);
-        else
-          {
-            c_parser_error (parser, "not enough perfectly nested loops");
-            if (bracecount)
-              {
-                open_brace_parsed = true;
-                bracecount--;
-              }
-            fail = true;
-            collapse = 0;
-            break;
-          }
-      }
-    while (1);
-
-    nbraces += bracecount;
-  }
-
-  save_break = c_break_label;
-  c_break_label = size_one_node;
-  save_cont = c_cont_label;
-  c_cont_label = NULL_TREE;
-  body = push_stmt_list ();
-
-  if (open_brace_parsed)
-    {
-      location_t here = c_parser_peek_token (parser)->location;
-      stmt = c_begin_compound_stmt (true);
-      c_parser_compound_statement_nostart (parser);
-      add_stmt (c_end_compound_stmt (here, stmt, true));
-    }
-  else
-    add_stmt (c_parser_c99_block_statement (parser));
-  if (c_cont_label)
-    {
-      tree t = build1 (LABEL_EXPR, void_type_node, c_cont_label);
-      SET_EXPR_LOCATION (t, loc);
-      add_stmt (t);
-    }
+  tree stmt;
 
-  body = pop_stmt_list (body);
-  c_break_label = save_break;
-  c_cont_label = save_cont;
+  block = c_end_compound_stmt (loc, block, true);
 
-  while (nbraces)
-    {
-      if (c_parser_next_token_is (parser, CPP_CLOSE_BRACE))
-        {
-          c_parser_consume_token (parser);
-          nbraces--;
-        }
-      else if (c_parser_next_token_is (parser, CPP_SEMICOLON))
-        c_parser_consume_token (parser);
-      else
-        {
-          c_parser_error (parser, "collapsed loops not perfectly nested");
-          while (nbraces)
-            {
-              location_t here = c_parser_peek_token (parser)->location;
-              stmt = c_begin_compound_stmt (true);
-              add_stmt (body);
-              c_parser_compound_statement_nostart (parser);
-              body = c_end_compound_stmt (here, stmt, true);
-              nbraces--;
-            }
-          goto pop_scopes;
-        }
-    }
+  stmt = make_node (ACC_LOOP);
+  TREE_TYPE (stmt) = void_type_node;
+  ACC_LOOP_CLAUSES (stmt) = clauses;
+  ACC_LOOP_BODY (stmt) = block;
+  SET_EXPR_LOCATION (stmt, loc);
 
-  /* Only bother calling c_finish_omp_for if we haven't already generated
-     an error from the initialization parsing.  */
-  if (!fail)
-  {
-    stmt = c_finish_acc_loop (loc, declv, initv, condv, incrv, body, NULL);
-    if (stmt)
-    {
-      if (par_clauses != NULL)
-      {
-      }
-      ACC_LOOP_CLAUSES (stmt) = clauses;
-    }
-    ret = stmt;
-  }
-pop_scopes:
-  while (!for_block->is_empty ())
-    {
-      stmt = c_end_compound_stmt (loc, for_block->pop (), true);
-      add_stmt (stmt);
-    }
-  release_tree_vector (for_block);
-  return ret;
+  return add_stmt (stmt);
 }
 
 /* OpenACC 1.0:
@@ -12531,18 +12003,16 @@ pop_scopes:
 static tree
 c_parser_acc_loop (location_t loc, c_parser *parser)
 {
-  tree block, clauses, ret;
-
-  clauses = c_parser_acc_all_clauses (parser,
-                                      ACC_LOOP_CLAUSE_MASK,
-                                      "#pragma acc loop");
+  const char *p_name = "#pragma acc loop";
+  tree clauses, stmt, block;
+  unsigned int mask = ACC_LOOP_CLAUSE_MASK;
 
-  block = c_begin_compound_stmt (true);
-  ret = c_parser_acc_loop_1 (loc, parser, clauses, NULL);
-  block = c_end_compound_stmt (loc, block, true);
-  add_stmt (block);
+  clauses = c_parser_acc_all_clauses (parser, mask, p_name);
+  block = c_begin_acc_loop ();
+  c_parser_statement (parser);
+  stmt = c_finish_acc_loop (loc, clauses, block);
 
-  return ret;
+  return stmt;
 }
 
 /* OpenACC 1.0:
@@ -12556,7 +12026,7 @@ c_parser_acc_parallel (location_t loc, c_parser *parser)
 {
   enum pragma_kind p_kind = PRAGMA_ACC_PARALLEL;
   const char *p_name = "#pragma acc parallel";
-  tree stmt, clauses, block;
+  tree stmt, clauses, block, loop_block;
   unsigned int mask = ACC_PARALLEL_CLAUSE_MASK;
   struct acc_context ctx;
 
@@ -12587,10 +12057,11 @@ c_parser_acc_parallel (location_t loc, c_parser *parser)
       stmt = c_finish_acc_parallel (loc, clauses, block);
       break;
     case PRAGMA_ACC_PARALLEL_LOOP:
-      block = c_begin_compound_stmt (true);
-      stmt = c_parser_acc_loop_1 (loc, parser, clauses, NULL);
-      block = c_end_compound_stmt (loc, block, true);
-      add_stmt (block);
+      block = c_begin_acc_parallel ();
+      loop_block = c_begin_acc_loop ();
+      c_parser_statement (parser);
+      c_finish_acc_loop (loc, clauses, loop_block);
+      stmt = c_finish_acc_parallel (loc, clauses, block);
       break;
     default:
       gcc_unreachable ();
@@ -12610,7 +12081,7 @@ c_parser_acc_kernels (location_t loc, c_parser *parser)
 {
   enum pragma_kind p_kind = PRAGMA_ACC_KERNELS;
   const char *p_name = "#pragma acc kernels";
-  tree stmt, clauses, block;
+  tree stmt, clauses, block, loop_block;
   unsigned int mask = ACC_KERNELS_CLAUSE_MASK;
   struct acc_context ctx;
 
@@ -12634,21 +12105,22 @@ c_parser_acc_kernels (location_t loc, c_parser *parser)
   clauses = c_parser_acc_all_clauses (parser, mask, p_name);
 
   switch (p_kind)
-  {
-  case PRAGMA_ACC_KERNELS:
-    block = c_begin_acc_kernels ();
-    c_parser_statement (parser);
-    stmt = c_finish_acc_kernels (loc, clauses, block);
-    break;
-  case PRAGMA_ACC_KERNELS_LOOP:
-    block = c_begin_compound_stmt (true);
-    stmt = c_parser_acc_loop_1 (loc, parser, clauses, NULL);
-    block = c_end_compound_stmt (loc, block, true);
-    add_stmt (block);
-    break;
-  default:
-    gcc_unreachable ();
-  }
+    {
+    case PRAGMA_ACC_KERNELS:
+      block = c_begin_acc_kernels ();
+      c_parser_statement (parser);
+      stmt = c_finish_acc_kernels (loc, clauses, block);
+      break;
+    case PRAGMA_ACC_KERNELS_LOOP:
+      block = c_begin_acc_kernels ();
+      loop_block = c_begin_acc_loop ();
+      c_parser_statement (parser);
+      c_finish_acc_loop (loc, clauses, loop_block);
+      stmt = c_finish_acc_kernels (loc, clauses, block);
+      break;
+    default:
+      gcc_unreachable ();
+    }
 
   acc_current_ctx = ctx.previous;
 
diff --git a/gcc/fortran/trans-openacc.c b/gcc/fortran/trans-openacc.c
index cdc0af646a9..c4394289154 100644
--- a/gcc/fortran/trans-openacc.c
+++ b/gcc/fortran/trans-openacc.c
@@ -742,178 +742,18 @@ typedef struct dovar_init_d {
 } dovar_init;
 
 static tree
-trans_acc_loop (gfc_code *code, stmtblock_t *pblock,
-                    gfc_acc_clauses *do_clauses)
+trans_acc_loop (gfc_code *code, gfc_acc_clauses *do_clauses)
 {
-  gfc_se se;
-  tree dovar, stmt, from, to, step, type, init, cond, incr;
-  tree count = NULL_TREE, cycle_label, tmp, acc_clauses;
   stmtblock_t block;
-  stmtblock_t body;
-  gfc_acc_clauses *clauses = code->ext.omp_clauses;
-  int i, collapse = clauses->collapse;
-  vec<dovar_init> inits = vNULL;
-  dovar_init *di;
-  unsigned ix;
-
-  if (collapse <= 0)
-    collapse = 1;
-
-  code = code->block->next;
-  gcc_assert (code->op == EXEC_DO);
-
-  init = make_tree_vec (collapse);
-  cond = make_tree_vec (collapse);
-  incr = make_tree_vec (collapse);
-
-  if (pblock == NULL)
-    {
-      gfc_start_block (&block);
-      pblock = &block;
-    }
-
-  acc_clauses = trans_acc_clauses (pblock, do_clauses, code->loc);
-
-  for (i = 0; i < collapse; i++)
-    {
-      int simple = 0;
-      int dovar_found = 0;
-      tree dovar_decl;
-
-      /* Evaluate all the expressions in the iterator.  */
-      gfc_init_se (&se, NULL);
-      gfc_conv_expr_lhs (&se, code->ext.iterator->var);
-      gfc_add_block_to_block (pblock, &se.pre);
-      dovar = se.expr;
-      type = TREE_TYPE (dovar);
-      gcc_assert (TREE_CODE (type) == INTEGER_TYPE);
-
-      gfc_init_se (&se, NULL);
-      gfc_conv_expr_val (&se, code->ext.iterator->start);
-      gfc_add_block_to_block (pblock, &se.pre);
-      from = gfc_evaluate_now (se.expr, pblock);
-
-      gfc_init_se (&se, NULL);
-      gfc_conv_expr_val (&se, code->ext.iterator->end);
-      gfc_add_block_to_block (pblock, &se.pre);
-      to = gfc_evaluate_now (se.expr, pblock);
-
-      gfc_init_se (&se, NULL);
-      gfc_conv_expr_val (&se, code->ext.iterator->step);
-      gfc_add_block_to_block (pblock, &se.pre);
-      step = gfc_evaluate_now (se.expr, pblock);
-      dovar_decl = dovar;
-
-      /* Special case simple loops.  */
-      if (TREE_CODE (dovar) == VAR_DECL)
-        {
-          if (integer_onep (step))
-            simple = 1;
-          else if (tree_int_cst_equal (step, integer_minus_one_node))
-            simple = -1;
-        }
-      else
-        dovar_decl
-          = trans_acc_variable (code->ext.iterator->var->symtree->n.sym);
-
-      /* Loop body.  */
-      if (simple)
-        {
-          TREE_VEC_ELT (init, i) = build2_v (MODIFY_EXPR, dovar, from);
-          /* The condition should not be folded.  */
-          TREE_VEC_ELT (cond, i) = build2_loc (input_location, simple > 0
-                                               ? LE_EXPR : GE_EXPR,
-                                               boolean_type_node, dovar, to);
-          TREE_VEC_ELT (incr, i) = fold_build2_loc (input_location, PLUS_EXPR,
-                                                    type, dovar, step);
-          TREE_VEC_ELT (incr, i) = fold_build2_loc (input_location,
-                                                    MODIFY_EXPR,
-                                                    type, dovar,
-                                                    TREE_VEC_ELT (incr, i));
-        }
-      else
-        {
-          /* STEP is not 1 or -1.  Use:
-             for (count = 0; count < (to + step - from) / step; count++)
-               {
-                 dovar = from + count * step;
-                 body;
-               cycle_label:;
-               }  */
-          tmp = fold_build2_loc (input_location, MINUS_EXPR, type, step, from);
-          tmp = fold_build2_loc (input_location, PLUS_EXPR, type, to, tmp);
-          tmp = fold_build2_loc (input_location, TRUNC_DIV_EXPR, type, tmp,
-                                 step);
-          tmp = gfc_evaluate_now (tmp, pblock);
-          count = gfc_create_var (type, "count");
-          TREE_VEC_ELT (init, i) = build2_v (MODIFY_EXPR, count,
-                                             build_int_cst (type, 0));
-          /* The condition should not be folded.  */
-          TREE_VEC_ELT (cond, i) = build2_loc (input_location, LT_EXPR,
-                                               boolean_type_node,
-                                               count, tmp);
-          TREE_VEC_ELT (incr, i) = fold_build2_loc (input_location, PLUS_EXPR,
-                                                    type, count,
-                                                    build_int_cst (type, 1));
-          TREE_VEC_ELT (incr, i) = fold_build2_loc (input_location,
-                                                    MODIFY_EXPR, type, count,
-                                                    TREE_VEC_ELT (incr, i));
-
-          /* Initialize DOVAR.  */
-          tmp = fold_build2_loc (input_location, MULT_EXPR, type, count, step);
-          tmp = fold_build2_loc (input_location, PLUS_EXPR, type, from, tmp);
-          dovar_init e = {dovar, tmp};
-          inits.safe_push (e);
-        }
-      gcc_assert (dovar_found != 2);
-
-      if (i + 1 < collapse)
-        code = code->block->next;
-    }
-
-  if (pblock != &block)
-    {
-      pushlevel ();
-      gfc_start_block (&block);
-    }
-
-  gfc_start_block (&body);
-
-  FOR_EACH_VEC_ELT (inits, ix, di)
-    gfc_add_modify (&body, di->var, di->init);
-  inits.release ();
-
-  /* Cycle statement is implemented with a goto.  Exit statement must not be
-     present for this loop.  */
-  cycle_label = gfc_build_label_decl (NULL_TREE);
-
-  /* Put these labels where they can be found later.  */
-
-  code->cycle_label = cycle_label;
-  code->exit_label = NULL_TREE;
-
-  /* Main loop body.  */
-  tmp = trans_acc_code (code->block->next, true);
-  gfc_add_expr_to_block (&body, tmp);
-
-  /* Label for cycle statements (if needed).  */
-  if (TREE_USED (cycle_label))
-    {
-      tmp = build1_v (LABEL_EXPR, cycle_label);
-      gfc_add_expr_to_block (&body, tmp);
-    }
-
-  /* End of loop body.  */
-  stmt = make_node (ACC_LOOP);
+  tree stmt, acc_clauses;
 
-  TREE_TYPE (stmt) = void_type_node;
-  ACC_LOOP_BODY (stmt) = gfc_finish_block (&body);
-  ACC_LOOP_CLAUSES (stmt) = acc_clauses;
-  ACC_LOOP_INIT (stmt) = init;
-  ACC_LOOP_COND (stmt) = cond;
-  ACC_LOOP_INCR (stmt) = incr;
+  gfc_start_block (&block);
+  acc_clauses = trans_acc_clauses (&block, do_clauses,
+                                       code->loc);
+  stmt = trans_acc_code (code->block->next, true);
+  stmt = build2_loc (input_location, ACC_LOOP, void_type_node, stmt,
+                     acc_clauses);
   gfc_add_expr_to_block (&block, stmt);
-
   return gfc_finish_block (&block);
 }
 
@@ -1020,7 +860,7 @@ trans_acc_parallel_loop (gfc_code *code)
     pblock = &block;
   else
     pushlevel ();
-  stmt = trans_acc_loop (code, pblock, &loop_clauses);
+  stmt = trans_acc_loop (code, &loop_clauses);
   if (TREE_CODE (stmt) != BIND_EXPR)
     stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
   else
@@ -1064,7 +904,7 @@ trans_acc_kernels_loop (gfc_code *code)
     pblock = &block;
   else
     pushlevel ();
-  stmt = trans_acc_loop (code, pblock, &loop_clauses);
+  stmt = trans_acc_loop (code, &loop_clauses);
   if (TREE_CODE (stmt) != BIND_EXPR)
     stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
   else
@@ -1103,7 +943,7 @@ gfc_trans_acc_directive (gfc_code *code)
     case EXEC_ACC_HOST_DATA:
       return trans_acc_host_data (code);
     case EXEC_ACC_LOOP:
-      return trans_acc_loop (code, NULL, code->ext.omp_clauses);
+      return trans_acc_loop (code, code->ext.omp_clauses);
     case EXEC_ACC_UPDATE:
       return trans_acc_update (code);
     case EXEC_ACC_WAIT:
diff --git a/gcc/gimple-oacc.c b/gcc/gimple-oacc.c
index 1184ebc5cf0..5044934d5ca 100644
--- a/gcc/gimple-oacc.c
+++ b/gcc/gimple-oacc.c
@@ -486,19 +486,15 @@ gimple_build_acc_host_data (gimple_seq body, tree clauses, tree child_fn,
 }
 
 gimple
-gimple_build_acc_loop (gimple_seq body, tree clauses, size_t collapse,
-                       gimple_seq pre_body)
+gimple_build_acc_loop (gimple_seq body, tree clauses, tree child_fn,
+                       tree data_arg)
 {
   gimple p = gimple_alloc (GIMPLE_ACC_LOOP, 0);
+
   if (body)
     gimple_acc_set_body (p, body);
-  gimple_acc_loop_set_clauses (p, clauses);
-  p->gimple_acc_loop.collapse = collapse;
-  p->gimple_acc_loop.iter
-      = ggc_alloc_cleared_vec_gimple_acc_loop_iter (collapse);
-  if (pre_body)
-    gimple_acc_loop_set_pre_body (p, pre_body);
 
+  gimple_acc_loop_set_clauses (p, clauses);
   return p;
 }
 
diff --git a/gcc/gimple-oacc.h b/gcc/gimple-oacc.h
index 6ed3668c4d4..7c8d2714fc4 100644
--- a/gcc/gimple-oacc.h
+++ b/gcc/gimple-oacc.h
@@ -135,7 +135,7 @@ extern gimple gimple_build_acc_data (gimple_seq, tree, tree, tree);
 extern gimple gimple_build_acc_cache (gimple_seq, tree, tree, tree);
 extern gimple gimple_build_acc_wait (gimple_seq, tree, tree, tree);
 extern gimple gimple_build_acc_host_data (gimple_seq, tree, tree, tree);
-extern gimple gimple_build_acc_loop (gimple_seq, tree, size_t, gimple_seq);
+extern gimple gimple_build_acc_loop (gimple_seq, tree, tree, tree);
 extern gimple gimple_build_acc_declare (gimple_seq, tree, tree, tree);
 extern gimple gimple_build_acc_update (gimple_seq, tree, tree, tree);
 extern gimple gimple_build_acc_compute_region_end ();
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 25a259ff245..495a3c494e4 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7576,152 +7576,28 @@ gimplify_acc_host_data (tree *expr_p, gimple_seq *pre_p)
 static enum gimplify_status
 gimplify_acc_loop (tree *expr_p, gimple_seq *pre_p)
 {
-  tree for_stmt, decl, var, t;
-  enum gimplify_status ret = GS_ALL_DONE;
-  enum gimplify_status tret;
-  gimple gfor;
-  gimple_seq for_body, for_pre_body;
-  int i;
-  bitmap has_decl_expr = NULL;
-
-  for_stmt = *expr_p;
-
-  gimplify_scan_acc_clauses (&ACC_LOOP_CLAUSES (for_stmt), pre_p, ART_LOOP);
-
-  /* Handle ACC_LOOP_INIT.  */
-  for_pre_body = NULL;
-
-  gimplify_and_add (ACC_LOOP_PRE_BODY (for_stmt), &for_pre_body);
-  ACC_LOOP_PRE_BODY (for_stmt) = NULL_TREE;
-
-  for_body = NULL;
-  gcc_assert (TREE_VEC_LENGTH (ACC_LOOP_INIT (for_stmt))
-              == TREE_VEC_LENGTH (ACC_LOOP_COND (for_stmt)));
-  gcc_assert (TREE_VEC_LENGTH (ACC_LOOP_INIT (for_stmt))
-              == TREE_VEC_LENGTH (ACC_LOOP_INCR (for_stmt)));
-  for (i = 0; i < TREE_VEC_LENGTH (ACC_LOOP_INIT (for_stmt)); i++)
-    {
-      t = TREE_VEC_ELT (ACC_LOOP_INIT (for_stmt), i);
-      gcc_assert (TREE_CODE (t) == MODIFY_EXPR);
-      decl = TREE_OPERAND (t, 0);
-      gcc_assert (DECL_P (decl));
-      gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
-                  || POINTER_TYPE_P (TREE_TYPE (decl)));
-
-      acc_add_variable (gimplify_acc_ctxp, decl, GAVD_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
-         modified in the body of the loop.  */
-      if (!is_gimple_reg (decl))
-        {
-          var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
-          TREE_OPERAND (t, 0) = var;
-
-          gimplify_seq_add_stmt (&for_body, gimple_build_assign (decl, var));
-
-          acc_add_variable (gimplify_acc_ctxp, var, GAVD_SEEN);
-        }
-      else
-        var = decl;
-
-      tret = gimplify_expr (&TREE_OPERAND (t, 1), &for_pre_body, NULL,
-                            is_gimple_val, fb_rvalue);
-      ret = MIN (ret, tret);
-      if (ret == GS_ERROR)
-        return ret;
-
-      /* Handle ACC_LOOP_COND.  */
-      t = TREE_VEC_ELT (ACC_LOOP_COND (for_stmt), i);
-      gcc_assert (COMPARISON_CLASS_P (t));
-      gcc_assert (TREE_OPERAND (t, 0) == decl);
-
-      tret = gimplify_expr (&TREE_OPERAND (t, 1), &for_pre_body, NULL,
-                            is_gimple_val, fb_rvalue);
-      ret = MIN (ret, tret);
-
-      /* Handle ACC_LOOP_INCR.  */
-      t = TREE_VEC_ELT (ACC_LOOP_INCR (for_stmt), i);
-      switch (TREE_CODE (t))
-        {
-        case PREINCREMENT_EXPR:
-        case POSTINCREMENT_EXPR:
-          t = build_int_cst (TREE_TYPE (decl), 1);
-          t = build2 (PLUS_EXPR, TREE_TYPE (decl), var, t);
-          t = build2 (MODIFY_EXPR, TREE_TYPE (var), var, t);
-          TREE_VEC_ELT (ACC_LOOP_INCR (for_stmt), i) = t;
-          break;
-
-        case PREDECREMENT_EXPR:
-        case POSTDECREMENT_EXPR:
-          t = build_int_cst (TREE_TYPE (decl), -1);
-          t = build2 (PLUS_EXPR, TREE_TYPE (decl), var, t);
-          t = build2 (MODIFY_EXPR, TREE_TYPE (var), var, t);
-          TREE_VEC_ELT (ACC_LOOP_INCR (for_stmt), i) = t;
-          break;
-
-        case MODIFY_EXPR:
-          gcc_assert (TREE_OPERAND (t, 0) == decl);
-          TREE_OPERAND (t, 0) = var;
-
-          t = TREE_OPERAND (t, 1);
-          switch (TREE_CODE (t))
-            {
-            case PLUS_EXPR:
-              if (TREE_OPERAND (t, 1) == decl)
-                {
-                  TREE_OPERAND (t, 1) = TREE_OPERAND (t, 0);
-                  TREE_OPERAND (t, 0) = var;
-                  break;
-                }
-
-              /* Fallthru.  */
-            case MINUS_EXPR:
-            case POINTER_PLUS_EXPR:
-              gcc_assert (TREE_OPERAND (t, 0) == decl);
-              TREE_OPERAND (t, 0) = var;
-              break;
-            default:
-              gcc_unreachable ();
-            }
-
-          tret = gimplify_expr (&TREE_OPERAND (t, 1), &for_pre_body, NULL,
-                                is_gimple_val, fb_rvalue);
-          ret = MIN (ret, tret);
-          break;
-
-        default:
-          gcc_unreachable ();
-        }
-    }
+  tree expr = *expr_p;
+  gimple g;
+  gimple_seq body = NULL;
+  struct gimplify_ctx gctx;
 
-  BITMAP_FREE (has_decl_expr);
+  gimplify_scan_acc_clauses (&ACC_LOOP_CLAUSES (expr), pre_p, ART_LOOP);
 
-  gimplify_and_add (ACC_LOOP_BODY (for_stmt), &for_body);
+  push_gimplify_context (&gctx);
 
-  gimplify_adjust_acc_clauses (&ACC_LOOP_CLAUSES (for_stmt));
+  g = gimplify_and_return_first (ACC_LOOP_BODY (expr), &body);
+  if (gimple_code (g) == GIMPLE_BIND)
+    pop_gimplify_context (g);
+  else
+    pop_gimplify_context (NULL);
 
-  gfor = gimple_build_acc_loop (for_body, ACC_LOOP_CLAUSES (for_stmt),
-                                TREE_VEC_LENGTH (ACC_LOOP_INIT (for_stmt)),
-                                for_pre_body);
+  gimplify_adjust_acc_clauses (&ACC_LOOP_CLAUSES (expr));
 
-  for (i = 0; i < TREE_VEC_LENGTH (ACC_LOOP_INIT (for_stmt)); i++)
-    {
-      t = TREE_VEC_ELT (ACC_LOOP_INIT (for_stmt), i);
-      gimple_acc_loop_set_index (gfor, i, TREE_OPERAND (t, 0));
-      gimple_acc_loop_set_initial (gfor, i, TREE_OPERAND (t, 1));
-      t = TREE_VEC_ELT (ACC_LOOP_COND (for_stmt), i);
-      gimple_acc_loop_set_cond (gfor, i, TREE_CODE (t));
-      gimple_acc_loop_set_final (gfor, i, TREE_OPERAND (t, 1));
-      t = TREE_VEC_ELT (ACC_LOOP_INCR (for_stmt), i);
-      gimple_acc_loop_set_incr (gfor, i, TREE_OPERAND (t, 1));
-    }
+  g = gimple_build_acc_loop (body, ACC_LOOP_CLAUSES (expr),
+                             NULL_TREE, NULL_TREE);
 
-  gimplify_seq_add_stmt (pre_p, gfor);
-  if (ret != GS_ALL_DONE)
-    return GS_ERROR;
+  gimplify_seq_add_stmt (pre_p, g);
   *expr_p = NULL_TREE;
-  return GS_ALL_DONE;
 }
 
 /* Gimplify the contents of an ACC_CACHE statement. */
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 0848bb15b9e..bdc8c279962 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -2625,54 +2625,7 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
     case ACC_LOOP:
       pp_string (buffer, "#pragma acc loop");
       dump_acc_clauses (buffer, ACC_LOOP_CLAUSES(node), spc, flags);
-      if (!(flags & TDF_SLIM))
-        {
-          int i;
-
-          if (ACC_LOOP_PRE_BODY (node))
-            {
-              newline_and_indent (buffer, spc + 2);
-              pp_character (buffer, '{');
-              spc += 4;
-              newline_and_indent (buffer, spc);
-              dump_generic_node (buffer, ACC_LOOP_PRE_BODY (node),
-                  spc, flags, false);
-            }
-          spc -= 2;
-          for (i = 0; i < TREE_VEC_LENGTH (ACC_LOOP_INIT (node)); i++)
-            {
-              spc += 2;
-              newline_and_indent (buffer, spc);
-              pp_string (buffer, "for (");
-              dump_generic_node (buffer, TREE_VEC_ELT (ACC_LOOP_INIT (node), i),
-                                 spc, flags, false);
-              pp_string (buffer, "; ");
-              dump_generic_node (buffer, TREE_VEC_ELT (ACC_LOOP_COND (node), i),
-                                 spc, flags, false);
-              pp_string (buffer, "; ");
-              dump_generic_node (buffer, TREE_VEC_ELT (ACC_LOOP_INCR (node), i),
-                                 spc, flags, false);
-              pp_string (buffer, ")");
-            }
-          if (ACC_LOOP_BODY (node))
-            {
-              newline_and_indent (buffer, spc + 2);
-              pp_character (buffer, '{');
-              newline_and_indent (buffer, spc + 4);
-              dump_generic_node (buffer, ACC_LOOP_BODY (node), spc + 4, flags,
-                  false);
-              newline_and_indent (buffer, spc + 2);
-              pp_character (buffer, '}');
-            }
-          spc -= 2 * TREE_VEC_LENGTH (ACC_LOOP_INIT (node)) - 2;
-          if (ACC_LOOP_PRE_BODY (node))
-            {
-              spc -= 4;
-              newline_and_indent (buffer, spc + 2);
-              pp_character (buffer, '}');
-            }
-        }
-      is_expr = false;
+      is_expr = dump_acc_body(flags, node, spc, is_expr, buffer);
 
       break;
     case ACC_HOST_DATA:
diff --git a/gcc/tree.def b/gcc/tree.def
index c7ad92b4e5a..5c2c2a3fd89 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1017,14 +1017,11 @@ DEFTREECODE (ACC_KERNELS, "acc_kernels", tcc_statement, 2)
 DEFTREECODE (ACC_DATA, "acc_data", tcc_statement, 2)
 
 /* #pragma acc loop */
+/* Consider #pragma acc loop as construct */
 /* Operand 0: BODY
    Operand 1: CLAUSES
-   Operand 2: INIT
-   Operand 3: COND
-   Operand 4: INCR
-   Operand 5: PRE_BODY
 */
-DEFTREECODE (ACC_LOOP, "acc_loop", tcc_statement, 6)
+DEFTREECODE (ACC_LOOP, "acc_loop", tcc_statement, 2)
 
 /* #pragma acc host_data */
 /* Operand 0: BODY
diff --git a/gcc/tree.h b/gcc/tree.h
index 691a01a05ba..5fb3a78e658 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1195,10 +1195,6 @@ extern void protected_set_expr_location (tree, location_t);
 
 #define ACC_LOOP_BODY(NODE)          TREE_OPERAND (ACC_LOOP_CHECK (NODE), 0)
 #define ACC_LOOP_CLAUSES(NODE)       TREE_OPERAND (ACC_LOOP_CHECK (NODE), 1)
-#define ACC_LOOP_INIT(NODE)          TREE_OPERAND (ACC_LOOP_CHECK (NODE), 2)
-#define ACC_LOOP_COND(NODE)          TREE_OPERAND (ACC_LOOP_CHECK (NODE), 3)
-#define ACC_LOOP_INCR(NODE)          TREE_OPERAND (ACC_LOOP_CHECK (NODE), 4)
-#define ACC_LOOP_PRE_BODY(NODE)      TREE_OPERAND (ACC_LOOP_CHECK (NODE), 5)
 
 /* OpenACC clauses */
 #define ACC_CLAUSE_NUM_GANGS_EXPR(NODE)          ACC_CLAUSE_DECL (NODE)