[PATCH 5/7] openmp: Add C/C++ support for "omp tile"

[Frederik Harwath <frederik@codesourcery.com>]

This commit adds the C and C++ front end support for the "omp tile"
directive.

gcc/c-family/ChangeLog:

        * c-omp.cc (c_omp_directives): Add PRAGMA_OMP_TILE.
        * c-pragma.cc (omp_pragmas_simd): Likewise.
        * c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_TILE.
        (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_TILE

gcc/c/ChangeLog:

        * c-parser.cc (c_parser_nested_omp_unroll_clauses): Rename and
        generalize ...
        (c_parser_omp_nested_loop_transform_clauses): ... to this.
        (c_parser_omp_for_loop): Handle "omp tile" parsing in loop nests.
        (c_parser_omp_tile_sizes): Parse single "sizes" clause.
        (c_parser_omp_loop_transform_clause): New function.
        (c_parser_omp_tile): New function for parsing "omp tile"
        (c_parser_omp_unroll): Adjust to renaming.
        (c_parser_omp_construct): Handle PRAGMA_OMP_TILE.

gcc/cp/ChangeLog:

        * parser.cc (cp_parser_omp_clause_unroll_partial): Adjust.
        (cp_parser_nested_omp_unroll_clauses): Rename ...
        (cp_parser_omp_nested_loop_transform_clauses): ... to this.
        (cp_parser_omp_for_loop): Handle "omp tile" parsing in loop nests.
        (cp_parser_omp_tile_sizes): New function, parses single "sizes" clause
        (cp_parser_omp_tile): New function for parsing "omp tile".
        (cp_parser_omp_loop_transform_clause): New  function.
        (cp_parser_omp_unroll): Adjust to renaming.
        (cp_parser_omp_construct): Handle PRAGMA_OMP_TILE.
        (cp_parser_pragma): Likewise.
        * pt.cc (tsubst_omp_clauses): Handle OMP_CLAUSE_TILE.
        * semantics.cc (finish_omp_clauses): Likewise.

gcc/ChangeLog:

        * gimplify.cc (omp_for_drop_tile_clauses): New function, ...
        (gimplify_omp_for): ... used here.

libgomp/ChangeLog:

        * testsuite/libgomp.c++/loop-transforms/tile-1.C: New test.
        * testsuite/libgomp.c++/loop-transforms/tile-2.C: New test.
        * testsuite/libgomp.c++/loop-transforms/tile-3.C: New test.

gcc/testsuite/ChangeLog:

        * c-c++-common/gomp/loop-transforms/tile-1.c: New test.
        * c-c++-common/gomp/loop-transforms/tile-2.c: New test.
        * c-c++-common/gomp/loop-transforms/tile-3.c: New test.
        * c-c++-common/gomp/loop-transforms/tile-4.c: New test.
        * c-c++-common/gomp/loop-transforms/tile-5.c: New test.
        * c-c++-common/gomp/loop-transforms/tile-6.c: New test.
        * c-c++-common/gomp/loop-transforms/tile-7.c: New test.
        * c-c++-common/gomp/loop-transforms/tile-8.c: New test.
        * c-c++-common/gomp/loop-transforms/unroll-2.c: Adapt.
        * g++.dg/gomp/loop-transforms/tile-1.h: New test.
        * g++.dg/gomp/loop-transforms/tile-1a.C: New test.
        * g++.dg/gomp/loop-transforms/tile-1b.C: New test.
---
 gcc/c-family/c-omp.cc                         |   4 +-
 gcc/c-family/c-pragma.cc                      |   1 +
 gcc/c-family/c-pragma.h                       |   2 +
 gcc/c/c-parser.cc                             | 277 ++++++++++++---
 gcc/cp/parser.cc                              | 289 +++++++++++++---
 gcc/cp/pt.cc                                  |   1 +
 gcc/cp/semantics.cc                           |  40 +++
 gcc/gimplify.cc                               |  28 ++
 .../gomp/loop-transforms/tile-1.c             | 164 +++++++++
 .../gomp/loop-transforms/tile-2.c             | 183 ++++++++++
 .../gomp/loop-transforms/tile-3.c             | 117 +++++++
 .../gomp/loop-transforms/tile-4.c             | 322 ++++++++++++++++++
 .../gomp/loop-transforms/tile-5.c             | 150 ++++++++
 .../gomp/loop-transforms/tile-6.c             |  34 ++
 .../gomp/loop-transforms/tile-7.c             |  31 ++
 .../gomp/loop-transforms/tile-8.c             |  40 +++
 .../gomp/loop-transforms/unroll-2.c           |  12 +-
 .../g++.dg/gomp/loop-transforms/tile-1.h      |  27 ++
 .../g++.dg/gomp/loop-transforms/tile-1a.C     |  27 ++
 .../g++.dg/gomp/loop-transforms/tile-1b.C     |  27 ++
 .../libgomp.c++/loop-transforms/tile-1.C      |  52 +++
 .../libgomp.c++/loop-transforms/tile-2.C      |  69 ++++
 .../libgomp.c++/loop-transforms/tile-3.C      |  28 ++
 23 files changed, 1823 insertions(+), 102 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-1.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-2.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-3.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-4.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-5.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-6.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-7.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-8.c
 create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1.h
 create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1a.C
 create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1b.C
 create mode 100644 libgomp/testsuite/libgomp.c++/loop-transforms/tile-1.C
 create mode 100644 libgomp/testsuite/libgomp.c++/loop-transforms/tile-2.C
 create mode 100644 libgomp/testsuite/libgomp.c++/loop-transforms/tile-3.C

diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index fec7f337772..2ab7faea2cc 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3207,8 +3207,8 @@ const struct c_omp_directive c_omp_directives[] = {
     C_OMP_DIR_STANDALONE, false },
   { "taskyield", nullptr, nullptr, PRAGMA_OMP_TASKYIELD,
     C_OMP_DIR_STANDALONE, false },
-  /* { "tile", nullptr, nullptr, PRAGMA_OMP_TILE,
-    C_OMP_DIR_CONSTRUCT, false },  */
+  { "tile", nullptr, nullptr, PRAGMA_OMP_TILE,
+    C_OMP_DIR_CONSTRUCT, false },
   { "teams", nullptr, nullptr, PRAGMA_OMP_TEAMS,
     C_OMP_DIR_CONSTRUCT, true },
   { "threadprivate", nullptr, nullptr, PRAGMA_OMP_THREADPRIVATE,
diff --git a/gcc/c-family/c-pragma.cc b/gcc/c-family/c-pragma.cc
index 96a28ac1b0c..75d5cabbafd 100644
--- a/gcc/c-family/c-pragma.cc
+++ b/gcc/c-family/c-pragma.cc
@@ -1593,6 +1593,7 @@ static const struct omp_pragma_def omp_pragmas_simd[] = {
   { "target", PRAGMA_OMP_TARGET },
   { "taskloop", PRAGMA_OMP_TASKLOOP },
   { "teams", PRAGMA_OMP_TEAMS },
+  { "tile", PRAGMA_OMP_TILE },
   { "unroll", PRAGMA_OMP_UNROLL },
 };

diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 6686abdc94d..c0476f74441 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -81,6 +81,7 @@ enum pragma_kind {
   PRAGMA_OMP_TASKYIELD,
   PRAGMA_OMP_THREADPRIVATE,
   PRAGMA_OMP_TEAMS,
+  PRAGMA_OMP_TILE,
   PRAGMA_OMP_UNROLL,
   /* PRAGMA_OMP__LAST_ should be equal to the last PRAGMA_OMP_* code.  */
   PRAGMA_OMP__LAST_ = PRAGMA_OMP_UNROLL,
@@ -157,6 +158,7 @@ enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_TASKGROUP,
   PRAGMA_OMP_CLAUSE_THREAD_LIMIT,
   PRAGMA_OMP_CLAUSE_THREADS,
+  PRAGMA_OMP_CLAUSE_TILE,
   PRAGMA_OMP_CLAUSE_TO,
   PRAGMA_OMP_CLAUSE_UNIFORM,
   PRAGMA_OMP_CLAUSE_UNTIED,
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index e7c9da99552..aac23dec9c0 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -20243,7 +20243,8 @@ c_parser_omp_scan_loop_body (c_parser *parser, bool open_brace_parsed)
                             "expected %<}%>");
 }

-static bool c_parser_nested_omp_unroll_clauses (c_parser *, tree &);
+static int c_parser_omp_nested_loop_transform_clauses (c_parser *, tree &, int,
+                                                      const char *);

 /* Parse the restricted form of loop statements allowed by OpenACC and OpenMP.
    The real trick here is to determine the loop control variable early
@@ -20263,16 +20264,18 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
   bool fail = false, open_brace_parsed = false;
   int i, collapse = 1, ordered = 0, count, nbraces = 0;
   location_t for_loc;
-  bool tiling = false;
+  bool oacc_tiling = false;
   bool inscan = false;
   vec<tree, va_gc> *for_block = make_tree_vector ();

   for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
     if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE)
-      collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl));
+      {
+       collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl));
+      }
     else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_OACC_TILE)
       {
-       tiling = true;
+       oacc_tiling = true;
        collapse = list_length (OMP_CLAUSE_OACC_TILE_LIST (cl));
       }
     else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_ORDERED
@@ -20295,21 +20298,31 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
       ordered = collapse;
     }

-  gcc_assert (tiling || (collapse >= 1 && ordered >= 0));
+  c_parser_omp_nested_loop_transform_clauses (parser, clauses, collapse,
+                                             "loop collapse");
+
+  /* Find the depth of the loop nest affected by "omp tile"
+     directives. There can be several such directives, but the tiling
+     depth of the outer ones may not be larger than the depth of the
+     innermost directive. */
+  int omp_tile_depth = 0;
+  for (tree c = clauses; c; c = TREE_CHAIN (c))
+    {
+      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TILE)
+       continue;
+
+      omp_tile_depth = list_length (OMP_CLAUSE_TILE_SIZES (c));
+    }
+
+  gcc_assert (oacc_tiling || (collapse >= 1 && ordered >= 0));
   count = ordered ? ordered : collapse;
+  count = MAX (count, omp_tile_depth);

   declv = make_tree_vec (count);
   initv = make_tree_vec (count);
   condv = make_tree_vec (count);
   incrv = make_tree_vec (count);

-  if (c_parser_nested_omp_unroll_clauses (parser, clauses)
-      && count > 1)
-    {
-      error_at (loc, "collapse cannot be larger than 1 on an unrolled loop");
-      return NULL;
-    }
-
   if (!c_parser_next_token_is_keyword (parser, RID_FOR))
     {
       c_parser_error (parser, "for statement expected");
@@ -23945,47 +23958,224 @@ c_parser_omp_taskloop (location_t loc, c_parser *parser,
        ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PARTIAL)      \
          | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FULL) )

-/* Parse zero or more '#pragma omp unroll' that follow
-   another directive that requires a canonical loop nest. */
+/* OpenMP 5.1: Parse sizes list for "omp tile sizes"
+   sizes ( size-expr-list ) */
+static tree
+c_parser_omp_tile_sizes (c_parser *parser, location_t loc)
+{
+  tree sizes = NULL_TREE;

-static bool
-c_parser_nested_omp_unroll_clauses (c_parser *parser, tree &clauses)
+  c_token *tok = c_parser_peek_token (parser);
+  if (tok->type != CPP_NAME
+      || strcmp ("sizes", IDENTIFIER_POINTER (tok->value)))
+    {
+      c_parser_error (parser, "expected %<sizes%>");
+      return error_mark_node;
+    }
+  c_parser_consume_token (parser);
+
+  if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    return error_mark_node;
+
+  do
+    {
+      if (sizes && !c_parser_require (parser, CPP_COMMA, "expected %<,%>"))
+       return error_mark_node;
+
+      location_t expr_loc = c_parser_peek_token (parser)->location;
+      c_expr cexpr = c_parser_expr_no_commas (parser, NULL);
+      cexpr = convert_lvalue_to_rvalue (expr_loc, cexpr, false, true);
+      tree expr = cexpr.value;
+
+      if (expr == error_mark_node)
+       {
+         c_parser_skip_until_found (parser, CPP_CLOSE_PAREN,
+                                    "expected %<)%>");
+         return error_mark_node;
+       }
+
+      expr = c_fully_fold (expr, false, NULL);
+
+      if (!INTEGRAL_TYPE_P (TREE_TYPE (expr)) || !tree_fits_shwi_p (expr)
+         || tree_to_shwi (expr) <= 0)
+       {
+         c_parser_error (parser, "%<tile sizes%> argument needs positive"
+                                 " integral constant");
+         expr = integer_zero_node;
+       }
+
+      sizes = tree_cons (NULL_TREE, expr, sizes);
+    }
+  while (c_parser_next_token_is_not (parser, CPP_CLOSE_PAREN));
+  c_parser_consume_token (parser);
+
+  gcc_assert (sizes);
+  tree c  = build_omp_clause (loc, OMP_CLAUSE_TILE);
+  OMP_CLAUSE_TILE_SIZES (c) = sizes;
+
+  return c;
+}
+
+/* Parse a single OpenMP loop transformation directive and return the
+   clause that is used internally to represent the directive. */
+
+static tree
+c_parser_omp_loop_transform_clause (c_parser *parser)
 {
-  static const char *p_name = "#pragma omp unroll";
-  c_token *tok;
-  bool found_unroll = false;
-  while (c_parser_next_token_is (parser, CPP_PRAGMA)
-        && (tok = c_parser_peek_token (parser),
-            tok->pragma_kind == PRAGMA_OMP_UNROLL))
+  c_token *tok = c_parser_peek_token (parser);
+  if (tok->type != CPP_PRAGMA)
+    return NULL_TREE;
+
+  tree c;
+  switch (tok->pragma_kind)
     {
+    case PRAGMA_OMP_UNROLL:
       c_parser_consume_pragma (parser);
-      tree c = c_parser_omp_all_clauses (parser, OMP_UNROLL_CLAUSE_MASK,
-                                        p_name, true);
-      if (c)
+      c = c_parser_omp_all_clauses (parser, OMP_UNROLL_CLAUSE_MASK,
+                                   "#pragma omp unroll", false, true);
+      if (!c)
        {
-         gcc_assert (!TREE_CHAIN (c));
-         found_unroll = true;
-         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_UNROLL_FULL)
-           {
-             error_at (tok->location, "%<full%> clause is invalid here; "
-                       "turns loop into non-loop");
-             continue;
-           }
+         if (c_parser_next_token_is (parser, CPP_PRAGMA_EOL))
+           c = build_omp_clause (tok->location, OMP_CLAUSE_UNROLL_NONE);
+         else
+           c = error_mark_node;
        }
-      else
+      c_parser_skip_to_pragma_eol (parser);
+      break;
+
+    case PRAGMA_OMP_TILE:
+      c_parser_consume_pragma (parser);
+      c = c_parser_omp_tile_sizes (parser, tok->location);
+      c_parser_skip_to_pragma_eol (parser);
+      break;
+
+    default:
+      c = NULL_TREE;
+      break;
+    }
+
+  gcc_assert (!c || !TREE_CHAIN (c));
+  return c;
+}
+
+/* Parse zero or more OpenMP loop transformation directives that
+   follow another directive that requires a canonical loop nest and
+   append all to CLAUSES.  Return the nesting depth
+   of the transformed loop nest.
+
+   REQUIRED_DEPTH is the nesting depth of the loop nest required by
+   the preceding directive.  OUTER_DESCR is a description of the
+   language construct that requires the loop nest depth (e.g. "loop
+   collpase", "outer transformation") that is used for error
+   messages. */
+
+static int
+c_parser_omp_nested_loop_transform_clauses (c_parser *parser, tree &clauses,
+                                           int required_depth,
+                                           const char *outer_descr)
+{
+  tree c = NULL_TREE;
+  tree last_c = tree_last (clauses);
+
+  /* The depth of the loop nest, counting from LEVEL, after the
+     transformations. That is, the nesting depth left by the outermost
+     transformation which is the first to be parsed, but the last to be
+     executed. */
+  int transformed_depth = 0;
+
+  /* The minimum nesting depth required by the last parsed transformation. */
+  int last_depth = required_depth;
+  while ((c = c_parser_omp_loop_transform_clause (parser)))
+    {
+      /* The nesting depth left after the current transformation */
+      int depth = 1;
+      if (TREE_CODE (c) == ERROR_MARK)
+       goto error;
+
+      gcc_assert (!TREE_CHAIN (c));
+      switch (OMP_CLAUSE_CODE (c))
        {
-         error_at (tok->location, "%<#pragma omp unroll%> without "
-                                  "%<partial%> clause is invalid here; "
-                                  "turns loop into non-loop");
-         continue;
+       case OMP_CLAUSE_UNROLL_FULL:
+         error_at (OMP_CLAUSE_LOCATION (c),
+                   "%<full%> clause is invalid here; "
+                   "turns loop into non-loop");
+         goto error;
+       case OMP_CLAUSE_UNROLL_NONE:
+         error_at (OMP_CLAUSE_LOCATION (c),
+                   "%<#pragma omp unroll%> without "
+                   "%<partial%> clause is invalid here; "
+                   "turns loop into non-loop");
+         goto error;
+       case OMP_CLAUSE_UNROLL_PARTIAL:
+         depth = 1;
+         break;
+       case OMP_CLAUSE_TILE:
+         depth = list_length (OMP_CLAUSE_TILE_SIZES (c));
+         break;
+       default:
+         gcc_unreachable ();
+       }
+
+      if (depth < last_depth)
+       {
+         bool is_outermost_clause = !transformed_depth;
+         error_at (OMP_CLAUSE_LOCATION (c),
+                   "nesting depth left after this transformation too low "
+                   "for %s",
+                   is_outermost_clause ? outer_descr
+                                       : "outer transformation");
+         goto error;
        }

-      clauses = chainon (clauses, c);
+      last_depth = depth;
+
+      if (!transformed_depth)
+       transformed_depth = last_depth;
+
+      if (!clauses)
+       clauses = c;
+      else if (last_c)
+       TREE_CHAIN (last_c) = c;
+
+      last_c = c;
     }

-  return found_unroll;
+  return transformed_depth;
+
+error:
+  while (c_parser_omp_loop_transform_clause (parser))
+    ;
+  clauses = NULL_TREE;
+  return -1;
 }

+/* OpenMP 5.1:
+   tile sizes ( size-expr-list ) */
+
+static tree
+c_parser_omp_tile (location_t loc, c_parser *parser, bool *if_p)
+{
+  tree block;
+  tree ret = error_mark_node;
+
+  tree clauses = c_parser_omp_tile_sizes (parser, loc);
+  c_parser_skip_to_pragma_eol (parser);
+
+  if (!clauses || clauses == error_mark_node)
+    return error_mark_node;
+
+  int required_depth = list_length (OMP_CLAUSE_TILE_SIZES (clauses));
+  c_parser_omp_nested_loop_transform_clauses (parser, clauses, required_depth,
+                                             "outer transformation");
+
+  block = c_begin_compound_stmt (true);
+  ret = c_parser_omp_for_loop (loc, parser, OMP_LOOP_TRANS, clauses, NULL, if_p);
+  block = c_end_compound_stmt (loc, block, true);
+  add_stmt (block);
+
+  return ret;
+ }
+
 static tree
 c_parser_omp_unroll (location_t loc, c_parser *parser, bool *if_p)
 {
@@ -23994,7 +24184,9 @@ c_parser_omp_unroll (location_t loc, c_parser *parser, bool *if_p)
   omp_clause_mask mask = OMP_UNROLL_CLAUSE_MASK;

   tree clauses = c_parser_omp_all_clauses (parser, mask, p_name, false);
-  c_parser_nested_omp_unroll_clauses (parser, clauses);
+  int required_depth = 1;
+  c_parser_omp_nested_loop_transform_clauses (parser, clauses, required_depth,
+                                             "outer transformation");

   if (!clauses)
     {
@@ -24496,6 +24688,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p)
     case PRAGMA_OMP_ASSUME:
       c_parser_omp_assume (parser, if_p);
       return;
+    case PRAGMA_OMP_TILE:
+      stmt = c_parser_omp_tile (loc, parser, if_p);
+      break;
     case PRAGMA_OMP_UNROLL:
       stmt = c_parser_omp_unroll (loc, parser, if_p);
       break;
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 90af40c4dbc..084ecd3ada5 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -43631,7 +43631,8 @@ cp_parser_omp_scan_loop_body (cp_parser *parser)
   braces.require_close (parser);
 }

-static bool cp_parser_nested_omp_unroll_clauses (cp_parser *, tree &);
+static int cp_parser_omp_nested_loop_transform_clauses (cp_parser *, tree &,
+                                                       int, const char *);

 /* Parse the restricted form of the for statement allowed by OpenMP.  */

@@ -43643,20 +43644,20 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tree_code code, tree clauses,
   tree orig_decl;
   tree real_decl, initv, condv, incrv, declv, orig_declv;
   tree this_pre_body, cl, ordered_cl = NULL_TREE;
-  location_t loc_first;
   bool collapse_err = false;
   int i, collapse = 1, ordered = 0, count, nbraces = 0;
   releasing_vec for_block;
   auto_vec<tree, 4> orig_inits;
-  bool tiling = false;
+  bool oacc_tiling = false;
   bool inscan = false;
+  location_t loc_first = cp_lexer_peek_token (parser->lexer)->location;

   for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
     if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE)
       collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl));
     else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_OACC_TILE)
       {
-       tiling = true;
+       oacc_tiling = true;
        collapse = list_length (OMP_CLAUSE_OACC_TILE_LIST (cl));
       }
     else if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_ORDERED
@@ -43679,26 +43680,33 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tree_code code, tree clauses,
       ordered = collapse;
     }

-  gcc_assert (tiling || (collapse >= 1 && ordered >= 0));
+
+  gcc_assert (oacc_tiling || (collapse >= 1 && ordered >= 0));
   count = ordered ? ordered : collapse;

+  cp_parser_omp_nested_loop_transform_clauses (parser, clauses, count,
+                                              "loop collapse");
+
+  /* Find the depth of the loop nest affected by "omp tile"
+     directives. There can be several such directives, but the tiling
+     depth of the outer ones may not be larger than the depth of the
+     innermost directive. */
+  int omp_tile_depth = 0;
+  for (tree c = clauses; c; c = TREE_CHAIN (c))
+    {
+      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TILE)
+       continue;
+
+      omp_tile_depth = list_length (OMP_CLAUSE_TILE_SIZES (c));
+    }
+  count = MAX (count, omp_tile_depth);
+
   declv = make_tree_vec (count);
   initv = make_tree_vec (count);
   condv = make_tree_vec (count);
   incrv = make_tree_vec (count);
   orig_declv = NULL_TREE;

-  loc_first = cp_lexer_peek_token (parser->lexer)->location;
-
-  if (cp_parser_nested_omp_unroll_clauses (parser, clauses)
-      && count > 1)
-    {
-      error_at (loc_first,
-               "collapse cannot be larger than 1 on an unrolled loop");
-      return NULL;
-    }
-
-
   for (i = 0; i < count; i++)
     {
       int bracecount = 0;
@@ -45734,51 +45742,224 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
   return true;
 }

+/* OpenMP 5.1: Parse sizes list for "omp tile sizes"
+   sizes ( size-expr-list ) */
+static tree
+cp_parser_omp_tile_sizes (cp_parser *parser, location_t loc)
+{
+  tree sizes = NULL_TREE;
+  cp_lexer *lexer = parser->lexer;
+
+  cp_token *tok = cp_lexer_peek_token (lexer);
+  if (tok->type != CPP_NAME
+      || strcmp ("sizes", IDENTIFIER_POINTER (tok->u.value)))
+    {
+      cp_parser_error (parser, "expected %<sizes%>");
+      return error_mark_node;
+    }
+  cp_lexer_consume_token (lexer);
+
+  if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+    return error_mark_node;
+
+  do
+    {
+      if (sizes && !cp_parser_require (parser, CPP_COMMA, RT_COMMA))
+       return error_mark_node;
+
+      tree expr = cp_parser_constant_expression (parser);
+      if (expr == error_mark_node)
+       {
+         cp_parser_skip_to_closing_parenthesis (parser,
+                                                /*recovering=*/true,
+                                                /*or_comma=*/false,
+                                                /*consume_paren=*/
+                                                true);
+         return error_mark_node;
+       }
+
+      sizes = tree_cons (NULL_TREE, expr, sizes);
+    }
+  while (cp_lexer_next_token_is_not (lexer, CPP_CLOSE_PAREN));
+  cp_lexer_consume_token (lexer);
+
+  gcc_assert (sizes);
+  tree c  = build_omp_clause (loc, OMP_CLAUSE_TILE);
+  OMP_CLAUSE_TILE_SIZES (c) = sizes;
+
+  return c;
+}
+
+/* OpenMP 5.1:
+   tile sizes ( size-expr-list ) */
+
+static tree
+cp_parser_omp_tile (cp_parser *parser, cp_token *tok, bool *if_p)
+{
+  tree block;
+  tree ret = error_mark_node;
+
+  tree clauses = cp_parser_omp_tile_sizes (parser, tok->location);
+  cp_parser_require_pragma_eol (parser, tok);
+
+  if (!clauses || clauses == error_mark_node)
+    return error_mark_node;
+
+  int required_depth = list_length (OMP_CLAUSE_TILE_SIZES (clauses));
+  cp_parser_omp_nested_loop_transform_clauses (
+      parser, clauses, required_depth, "outer transformation");
+
+  block = begin_omp_structured_block ();
+  clauses = finish_omp_clauses (clauses, C_ORT_OMP);
+
+  ret = cp_parser_omp_for_loop (parser, OMP_LOOP_TRANS, clauses, NULL, if_p);
+  block = finish_omp_structured_block (block);
+  add_stmt (block);
+
+  return ret;
+}
+
 #define OMP_UNROLL_CLAUSE_MASK                                 \
        ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PARTIAL)      \
          | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FULL) )

-/* Parse zero or more '#pragma omp unroll' that follow
-   another directive that requires a canonical loop nest. */
+/* Parse a single OpenMP loop transformation directive and return the
+   clause that is used internally to represent the directive. */

-static bool
-cp_parser_nested_omp_unroll_clauses (cp_parser *parser, tree &clauses)
+static tree
+cp_parser_omp_loop_transform_clause (cp_parser *parser)
 {
-  static const char *p_name = "#pragma omp unroll";
-  cp_token *tok;
-  bool unroll_found = false;
-  while (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA)
-        && (tok = cp_lexer_peek_token (parser->lexer),
-            cp_parser_pragma_kind (tok) == PRAGMA_OMP_UNROLL))
+  cp_lexer *lexer = parser->lexer;
+  cp_token *tok = cp_lexer_peek_token (lexer);
+  if (tok->type != CPP_PRAGMA)
+    return NULL_TREE;
+
+  tree c;
+  switch (cp_parser_pragma_kind (tok))
     {
-      cp_lexer_consume_token (parser->lexer);
-      gcc_assert (tok->type == CPP_PRAGMA);
-      parser->lexer->in_pragma = true;
-      tree c = cp_parser_omp_all_clauses (parser, OMP_UNROLL_CLAUSE_MASK,
-                                         p_name, tok);
-      if (c)
-       {
-         gcc_assert (!TREE_CHAIN (c));
-         unroll_found = true;
-         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_UNROLL_FULL)
-           {
-             error_at (tok->location, "%<full%> clause is invalid here; "
-                       "turns loop into non-loop");
-             continue;
-           }
+    case PRAGMA_OMP_UNROLL:
+      cp_lexer_consume_token (lexer);
+      lexer->in_pragma = true;
+      c = cp_parser_omp_all_clauses (parser, OMP_UNROLL_CLAUSE_MASK,
+                                    "#pragma omp unroll", tok,
+                                    false, true);
+      if (!c)
+       {
+         if (cp_lexer_next_token_is (lexer, CPP_PRAGMA_EOL))
+           c = build_omp_clause (tok->location, OMP_CLAUSE_UNROLL_NONE);
+         else
+           c = error_mark_node;
+       }
+      cp_parser_skip_to_pragma_eol (parser, tok);
+      break;

-         c = finish_omp_clauses (c, C_ORT_OMP);
+    case PRAGMA_OMP_TILE:
+      cp_lexer_consume_token (lexer);
+      lexer->in_pragma = true;
+      c = cp_parser_omp_tile_sizes (parser, tok->location);
+      cp_parser_require_pragma_eol (parser, tok);
+      break;
+
+    default:
+      c = NULL_TREE;
+      break;
+    }
+
+  gcc_assert (!c || !TREE_CHAIN (c));
+  return c;
+}
+
+/* Parse zero or more OpenMP loop transformation directives that
+   follow another directive that requires a canonical loop nest and
+   append all to CLAUSES.  Return the nesting depth
+   of the transformed loop nest.
+
+   REQUIRED_DEPTH is the nesting depth of the loop nest required by
+   the preceding directive.  OUTER_DESCR is a description of the
+   language construct that requires the loop nest depth (e.g. "loop
+   collpase", "outer transformation") that is used for error
+   messages. */
+
+static int
+cp_parser_omp_nested_loop_transform_clauses (cp_parser *parser, tree &clauses,
+                                            int required_depth,
+                                            const char *outer_descr)
+{
+  tree c = NULL_TREE;
+  tree last_c = tree_last (clauses);
+
+  /* The depth of the loop nest after the transformations. That is,
+     the nesting depth left by the outermost transformation which is
+     the first to be parsed, but the last to be executed. */
+  int transformed_depth = 0;
+
+  /* The minimum nesting depth required by the last parsed transformation. */
+  int last_depth = required_depth;
+
+  while ((c = cp_parser_omp_loop_transform_clause (parser)))
+    {
+      /* The nesting depth left after the current transformation */
+      int depth = 1;
+      if (TREE_CODE (c) == ERROR_MARK)
+       goto error;
+
+      gcc_assert (!TREE_CHAIN (c));
+      switch (OMP_CLAUSE_CODE (c))
+       {
+       case OMP_CLAUSE_UNROLL_FULL:
+         error_at (OMP_CLAUSE_LOCATION (c),
+                   "%<full%> clause is invalid here; "
+                   "turns loop into non-loop");
+         goto error;
+       case OMP_CLAUSE_UNROLL_NONE:
+         error_at (OMP_CLAUSE_LOCATION (c),
+                   "%<#pragma omp unroll%> without "
+                   "%<partial%> clause is invalid here; "
+                   "turns loop into non-loop");
+         goto error;
+       case OMP_CLAUSE_UNROLL_PARTIAL:
+         depth = 1;
+         break;
+       case OMP_CLAUSE_TILE:
+         depth = list_length (OMP_CLAUSE_TILE_SIZES (c));
+         break;
+       default:
+         gcc_unreachable ();
        }
-      else
+
+      if (depth < last_depth)
        {
-         error_at (tok->location, "%<#pragma omp unroll%> without "
-                                  "%<partial%> clause is invalid here; "
-                                  "turns loop into non-loop");
-         continue;
+         bool is_outermost_clause = !transformed_depth;
+         error_at (OMP_CLAUSE_LOCATION (c),
+                   "nesting depth left after this transformation too low "
+                   "for %s",
+                   is_outermost_clause ? outer_descr
+                                       : "outer transformation");
+         goto error;
        }
-      clauses = chainon (clauses, c);
+
+      last_depth = depth;
+
+      if (!transformed_depth)
+       transformed_depth = last_depth;
+
+      c = finish_omp_clauses (c, C_ORT_OMP);
+
+      if (!clauses)
+       clauses = c;
+      else if (last_c)
+       TREE_CHAIN (last_c) = c;
+
+      last_c = c;
     }
-  return unroll_found;
+
+  return transformed_depth;
+
+error:
+  while (cp_parser_omp_loop_transform_clause (parser))
+    ;
+  clauses = NULL_TREE;
+  return -1;
 }

 static tree
@@ -45788,7 +45969,7 @@ cp_parser_omp_unroll (cp_parser *parser, cp_token *tok, bool *if_p)
   static const char *p_name = "#pragma omp unroll";
   omp_clause_mask mask = OMP_UNROLL_CLAUSE_MASK;

-  tree clauses = cp_parser_omp_all_clauses (parser, mask, p_name, tok, false);
+  tree clauses = cp_parser_omp_all_clauses (parser, mask, p_name, tok, true);

   if (!clauses)
     {
@@ -45797,7 +45978,9 @@ cp_parser_omp_unroll (cp_parser *parser, cp_token *tok, bool *if_p)
       clauses = c;
     }

-  cp_parser_nested_omp_unroll_clauses (parser, clauses);
+  int required_depth = 1;
+  cp_parser_omp_nested_loop_transform_clauses (
+      parser, clauses, required_depth, "outer transformation");

   block = begin_omp_structured_block ();
   ret = cp_parser_omp_for_loop (parser, OMP_LOOP_TRANS, clauses, NULL, if_p);
@@ -48900,6 +49083,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
     case PRAGMA_OMP_ASSUME:
       cp_parser_omp_assume (parser, pragma_tok, if_p);
       return;
+    case PRAGMA_OMP_TILE:
+      stmt = cp_parser_omp_tile (parser, pragma_tok, if_p);
+      break;
     case PRAGMA_OMP_UNROLL:
       stmt = cp_parser_omp_unroll (parser, pragma_tok, if_p);
       break;
@@ -49529,6 +49715,7 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p)
       cp_parser_omp_construct (parser, pragma_tok, if_p);
       pop_omp_privatization_clauses (stmt);
       return true;
+    case PRAGMA_OMP_TILE:
     case PRAGMA_OMP_UNROLL:
       if (context != pragma_stmt && context != pragma_compound)
        goto bad_stmt;
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index 16197b17e5a..a9d36d66caf 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -18087,6 +18087,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
        case OMP_CLAUSE_WAIT:
        case OMP_CLAUSE_DETACH:
        case OMP_CLAUSE_UNROLL_PARTIAL:
+       case OMP_CLAUSE_TILE:
          OMP_CLAUSE_OPERAND (nc, 0)
            = tsubst_expr (OMP_CLAUSE_OPERAND (oc, 0), args, complain, in_decl);
          break;
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index c87e252ff06..15f7c7e6dc4 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8769,6 +8769,46 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
            }
          break;

+       case OMP_CLAUSE_TILE:
+         for (tree list = OMP_CLAUSE_TILE_SIZES (c); !remove && list;
+              list = TREE_CHAIN (list))
+           {
+             t = TREE_VALUE (list);
+
+             if (t == error_mark_node)
+               remove = true;
+             else if (!type_dependent_expression_p (t)
+                      && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+               {
+                 error_at (OMP_CLAUSE_LOCATION (c),
+                           "%<tile sizes%> argument needs integral type");
+                 remove = true;
+               }
+             else
+               {
+                 t = mark_rvalue_use (t);
+                 if (!processing_template_decl)
+                   {
+                     t = maybe_constant_value (t);
+                     int n;
+                     if (!tree_fits_shwi_p (t)
+                         || !INTEGRAL_TYPE_P (TREE_TYPE (t))
+                         || (n = tree_to_shwi (t)) <= 0 || (int)n != n)
+                       {
+                         error_at (OMP_CLAUSE_LOCATION (c),
+                                   "%<tile sizes%> argument needs positive "
+                                   "integral constant");
+                         remove = true;
+                       }
+                     t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+                   }
+               }
+
+             /* Update list item.  */
+             TREE_VALUE (list) = t;
+           }
+         break;
+
        case OMP_CLAUSE_ORDERED:
          ordered_seen = true;
          break;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 4d504a12451..365897afb61 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -13572,6 +13572,29 @@ find_standalone_omp_ordered (tree *tp, int *walk_subtrees, void *)
   return NULL_TREE;
 }

+static void omp_for_drop_tile_clauses (tree for_stmt)
+{
+  /* Drop erroneous loop transformation clauses to avoid follow up errors
+     in pass-omp_transform_loops. */
+  tree last_c = NULL_TREE;
+  for (tree c = OMP_FOR_CLAUSES (for_stmt); c;
+       c = OMP_CLAUSE_CHAIN (c))
+    {
+
+      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TILE)
+       continue;
+
+      if (last_c)
+       TREE_CHAIN (last_c) = TREE_CHAIN (c);
+      else
+       OMP_FOR_CLAUSES (for_stmt) = TREE_CHAIN (c);
+
+      error_at (OMP_CLAUSE_LOCATION (c),
+               "'tile' loop transformation may not appear on "
+               "non-rectangular for");
+    }
+}
+
 /* Gimplify the gross structure of an OMP_FOR statement.  */

 static enum gimplify_status
@@ -13763,6 +13786,8 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
     case OMP_FOR:
       if (OMP_FOR_NON_RECTANGULAR (inner_for_stmt ? inner_for_stmt : for_stmt))
        {
+         omp_for_drop_tile_clauses (for_stmt);
+
          if (omp_find_clause (OMP_FOR_CLAUSES (for_stmt),
                               OMP_CLAUSE_SCHEDULE))
            error_at (EXPR_LOCATION (for_stmt),
@@ -13808,6 +13833,8 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       ort = ORT_SIMD;
       break;
     case OMP_LOOP_TRANS:
+      if (OMP_FOR_NON_RECTANGULAR (inner_for_stmt ? inner_for_stmt : for_stmt))
+       omp_for_drop_tile_clauses (for_stmt);
       break;
     default:
       gcc_unreachable ();
@@ -14693,6 +14720,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
          case OMP_CLAUSE_UNROLL_FULL:
          case OMP_CLAUSE_UNROLL_NONE:
          case OMP_CLAUSE_UNROLL_PARTIAL:
+         case OMP_CLAUSE_TILE:
            *gfor_clauses_ptr = c;
            gfor_clauses_ptr = &OMP_CLAUSE_CHAIN (c);
            break;
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-1.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-1.c
new file mode 100644
index 00000000000..8a2f2126af4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-1.c
@@ -0,0 +1,164 @@
+extern void dummy (int);
+
+void
+test ()
+{
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(0) /* { dg-error {'tile sizes' argument needs positive integral constant} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(-1) /* { dg-error {'tile sizes' argument needs positive integral constant} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes() /* { dg-error {expected expression before} "" { target c} } */
+    /* { dg-error {expected primary-expression before} "" { target c++ } .-1 } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(,) /* { dg-error {expected expression before} "" { target c } } */
+    /* { dg-error {expected primary-expression before} "" { target c++ } .-1 } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(1,2 /* { dg-error {expected '\,' before end of line} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes /* { dg-error {expected '\(' before end of line} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(1) sizes(1) /* { dg-error {expected end of line before 'sizes'} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(1)
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(1, 2)
+    #pragma omp tile sizes(1) /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp tile sizes(1, 2)
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp tile sizes(5, 6)
+    #pragma omp tile sizes(1, 2, 3)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+    for (int k = 0; k < 100; ++k)
+       dummy (i);
+
+    #pragma omp tile sizes(1)
+    #pragma omp unroll partia /* { dg-error {expected '#pragma omp' clause before 'partia'} } */
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(1)
+    #pragma omp unroll /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(1)
+    #pragma omp unroll full /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(1)
+    #pragma omp unroll partial
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(8,8)
+    #pragma omp unroll partial /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(8,8)
+    #pragma omp unroll partial /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp tile sizes(1, 2) /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = i; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp tile sizes(1, 2) /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = 2; j < i; ++j)
+       dummy (i);
+
+    #pragma omp tile sizes(1, 2, 3)
+    for (int i = 0; i < 100; ++i)
+      for (int j = 0; j < 100; ++j)
+        dummy (i); /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } } */
+    /* { dg-error {not enough for loops to collapse} "" { target c++ } .-1 } */
+    /* { dg-error {'i' was not declared in this scope} "" { target c++ } .-2 } */
+
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+      {
+       dummy (i);
+        for (int j = 0; j < 100; ++j)
+          dummy (i);
+      }
+
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+      {
+        for (int j = 0; j < 100; ++j)
+           dummy (j);
+       dummy (i);
+      }
+
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+      {
+        dummy (i); /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } } */
+       /* { dg-error {not enough for loops to collapse} "" { target c++ } .-1 } */
+        for (int j = 0; j < 100; ++j)
+          dummy (j);
+      }
+
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+      {
+        for (int j = 0; j < 100; ++j)
+         dummy (j);
+       dummy (i); /* { dg-error {collapsed loops not perfectly nested before 'dummy'} "" { target c} } */
+       /* { dg-error {collapsed loops not perfectly nested} "" { target c++ } .-1 } */
+      }
+
+    int s;
+    #pragma omp tile sizes(s) /* { dg-error {'tile sizes' argument needs positive integral constant} "" { target { ! c++98_only } } } */
+    /* { dg-error {the value of 's' is not usable in a constant expression} "" { target { c++ && { ! c++98_only } } } .-1 } */
+    /* { dg-error {'s' cannot appear in a constant-expression} "" { target c++98_only } .-2 } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp tile sizes(42.0) /* { dg-error {'tile sizes' argument needs positive integral constant} "" { target c } } */
+    /* { dg-error {'tile sizes' argument needs integral type} "" { target c++ } .-1 } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-2.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-2.c
new file mode 100644
index 00000000000..51d62552945
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-2.c
@@ -0,0 +1,183 @@
+extern void dummy (int);
+
+void
+test ()
+{
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(0) /* { dg-error {'tile sizes' argument needs positive integral constant} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(-1) /* { dg-error {'tile sizes' argument needs positive integral constant} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes() /* { dg-error {expected expression before} "" { target c} } */
+    /* { dg-error {expected primary-expression before} "" { target c++ } .-1 } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(,) /* { dg-error {expected expression before} "" { target c } } */
+    /* { dg-error {expected primary-expression before} "" { target c++ } .-1 } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1,2 /* { dg-error {expected '\,' before end of line} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes /* { dg-error {expected '\(' before end of line} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1) sizes(1) /* { dg-error {expected end of line before 'sizes'} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1, 2)
+    #pragma omp tile sizes(1) /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1, 2)
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(5, 6)
+    #pragma omp tile sizes(1, 2, 3)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+    for (int k = 0; k < 100; ++k)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    #pragma omp unroll partia /* { dg-error {expected '#pragma omp' clause before 'partia'} } */
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    #pragma omp unroll /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    #pragma omp unroll full /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    #pragma omp unroll partial
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(8,8)
+    #pragma omp unroll partial /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(8,8)
+    #pragma omp unroll partial /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1, 2) /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = i; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1, 2) /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = 2; j < i; ++j)
+       dummy (i);
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1, 2, 3)
+    for (int i = 0; i < 100; ++i)
+      for (int j = 0; j < 100; ++j)
+        dummy (i); /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } } */
+    /* { dg-error {not enough for loops to collapse} "" { target c++ } .-1 } */
+    /* { dg-error {'i' was not declared in this scope} "" { target c++ } .-2 } */
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+      {
+       dummy (i);
+        for (int j = 0; j < 100; ++j)
+          dummy (i);
+      }
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+      {
+        for (int j = 0; j < 100; ++j)
+           dummy (j);
+       dummy (i);
+      }
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+      {
+        dummy (i); /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } } */
+       /* { dg-error {not enough for loops to collapse} "" { target c++ } .-1 } */
+        for (int j = 0; j < 100; ++j)
+          dummy (j);
+      }
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+      {
+        for (int j = 0; j < 100; ++j)
+         dummy (j);
+       dummy (i); /* { dg-error {collapsed loops not perfectly nested before 'dummy'} "" { target c} } */
+       /* { dg-error {collapsed loops not perfectly nested} "" { target c++ } .-1 } */
+      }
+
+    #pragma omp parallel for
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-3.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-3.c
new file mode 100644
index 00000000000..7fffc72b335
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-3.c
@@ -0,0 +1,117 @@
+extern void dummy (int);
+
+void
+test ()
+{
+    #pragma omp for
+    #pragma omp tile sizes(1, 2) /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = i; j < 100; ++j)
+       dummy (i);
+
+    #pragma omp for
+    #pragma omp tile sizes(1, 2) /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < i; ++j)
+       dummy (i);
+
+
+#pragma omp for collapse(1)
+    #pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+#pragma omp for collapse(2)
+    #pragma omp tile sizes(1) /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+#pragma omp for collapse(2)
+    #pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+#pragma omp for collapse(3)
+    #pragma omp tile sizes(1, 2) /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+    /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } .-1 } */
+    /* { dg-error {not enough for loops to collapse} "" { target c++ } .-2 } */
+    /* { dg-error {'i' was not declared in this scope} "" { target c++ } .-3 } */
+
+#pragma omp for collapse(1)
+#pragma omp tile sizes(1)
+#pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+#pragma omp for collapse(2)
+#pragma omp tile sizes(1, 2)
+#pragma omp tile sizes(1) /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+    for (int i = 0; i < 100; ++i)
+       dummy (i); /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } } */
+    /* { dg-error {not enough for loops to collapse} "" { target c++ } .-1 } */
+
+#pragma omp for collapse(2)
+#pragma omp tile sizes(1, 2)
+#pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+       dummy (i); /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } } */
+    /* { dg-error {not enough for loops to collapse} "" { target c++ } .-1 } */
+
+#pragma omp for collapse(2)
+#pragma omp tile sizes(5, 6)
+#pragma omp tile sizes(1, 2, 3)
+    for (int i = 0; i < 100; ++i)
+       dummy (i); /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } } */
+    /* { dg-error {not enough for loops to collapse} "" { target c++ } .-1 } */
+
+
+#pragma omp for collapse(1)
+#pragma omp tile sizes(1)
+#pragma omp tile sizes(1)
+    for (int i = 0; i < 100; ++i)
+       dummy (i);
+
+#pragma omp for collapse(2)
+#pragma omp tile sizes(1, 2)
+#pragma omp tile sizes(1) /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+#pragma omp for collapse(2)
+#pragma omp tile sizes(1, 2)
+#pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i);
+
+#pragma omp for collapse(2)
+#pragma omp tile sizes(5, 6)
+#pragma omp tile sizes(1, 2, 3)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+    for (int k = 0; k < 100; ++k)
+       dummy (i);
+
+#pragma omp for collapse(3)
+#pragma omp tile sizes(1, 2) /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+#pragma omp tile sizes(1, 2)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+       dummy (i); /* { dg-error {not enough perfectly nested loops before 'dummy'} "" { target c } } */
+    /* { dg-error {not enough for loops to collapse} "" { target c++ } .-1 } */
+
+#pragma omp for collapse(3)
+#pragma omp tile sizes(5, 6) /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+#pragma omp tile sizes(1, 2, 3)
+    for (int i = 0; i < 100; ++i)
+    for (int j = 0; j < 100; ++j)
+    for (int k = 0; k < 100; ++k)
+       dummy (i);
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-4.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-4.c
new file mode 100644
index 00000000000..d46bb0cb642
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-4.c
@@ -0,0 +1,322 @@
+/* { dg-do run } */
+/* { dg-options "-O0 -fopenmp-simd" } */
+
+#include <stdio.h>
+
+#define ASSERT_EQ(var, val) if (var != val) { fprintf (stderr, "%s:%d: Unexpected value %d, expected %d\n", __FILE__, __LINE__, var, val); \
+    __builtin_abort (); }
+
+int
+test1 ()
+{
+  int iter = 0;
+  int i;
+#pragma omp tile sizes(3)
+  for (i = 0; i < 10; i=i+2)
+       {
+         ASSERT_EQ (i, iter)
+         iter = iter + 2;
+       }
+
+  ASSERT_EQ (i, 10)
+  return iter;
+}
+
+int
+test2 ()
+{
+  int iter = 0;
+  int i;
+#pragma omp tile sizes(3)
+  for (i = 0; i < 10; i=i+2)
+       {
+         ASSERT_EQ (i, iter)
+         iter = iter + 2;
+       }
+
+  ASSERT_EQ (i, 10)
+  return iter;
+}
+
+int
+test3 ()
+{
+  int iter = 0;
+  int i;
+#pragma omp tile sizes(8)
+  for (i = 0; i < 10; i=i+2)
+       {
+         ASSERT_EQ (i, iter)
+         iter = iter + 2;
+       }
+
+  ASSERT_EQ (i, 10)
+  return iter;
+}
+
+int
+test4 ()
+{
+  int iter = 10;
+  int i;
+#pragma omp tile sizes(8)
+  for (i = 10; i > 0; i=i-2)
+       {
+         ASSERT_EQ (i, iter)
+         iter = iter - 2;
+       }
+  ASSERT_EQ (i, 0)
+  return iter;
+}
+
+int
+test5 ()
+{
+  int iter = 10;
+  int i;
+#pragma omp tile sizes(71)
+  for (i = 10; i > 0; i=i-2)
+       {
+         ASSERT_EQ (i, iter)
+         iter = iter - 2;
+       }
+
+  ASSERT_EQ (i, 0)
+  return iter;
+}
+
+int
+test6 ()
+{
+  int iter = 10;
+  int i;
+#pragma omp tile sizes(1)
+  for (i = 10; i > 0; i=i-2)
+       {
+         ASSERT_EQ (i, iter)
+         iter = iter - 2;
+       }
+  ASSERT_EQ (i, 0)
+  return iter;
+}
+
+int
+test7 ()
+{
+  int iter = 5;
+  int i;
+#pragma omp tile sizes(2)
+  for (i = 5; i < -5; i=i-3)
+       {
+         fprintf (stderr, "%d\n", i);
+         __builtin_abort ();
+         iter = iter - 3;
+       }
+
+  ASSERT_EQ (i, 5)
+
+  /* No iteration expected */
+  return iter;
+}
+
+int
+test8 ()
+{
+  int iter = 5;
+  int i;
+#pragma omp tile sizes(2)
+  for (i = 5; i > -5; i=i-3)
+       {
+         ASSERT_EQ (i, iter)
+         /* Expect only first iteration of the last tile to execute */
+         if (iter != -4)
+           iter = iter - 3;
+       }
+
+  ASSERT_EQ (i, -7)
+  return iter;
+}
+
+
+int
+test9 ()
+{
+  int iter = 5;
+  int i;
+#pragma omp tile sizes(5)
+  for (i = 5; i >= -5; i=i-4)
+       {
+         ASSERT_EQ (i, iter)
+         /* Expect only first iteration of the last tile to execute */
+         if (iter != - 3)
+           iter = iter - 4;
+       }
+
+  ASSERT_EQ (i, -7)
+  return iter;
+}
+
+int
+test10 ()
+{
+  int iter = 5;
+  int i;
+#pragma omp tile sizes(5)
+  for (i = 5; i >= -5; i--)
+       {
+         ASSERT_EQ (i, iter)
+         iter--;
+       }
+
+  ASSERT_EQ (i, -6)
+  return iter;
+}
+
+int
+test11 ()
+{
+  int iter = 5;
+  int i;
+#pragma omp tile sizes(15)
+  for (i = 5; i != -5; i--)
+       {
+         ASSERT_EQ (i, iter)
+         iter--;
+       }
+  ASSERT_EQ (i, -5)
+  return iter;
+}
+
+int
+test12 ()
+{
+  int iter = 0;
+  unsigned i;
+#pragma omp tile sizes(3)
+  for (i = 0; i != 5; i++)
+       {
+         ASSERT_EQ (i, iter)
+         iter++;
+       }
+
+  ASSERT_EQ (i, 5)
+  return iter;
+}
+
+int
+test13 ()
+{
+  int iter = -5;
+  long long unsigned int i;
+#pragma omp tile sizes(15)
+  for (int i = -5; i < 5; i=i+3)
+       {
+         ASSERT_EQ (i, iter)
+         iter++;
+       }
+
+  ASSERT_EQ (i, 5)
+  return iter;
+}
+
+int
+test14 (unsigned init, int step)
+{
+  int iter = init;
+  long long unsigned int i;
+#pragma omp tile sizes(8)
+  for (i = init; i < 2*init; i=i+step)
+    iter++;
+
+  ASSERT_EQ (i, 2*init)
+  return iter;
+}
+
+int
+test15 (unsigned init, int step)
+{
+  int iter = init;
+  int i;
+#pragma omp tile sizes(8)
+  for (unsigned i = init; i > 2* init; i=i+step)
+    iter++;
+
+  return iter;
+}
+
+int
+main ()
+{
+  int last_iter;
+
+  last_iter = test1 ();
+  ASSERT_EQ (last_iter, 10);
+
+  last_iter = test2 ();
+  ASSERT_EQ (last_iter, 10);
+
+  last_iter = test3 ();
+  ASSERT_EQ (last_iter, 10);
+
+  last_iter = test4 ();
+  ASSERT_EQ (last_iter, 0);
+
+  last_iter = test5 ();
+  ASSERT_EQ (last_iter, 0);
+
+  last_iter = test6 ();
+  ASSERT_EQ (last_iter, 0);
+
+  last_iter = test7 ();
+  ASSERT_EQ (last_iter, 5);
+
+  last_iter = test8 ();
+  ASSERT_EQ (last_iter, -4);
+
+  last_iter = test9 ();
+  ASSERT_EQ (last_iter, -3);
+
+  last_iter = test10 ();
+  ASSERT_EQ (last_iter, -6);
+  return 0;
+
+  last_iter = test11 ();
+  ASSERT_EQ (last_iter, -4);
+  return 0;
+
+  last_iter = test12 ();
+  ASSERT_EQ (last_iter, 5);
+  return 0;
+
+  last_iter = test13 ();
+  ASSERT_EQ (last_iter, 4);
+  return 0;
+
+  last_iter = test14 (0, 1);
+  ASSERT_EQ (last_iter, 0);
+  return 0;
+
+  last_iter = test14 (0, -1);
+  ASSERT_EQ (last_iter, 0);
+  return 0;
+
+  last_iter = test14 (8, 2);
+  ASSERT_EQ (last_iter, 16);
+  return 0;
+
+  last_iter = test14 (5, 3);
+  ASSERT_EQ (last_iter, 9);
+  return 0;
+
+  last_iter = test15 (8, -1);
+  ASSERT_EQ (last_iter, 9);
+  return 0;
+
+  last_iter = test15 (8, -2);
+  ASSERT_EQ (last_iter, 10);
+  return 0;
+
+  last_iter = test15 (5, -3);
+  ASSERT_EQ (last_iter, 6);
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-5.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-5.c
new file mode 100644
index 00000000000..815318ab27a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-5.c
@@ -0,0 +1,150 @@
+/* { dg-do run } */
+/* { dg-options "-O0 -fopenmp-simd" } */
+
+#include <stdio.h>
+
+#define ASSERT_EQ(var, val) if (var != val) { fprintf (stderr, "%s:%d: Unexpected value %d\n", __FILE__, __LINE__, var); \
+    __builtin_abort (); }
+
+#define ASSERT_EQ_PTR(var, ptr) if (var != ptr) { fprintf (stderr, "%s:%d: Unexpected value %p\n", __FILE__, __LINE__, var); \
+    __builtin_abort (); }
+
+int
+test1 (int data[10])
+{
+  int iter = 0;
+  int *i;
+  #pragma omp tile sizes(5)
+  for (i = data; i < data + 10 ; i++)
+       {
+         ASSERT_EQ (*i, data[iter]);
+         ASSERT_EQ_PTR (i, data + iter);
+         iter++;
+       }
+
+  ASSERT_EQ_PTR (i, data + 10)
+  return iter;
+}
+
+int
+test2 (int data[10])
+{
+  int iter = 0;
+  int *i;
+  #pragma omp tile sizes(5)
+  for (i = data; i < data + 10 ; i=i+2)
+       {
+         ASSERT_EQ_PTR (i, data + 2 * iter);
+         ASSERT_EQ (*i, data[2 * iter]);
+         iter++;
+       }
+
+  ASSERT_EQ_PTR (i, data + 10)
+  return iter;
+}
+
+int
+test3 (int data[10])
+{
+  int iter = 0;
+  int *i;
+  #pragma omp tile sizes(5)
+  for (i = data; i <= data + 9 ; i=i+2)
+       {
+         ASSERT_EQ (*i, data[2 * iter]);
+         iter++;
+       }
+
+  ASSERT_EQ_PTR (i, data + 10)
+  return iter;
+}
+
+int
+test4 (int data[10])
+{
+  int iter = 0;
+  int *i;
+  #pragma omp tile sizes(5)
+  for (i = data; i != data + 10 ; i=i+1)
+       {
+         ASSERT_EQ (*i, data[iter]);
+         iter++;
+       }
+
+  ASSERT_EQ_PTR (i, data + 10)
+  return iter;
+}
+
+int
+test5 (int data[10])
+{
+  int iter = 0;
+  int *i;
+  #pragma omp tile sizes(3)
+  for (i = data + 9; i >= data ; i--)
+       {
+         ASSERT_EQ (*i, data[9 - iter]);
+         iter++;
+       }
+
+  ASSERT_EQ_PTR (i, data - 1)
+  return iter;
+}
+
+int
+test6 (int data[10])
+{
+  int iter = 0;
+  int *i;
+  #pragma omp tile sizes(3)
+  for (i = data + 9; i > data - 1 ; i--)
+       {
+         ASSERT_EQ (*i, data[9 - iter]);
+         iter++;
+       }
+
+  ASSERT_EQ_PTR (i, data - 1)
+  return iter;
+}
+
+int
+test7 (int data[10])
+{
+  int iter = 0;
+  #pragma omp tile sizes(1)
+  for (int *i = data + 9; i != data - 1 ; i--)
+       {
+         ASSERT_EQ (*i, data[9 - iter]);
+         iter++;
+       }
+
+  return iter;
+}
+
+int
+main ()
+{
+  int iter_count;
+  int data[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
+
+  iter_count = test1 (data);
+  ASSERT_EQ (iter_count, 10);
+
+  iter_count = test2 (data);
+  ASSERT_EQ (iter_count, 5);
+
+  iter_count = test3 (data);
+  ASSERT_EQ (iter_count, 5);
+
+  iter_count = test4 (data);
+  ASSERT_EQ (iter_count, 10);
+
+  iter_count = test5 (data);
+  ASSERT_EQ (iter_count, 10);
+
+  iter_count = test6 (data);
+  ASSERT_EQ (iter_count, 10);
+
+  iter_count = test7 (data);
+  ASSERT_EQ (iter_count, 10);
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-6.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-6.c
new file mode 100644
index 00000000000..8132128a5a8
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-6.c
@@ -0,0 +1,34 @@
+/* { dg-do run } */
+/* { dg-options "-O0 -fopenmp-simd" } */
+
+#include <stdio.h>
+
+int
+test1 ()
+{
+   int sum = 0;
+for (int k = 0; k < 10; k++)
+  {
+#pragma omp tile sizes(5,7)
+  for (int i = 0; i < 10; i++)
+  for (int j = 0; j < 10; j=j+2)
+       {
+         sum = sum + 1;
+       }
+  }
+
+  return sum;
+}
+
+int
+main ()
+{
+  int result = test1 ();
+
+  if (result != 500)
+    {
+      fprintf (stderr, "Wrong result: %d\n", result);
+    __builtin_abort ();
+    }
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-7.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-7.c
new file mode 100644
index 00000000000..cd25a62c5c0
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-7.c
@@ -0,0 +1,31 @@
+/* { dg-do run } */
+/* { dg-options "-O0 -fopenmp-simd" } */
+
+#include <stdio.h>
+#define ASSERT_EQ(var, val) if (var != val) { fprintf (stderr, "%s:%d: Unexpected value %d\n", __FILE__, __LINE__, var); \
+    __builtin_abort (); }
+
+#define ASSERT_EQ_PTR(var, ptr) if (var != ptr) { fprintf (stderr, "%s:%d: Unexpected value %p\n", __FILE__, __LINE__, var); \
+    __builtin_abort (); }
+
+int
+main ()
+{
+  int iter_count;
+  int data[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
+
+  int iter = 0;
+  int *i;
+  #pragma omp tile sizes(1)
+  for (i = data; i < data + 10; i=i+2)
+       {
+         ASSERT_EQ_PTR (i, data + 2 * iter);
+         ASSERT_EQ (*i, data[2 * iter]);
+         iter++;
+       }
+
+  unsigned long real_iter_count = ((unsigned long)i - (unsigned long)data) / (sizeof (int) * 2);
+  ASSERT_EQ (real_iter_count, 5);
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-8.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-8.c
new file mode 100644
index 00000000000..c26e03d7e74
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/tile-8.c
@@ -0,0 +1,40 @@
+/* { dg-do run } */
+/* { dg-options "-O0 -fopenmp-simd" } */
+
+#include <stdio.h>
+
+#define ASSERT_EQ(var, val) if (var != val) { fprintf (stderr, "%s:%d: Unexpected value %d, expected %d\n", __FILE__, __LINE__, var, val); \
+    __builtin_abort (); }
+
+int
+main ()
+{
+  int iter_j = 0, iter_k = 0;
+  unsigned i, j, k;
+#pragma omp tile sizes(3,5,8)
+  for (i = 0; i < 2; i=i+2)
+  for (j = 0; j < 3; j=j+1)
+  for (k = 0; k < 5; k=k+3)
+       {
+         /* fprintf (stderr, "i=%d j=%d k=%d\n", i, j, k);
+          * fprintf (stderr, "iter_j=%d iter_k=%d\n", iter_j, iter_k); */
+         ASSERT_EQ (i, 0);
+         if (k == 0)
+           {
+             ASSERT_EQ (j, iter_j);
+             iter_k = 0;
+           }
+
+         ASSERT_EQ (k, iter_k);
+
+         iter_k = iter_k + 3;
+         if (k == 3)
+           iter_j++;
+       }
+
+  ASSERT_EQ (i, 2);
+  ASSERT_EQ (j, 3);
+  ASSERT_EQ (k, 6);
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/loop-transforms/unroll-2.c b/gcc/testsuite/c-c++-common/gomp/loop-transforms/unroll-2.c
index 8f7c3088a2e..e4fee72c04d 100644
--- a/gcc/testsuite/c-c++-common/gomp/loop-transforms/unroll-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/loop-transforms/unroll-2.c
@@ -19,7 +19,7 @@ test ()

 #pragma omp for
 #pragma omp unroll full /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
-#pragma omp unroll full /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
+#pragma omp unroll full
   for (int i = -300; i != 100; ++i)
     dummy (i);

@@ -45,13 +45,11 @@ test ()
   int i;
 #pragma omp for
 #pragma omp unroll( /* { dg-error {expected '#pragma omp' clause before '\(' token} } */
-  /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} "" { target *-*-* } .-1 } */
   for (int i = -300; i != 100; ++i)
     dummy (i);

 #pragma omp for
 #pragma omp unroll foo /* { dg-error {expected '#pragma omp' clause before 'foo'} } */
-  /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} "" { target *-*-* } .-1 } */
   for (int i = -300; i != 100; ++i)
     dummy (i);

@@ -67,7 +65,7 @@ test ()

 #pragma omp unroll partial(i)
  /* { dg-error {the value of 'i' is not usable in a constant expression} "" { target c++ } .-1 } */
- /* { dg-error {partial argument needs positive constant integer expression} "" { target c } .-2 } */
+ /* { dg-error {partial argument needs positive constant integer expression} "" { target *-*-* } .-2 } */
   for (int i = -300; i != 100; ++i)
     dummy (i);

@@ -78,20 +76,18 @@ test ()
 #pragma omp for
 #pragma omp unroll partial(1)
 #pragma omp unroll parti /* { dg-error {expected '#pragma omp' clause before 'parti'} } */
-  /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} "" { target *-*-* } .-1 } */
   for (int i = -300; i != 100; ++i)
     dummy (i);

 #pragma omp for
 #pragma omp unroll partial(1)
 #pragma omp unroll parti /* { dg-error {expected '#pragma omp' clause before 'parti'} } */
-  /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} "" { target *-*-* } .-1 } */
   for (int i = -300; i != 100; ++i)
     dummy (i);

 int sum = 0;
-#pragma omp parallel for reduction(+ : sum) collapse(2) /* { dg-error {collapse cannot be larger than 1 on an unrolled loop} "" { target c } } */
-#pragma omp unroll partial(1) /* { dg-error {collapse cannot be larger than 1 on an unrolled loop} "" { target c++ } } */
+#pragma omp parallel for reduction(+ : sum) collapse(2)
+#pragma omp unroll partial(1) /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
   for (int i = 3; i < 10; ++i)
     for (int j = -2; j < 7; ++j)
       sum++;
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1.h b/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1.h
new file mode 100644
index 00000000000..166d1d48677
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1.h
@@ -0,0 +1,27 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++11" }
+
+#include <vector>
+
+extern void dummy (int);
+
+template<class T, int U, unsigned V> void
+test1_template ()
+{
+  std::vector<int> v;
+
+  for (unsigned i = 0; i < 10; i++)
+    v.push_back (i);
+
+#pragma omp for
+  for (int i : v)
+    dummy (i);
+
+#pragma omp tile sizes (U, 10, V)
+  for (T i : v)
+  for (T j : v)
+  for (T k : v)
+    dummy (i);
+}
+
+void test () { test1_template <long, 5, 3> (); };
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1a.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1a.C
new file mode 100644
index 00000000000..1ee76da3d4a
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1a.C
@@ -0,0 +1,27 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++11" }
+
+#include <vector>
+
+extern void dummy (int);
+
+template<class T, int U, unsigned V> void
+test1_template ()
+{
+  std::vector<int> v;
+
+  for (unsigned i = 0; i < 10; i++)
+    v.push_back (i);
+
+#pragma omp teams distribute parallel for num_teams(V)
+  for (int i : v)
+    dummy (i);
+
+#pragma omp tile sizes (V, U)
+  for (T i : v)
+  for (T j : v)
+  for (T k : v)
+    dummy (i);
+}
+
+void test () { test1_template <long, 5, 3> (); };
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1b.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1b.C
new file mode 100644
index 00000000000..263c9b301c6
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/tile-1b.C
@@ -0,0 +1,27 @@
+// { dg-do compile }
+// { dg-additional-options "-std=c++11" }
+
+#include <vector>
+
+extern void dummy (int);
+
+template<class T, int U, unsigned V> void
+test1_template ()
+{
+  std::vector<int> v;
+
+  for (unsigned i = 0; i < 10; i++)
+    v.push_back (i);
+
+#pragma omp for
+  for (int i : v)
+    dummy (i);
+
+#pragma omp tile sizes (U, 10, V) // { dg-error {'tile sizes' argument needs positive integral constant} }
+  for (T i : v)
+  for (T j : v)
+  for (T k : v)
+    dummy (i);
+}
+
+void test () { test1_template <long, 5, 0> (); };
diff --git a/libgomp/testsuite/libgomp.c++/loop-transforms/tile-1.C b/libgomp/testsuite/libgomp.c++/loop-transforms/tile-1.C
new file mode 100644
index 00000000000..2a4d760720d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/loop-transforms/tile-1.C
@@ -0,0 +1,52 @@
+#include <string.h>
+#include <stdio.h>
+#include <math.h>
+
+void
+mult (float *matrix1, float *matrix2, float *result, unsigned dim0,
+      unsigned dim1)
+{
+  memset (result, 0, sizeof (float) * dim0 * dim1);
+#pragma omp target parallel for collapse(3) map(tofrom:result[0:dim0*dim1]) map(to:matrix1[0:dim0*dim1], matrix2[0:dim0*dim1])
+#pragma omp tile sizes(8, 16, 4)
+  for (unsigned i = 0; i < dim0; i++)
+    for (unsigned j = 0; j < dim1; j++)
+      for (unsigned k = 0; k < dim1; k++)
+       result[i * dim1 + j] += matrix1[i * dim1 + k] * matrix2[k * dim0 + j];
+}
+
+int
+main ()
+{
+  unsigned dim0 = 20;
+  unsigned dim1 = 20;
+
+  float *result = (float *)malloc (sizeof (float) * dim0 * dim1);
+  float *matrix1 = (float *)malloc (sizeof (float) * dim0 * dim1);
+  float *matrix2 = (float *)malloc (sizeof (float) * dim0 * dim1);
+
+  for (unsigned i = 0; i < dim0; i++)
+    for (unsigned j = 0; j < dim1; j++)
+      matrix1[i * dim1 + j] = j;
+
+  for (unsigned i = 0; i < dim1; i++)
+    for (unsigned j = 0; j < dim0; j++)
+      if (i == j)
+       matrix2[i * dim0 + j] = 1;
+      else
+       matrix2[i * dim0 + j] = 0;
+
+  mult (matrix1, matrix2, result, dim0, dim1);
+
+  for (unsigned i = 0; i < dim0; i++)
+    for (unsigned j = 0; j < dim1; j++)
+      {
+       if (matrix1[i * dim1 + j] != result[i * dim1 + j])
+         {
+           printf ("ERROR at %d, %d\n", i, j);
+           __builtin_abort ();
+         }
+      }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/loop-transforms/tile-2.C b/libgomp/testsuite/libgomp.c++/loop-transforms/tile-2.C
new file mode 100644
index 00000000000..780421fa4c7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/loop-transforms/tile-2.C
@@ -0,0 +1,69 @@
+// { dg-additional-options "-std=c++11" }
+// { dg-additional-options "-O0" }
+
+#include <vector>
+#include <stdio.h>
+
+constexpr unsigned fib (unsigned n)
+{
+  return n <= 2 ? 1 : fib (n-1) + fib (n-2);
+}
+
+int
+test1 ()
+{
+  std::vector<int> v;
+
+  for (unsigned i = 0; i <= 9; i++)
+    v.push_back (1);
+
+  int sum = 0;
+  for (int k = 0; k < 10; k++)
+    #pragma omp tile sizes(fib(4))
+    for (int i : v) {
+      for (int j = 8; j != -2; --j)
+       sum = sum + i;
+    }
+
+  return sum;
+}
+
+int
+test2 ()
+{
+  std::vector<int> v;
+
+  for (unsigned i = 0; i <= 10; i++)
+    v.push_back (i);
+
+  int sum = 0;
+  for (int k = 0; k < 10; k++)
+#pragma omp parallel for collapse(2) reduction(+:sum)
+#pragma omp tile sizes(fib(4), 1)
+  for (int i : v)
+    for (int j = 8; j > -2; --j)
+       sum = sum + i;
+
+  return sum;
+}
+
+int
+main ()
+{
+  int result = test1 ();
+
+  if (result != 1000)
+    {
+      fprintf (stderr, "%d: Wrong result: %d\n", __LINE__, result);
+      __builtin_abort ();
+    }
+
+  result = test2 ();
+  if (result != 5500)
+    {
+      fprintf (stderr, "%d: Wrong result: %d\n", __LINE__, result);
+    __builtin_abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/loop-transforms/tile-3.C b/libgomp/testsuite/libgomp.c++/loop-transforms/tile-3.C
new file mode 100644
index 00000000000..91ec8f5c137
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/loop-transforms/tile-3.C
@@ -0,0 +1,28 @@
+// { dg-additional-options "-std=c++11" }
+// { dg-additional-options "-O0" }
+
+#include <vector>
+
+int
+main ()
+{
+  std::vector<int> v;
+  std::vector<int> w;
+
+  for (unsigned i = 0; i <= 9; i++)
+    v.push_back (i);
+
+  int iter = 0;
+#pragma omp for
+#pragma omp tile sizes(5)
+  for (int i : v)
+    {
+      w.push_back (iter);
+      iter++;
+    }
+
+  for (int i = 0; i < w.size (); i++)
+    if (w[i] != i)
+      __builtin_abort ();
+  return 0;
+}
--
2.36.1

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955