diff --git a/gcc/gimple.c b/gcc/gimple.c
index d876e90..4658f29 100644
--- a/gcc/gimple.c
+++ b/gcc/gimple.c
@@ -1098,17 +1098,6 @@ gimple_build_omp_target (gimple_seq body, int kind, tree clauses)
   return p;
 }
 
-/* Set dimensions of TARGET to NUM and allocate kernel_dim array of the
-   statement with the appropriate number of elements.  */
-
-void
-gimple_omp_target_init_dimensions (gomp_target *target, size_t num)
-{
-  gcc_assert (num > 0);
-  target->dimensions = num;
-  target->kernel_dim = ggc_cleared_vec_alloc<gimple_omp_target_grid_dim> (num);
-}
-
 /* Build a GIMPLE_OMP_TEAMS statement.
 
    BODY is the sequence of statements that will be executed.
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 14e6cf6..4c4c799 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -661,21 +661,7 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
      Shared data argument.  */
   tree data_arg;
 
-  /* TODO: Revisit placement of the following two fields.  On one hand, we
-     currently only use them on target construct.  On the other, use on
-     parallel construct is also possible in the future.  */
-
   /* [ WORD 11 ] */
-  /* Number of elements in kernel_iter array.  */
-  size_t dimensions;
-
-  /* [ WORD 12 ] */
-  /* If target also contains a GPU kernel, it should be run with the
-     following grid sizes.  */
-  struct gimple_omp_target_grid_dim
-    * GTY((length ("%h.dimensions"))) kernel_dim;
-
-  /* [ WORD 13 ] */
   /* If set, this statement is part of a gridified kernel, its clauses need to
      be scanned and lowered but the statement should be discarded after
      lowering.  */
@@ -1504,7 +1490,6 @@ gomp_sections *gimple_build_omp_sections (gimple_seq, tree);
 gimple *gimple_build_omp_sections_switch (void);
 gomp_single *gimple_build_omp_single (gimple_seq, tree);
 gomp_target *gimple_build_omp_target (gimple_seq, int, tree);
-void gimple_omp_target_init_dimensions (gomp_target *, size_t);
 gomp_teams *gimple_build_omp_teams (gimple_seq, tree);
 gomp_atomic_load *gimple_build_omp_atomic_load (tree, tree);
 gomp_atomic_store *gimple_build_omp_atomic_store (tree);
@@ -5683,73 +5668,6 @@ gimple_omp_target_set_data_arg (gomp_target *omp_target_stmt,
   omp_target_stmt->data_arg = data_arg;
 }
 
-/* Return the number of dimensions of kernel grid.  */
-
-static inline size_t
-gimple_omp_target_dimensions (gomp_target *omp_target_stmt)
-{
-  return omp_target_stmt->dimensions;
-}
-
-/* Return the size of kernel grid of OMP_TARGET_STMT along dimension N.  */
-
-static inline tree
-gimple_omp_target_grid_size (gomp_target *omp_target_stmt, unsigned n)
-{
-  gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n);
-  return omp_target_stmt->kernel_dim[n].grid_size;
-}
-
-/* Return pointer to tree specifying the size of kernel grid of OMP_TARGET_STMT
-   along dimension N.  */
-
-static inline tree *
-gimple_omp_target_grid_size_ptr (gomp_target *omp_target_stmt, unsigned n)
-{
-  gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n);
-  return &omp_target_stmt->kernel_dim[n].grid_size;
-}
-
-/* Set the size of kernel grid of OMP_TARGET_STMT along dimension N to V  */
-
-static inline void
-gimple_omp_target_set_grid_size (gomp_target *omp_target_stmt, unsigned n,
-				 tree v)
-{
-  gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n);
-  omp_target_stmt->kernel_dim[n].grid_size = v;
-}
-
-/* Return the size of kernel work group of OMP_TARGET_STMT along dimension N.  */
-
-static inline tree
-gimple_omp_target_workgroup_size (gomp_target *omp_target_stmt, unsigned n)
-{
-  gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n);
-  return omp_target_stmt->kernel_dim[n].workgroup_size;
-}
-
-/* Return pointer to tree specifying the size of kernel work group of
-   OMP_TARGET_STMT along dimension N.  */
-
-static inline tree *
-gimple_omp_target_workgroup_size_ptr (gomp_target *omp_target_stmt, unsigned n)
-{
-  gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n);
-  return &omp_target_stmt->kernel_dim[n].workgroup_size;
-}
-
-/* Set the size of kernel workgroup of OMP_TARGET_STMT along dimension N to
-   V */
-
-static inline void
-gimple_omp_target_set_workgroup_size (gomp_target *omp_target_stmt, unsigned n,
-				      tree v)
-{
-  gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n);
-  omp_target_stmt->kernel_dim[n].workgroup_size = v;
-}
-
 /* Return the clauses associated with OMP_TEAMS GS.  */
 
 static inline tree
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index f1d10a2..5933c60 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2140,6 +2140,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	    }
 	  break;
 
+	case OMP_CLAUSE__GRIDDIM_:
+	  if (ctx->outer)
+	    {
+	      scan_omp_op (&OMP_CLAUSE_GRIDDIM_SIZE (c), ctx->outer);
+	      scan_omp_op (&OMP_CLAUSE_GRIDDIM_GROUP (c), ctx->outer);
+	    }
+	  break;
+
 	case OMP_CLAUSE_NOWAIT:
 	case OMP_CLAUSE_ORDERED:
 	case OMP_CLAUSE_COLLAPSE:
@@ -2336,6 +2344,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE__GRIDDIM_:
 	  break;
 
 	case OMP_CLAUSE_DEVICE_RESIDENT:
@@ -3088,12 +3097,6 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
   TYPE_NAME (ctx->record_type) = name;
   TYPE_ARTIFICIAL (ctx->record_type) = 1;
 
-  for (size_t i = 0; i < gimple_omp_target_dimensions (stmt); i++)
-    {
-      scan_omp_op (gimple_omp_target_grid_size_ptr (stmt, i), ctx);
-      scan_omp_op (gimple_omp_target_workgroup_size_ptr (stmt, i), ctx);
-    }
-
   if (offloaded)
     {
       create_omp_child_function (ctx, false);
@@ -6310,7 +6313,9 @@ region_needs_kernel_p (struct omp_region *region)
 	{
 	  gomp_target *tgt_stmt;
 	  tgt_stmt = as_a <gomp_target *> (last_stmt (region->entry));
-	  if (gimple_omp_target_dimensions (tgt_stmt))
+
+	  if (find_omp_clause (gimple_omp_target_clauses (tgt_stmt),
+			       OMP_CLAUSE__GRIDDIM_))
 	    return indirect;
 	  else
 	    return true;
@@ -12624,26 +12629,30 @@ get_kernel_launch_attributes (gimple_stmt_iterator *gsi, gomp_target *tgt_stmt)
   tree u32_one = build_one_cst (uint32_type_node);
   tree lattrs = create_tmp_var (kernel_launch_attributes_type,
 				"__kernel_launch_attrs");
+
+  unsigned max_dim = 0;
+  for (tree clause = gimple_omp_target_clauses (tgt_stmt);
+       clause;
+       clause = OMP_CLAUSE_CHAIN (clause))
+    {
+      if (OMP_CLAUSE_CODE (clause) != OMP_CLAUSE__GRIDDIM_)
+	continue;
+
+      unsigned dim = OMP_CLAUSE_GRIDDIM_DIMENSION (clause);
+      max_dim = MAX (dim, max_dim);
+
+      insert_store_range_dim (gsi, lattrs, kernel_lattrs_grid_decl, dim,
+			      OMP_CLAUSE_GRIDDIM_SIZE (clause));
+      insert_store_range_dim (gsi, lattrs, kernel_lattrs_group_decl, dim,
+			      OMP_CLAUSE_GRIDDIM_GROUP (clause));
+    }
+
   tree dimref = build3 (COMPONENT_REF, uint32_type_node,
 			lattrs, kernel_lattrs_dimnum_decl, NULL_TREE);
   /* At this moment we cannot gridify a loop with a collapse clause.  */
   /* TODO: Adjust when we support bigger collapse.  */
-  gcc_assert (gimple_omp_target_dimensions (tgt_stmt) == 1);
+  gcc_assert (max_dim == 0);
   gsi_insert_before (gsi, gimple_build_assign (dimref, u32_one), GSI_SAME_STMT);
-
-  /* Calculation of grid size: */
-  insert_store_range_dim (gsi, lattrs, kernel_lattrs_grid_decl, 0,
-			  gimple_omp_target_grid_size (tgt_stmt, 0));
-  insert_store_range_dim (gsi, lattrs, kernel_lattrs_group_decl, 0,
-			  gimple_omp_target_workgroup_size (tgt_stmt, 0));
-  insert_store_range_dim (gsi, lattrs, kernel_lattrs_grid_decl, 1,
-			  u32_one);
-  insert_store_range_dim (gsi, lattrs, kernel_lattrs_group_decl, 2,
-			  u32_one);
-  insert_store_range_dim (gsi, lattrs, kernel_lattrs_grid_decl, 2,
-			  u32_one);
-  insert_store_range_dim (gsi, lattrs, kernel_lattrs_group_decl, 1,
-			  u32_one);
   TREE_ADDRESSABLE (lattrs) = 1;
   return build_fold_addr_expr (lattrs);
 }
@@ -12717,7 +12726,8 @@ get_target_arguments (gimple_stmt_iterator *gsi, gomp_target *tgt_stmt)
   args.quick_push (t);
 
   /* Add HSA-specific grid sizes, if available.  */
-  if (gimple_omp_target_dimensions (tgt_stmt))
+  if (find_omp_clause (gimple_omp_target_clauses (tgt_stmt),
+		       OMP_CLAUSE__GRIDDIM_))
     {
       t = get_target_argument_identifier (GOMP_DEVICE_HSA, true,
 					  GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES);
@@ -13392,14 +13402,16 @@ expand_target_kernel_body (struct omp_region *target)
       if (gimple_omp_target_kind (tgt_stmt) != GF_OMP_TARGET_KIND_REGION)
 	return;
       gcc_checking_assert (orig_child_fndecl);
-      gcc_assert (!gimple_omp_target_dimensions (tgt_stmt));
+      gcc_assert (!find_omp_clause (gimple_omp_target_clauses (tgt_stmt),
+				    OMP_CLAUSE__GRIDDIM_));
       cgraph_node *n = cgraph_node::get (orig_child_fndecl);
 
       hsa_register_kernel (n);
       return;
     }
 
-  gcc_assert (gimple_omp_target_dimensions (tgt_stmt));
+  gcc_assert (find_omp_clause (gimple_omp_target_clauses (tgt_stmt),
+			       OMP_CLAUSE__GRIDDIM_));
   tree inside_block = gimple_block (first_stmt (single_succ (gpukernel->entry)));
   *pp = gpukernel->next;
   for (pp = &gpukernel->inner; *pp; pp = &(*pp)->next)
@@ -17470,7 +17482,6 @@ attempt_target_gridification (gomp_target *target, gimple_stmt_iterator *gsi,
 
   walk_tree (&group_size, remap_prebody_decls, &wi, NULL);
   size_t collapse = gimple_omp_for_collapse (inner_loop);
-  gimple_omp_target_init_dimensions (target, collapse);
   for (size_t i = 0; i < collapse; i++)
     {
       gimple_omp_for_iter iter = inner_loop->iter[i];
@@ -17506,7 +17517,6 @@ attempt_target_gridification (gomp_target *target, gimple_stmt_iterator *gsi,
       t = fold_convert (uint32_type_node, t);
       tree gs = force_gimple_operand_gsi (gsi, t, true, NULL_TREE, true,
 					  GSI_SAME_STMT);
-      gimple_omp_target_set_grid_size (target, i, gs);
       tree ws;
       if (i == 0 && group_size)
 	{
@@ -17516,7 +17526,13 @@ attempt_target_gridification (gomp_target *target, gimple_stmt_iterator *gsi,
 	}
       else
 	ws = build_zero_cst (uint32_type_node);
-      gimple_omp_target_set_workgroup_size (target, i, ws);
+
+      tree c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE__GRIDDIM_);
+      OMP_CLAUSE_SET_GRIDDIM_DIMENSION (c, (unsigned int) i);
+      OMP_CLAUSE_GRIDDIM_SIZE (c) = gs;
+      OMP_CLAUSE_GRIDDIM_GROUP (c) = ws;
+      OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (target);
+      gimple_omp_target_set_clauses (target, c);
     }
 
   delete declmap;
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 9cc64d9..858f220 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -460,7 +460,11 @@ enum omp_clause_code {
   OMP_CLAUSE_VECTOR_LENGTH,
 
   /* OpenACC clause: tile ( size-expr-list ).  */
-  OMP_CLAUSE_TILE
+  OMP_CLAUSE_TILE,
+
+  /* OpenMP internal-only clause to specify grid dimensions of a gridified
+     kernel.  */
+  OMP_CLAUSE__GRIDDIM_
 };
 
 #undef DEFTREESTRUCT
@@ -1377,6 +1381,9 @@ struct GTY(()) tree_omp_clause {
     enum tree_code                 reduction_code;
     enum omp_clause_linear_kind    linear_kind;
     enum tree_code                 if_modifier;
+    /* The dimension a OMP_CLAUSE__GRIDDIM_ clause of a gridified target
+       construct describes.  */
+    unsigned int		   dimension;
   } GTY ((skip)) subcode;
 
   /* The gimplification of OMP_CLAUSE_REDUCTION_{INIT,MERGE} for omp-low's
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index caec760..ad5cfdb 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -945,6 +945,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
       pp_right_paren (pp);
       break;
 
+    case OMP_CLAUSE__GRIDDIM_:
+      pp_string (pp, "_griddim_(");
+      pp_unsigned_wide_integer (pp, OMP_CLAUSE_GRIDDIM_DIMENSION (clause));
+      pp_colon (pp);
+      dump_generic_node (pp, OMP_CLAUSE_GRIDDIM_SIZE (clause), spc, flags,
+			 false);
+      pp_comma (pp);
+      dump_generic_node (pp, OMP_CLAUSE_GRIDDIM_GROUP (clause), spc, flags,
+			 false);
+      pp_right_paren (pp);
+      break;
+
     default:
       /* Should never happen.  */
       dump_generic_node (pp, clause, spc, flags, false);
diff --git a/gcc/tree.c b/gcc/tree.c
index 2387deb..3a74982 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -329,6 +329,7 @@ unsigned const char omp_clause_num_ops[] =
   1, /* OMP_CLAUSE_NUM_WORKERS  */
   1, /* OMP_CLAUSE_VECTOR_LENGTH  */
   1, /* OMP_CLAUSE_TILE  */
+  2, /* OMP_CLAUSE__GRIDDIM_  */
 };
 
 const char * const omp_clause_code_name[] =
@@ -400,7 +401,8 @@ const char * const omp_clause_code_name[] =
   "num_gangs",
   "num_workers",
   "vector_length",
-  "tile"
+  "tile",
+  "griddim"
 };
 
 
@@ -11603,6 +11605,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
       switch (OMP_CLAUSE_CODE (*tp))
 	{
 	case OMP_CLAUSE_GANG:
+	case OMP_CLAUSE__GRIDDIM_:
 	  WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 1));
 	  /* FALLTHRU */
 
diff --git a/gcc/tree.h b/gcc/tree.h
index 0c1602e..7b9bcb3 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1636,6 +1636,17 @@ extern void protected_set_expr_location (tree, location_t);
 #define OMP_CLAUSE_TILE_LIST(NODE) \
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_TILE), 0)
 
+#define OMP_CLAUSE_GRIDDIM_DIMENSION(NODE) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__GRIDDIM_)\
+   ->omp_clause.subcode.dimension)
+#define OMP_CLAUSE_SET_GRIDDIM_DIMENSION(NODE, DIMENSION) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__GRIDDIM_)\
+   ->omp_clause.subcode.dimension = (DIMENSION))
+#define OMP_CLAUSE_GRIDDIM_SIZE(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__GRIDDIM_), 0)
+#define OMP_CLAUSE_GRIDDIM_GROUP(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__GRIDDIM_), 1)
+
 /* SSA_NAME accessors.  */
 
 /* Returns the IDENTIFIER_NODE giving the SSA name a name or NULL_TREE
