@@ -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.
@@ -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
@@ -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;
@@ -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
@@ -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);
@@ -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 */
@@ -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