@@ -1,3 +1,8 @@
+2010-12-27 Sebastian Pop <sebastian.pop@amd.com>
+
+ * graphite-opencl-codegen.c: Sort topologically static functions.
+ * graphite-opencl.c: Same.
+
2010-12-25 Sebastian Pop <sebastian.pop@amd.com>
* graphite-cloog-compat.h (cloog_names_nb_scattering): New.
@@ -61,43 +61,6 @@
#include "dyn-string.h"
#include "graphite-opencl.h"
-
-/* These functions implement code generation from different clast
- structures. */
-static void opencl_print_stmt_list (struct clast_stmt *, opencl_main, int);
-static void opencl_print_for (struct clast_for *, opencl_main, int);
-static void opencl_print_guard (struct clast_guard *, opencl_main, int);
-static void opencl_print_equation (struct clast_equation *, opencl_main);
-static void opencl_print_expr (struct clast_expr *, opencl_main);
-static void opencl_add_variable (const char *, tree, opencl_main);
-static void opencl_print_term (struct clast_term *, opencl_main);
-static void opencl_print_reduction (struct clast_reduction *, opencl_main);
-static void opencl_print_sum (struct clast_reduction *, opencl_main);
-static void opencl_print_binary (struct clast_binary *, opencl_main);
-static void opencl_print_minmax_c (struct clast_reduction *, opencl_main);
-
-/* These function implement code generation from different gimple
- objects. */
-static void opencl_print_bb (basic_block, opencl_main);
-static void opencl_print_gimple_assign_operation (gimple, opencl_main);
-static void opencl_print_gimple_assign (gimple, opencl_main);
-static void opencl_print_gimple (gimple, opencl_main);
-static int opencl_print_operand (tree, bool, opencl_main);
-
-
-static void opencl_print_local_vars (const char *, const char *, const char *,
- opencl_main);
-static void opencl_try_variable (opencl_main, tree);
-static const char *opencl_get_var_name (tree);
-static void opencl_build_defines (tree, opencl_main);
-static void opencl_expand_scalar_vars (opencl_main, gimple);
-static void opencl_add_function_arg (opencl_main, tree, const char *);
-static void opencl_add_data_refs_pbb (poly_bb_p, opencl_main);
-static void opencl_add_non_scalar_type_decl (tree, dyn_string_t, const char *);
-static const char *opencl_print_function_arg_with_type (const char *, tree);
-static bool check_and_mark_arg (opencl_main, const char *, bool);
-
-
/* Compare two clast names based on their indexes. */
static int
@@ -203,7 +166,6 @@ opencl_get_main_type (tree type)
return build_pointer_type (type);
}
-
/* Create the base part of FUNCTION declaration, similar to this:
"__global void __opencl_function_0". */
@@ -455,6 +417,94 @@ gen_type_with_name (const char *name, tree t)
return concat (data_type, " ", type_part, NULL);
}
+/* Get name of the variable, represented by tree NODE. If variable is
+ temporary, generate name for it. */
+
+static const char *
+opencl_get_var_name (tree node)
+{
+ bool ssa_name = TREE_CODE (node) == SSA_NAME;
+ tree name;
+ int num = 0;
+ if (ssa_name)
+ {
+ num = SSA_NAME_VERSION (node);
+ node = SSA_NAME_VAR (node);
+ }
+ name = DECL_NAME (node);
+ if (name)
+ {
+ if (!ssa_name)
+ return identifier_to_locale (IDENTIFIER_POINTER (name));
+ else
+ {
+ const char *base = identifier_to_locale (IDENTIFIER_POINTER (name));
+ char *buff = XNEWVEC (char, strlen (base) + 5);
+ sprintf (buff, "%s_%d", base, num);
+ return buff;
+ }
+ }
+ else
+ {
+ int tmp_var_uid = DECL_UID (node);
+ char *tmp = XNEWVEC (char, 30);
+ sprintf (tmp, "opencl_var_%d_%d", tmp_var_uid, num);
+ return tmp;
+ }
+}
+
+/* Replace all dots to underscores in string pointed to by P. Return P. */
+
+static char *
+filter_dots (char *p)
+{
+ char *s;
+ for (s = p; *s; s++)
+ if (*s == '.')
+ *s = '_';
+ return p;
+}
+
+/* Return string with varibale definition. ARG_NAME is the name of
+ the variable and TYPE is it's type. */
+
+static const char *
+opencl_print_function_arg_with_type (const char *arg_name, tree type)
+{
+ const char *decl = gen_type_with_name (arg_name, type);
+ char *ddecl;
+ ddecl = xstrdup (decl);
+ return filter_dots (ddecl);
+}
+
+/* Check whether variable with name NAME has been defined as global or
+ local variable and mark it as defined. This function returns false
+ if variable has already been defined, otherwise it returns true. */
+
+static bool
+check_and_mark_arg (opencl_main code_gen, const char *name, bool local)
+{
+ const char **slot;
+ gcc_assert (code_gen->defined_vars || !local);
+ if (code_gen->defined_vars)
+ {
+ slot = (const char **)htab_find_slot (code_gen->defined_vars,
+ name, INSERT);
+ if (*slot)
+ return false;
+ if (local)
+ *slot = name;
+ }
+
+ slot = (const char **)htab_find_slot (code_gen->global_defined_vars,
+ name, INSERT);
+ if (*slot)
+ return false;
+ if (!local)
+ *slot = name;
+ return true;
+}
+
/* Replace perfect nested loop nest represented by F with opencl kernel.
For example, loop nest like this
@@ -611,61 +661,40 @@ opencl_perfect_nested_to_kernel (opencl_main code_gen, struct clast_for *f,
VEC_free (tree, heap, mod);
}
-/* Generate code for loop statement F. DEPTH is the depth of F in
- current loop nest. CODE_GEN holds information related to OpenCL
- code generation. */
+/* Append list of names of loop iterators from CODE_GEN with same type
+ TYPE to current kernel. FIRST and LAST define outermost and
+ innermost iterators to append respectively. */
-static opencl_body
-opencl_print_loop (struct clast_for *f, opencl_main code_gen, int depth)
+static void
+opencl_print_local_vars (const char *fist, const char *last,
+ const char *type, opencl_main code_gen)
{
- opencl_body current_body = code_gen->current_body;
-
- code_gen->global_defined_vars
- = htab_create (10, htab_hash_string, opencl_cmp_str, NULL);
-
- opencl_perfect_nested_to_kernel (code_gen, f, current_body, depth);
-
- /* Define local loop iterators. */
- opencl_print_local_vars (current_body->first_iter,
- current_body->last_iter,
- "unsigned int", code_gen);
-
- /* Generate code for kernel body. */
- opencl_print_stmt_list (current_body->clast_body, code_gen, depth + 1);
- opencl_append_string_to_body ("}\n", code_gen);
-
- if (current_body->num_of_data_writes)
+ char **names = cloog_names_scattering (code_gen->root_names);
+ int len = cloog_names_nb_scattering (code_gen->root_names);
+ int i;
+ for (i = 0; i < len; i++)
{
- dyn_string_t header = current_body->header;
- int offset;
-
- dyn_string_append (header, current_body->non_scalar_args);
- offset = dyn_string_length (header) - 2;
-
- if (*(dyn_string_buf (header) + offset) == ',')
- *(dyn_string_buf (header) + offset) = ' ';
-
- opencl_append_string_to_header (")\n{\n", code_gen);
- }
-
- return current_body;
-}
+ const char *tmp = names[i];
+ if (opencl_cmp_scat (fist, tmp) <= 0
+ && opencl_cmp_scat (last, tmp) >= 0)
+ {
+ const char **slot =
+ (const char **) htab_find_slot (code_gen->global_defined_vars,
+ tmp, INSERT);
+ *slot = tmp;
+ continue;
+ }
-/* Generate OpenCL code for clast_assignment A.
- CODE_GEN holds information related to OpenCL code generation. */
+ if (opencl_cmp_scat (fist, tmp) > 0)
+ continue;
-static void
-opencl_print_assignment (struct clast_assignment *a, opencl_main code_gen)
-{
- /* Real assignment. */
- if (a->LHS)
- {
- opencl_append_string_to_body (a->LHS, code_gen);
- opencl_append_string_to_body (" = ", code_gen);
+ opencl_append_string_to_body (type, code_gen);
+ opencl_append_string_to_body (" ", code_gen);
+ opencl_append_string_to_body (tmp, code_gen);
+ opencl_append_string_to_body (";\n", code_gen);
+ *((const char **)htab_find_slot (code_gen->global_defined_vars,
+ tmp, INSERT)) = tmp;
}
-
- /* Just expression. */
- opencl_print_expr (a->RHS, code_gen);
}
/* Return tree with variable, corresponging to given clast name NAME.
@@ -695,6 +724,24 @@ opencl_get_scat_real_name (opencl_main code_gen, clast_name_p name)
return opencl_get_var_name (opencl_clast_name_to_tree (code_gen, name));
}
+/* Add variable VAR with name NAME as function argument. Append it's
+ declaration in finction header and add it as function parameter.
+ CODE_GEN holds information related to OpenCL code generation. */
+
+static void
+opencl_add_function_arg (opencl_main code_gen, tree var, const char *name)
+{
+ opencl_body body;
+ const char *decl;
+ tree type;
+ type = TREE_TYPE (var);
+ body = code_gen->current_body;
+ decl = opencl_print_function_arg_with_type (name, type);
+ dyn_string_append_cstr (body->header, decl);
+ dyn_string_append_cstr (body->header, ", ");
+ VEC_safe_push (tree, heap, body->function_args, var);
+}
+
/* Add clast variable (scat_i) as kernel argument. NAME is a new name
of loop iterator (scat_*), REAL_NAME is an old (origin) name of
loop iterator. CODE_GEN holds information related to OpenCL code
@@ -713,514 +760,311 @@ opencl_add_scat_as_arg (opencl_main code_gen, clast_name_p name,
opencl_add_function_arg (code_gen, var, real_name);
}
-/* Generate OpenCL code for user statement U. Code will be generated
- from basic block, related to U. Also induction variables mapping
- to old variables must be calculated to process basic block.
- CODE_GEN holds information related to OpenCL code generation. */
+/* Append variable name NAME to function body. Differs from appending
+ string by replacing `.' by `_'. CODE_GEN holds information related
+ to OpenCL code generation. */
static void
-opencl_print_user_stmt (struct clast_user_stmt *u, opencl_main code_gen)
+opencl_append_var_name (const char *name, opencl_main code_gen)
{
- CloogStatement * cs;
- poly_bb_p pbb;
- gimple_bb_p gbbp;
- basic_block bb;
+ int len = strlen (name);
+ char *tmp = XNEWVEC (char, len + 1);
int i;
- int nb_loops = number_of_loops ();
- code_gen->iv_map = VEC_alloc (tree, heap, nb_loops);
-
- for (i = 0; i < nb_loops; i++)
- VEC_safe_push (tree, heap, code_gen->iv_map, NULL_TREE);
- build_iv_mapping (code_gen->iv_map, code_gen->region,
- code_gen->newivs,
- code_gen->newivs_index, u,
- code_gen->params_index);
-
- code_gen->defined_vars
- = htab_create (10, htab_hash_string, opencl_cmp_str, NULL);
- opencl_append_string_to_body ("{\n", code_gen);
-
- cs = u->statement;
- pbb = (poly_bb_p) cloog_statement_usr (cs);
- gbbp = PBB_BLACK_BOX (pbb);
- bb = GBB_BB (gbbp);
- code_gen->context_loop = bb->loop_father;
-
- opencl_add_data_refs_pbb (pbb, code_gen);
- opencl_print_bb (bb, code_gen);
- opencl_append_string_to_body ("}\n", code_gen);
- htab_delete (code_gen->defined_vars);
- code_gen->defined_vars = NULL;
- VEC_free (tree, heap, code_gen->iv_map);
+ for (i = 0; i <= len; i++)
+ {
+ char tt = name[i];
+ if (tt == '.')
+ tt = '_';
+ tmp[i] = tt;
+ }
+ opencl_append_string_to_body (tmp, code_gen);
+ free (tmp);
}
-/* If tree node NODE defined in current sese build and insert define
- statements for it, otherwise mark node as external (parameter for
- kernel). If tree defined in current sese, also recursively build
- defines for all trees in definition expression. */
+/* Generate code for clast term T. CODE_GEN holds information
+ related to OpenCL code generation. */
static void
-opencl_build_defines (tree node, opencl_main code_gen)
+opencl_print_term (struct clast_term *t, opencl_main code_gen)
{
- enum tree_code code = TREE_CODE (node);
- switch (code)
+ if (t->var)
{
- case SSA_NAME:
- {
- const char *tmp = opencl_get_var_name (node);
- gimple def_stmt;
-
- /* If name defined in other sese it is kernel's parameter. */
- if (!defined_in_sese_p (node, code_gen->region))
- return;
-
- /* Bail out if this name was defined earlier either in this
- or other region. */
- if (*(const char **)htab_find_slot (code_gen->defined_vars,
- tmp, INSERT))
- return;
+ const char *real_name = opencl_get_scat_real_name (code_gen, t->var);
- /* Get definition statement. */
- def_stmt = SSA_NAME_DEF_STMT (node);
- opencl_expand_scalar_vars (code_gen, def_stmt);
- opencl_print_gimple (def_stmt, code_gen);
- return;
- }
- case ARRAY_REF:
- {
- tree arr = TREE_OPERAND (node, 0);
- tree offset = TREE_OPERAND (node, 1);
- opencl_build_defines (arr, code_gen);
- opencl_build_defines (offset, code_gen);
- return;
- }
- default:
- gcc_unreachable ();
+ if (mpz_cmp_si (t->val, 1) == 0)
+ opencl_append_var_name (real_name, code_gen);
+ else if (mpz_cmp_si (t->val, -1) == 0)
+ {
+ opencl_append_string_to_body ("-", code_gen);
+ opencl_append_var_name (real_name, code_gen);
+ }
+ else
+ {
+ opencl_append_num_to_body (code_gen, mpz_get_si (t->val), "%d");
+ opencl_append_string_to_body ("*", code_gen);
+ opencl_append_var_name (real_name, code_gen);
+ }
+ opencl_add_scat_as_arg (code_gen, t->var, real_name);
}
+ else
+ opencl_append_num_to_body (code_gen, mpz_get_si (t->val), "%d");
}
-/* For a given gimple statement STMT build definition for all names,
- used in this stament. If name has been defined in other sese, mark
- it as kernel parameter. CODE_GEN holds information related to
- OpenCL code generation. */
+/* Generate code for clast sum statement R. CODE_GEN holds information
+ related to OpenCL code generation. */
static void
-opencl_expand_scalar_vars (opencl_main code_gen, gimple stmt)
+opencl_print_sum (struct clast_reduction *r, opencl_main code_gen)
{
- ssa_op_iter iter;
- use_operand_p use_p;
- FOR_EACH_SSA_USE_OPERAND (use_p, stmt, iter, SSA_OP_ALL_USES)
- {
- tree use = USE_FROM_PTR (use_p);
- if (!is_gimple_reg (use))
- continue;
- opencl_build_defines (use, code_gen);
- }
-}
+ int i;
+ struct clast_term *t;
-/* Generate code for a single basic block BB. CODE_GEN holds
- information related to OpenCL code generation. */
+ gcc_assert (r->n >= 1 && r->elts[0]->type == clast_expr_term);
+ t = (struct clast_term *) r->elts[0];
+ opencl_print_term (t, code_gen);
-static void
-opencl_print_bb (basic_block bb, opencl_main code_gen)
-{
- gimple_stmt_iterator gsi;
- for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
+ for (i = 1; i < r->n; ++i)
{
- gimple stmt = gsi_stmt (gsi);
- opencl_expand_scalar_vars (code_gen, stmt);
- opencl_print_gimple (stmt, code_gen);
+ gcc_assert (r->elts[i]->type == clast_expr_term);
+ t = (struct clast_term *) r->elts[i];
+ if (mpz_sgn (t->val) > 0)
+ opencl_append_string_to_body ("+", code_gen);
+ opencl_print_term (t, code_gen);
}
}
-/* Print operation simbol (`+' `-' `*') for assignment operation GMA.
- CODE_GEN holds information related to OpenCL code generation. */
-
-static void
-opencl_print_gimple_assign_operation (gimple gmp, opencl_main code_gen)
-{
- opencl_append_string_to_body
- (op_symbol_code (gimple_assign_rhs_code (gmp)), code_gen);
-}
+static void opencl_print_expr (struct clast_expr *, opencl_main);
-/* Print pointer expression represented by EXPR. TYPE_SIZE represents
- size of the base type for EXPR. CODE_GEN holds information related
- to OpenCL code generation. */
+/* Generate code for clast min/max operation R. CODE_GEN holds
+ information related to OpenCL code generation. */
static void
-opencl_print_addr_operand (tree expr, tree type_size, opencl_main code_gen)
+opencl_print_minmax_c ( struct clast_reduction *r, opencl_main code_gen)
{
- if (TREE_CODE (TREE_TYPE (expr)) != POINTER_TYPE)
+ int i;
+ for (i = 1; i < r->n; ++i)
+ opencl_append_string_to_body (r->type == clast_red_max ? "max (" : "min (",
+ code_gen);
+ if (r->n > 0)
{
- opencl_append_string_to_body ("(", code_gen);
- opencl_print_operand (expr, false, code_gen);
- opencl_append_string_to_body ("/", code_gen);
- opencl_print_operand (type_size, false, code_gen);
+ opencl_append_string_to_body ("(unsigned int)(", code_gen);
+ opencl_print_expr (r->elts[0], code_gen);
opencl_append_string_to_body (")", code_gen);
}
- else
- opencl_print_operand (expr, false, code_gen);
-
-}
-
-/* Print unary gimple operation GMP. CODE_GEN holds information
- related to OpenCL code generation. */
-
-static void
-opencl_print_unary (gimple gmp, opencl_main code_gen)
-{
- switch (gimple_assign_rhs_code (gmp))
+ for (i = 1; i < r->n; ++i)
{
- case BIT_NOT_EXPR:
- opencl_append_string_to_body ("~", code_gen);
- return;
- case TRUTH_NOT_EXPR:
- opencl_append_string_to_body ("!", code_gen);
- return;
- case NEGATE_EXPR:
- opencl_append_string_to_body ("-", code_gen);
- return;
- case MODIFY_EXPR:
- default:
- return;
+ opencl_append_string_to_body (",", code_gen);
+ opencl_append_string_to_body ("(unsigned int)(", code_gen);
+ opencl_print_expr (r->elts[i], code_gen);
+ opencl_append_string_to_body ("))", code_gen);
}
}
-/* Generate code for min or max gimple operand GMP. CODE_GEN holds
+/* Generate code for clast reduction statement R. CODE_GEN holds
information related to OpenCL code generation. */
static void
-opencl_print_max_min_assign (gimple gmp, opencl_main code_gen)
+opencl_print_reduction (struct clast_reduction *r, opencl_main code_gen)
{
- tree lhs = gimple_assign_lhs (gmp);
- tree rhs1 = gimple_assign_rhs1 (gmp);
- tree rhs2 = gimple_assign_rhs2 (gmp);
- bool max = gimple_assign_rhs_code (gmp) == MAX_EXPR;
-
- opencl_print_operand (lhs, true, code_gen);
- opencl_append_string_to_body (max?" = fmax (":"= fmin (", code_gen);
- opencl_print_operand (rhs1, false, code_gen);
- opencl_append_string_to_body (",", code_gen);
- opencl_print_operand (rhs2, false, code_gen);
- opencl_append_string_to_body (");\n", code_gen);
-
+ switch (r->type)
+ {
+ case clast_red_sum:
+ opencl_print_sum (r, code_gen);
+ break;
+ case clast_red_min:
+ case clast_red_max:
+ if (r->n == 1)
+ {
+ opencl_print_expr (r->elts[0], code_gen);
+ break;
+ }
+ opencl_print_minmax_c (r, code_gen);
+ break;
+ default:
+ gcc_unreachable ();
+ }
}
-/* Generate code for gimple assignment statement GMP. CODE_GEN holds
+/* Generate code for clast binary operation B. CODE_GEN holds
information related to OpenCL code generation. */
static void
-opencl_print_gimple_assign (gimple gmp, opencl_main code_gen)
+opencl_print_binary (struct clast_binary *b, opencl_main code_gen)
{
- int num_of_ops = gimple_num_ops (gmp);
- tree lhs;
- tree rhs1;
- tree rhs2;
- bool addr_expr;
- int result;
- tree result_size = NULL;
-
- if (gimple_assign_rhs_code (gmp) == MAX_EXPR
- || gimple_assign_rhs_code (gmp) == MIN_EXPR)
- {
- opencl_print_max_min_assign (gmp, code_gen);
- return;
- }
- gcc_assert (num_of_ops == 2 || num_of_ops == 3);
- lhs = gimple_assign_lhs (gmp);
-
- addr_expr = (TREE_CODE (TREE_TYPE (lhs)) == POINTER_TYPE);
- if (addr_expr)
- result_size = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (lhs)));
-
- rhs1 = gimple_assign_rhs1 (gmp);
- rhs2 = gimple_assign_rhs2 (gmp);
- result = opencl_print_operand (lhs, true, code_gen);
- if (result != 0)
- return;
- opencl_append_string_to_body (" = ", code_gen);
+ const char *s1 = NULL, *s2 = NULL, *s3 = NULL;
+ bool group = (b->LHS->type == clast_expr_red
+ && ((struct clast_reduction*) b->LHS)->n > 1);
- if (addr_expr)
- opencl_print_addr_operand (rhs1, result_size, code_gen);
- else
- {
- if (rhs2 == NULL)
- opencl_print_unary (gmp, code_gen);
- opencl_print_operand (rhs1, false, code_gen);
- }
- if (rhs2 != NULL_TREE)
+ switch (b->type)
{
- opencl_print_gimple_assign_operation (gmp, code_gen);
- if (addr_expr)
- opencl_print_addr_operand (rhs2, result_size, code_gen);
+ case clast_bin_fdiv:
+ s1 = "floor ((", s2 = ")/(", s3 = "))";
+ break;
+ case clast_bin_cdiv:
+ s1 = "ceil ((", s2 = ")/(", s3 = "))";
+ break;
+ case clast_bin_div:
+ if (group)
+ s1 = "(", s2 = ")/", s3 = "";
else
- opencl_print_operand (rhs2, false, code_gen);
- }
- opencl_append_string_to_body (";\n",code_gen);
-}
-
-/* Generate code for arguments for gimple call statement GMP.
- CODE_GEN hold information related to OpenCL code generation. */
-
-static void
-opencl_print_gimple_call_args (opencl_main code_gen, gimple gmp)
-{
- size_t len = gimple_call_num_args (gmp);
- size_t i;
- opencl_append_string_to_body (" (",code_gen);
- for (i = 0; i < len; i++)
- {
- opencl_print_operand (gimple_call_arg (gmp, i), false, code_gen);
- if (i < len - 1)
- opencl_append_string_to_body (", ",code_gen);
+ s1 = "", s2 = "/", s3 = "";
+ break;
+ case clast_bin_mod:
+ if (group)
+ s1 = "(", s2 = ")%", s3 = "";
+ else
+ s1 = "", s2 = "%", s3 = "";
+ break;
}
- opencl_append_string_to_body (")",code_gen);
-}
-/* Replace some function names. */
-
-static const char *
-opencl_get_function_name (tree function)
-{
- const char *gimple_name = IDENTIFIER_POINTER (DECL_NAME (function));
- if (!strcmp (gimple_name, "__builtin_powf"))
- return "pow";
- return gimple_name;
+ opencl_append_string_to_body (s1, code_gen);
+ opencl_print_expr (b->LHS, code_gen);
+ opencl_append_string_to_body (s2, code_gen);
+ opencl_append_num_to_body (code_gen, mpz_get_si (b->RHS), "%d");
+ opencl_append_string_to_body (s3, code_gen);
}
-/* Generate code for gimple call statement GMP. CODE_GEN holds information
+/* Generate code for clast expression E. CODE_GEN holds information
related to OpenCL code generation. */
static void
-opencl_print_gimple_call (opencl_main code_gen, gimple gmp)
-{
- tree lhs = gimple_call_lhs (gmp);
- tree function = gimple_call_fn (gmp);
- opencl_print_operand (lhs, true, code_gen);
- opencl_append_string_to_body (" = ", code_gen);
-
- while (TREE_CODE (function) == ADDR_EXPR
- || TREE_CODE (function) == INDIRECT_REF)
- function = TREE_OPERAND (function, 0);
- opencl_append_string_to_body (opencl_get_function_name (function), code_gen);
- opencl_print_gimple_call_args (code_gen, gmp);
- opencl_append_string_to_body (";\n",code_gen);
-}
-
-/* Generate code for gimple statment SMP. Now only assignment
- operation are supported, but it seems enough for clast translation.
- GIMPLE_COND statements are loop bound conditions and can be safely
- ignored. CODE_GEN holds information related to OpenCL code
- generation. */
-
-static void
-opencl_print_gimple (gimple gmp, opencl_main code_gen)
+opencl_print_expr (struct clast_expr *e, opencl_main code_gen)
{
- if (!gmp)
+ if (!e)
return;
-
- switch (gimple_code (gmp))
+ switch (e->type)
{
- case GIMPLE_ASSIGN:
- opencl_print_gimple_assign (gmp, code_gen);
- break;
- case GIMPLE_COND:
- break;
- case GIMPLE_PHI:
- break;
- case GIMPLE_CALL:
- opencl_print_gimple_call (code_gen, gmp);
+ case clast_expr_term:
+ opencl_print_term ((struct clast_term*) e, code_gen);
break;
- case GIMPLE_DEBUG:
+ case clast_expr_red:
+ opencl_print_reduction ((struct clast_reduction*) e, code_gen);
break;
- case GIMPLE_LABEL:
- {
- tree label = gimple_label_label (gmp);
- opencl_print_operand (label, false, code_gen);
- opencl_append_string_to_body (": ", code_gen);
- }
+ case clast_expr_bin:
+ opencl_print_binary ((struct clast_binary*) e, code_gen);
break;
default:
- debug_gimple_stmt (gmp);
gcc_unreachable ();
}
}
-/* Get name of the variable, represented by tree NODE. If variable is
- temporary, generate name for it. */
-
-static const char *
-opencl_get_var_name (tree node)
-{
- bool ssa_name = TREE_CODE (node) == SSA_NAME;
- tree name;
- int num = 0;
- if (ssa_name)
- {
- num = SSA_NAME_VERSION (node);
- node = SSA_NAME_VAR (node);
- }
- name = DECL_NAME (node);
- if (name)
- {
- if (!ssa_name)
- return identifier_to_locale (IDENTIFIER_POINTER (name));
- else
- {
- const char *base = identifier_to_locale (IDENTIFIER_POINTER (name));
- char *buff = XNEWVEC (char, strlen (base) + 5);
- sprintf (buff, "%s_%d", base, num);
- return buff;
- }
- }
- else
- {
- int tmp_var_uid = DECL_UID (node);
- char *tmp = XNEWVEC (char, 30);
- sprintf (tmp, "opencl_var_%d_%d", tmp_var_uid, num);
- return tmp;
- }
-}
-
-/* Append variable name NAME to function body. Differs from appending
- string by replacing `.' by `_'. CODE_GEN holds information related
- to OpenCL code generation. */
+/* Generate OpenCL code for clast_assignment A.
+ CODE_GEN holds information related to OpenCL code generation. */
static void
-opencl_append_var_name (const char *name, opencl_main code_gen)
+opencl_print_assignment (struct clast_assignment *a, opencl_main code_gen)
{
- int len = strlen (name);
- char *tmp = XNEWVEC (char, len + 1);
- int i;
- for (i = 0; i <= len; i++)
+ /* Real assignment. */
+ if (a->LHS)
{
- char tt = name[i];
- if (tt == '.')
- tt = '_';
- tmp[i] = tt;
+ opencl_append_string_to_body (a->LHS, code_gen);
+ opencl_append_string_to_body (" = ", code_gen);
}
- opencl_append_string_to_body (tmp, code_gen);
- free (tmp);
+
+ /* Just expression. */
+ opencl_print_expr (a->RHS, code_gen);
}
-/* If variable VAR_DECL is not defined and it is not marked as a
- parameter, mark it as a parameter and add it to parameters list.
+/* Print operation simbol (`+' `-' `*') for assignment operation GMA.
CODE_GEN holds information related to OpenCL code generation. */
static void
-opencl_try_variable (opencl_main code_gen, tree var_decl)
+opencl_print_gimple_assign_operation (gimple gmp, opencl_main code_gen)
{
- const char *name = opencl_get_var_name (var_decl);
- gcc_assert (code_gen->defined_vars);
-
- if (check_and_mark_arg (code_gen, name, false))
- opencl_add_function_arg (code_gen, var_decl, name);
+ opencl_append_string_to_body
+ (op_symbol_code (gimple_assign_rhs_code (gmp)), code_gen);
}
-/* Define non scalar variable, represented be DATA as either local
- variable or kernel argument. CODE_GEN holds information related to
- OpenCL code generation. */
+/* Generate definition for non scalar variable VAR and place it to
+ string DEST. Use DECL_NAME as variable name. */
static void
-opencl_add_non_scalar_function_arg (opencl_main code_gen,
- opencl_data data)
+opencl_add_non_scalar_type_decl (tree var, dyn_string_t dest,
+ const char *decl_name)
{
- const char *decl;
- static int counter = 0;
- opencl_body body = code_gen->current_body;
- tree var = data->exact_object;
- const char *name = opencl_get_var_name (var);
tree type = TREE_TYPE (var);
+ const char *name = opencl_get_var_name (var);
+ static int counter = 0;
+ char type_name [30];
+ char *tmp_name = xstrdup (name);
+ const char *new_type;
+ tree inner_type = TREE_TYPE (type);
- /* Check whether given variable can be privatized. */
- if (data->privatized)
- {
- /* Define variable as local variable. */
- gcc_assert (TREE_CODE (type) == ARRAY_TYPE);
- decl = opencl_print_function_arg_with_type (name, type);
- dyn_string_append_cstr (body->pre_header, decl);
- dyn_string_append_cstr (body->pre_header, ";\n");
- return;
- }
- else
- {
- /* Define variable as kernel argument. */
- char decl_name [30];
- tree main_type = opencl_get_main_type (type);
- sprintf (decl_name, "oclFTmpArg%d", counter++);
- decl = opencl_print_function_arg_with_type (decl_name, main_type);
- dyn_string_append_cstr (body->non_scalar_args, "__global ");
- opencl_add_non_scalar_type_decl (var, body->pre_header, decl_name);
- dyn_string_append_cstr (body->non_scalar_args, decl);
- dyn_string_append_cstr (body->non_scalar_args, ", ");
- VEC_safe_push (opencl_data, heap, body->data_refs, data);
- }
-}
+ filter_dots (tmp_name);
-/* Register data reference REF to variable DATA. Do nothing, if it
- has already been registered. CODE_GEN holds information related to
- OpenCL code generation. */
+ sprintf (type_name, "oclFTmpType%d", counter++);
-static void
-opencl_try_data_ref (opencl_main code_gen, data_reference_p ref,
- opencl_data data)
-{
- tree var = dr_outermost_base_object (ref);
- const char *name = opencl_get_var_name (var);
- const char ** slot;
- gcc_assert (code_gen->defined_vars);
+ new_type = opencl_print_function_arg_with_type (type_name, inner_type);
- slot = (const char **)htab_find_slot (code_gen->global_defined_vars,
- name, INSERT);
- if (*slot)
- return;
- *slot = name;
- opencl_add_non_scalar_function_arg (code_gen, data);
+ dyn_string_append_cstr (dest, "typedef __global ");
+ dyn_string_append_cstr (dest, new_type);
+ dyn_string_append_cstr (dest, ";\n");
+
+ dyn_string_append_cstr (dest, type_name);
+ dyn_string_append_cstr (dest, " *");
+ dyn_string_append_cstr (dest, tmp_name);
+ if (decl_name != NULL)
+ {
+ dyn_string_append_cstr (dest, " = (");
+ dyn_string_append_cstr (dest, type_name);
+ dyn_string_append_cstr (dest, "*)");
+ dyn_string_append_cstr (dest, decl_name);
+ dyn_string_append_cstr (dest, ";\n");
+ }
+ free (tmp_name);
}
-/* Register data reference D_REF in current kernel. CODE_GEN hold
- information related to OpenCL code generation. */
+/* Append variable VAR with name VAR_NAME to current function body.
+ If variable has been defined in current scope, but definition for
+ it has not been generated - then generate it's definition and mark
+ variable as defined. CODE_GEN holds information related to OpenCL
+ code generation. */
static void
-opencl_add_data_ref (opencl_main code_gen, data_reference_p d_ref)
+opencl_add_variable (const char *var_name, tree var, opencl_main code_gen)
{
- opencl_data tmp = opencl_get_data_by_data_ref (code_gen, d_ref);
-
- gcc_assert (tmp);
- if (!DR_IS_READ (d_ref))
+ const char **slot;
+ if (htab_find (code_gen->global_defined_vars, var_name))
{
- bitmap_set_bit (code_gen->curr_meta->modified_on_device, tmp->id);
- tmp->written_in_current_body = true;
- tmp->ever_written_on_device = true;
- code_gen->current_body->num_of_data_writes ++;
+ opencl_append_var_name (var_name, code_gen);
+ return;
}
- else
+
+ slot = (const char **) htab_find_slot
+ (code_gen->defined_vars, var_name, INSERT);
+
+ if (! (*slot) && defined_in_sese_p (var, code_gen->region))
{
- tmp->read_in_current_body = true;
- tmp->ever_read_on_device = true;
+ const char *decl;
+ tree type = TREE_TYPE (var);
+ *slot = var_name;
+ if (TREE_CODE (type) == POINTER_TYPE
+ || TREE_CODE (type) == ARRAY_TYPE)
+ opencl_add_non_scalar_type_decl (var, code_gen->current_body->body,
+ NULL);
+ else
+ {
+ var = SSA_NAME_VAR (var);
+ decl = opencl_print_function_arg_with_type (var_name, type);
+ opencl_append_string_to_body (decl, code_gen);
+ }
+ return;
}
- if (!tmp->privatized)
- tmp->used_on_device = true;
-
- opencl_try_data_ref (code_gen, d_ref, tmp);
+ opencl_append_var_name (var_name, code_gen);
}
-/* Add base objects of all data references in PBB as arguments to
- current kernel. CODE_GEN holds information related to OpenCL code
- generation. */
+/* If variable VAR_DECL is not defined and it is not marked as a
+ parameter, mark it as a parameter and add it to parameters list.
+ CODE_GEN holds information related to OpenCL code generation. */
static void
-opencl_add_data_refs_pbb (poly_bb_p pbb, opencl_main code_gen)
+opencl_try_variable (opencl_main code_gen, tree var_decl)
{
- VEC (poly_dr_p, heap) *drs = PBB_DRS (pbb);
- int i;
- poly_dr_p curr;
+ const char *name = opencl_get_var_name (var_decl);
+ gcc_assert (code_gen->defined_vars);
- for (i = 0; VEC_iterate (poly_dr_p, drs, i, curr); i++)
- {
- data_reference_p d_ref = (data_reference_p) PDR_CDR (curr);
- opencl_add_data_ref (code_gen, d_ref);
- }
+ if (check_and_mark_arg (code_gen, name, false))
+ opencl_add_function_arg (code_gen, var_decl, name);
}
/* Generate operand for tree node NODE. If LSH is true, generated
@@ -1422,237 +1266,437 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
return 0;
}
-/* Append variable VAR with name VAR_NAME to current function body.
- If variable has been defined in current scope, but definition for
- it has not been generated - then generate it's definition and mark
- variable as defined. CODE_GEN holds information related to OpenCL
- code generation. */
+/* Generate code for min or max gimple operand GMP. CODE_GEN holds
+ information related to OpenCL code generation. */
static void
-opencl_add_variable (const char *var_name, tree var, opencl_main code_gen)
+opencl_print_max_min_assign (gimple gmp, opencl_main code_gen)
{
- const char ** slot;
- if (htab_find (code_gen->global_defined_vars, var_name))
+ tree lhs = gimple_assign_lhs (gmp);
+ tree rhs1 = gimple_assign_rhs1 (gmp);
+ tree rhs2 = gimple_assign_rhs2 (gmp);
+ bool max = gimple_assign_rhs_code (gmp) == MAX_EXPR;
+
+ opencl_print_operand (lhs, true, code_gen);
+ opencl_append_string_to_body (max?" = fmax (":"= fmin (", code_gen);
+ opencl_print_operand (rhs1, false, code_gen);
+ opencl_append_string_to_body (",", code_gen);
+ opencl_print_operand (rhs2, false, code_gen);
+ opencl_append_string_to_body (");\n", code_gen);
+}
+
+/* Print pointer expression represented by EXPR. TYPE_SIZE represents
+ size of the base type for EXPR. CODE_GEN holds information related
+ to OpenCL code generation. */
+
+static void
+opencl_print_addr_operand (tree expr, tree type_size, opencl_main code_gen)
+{
+ if (TREE_CODE (TREE_TYPE (expr)) != POINTER_TYPE)
{
- opencl_append_var_name (var_name, code_gen);
- return;
+ opencl_append_string_to_body ("(", code_gen);
+ opencl_print_operand (expr, false, code_gen);
+ opencl_append_string_to_body ("/", code_gen);
+ opencl_print_operand (type_size, false, code_gen);
+ opencl_append_string_to_body (")", code_gen);
}
+ else
+ opencl_print_operand (expr, false, code_gen);
+}
- slot = (const char **) htab_find_slot
- (code_gen->defined_vars, var_name, INSERT);
+/* Print unary gimple operation GMP. CODE_GEN holds information
+ related to OpenCL code generation. */
- if (! (*slot) && defined_in_sese_p (var, code_gen->region))
+static void
+opencl_print_unary (gimple gmp, opencl_main code_gen)
+{
+ switch (gimple_assign_rhs_code (gmp))
{
- const char *decl;
- tree type = TREE_TYPE (var);
- *slot = var_name;
- if (TREE_CODE (type) == POINTER_TYPE
- || TREE_CODE (type) == ARRAY_TYPE)
- {
- opencl_add_non_scalar_type_decl (var, code_gen->current_body->body,
- NULL);
- }
- else
- {
- var = SSA_NAME_VAR (var);
- decl = opencl_print_function_arg_with_type (var_name, type);
- opencl_append_string_to_body (decl, code_gen);
- }
+ case BIT_NOT_EXPR:
+ opencl_append_string_to_body ("~", code_gen);
+ return;
+ case TRUTH_NOT_EXPR:
+ opencl_append_string_to_body ("!", code_gen);
+ return;
+ case NEGATE_EXPR:
+ opencl_append_string_to_body ("-", code_gen);
+ return;
+ case MODIFY_EXPR:
+ default:
return;
}
- opencl_append_var_name (var_name, code_gen);
}
-/* Append list of names of loop iterators from CODE_GEN with same type
- TYPE to current kernel. FIRST and LAST define outermost and
- innermost iterators to append respectively. */
+/* Generate code for gimple assignment statement GMP. CODE_GEN holds
+ information related to OpenCL code generation. */
static void
-opencl_print_local_vars (const char *fist, const char *last,
- const char *type, opencl_main code_gen)
+opencl_print_gimple_assign (gimple gmp, opencl_main code_gen)
{
- char **names = cloog_names_scattering (code_gen->root_names);
- int len = cloog_names_nb_scattering (code_gen->root_names);
- int i;
- for (i = 0; i < len; i++)
+ int num_of_ops = gimple_num_ops (gmp);
+ tree lhs;
+ tree rhs1;
+ tree rhs2;
+ bool addr_expr;
+ int result;
+ tree result_size = NULL;
+
+ if (gimple_assign_rhs_code (gmp) == MAX_EXPR
+ || gimple_assign_rhs_code (gmp) == MIN_EXPR)
{
- const char *tmp = names[i];
- if (opencl_cmp_scat (fist, tmp) <= 0
- && opencl_cmp_scat (last, tmp) >= 0)
- {
- const char ** slot =
- (const char **) htab_find_slot (code_gen->global_defined_vars,
- tmp, INSERT);
- *slot = tmp;
- continue;
- }
+ opencl_print_max_min_assign (gmp, code_gen);
+ return;
+ }
+ gcc_assert (num_of_ops == 2 || num_of_ops == 3);
+ lhs = gimple_assign_lhs (gmp);
- if (opencl_cmp_scat (fist, tmp) > 0)
- continue;
+ addr_expr = (TREE_CODE (TREE_TYPE (lhs)) == POINTER_TYPE);
+ if (addr_expr)
+ result_size = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (lhs)));
- opencl_append_string_to_body (type, code_gen);
- opencl_append_string_to_body (" ", code_gen);
- opencl_append_string_to_body (tmp, code_gen);
- opencl_append_string_to_body (";\n", code_gen);
- *((const char **)htab_find_slot (code_gen->global_defined_vars,
- tmp, INSERT)) = tmp;
+ rhs1 = gimple_assign_rhs1 (gmp);
+ rhs2 = gimple_assign_rhs2 (gmp);
+ result = opencl_print_operand (lhs, true, code_gen);
+ if (result != 0)
+ return;
+ opencl_append_string_to_body (" = ", code_gen);
+
+ if (addr_expr)
+ opencl_print_addr_operand (rhs1, result_size, code_gen);
+ else
+ {
+ if (rhs2 == NULL)
+ opencl_print_unary (gmp, code_gen);
+ opencl_print_operand (rhs1, false, code_gen);
+ }
+ if (rhs2 != NULL_TREE)
+ {
+ opencl_print_gimple_assign_operation (gmp, code_gen);
+ if (addr_expr)
+ opencl_print_addr_operand (rhs2, result_size, code_gen);
+ else
+ opencl_print_operand (rhs2, false, code_gen);
}
+ opencl_append_string_to_body (";\n",code_gen);
}
-/* Replace all dots to underscores in string pointed to by P. Return P. */
+/* Generate code for arguments for gimple call statement GMP.
+ CODE_GEN hold information related to OpenCL code generation. */
-static char *
-filter_dots (char *p)
+static void
+opencl_print_gimple_call_args (opencl_main code_gen, gimple gmp)
{
- char *s;
- for (s = p; *s; s++)
- if (*s == '.')
- *s = '_';
- return p;
+ size_t len = gimple_call_num_args (gmp);
+ size_t i;
+ opencl_append_string_to_body (" (",code_gen);
+ for (i = 0; i < len; i++)
+ {
+ opencl_print_operand (gimple_call_arg (gmp, i), false, code_gen);
+ if (i < len - 1)
+ opencl_append_string_to_body (", ",code_gen);
+ }
+ opencl_append_string_to_body (")",code_gen);
}
-/* Return string with varibale definition. ARG_NAME is the name of
- the variable and TYPE is it's type. */
+/* Replace some function names. */
static const char *
-opencl_print_function_arg_with_type (const char *arg_name, tree type)
+opencl_get_function_name (tree function)
{
- const char *decl = gen_type_with_name (arg_name, type);
- char *ddecl;
- ddecl = xstrdup (decl);
- return filter_dots (ddecl);
+ const char *gimple_name = IDENTIFIER_POINTER (DECL_NAME (function));
+ if (!strcmp (gimple_name, "__builtin_powf"))
+ return "pow";
+ return gimple_name;
}
-/* Generate definition for non scalar variable VAR and place it to
- string DEST. Use DECL_NAME as variable name. */
+/* Generate code for gimple call statement GMP. CODE_GEN holds information
+ related to OpenCL code generation. */
static void
-opencl_add_non_scalar_type_decl (tree var, dyn_string_t dest,
- const char *decl_name)
+opencl_print_gimple_call (opencl_main code_gen, gimple gmp)
{
- tree type = TREE_TYPE (var);
- const char *name = opencl_get_var_name (var);
- static int counter = 0;
- char type_name [30];
- char *tmp_name = xstrdup (name);
- const char *new_type;
- tree inner_type = TREE_TYPE (type);
+ tree lhs = gimple_call_lhs (gmp);
+ tree function = gimple_call_fn (gmp);
+ opencl_print_operand (lhs, true, code_gen);
+ opencl_append_string_to_body (" = ", code_gen);
- filter_dots (tmp_name);
+ while (TREE_CODE (function) == ADDR_EXPR
+ || TREE_CODE (function) == INDIRECT_REF)
+ function = TREE_OPERAND (function, 0);
+ opencl_append_string_to_body (opencl_get_function_name (function), code_gen);
+ opencl_print_gimple_call_args (code_gen, gmp);
+ opencl_append_string_to_body (";\n",code_gen);
+}
- sprintf (type_name, "oclFTmpType%d", counter++);
+/* Generate code for gimple statment SMP. Now only assignment
+ operation are supported, but it seems enough for clast translation.
+ GIMPLE_COND statements are loop bound conditions and can be safely
+ ignored. CODE_GEN holds information related to OpenCL code
+ generation. */
- new_type = opencl_print_function_arg_with_type (type_name, inner_type);
+static void
+opencl_print_gimple (gimple gmp, opencl_main code_gen)
+{
+ if (!gmp)
+ return;
- dyn_string_append_cstr (dest, "typedef __global ");
- dyn_string_append_cstr (dest, new_type);
- dyn_string_append_cstr (dest, ";\n");
+ switch (gimple_code (gmp))
+ {
+ case GIMPLE_ASSIGN:
+ opencl_print_gimple_assign (gmp, code_gen);
+ break;
+ case GIMPLE_COND:
+ break;
+ case GIMPLE_PHI:
+ break;
+ case GIMPLE_CALL:
+ opencl_print_gimple_call (code_gen, gmp);
+ break;
+ case GIMPLE_DEBUG:
+ break;
+ case GIMPLE_LABEL:
+ {
+ tree label = gimple_label_label (gmp);
+ opencl_print_operand (label, false, code_gen);
+ opencl_append_string_to_body (": ", code_gen);
+ }
+ break;
+ default:
+ debug_gimple_stmt (gmp);
+ gcc_unreachable ();
+ }
+}
- dyn_string_append_cstr (dest, type_name);
- dyn_string_append_cstr (dest, " *");
- dyn_string_append_cstr (dest, tmp_name);
- if (decl_name != NULL)
+static void opencl_build_defines (tree, opencl_main);
+
+/* For a given gimple statement STMT build definition for all names,
+ used in this stament. If name has been defined in other sese, mark
+ it as kernel parameter. CODE_GEN holds information related to
+ OpenCL code generation. */
+
+static void
+opencl_expand_scalar_vars (opencl_main code_gen, gimple stmt)
+{
+ ssa_op_iter iter;
+ use_operand_p use_p;
+ FOR_EACH_SSA_USE_OPERAND (use_p, stmt, iter, SSA_OP_ALL_USES)
{
- dyn_string_append_cstr (dest, " = (");
- dyn_string_append_cstr (dest, type_name);
- dyn_string_append_cstr (dest, "*)");
- dyn_string_append_cstr (dest, decl_name);
- dyn_string_append_cstr (dest, ";\n");
+ tree use = USE_FROM_PTR (use_p);
+ if (!is_gimple_reg (use))
+ continue;
+ opencl_build_defines (use, code_gen);
}
- free (tmp_name);
+}
+
+/* If tree node NODE defined in current sese build and insert define
+ statements for it, otherwise mark node as external (parameter for
+ kernel). If tree defined in current sese, also recursively build
+ defines for all trees in definition expression. */
+
+static void
+opencl_build_defines (tree node, opencl_main code_gen)
+{
+ enum tree_code code = TREE_CODE (node);
+ switch (code)
+ {
+ case SSA_NAME:
+ {
+ const char *tmp = opencl_get_var_name (node);
+ gimple def_stmt;
+
+ /* If name defined in other sese it is kernel's parameter. */
+ if (!defined_in_sese_p (node, code_gen->region))
+ return;
+
+ /* Bail out if this name was defined earlier either in this
+ or other region. */
+ if (*(const char **)htab_find_slot (code_gen->defined_vars,
+ tmp, INSERT))
+ return;
+
+ /* Get definition statement. */
+ def_stmt = SSA_NAME_DEF_STMT (node);
+ opencl_expand_scalar_vars (code_gen, def_stmt);
+ opencl_print_gimple (def_stmt, code_gen);
+ return;
+ }
+ case ARRAY_REF:
+ {
+ tree arr = TREE_OPERAND (node, 0);
+ tree offset = TREE_OPERAND (node, 1);
+ opencl_build_defines (arr, code_gen);
+ opencl_build_defines (offset, code_gen);
+ return;
+ }
+ default:
+ gcc_unreachable ();
+ }
+}
+/* Generate code for a single basic block BB. CODE_GEN holds
+ information related to OpenCL code generation. */
+
+static void
+opencl_print_bb (basic_block bb, opencl_main code_gen)
+{
+ gimple_stmt_iterator gsi;
+ for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
+ {
+ gimple stmt = gsi_stmt (gsi);
+ opencl_expand_scalar_vars (code_gen, stmt);
+ opencl_print_gimple (stmt, code_gen);
+ }
}
-/* Check whether variable with name NAME has been defined as global or
- local variable and mark it as defined. This function returns false
- if variable has already been defined, otherwise it returns true. */
+/* Define non scalar variable, represented be DATA as either local
+ variable or kernel argument. CODE_GEN holds information related to
+ OpenCL code generation. */
-static bool
-check_and_mark_arg (opencl_main code_gen, const char *name, bool local)
+static void
+opencl_add_non_scalar_function_arg (opencl_main code_gen,
+ opencl_data data)
{
- const char ** slot;
- gcc_assert (code_gen->defined_vars || !local);
- if (code_gen->defined_vars)
+ const char *decl;
+ static int counter = 0;
+ opencl_body body = code_gen->current_body;
+ tree var = data->exact_object;
+ const char *name = opencl_get_var_name (var);
+ tree type = TREE_TYPE (var);
+
+ /* Check whether given variable can be privatized. */
+ if (data->privatized)
{
- slot = (const char **)htab_find_slot (code_gen->defined_vars,
- name, INSERT);
- if (*slot)
- return false;
- if (local)
- *slot = name;
+ /* Define variable as local variable. */
+ gcc_assert (TREE_CODE (type) == ARRAY_TYPE);
+ decl = opencl_print_function_arg_with_type (name, type);
+ dyn_string_append_cstr (body->pre_header, decl);
+ dyn_string_append_cstr (body->pre_header, ";\n");
+ return;
}
+ else
+ {
+ /* Define variable as kernel argument. */
+ char decl_name [30];
+ tree main_type = opencl_get_main_type (type);
+ sprintf (decl_name, "oclFTmpArg%d", counter++);
+ decl = opencl_print_function_arg_with_type (decl_name, main_type);
+ dyn_string_append_cstr (body->non_scalar_args, "__global ");
+ opencl_add_non_scalar_type_decl (var, body->pre_header, decl_name);
+ dyn_string_append_cstr (body->non_scalar_args, decl);
+ dyn_string_append_cstr (body->non_scalar_args, ", ");
+ VEC_safe_push (opencl_data, heap, body->data_refs, data);
+ }
+}
+
+/* Register data reference REF to variable DATA. Do nothing, if it
+ has already been registered. CODE_GEN holds information related to
+ OpenCL code generation. */
+
+static void
+opencl_try_data_ref (opencl_main code_gen, data_reference_p ref,
+ opencl_data data)
+{
+ tree var = dr_outermost_base_object (ref);
+ const char *name = opencl_get_var_name (var);
+ const char **slot;
+ gcc_assert (code_gen->defined_vars);
slot = (const char **)htab_find_slot (code_gen->global_defined_vars,
name, INSERT);
if (*slot)
- return false;
- if (!local)
- *slot = name;
- return true;
+ return;
+ *slot = name;
+ opencl_add_non_scalar_function_arg (code_gen, data);
}
-/* Add variable VAR with name NAME as function argument. Append it's
- declaration in finction header and add it as function parameter.
- CODE_GEN holds information related to OpenCL code generation. */
+/* Register data reference D_REF in current kernel. CODE_GEN hold
+ information related to OpenCL code generation. */
static void
-opencl_add_function_arg (opencl_main code_gen, tree var, const char *name)
+opencl_add_data_ref (opencl_main code_gen, data_reference_p d_ref)
{
- opencl_body body;
- const char *decl;
- tree type;
- type = TREE_TYPE (var);
- body = code_gen->current_body;
- decl = opencl_print_function_arg_with_type (name, type);
- dyn_string_append_cstr (body->header, decl);
- dyn_string_append_cstr (body->header, ", ");
- VEC_safe_push (tree, heap, body->function_args, var);
+ opencl_data tmp = opencl_get_data_by_data_ref (code_gen, d_ref);
+
+ gcc_assert (tmp);
+ if (!DR_IS_READ (d_ref))
+ {
+ bitmap_set_bit (code_gen->curr_meta->modified_on_device, tmp->id);
+ tmp->written_in_current_body = true;
+ tmp->ever_written_on_device = true;
+ code_gen->current_body->num_of_data_writes ++;
+ }
+ else
+ {
+ tmp->read_in_current_body = true;
+ tmp->ever_read_on_device = true;
+ }
+ if (!tmp->privatized)
+ tmp->used_on_device = true;
+
+ opencl_try_data_ref (code_gen, d_ref, tmp);
}
-/* Generate kernel function code for clast for statement F, located on
- depth DEPTH. CODE_GEN holds information related to OpenCL code
+/* Add base objects of all data references in PBB as arguments to
+ current kernel. CODE_GEN holds information related to OpenCL code
generation. */
-opencl_body
-opencl_clast_to_kernel (struct clast_for * f, opencl_main code_gen,
- int depth)
+static void
+opencl_add_data_refs_pbb (poly_bb_p pbb, opencl_main code_gen)
{
- opencl_body tmp = opencl_body_create ();
- code_gen->current_body = tmp;
- return opencl_print_loop (f, code_gen, depth);
+ VEC (poly_dr_p, heap) *drs = PBB_DRS (pbb);
+ int i;
+ poly_dr_p curr;
+
+ for (i = 0; VEC_iterate (poly_dr_p, drs, i, curr); i++)
+ {
+ data_reference_p d_ref = (data_reference_p) PDR_CDR (curr);
+ opencl_add_data_ref (code_gen, d_ref);
+ }
}
-/* Generate code for clast statement S, located on depth DEPTH.
+/* Generate OpenCL code for user statement U. Code will be generated
+ from basic block, related to U. Also induction variables mapping
+ to old variables must be calculated to process basic block.
CODE_GEN holds information related to OpenCL code generation. */
static void
-opencl_print_stmt_list (struct clast_stmt *s, opencl_main code_gen, int depth)
+opencl_print_user_stmt (struct clast_user_stmt *u, opencl_main code_gen)
{
- for ( ; s; s = s->next) {
- gcc_assert (!CLAST_STMT_IS_A (s, stmt_root));
- if (CLAST_STMT_IS_A (s, stmt_ass))
- {
- opencl_print_assignment ((struct clast_assignment *) s, code_gen);
- opencl_append_string_to_body (";\n", code_gen);
- }
- else if (CLAST_STMT_IS_A (s, stmt_user))
- opencl_print_user_stmt ((struct clast_user_stmt *) s, code_gen);
- else if (CLAST_STMT_IS_A (s, stmt_for))
- opencl_print_for ((struct clast_for *) s, code_gen, depth);
- else if (CLAST_STMT_IS_A (s, stmt_guard))
- opencl_print_guard ((struct clast_guard *) s, code_gen, depth);
- else if (CLAST_STMT_IS_A (s, stmt_block))
- {
- opencl_append_string_to_body ("{\n", code_gen);
- opencl_print_stmt_list (((struct clast_block *)s)->body, code_gen,
- depth);
- opencl_append_string_to_body ("}\n", code_gen);
- }
- else
- gcc_unreachable ();
- }
+ CloogStatement *cs;
+ poly_bb_p pbb;
+ gimple_bb_p gbbp;
+ basic_block bb;
+ int i;
+ int nb_loops = number_of_loops ();
+ code_gen->iv_map = VEC_alloc (tree, heap, nb_loops);
+
+ for (i = 0; i < nb_loops; i++)
+ VEC_safe_push (tree, heap, code_gen->iv_map, NULL_TREE);
+ build_iv_mapping (code_gen->iv_map, code_gen->region,
+ code_gen->newivs,
+ code_gen->newivs_index, u,
+ code_gen->params_index);
+
+ code_gen->defined_vars
+ = htab_create (10, htab_hash_string, opencl_cmp_str, NULL);
+ opencl_append_string_to_body ("{\n", code_gen);
+
+ cs = u->statement;
+ pbb = (poly_bb_p) cloog_statement_usr (cs);
+ gbbp = PBB_BLACK_BOX (pbb);
+ bb = GBB_BB (gbbp);
+ code_gen->context_loop = bb->loop_father;
+
+ opencl_add_data_refs_pbb (pbb, code_gen);
+ opencl_print_bb (bb, code_gen);
+ opencl_append_string_to_body ("}\n", code_gen);
+ htab_delete (code_gen->defined_vars);
+ code_gen->defined_vars = NULL;
+ VEC_free (tree, heap, code_gen->iv_map);
}
+static void opencl_print_stmt_list (struct clast_stmt *, opencl_main, int);
+
/* Generate code for clast for statement F, locate on depth LEVEL.
CODE_GEN holds information related to OpenCL code generation. */
@@ -1712,6 +1756,22 @@ opencl_print_for (struct clast_for *f, opencl_main code_gen, int level)
opencl_append_string_to_body ("}\n", code_gen);
}
+/* Generate code for clast equation EQ. CODE_GEN holds information
+ related to OpenCL code generation. */
+
+static void
+opencl_print_equation (struct clast_equation *eq, opencl_main code_gen)
+{
+ opencl_print_expr (eq->LHS, code_gen);
+ if (eq->sign == 0)
+ opencl_append_string_to_body (" == ", code_gen);
+ else if (eq->sign > 0)
+ opencl_append_string_to_body (" >= ", code_gen);
+ else
+ opencl_append_string_to_body (" <= ", code_gen);
+ opencl_print_expr (eq->RHS, code_gen);
+}
+
/* Generate code for clast conditional statement G, locate on depth DEPTH.
CODE_GEN holds information related to OpenCL code generation. */
@@ -1737,186 +1797,88 @@ opencl_print_guard (struct clast_guard *g, opencl_main code_gen, int depth)
opencl_append_string_to_body ("}\n", code_gen);
}
-
-/* Generate code for clast equation EQ. CODE_GEN holds information
- related to OpenCL code generation. */
-
-static void
-opencl_print_equation (struct clast_equation *eq, opencl_main code_gen)
-{
- opencl_print_expr (eq->LHS, code_gen);
- if (eq->sign == 0)
- opencl_append_string_to_body (" == ", code_gen);
- else if (eq->sign > 0)
- opencl_append_string_to_body (" >= ", code_gen);
- else
- opencl_append_string_to_body (" <= ", code_gen);
- opencl_print_expr (eq->RHS, code_gen);
-}
-
-/* Generate code for clast expression E. CODE_GEN holds information
- related to OpenCL code generation. */
+/* Generate code for clast statement S, located on depth DEPTH.
+ CODE_GEN holds information related to OpenCL code generation. */
static void
-opencl_print_expr (struct clast_expr *e, opencl_main code_gen)
+opencl_print_stmt_list (struct clast_stmt *s, opencl_main code_gen, int depth)
{
- if (!e)
- return;
- switch (e->type)
- {
- case clast_expr_term:
- opencl_print_term ((struct clast_term*) e, code_gen);
- break;
- case clast_expr_red:
- opencl_print_reduction ((struct clast_reduction*) e, code_gen);
- break;
- case clast_expr_bin:
- opencl_print_binary ((struct clast_binary*) e, code_gen);
- break;
- default:
+ for ( ; s; s = s->next) {
+ gcc_assert (!CLAST_STMT_IS_A (s, stmt_root));
+ if (CLAST_STMT_IS_A (s, stmt_ass))
+ {
+ opencl_print_assignment ((struct clast_assignment *) s, code_gen);
+ opencl_append_string_to_body (";\n", code_gen);
+ }
+ else if (CLAST_STMT_IS_A (s, stmt_user))
+ opencl_print_user_stmt ((struct clast_user_stmt *) s, code_gen);
+ else if (CLAST_STMT_IS_A (s, stmt_for))
+ opencl_print_for ((struct clast_for *) s, code_gen, depth);
+ else if (CLAST_STMT_IS_A (s, stmt_guard))
+ opencl_print_guard ((struct clast_guard *) s, code_gen, depth);
+ else if (CLAST_STMT_IS_A (s, stmt_block))
+ {
+ opencl_append_string_to_body ("{\n", code_gen);
+ opencl_print_stmt_list (((struct clast_block *)s)->body, code_gen,
+ depth);
+ opencl_append_string_to_body ("}\n", code_gen);
+ }
+ else
gcc_unreachable ();
- }
+ }
}
-/* Generate code for clast term T. CODE_GEN holds information
- related to OpenCL code generation. */
+/* Generate code for loop statement F. DEPTH is the depth of F in
+ current loop nest. CODE_GEN holds information related to OpenCL
+ code generation. */
-static void
-opencl_print_term (struct clast_term *t, opencl_main code_gen)
+static opencl_body
+opencl_print_loop (struct clast_for *f, opencl_main code_gen, int depth)
{
- if (t->var)
- {
- const char *real_name = opencl_get_scat_real_name (code_gen, t->var);
-
- if (mpz_cmp_si (t->val, 1) == 0)
- opencl_append_var_name (real_name, code_gen);
- else if (mpz_cmp_si (t->val, -1) == 0)
- {
- opencl_append_string_to_body ("-", code_gen);
- opencl_append_var_name (real_name, code_gen);
- }
- else
- {
- opencl_append_num_to_body (code_gen, mpz_get_si (t->val), "%d");
- opencl_append_string_to_body ("*", code_gen);
- opencl_append_var_name (real_name, code_gen);
- }
- opencl_add_scat_as_arg (code_gen, t->var, real_name);
- }
- else
- opencl_append_num_to_body (code_gen, mpz_get_si (t->val), "%d");
-}
-
-/* Generate code for clast reduction statement R. CODE_GEN holds
- information related to OpenCL code generation. */
+ opencl_body current_body = code_gen->current_body;
-static void
-opencl_print_reduction (struct clast_reduction *r, opencl_main code_gen)
-{
- switch (r->type)
- {
- case clast_red_sum:
- opencl_print_sum (r, code_gen);
- break;
- case clast_red_min:
- case clast_red_max:
- if (r->n == 1)
- {
- opencl_print_expr (r->elts[0], code_gen);
- break;
- }
- opencl_print_minmax_c (r, code_gen);
- break;
- default:
- gcc_unreachable ();
- }
-}
+ code_gen->global_defined_vars
+ = htab_create (10, htab_hash_string, opencl_cmp_str, NULL);
-/* Generate code for clast sum statement R. CODE_GEN holds information
- related to OpenCL code generation. */
+ opencl_perfect_nested_to_kernel (code_gen, f, current_body, depth);
-static void
-opencl_print_sum (struct clast_reduction *r, opencl_main code_gen)
-{
- int i;
- struct clast_term *t;
+ /* Define local loop iterators. */
+ opencl_print_local_vars (current_body->first_iter,
+ current_body->last_iter,
+ "unsigned int", code_gen);
- gcc_assert (r->n >= 1 && r->elts[0]->type == clast_expr_term);
- t = (struct clast_term *) r->elts[0];
- opencl_print_term (t, code_gen);
+ /* Generate code for kernel body. */
+ opencl_print_stmt_list (current_body->clast_body, code_gen, depth + 1);
+ opencl_append_string_to_body ("}\n", code_gen);
- for (i = 1; i < r->n; ++i)
+ if (current_body->num_of_data_writes)
{
- gcc_assert (r->elts[i]->type == clast_expr_term);
- t = (struct clast_term *) r->elts[i];
- if (mpz_sgn (t->val) > 0)
- opencl_append_string_to_body ("+", code_gen);
- opencl_print_term (t, code_gen);
- }
-}
+ dyn_string_t header = current_body->header;
+ int offset;
-/* Generate code for clast binary operation B. CODE_GEN holds
- information related to OpenCL code generation. */
+ dyn_string_append (header, current_body->non_scalar_args);
+ offset = dyn_string_length (header) - 2;
-static void
-opencl_print_binary (struct clast_binary *b, opencl_main code_gen)
-{
- const char *s1 = NULL, *s2 = NULL, *s3 = NULL;
- bool group = (b->LHS->type == clast_expr_red
- && ((struct clast_reduction*) b->LHS)->n > 1);
+ if (*(dyn_string_buf (header) + offset) == ',')
+ *(dyn_string_buf (header) + offset) = ' ';
- switch (b->type)
- {
- case clast_bin_fdiv:
- s1 = "floor ((", s2 = ")/(", s3 = "))";
- break;
- case clast_bin_cdiv:
- s1 = "ceil ((", s2 = ")/(", s3 = "))";
- break;
- case clast_bin_div:
- if (group)
- s1 = "(", s2 = ")/", s3 = "";
- else
- s1 = "", s2 = "/", s3 = "";
- break;
- case clast_bin_mod:
- if (group)
- s1 = "(", s2 = ")%", s3 = "";
- else
- s1 = "", s2 = "%", s3 = "";
- break;
+ opencl_append_string_to_header (")\n{\n", code_gen);
}
- opencl_append_string_to_body (s1, code_gen);
- opencl_print_expr (b->LHS, code_gen);
- opencl_append_string_to_body (s2, code_gen);
- opencl_append_num_to_body (code_gen, mpz_get_si (b->RHS), "%d");
- opencl_append_string_to_body (s3, code_gen);
+ return current_body;
}
-/* Generate code for clast min/max operation R. CODE_GEN holds
- information related to OpenCL code generation. */
+/* Generate kernel function code for clast for statement F, located on
+ depth DEPTH. CODE_GEN holds information related to OpenCL code
+ generation. */
-static void
-opencl_print_minmax_c ( struct clast_reduction *r, opencl_main code_gen)
+opencl_body
+opencl_clast_to_kernel (struct clast_for *f, opencl_main code_gen,
+ int depth)
{
- int i;
- for (i = 1; i < r->n; ++i)
- opencl_append_string_to_body (r->type == clast_red_max ? "max (" : "min (",
- code_gen);
- if (r->n > 0)
- {
- opencl_append_string_to_body ("(unsigned int)(", code_gen);
- opencl_print_expr (r->elts[0], code_gen);
- opencl_append_string_to_body (")", code_gen);
- }
- for (i = 1; i < r->n; ++i)
- {
- opencl_append_string_to_body (",", code_gen);
- opencl_append_string_to_body ("(unsigned int)(", code_gen);
- opencl_print_expr (r->elts[i], code_gen);
- opencl_append_string_to_body ("))", code_gen);
- }
+ opencl_body tmp = opencl_body_create ();
+ code_gen->current_body = tmp;
+ return opencl_print_loop (f, code_gen, depth);
}
#endif
@@ -308,7 +308,6 @@ opencl_supported_type_access_p (opencl_main code_gen, basic_block bb)
return true;
}
-
/* Mark variable, represented by tree OBJ as visited in bitmap VISITED.
If DEF is true and given variable can be privatized, mark it as
privatized in META. CODE_GEN holds information about non
@@ -111,27 +111,6 @@ enum OPENCL_FUNCTUONS
WAIT_FOR_EVENTS = 18
};
-/* Constructors and destructors. */
-static opencl_main opencl_main_create (CloogNames *, sese, edge, htab_t);
-static void opencl_main_delete (opencl_main);
-static void opencl_clast_meta_delete (opencl_clast_meta);
-static tree opencl_create_function_decl (enum OPENCL_FUNCTUONS);
-static edge opencl_create_function_call (edge);
-static void opencl_init_data (scop_p, opencl_main);
-static int opencl_get_non_scalar_type_depth (tree);
-static tree opencl_create_memory_for_pointer (opencl_data);
-static void opencl_init_basic_blocks (opencl_main);
-static edge opencl_set_context_properties (edge, tree);
-static tree opencl_create_clCreateContextFromType (tree);
-static tree opencl_create_clGetContextInfo_1 (tree);
-static void opencl_create_gimple_variables (void);
-static tree opencl_create_clCreateCommandQueue (tree);
-static tree opencl_create_malloc_call (tree);
-static edge opencl_create_init_context (edge);
-static void opencl_wait_for_event (opencl_main, tree);
-static void opencl_transform_stmt_list (struct clast_stmt *, opencl_main, int);
-static void opencl_create_gimple_for_body (opencl_body, opencl_main);
-
/* Data structure to be used in data_reference_p to opencl_data hash
table. */
struct map_ref_to_data_def
@@ -424,6 +403,21 @@ graphite_artificial_array_p (tree var)
return opencl_private_var_name_p (IDENTIFIER_POINTER (name));
}
+/* Get depth of type TYPE scalar (base) part. */
+
+static int
+opencl_get_non_scalar_type_depth (tree type)
+{
+ int count = 0;
+ while (TREE_CODE (type) == ARRAY_TYPE
+ || TREE_CODE (type) == POINTER_TYPE)
+ {
+ count++;
+ type = TREE_TYPE (type);
+ }
+ return count;
+}
+
/* Constructors & destructors.
<name>_create - creates a new object of such type and returns it.
<name>_delete - delete object (like destructor). */
@@ -510,21 +504,6 @@ opencl_main_delete (opencl_main data)
free (data);
}
-/* Get depth of type TYPE scalar (base) part. */
-
-static int
-opencl_get_non_scalar_type_depth (tree type)
-{
- int count = 0;
- while (TREE_CODE (type) == ARRAY_TYPE
- || TREE_CODE (type) == POINTER_TYPE)
- {
- count++;
- type = TREE_TYPE (type);
- }
- return count;
-}
-
/* Add function call CALL to edge SRC. If FLAG_GRAPHITE_OPENCL_DEBUG is
enabled, then add the following:
@@ -653,6 +632,385 @@ opencl_get_edge_for_init (opencl_main code_gen, int data_id, bool device)
return curr->init_edge;
}
+/* Return tree, which represents function selected by ID.
+ If ID is STATIC_INIT, init all required data. */
+
+static tree
+opencl_create_function_decl (enum OPENCL_FUNCTUONS id)
+{
+ static tree create_context_from_type_decl = NULL;
+ static tree get_context_info_decl = NULL;
+ static tree create_command_queue_decl = NULL;
+ static tree create_program_with_source_decl = NULL;
+ static tree build_program_decl = NULL;
+ static tree create_kernel_decl = NULL;
+ static tree create_buffer_decl = NULL;
+ static tree set_kernel_arg_decl = NULL;
+ static tree enqueue_nd_range_kernel_decl = NULL;
+ static tree enqueue_read_buffer_decl = NULL;
+ static tree enqueue_write_buffer_decl = NULL;
+ static tree release_memory_obj_decl = NULL;
+ static tree release_context_decl = NULL;
+ static tree release_command_queue_decl = NULL;
+ static tree release_program_decl = NULL;
+ static tree release_kernel_decl = NULL;
+ static tree get_platform_ids_decl = NULL;
+ static tree get_wait_for_events_decl = NULL;
+ switch (id)
+ {
+ case STATIC_INIT:
+ {
+ tree const_char_type = build_qualified_type (char_type_node,
+ TYPE_QUAL_CONST);
+ tree const_char_ptr = build_pointer_type (const_char_type);
+ tree const_char_ptr_ptr = build_pointer_type (const_char_ptr);
+
+ tree const_size_t = build_qualified_type (size_type_node,
+ TYPE_QUAL_CONST);
+ tree const_size_t_ptr = build_pointer_type (const_size_t);
+
+ tree size_t_ptr = build_pointer_type (size_type_node);
+
+ tree cl_device_type = integer_type_node;
+ tree cl_context_info = unsigned_type_node;
+ tree cl_command_queue_properties = long_unsigned_type_node;
+ tree cl_mem_flags = long_unsigned_type_node;
+
+ tree cl_context = ptr_type_node;
+ tree cl_context_properties = ptr_type_node;
+ tree cl_command_queue = ptr_type_node;
+ tree cl_device_id = ptr_type_node;
+ tree cl_program = ptr_type_node;
+ tree cl_kernel = ptr_type_node;
+ tree cl_event = ptr_type_node;
+ tree cl_mem = ptr_type_node;
+
+ tree const_cl_event = build_qualified_type (cl_event,
+ TYPE_QUAL_CONST);
+ tree cl_event_ptr = build_pointer_type (cl_event);
+ tree const_cl_event_ptr = build_pointer_type (const_cl_event);
+
+ tree const_cl_device_id = build_qualified_type (cl_device_id,
+ TYPE_QUAL_CONST);
+ tree const_cl_device_id_ptr = build_pointer_type (const_cl_device_id);
+
+ tree cl_platford_id = long_integer_type_node;
+ tree cl_platford_id_ptr = build_pointer_type (cl_platford_id);
+
+ tree function_type;
+ /* | cl_context
+ | clCreateContextFromType (cl_context_properties *properties,
+ | cl_device_type device_type,
+ | void (*pfn_notify) (const char *errinfo,
+ | const void *private_info, size_t cb,
+ | void *user_data),
+ | void *user_data,
+ | cl_int *errcode_ret) */
+ function_type
+ = build_function_type_list (cl_context,
+ cl_context_properties,
+ cl_device_type,
+ ptr_type_node,
+ ptr_type_node,
+ integer_ptr_type_node,
+ NULL_TREE);
+ create_context_from_type_decl
+ = build_fn_decl (opencl_function_names[0], function_type);
+
+ /* | cl_int clGetContextInfo (cl_context context,
+ | cl_context_info param_name,
+ | size_t param_value_size,
+ | void *param_value,
+ | size_t *param_value_size_ret) */
+ function_type
+ = build_function_type_list (integer_type_node,
+ cl_context,
+ cl_context_info,
+ size_type_node,
+ ptr_type_node,
+ size_t_ptr,
+ NULL_TREE);
+ get_context_info_decl
+ = build_fn_decl (opencl_function_names[1], function_type);
+
+ /* | cl_command_queue
+ | clCreateCommandQueue (cl_context context,
+ | cl_device_id device,
+ | cl_command_queue_properties properties,
+ | cl_int *errcode_ret) */
+ function_type
+ = build_function_type_list (cl_command_queue,
+ cl_context,
+ cl_device_id,
+ cl_command_queue_properties,
+ integer_ptr_type_node,
+ NULL_TREE);
+ create_command_queue_decl
+ = build_fn_decl (opencl_function_names[2], function_type);
+
+ /* | cl_program clCreateProgramWithSource (cl_context context,
+ | cl_uint count,
+ | const char **strings,
+ | const size_t *lengths,
+ | cl_int *errcode_ret) */
+ function_type
+ = build_function_type_list (cl_program,
+ cl_context,
+ unsigned_type_node,
+ const_char_ptr_ptr,
+ const_size_t_ptr,
+ integer_ptr_type_node,
+ NULL_TREE);
+ create_program_with_source_decl
+ = build_fn_decl (opencl_function_names[3], function_type);
+
+ /* | cl_int
+ | clBuildProgram (cl_program program,
+ | cl_uint num_devices,
+ | const cl_device_id *device_list,
+ | const char *options,
+ | void (*pfn_notify) (cl_program, void *user_data),
+ | void *user_data) */
+ function_type
+ = build_function_type_list (integer_type_node,
+ cl_program,
+ unsigned_type_node,
+ const_cl_device_id_ptr,
+ const_char_ptr,
+ ptr_type_node,
+ ptr_type_node,
+ NULL_TREE);
+ build_program_decl
+ = build_fn_decl (opencl_function_names[4], function_type);
+
+ /* | cl_kernel clCreateKernel (cl_program program,
+ | const char *kernel_name,
+ | cl_int *errcode_ret) */
+ function_type
+ = build_function_type_list (cl_kernel,
+ cl_program,
+ const_char_ptr,
+ integer_ptr_type_node,
+ NULL_TREE);
+
+ create_kernel_decl
+ = build_fn_decl (opencl_function_names[5], function_type);
+
+ /* | cl_mem clCreateBuffer (cl_context context,
+ | cl_mem_flags flags,
+ | size_t size,
+ | void *host_ptr,
+ | cl_int *errcode_ret) */
+
+ function_type
+ = build_function_type_list (cl_mem,
+ cl_context,
+ cl_mem_flags,
+ size_type_node,
+ ptr_type_node,
+ integer_ptr_type_node,
+ NULL_TREE);
+ create_buffer_decl
+ = build_fn_decl (opencl_function_names[6], function_type);
+
+
+ /* | cl_int clSetKernelArg (cl_kernel kernel,
+ | cl_uint arg_index,
+ | size_t arg_size,
+ | const void *arg_value) */
+
+ function_type
+ = build_function_type_list (integer_type_node,
+ cl_kernel,
+ unsigned_type_node,
+ size_type_node,
+ const_ptr_type_node,
+ NULL_TREE);
+ set_kernel_arg_decl
+ = build_fn_decl (opencl_function_names[7], function_type);
+
+ /* | cl_int clEnqueueNDRangeKernel (cl_command_queue command_queue,
+ | cl_kernel kernel,
+ | cl_uint work_dim,
+ | const size_t *global_work_offset,
+ | const size_t *global_work_size,
+ | const size_t *local_work_size,
+ | cl_uint num_events_in_wait_list,
+ | const cl_event *event_wait_list,
+ | cl_event *event) */
+
+ function_type
+ = build_function_type_list (integer_type_node,
+ cl_command_queue,
+ cl_kernel,
+ unsigned_type_node,
+ const_size_t_ptr,
+ const_size_t_ptr,
+ const_size_t_ptr,
+ unsigned_type_node,
+ const_cl_event_ptr,
+ cl_event_ptr,
+ NULL_TREE);
+
+ enqueue_nd_range_kernel_decl
+ = build_fn_decl (opencl_function_names[8], function_type);
+
+ /* | cl_int clEnqueueReadBuffer (cl_command_queue command_queue,
+ | cl_mem buffer,
+ | cl_bool blocking_read,
+ | size_t offset,
+ | size_t cb,
+ | void *ptr,
+ | cl_uint num_events_in_wait_list,
+ | const cl_event *event_wait_list,
+ | cl_event *event) */
+
+ function_type
+ = build_function_type_list (integer_type_node,
+ cl_command_queue,
+ cl_mem,
+ unsigned_type_node,
+ size_type_node,
+ size_type_node,
+ ptr_type_node,
+ unsigned_type_node,
+ const_cl_event_ptr,
+ cl_event_ptr,
+ NULL_TREE);
+
+ enqueue_read_buffer_decl
+ = build_fn_decl (opencl_function_names[9], function_type);
+
+ /* | cl_int clEnqueueWriteBuffer (cl_command_queue command_queue,
+ | cl_mem buffer,
+ | cl_bool blocking_write,
+ | size_t offset,
+ | size_t cb,
+ | const void *ptr,
+ | cl_uint num_events_in_wait_list,
+ | const cl_event *event_wait_list,
+ | cl_event *event) */
+
+ function_type
+ = build_function_type_list (integer_type_node,
+ cl_command_queue,
+ cl_mem,
+ unsigned_type_node,
+ size_type_node,
+ size_type_node,
+ const_ptr_type_node,
+ unsigned_type_node,
+ const_cl_event_ptr,
+ cl_event_ptr,
+ NULL_TREE);
+
+ enqueue_write_buffer_decl
+ = build_fn_decl (opencl_function_names[10], function_type);
+
+
+ /* cl_int clReleaseMemObject (cl_mem memobj) */
+
+ function_type
+ = build_function_type_list (integer_type_node, cl_mem, NULL_TREE);
+
+ release_memory_obj_decl
+ = build_fn_decl (opencl_function_names[11], function_type);
+
+
+ /* cl_int clReleaseContext (cl_context context) */
+ function_type
+ = build_function_type_list (integer_type_node, cl_context,
+ NULL_TREE);
+
+ release_context_decl
+ = build_fn_decl (opencl_function_names[12], function_type);
+
+ /* cl_int clReleaseCommandQueue (cl_command_queue command_queue) */
+ function_type
+ = build_function_type_list (integer_type_node, cl_command_queue,
+ NULL_TREE);
+
+ release_command_queue_decl
+ = build_fn_decl (opencl_function_names[13], function_type);
+
+ /* cl_int clReleaseProgram (cl_program program) */
+ function_type
+ = build_function_type_list (integer_type_node, cl_program,
+ NULL_TREE);
+
+ release_program_decl
+ = build_fn_decl (opencl_function_names[14], function_type);
+
+ /* cl_int clReleaseKernel (cl_kernel kernel) */
+ function_type
+ = build_function_type_list (integer_type_node, cl_kernel, NULL_TREE);
+
+ release_kernel_decl
+ = build_fn_decl (opencl_function_names[15], function_type);
+
+ /* | cl_int clGetPlatformIDs (cl_uint num_entries,
+ | cl_platform_id *platforms,
+ | cl_uint *num_platforms) */
+
+
+ function_type
+ = build_function_type_list (integer_type_node,
+ unsigned_type_node,
+ cl_platford_id_ptr,
+ build_pointer_type (unsigned_type_node),
+ NULL_TREE);
+ get_platform_ids_decl
+ = build_fn_decl (opencl_function_names [16], function_type);
+
+
+ /* | cl_int clWaitForEvents (cl_uint num_events,
+ | const cl_event *event_list) */
+
+ function_type
+ = build_function_type_list (integer_type_node,
+ unsigned_type_node,
+ const_cl_event_ptr,
+ NULL_TREE);
+
+ get_wait_for_events_decl
+ = build_fn_decl (opencl_function_names [17], function_type);
+
+ return NULL_TREE;
+ }
+ case CREATE_CONTEXT_FROM_TYPE: return create_context_from_type_decl;
+ case GET_CONTEXT_INFO: return get_context_info_decl;
+ case CREATE_COMMAND_QUEUE: return create_command_queue_decl;
+ case CREATE_PROGRAM_WITH_SOURCE: return create_program_with_source_decl;
+ case BUILD_PROGRAM: return build_program_decl;
+ case CREATE_KERNEL: return create_kernel_decl;
+ case CREATE_BUFFER: return create_buffer_decl;
+ case SET_KERNEL_ARG: return set_kernel_arg_decl;
+ case ENQUEUE_ND_RANGE_KERNEL: return enqueue_nd_range_kernel_decl;
+ case ENQUEUE_READ_BUFFER: return enqueue_read_buffer_decl;
+ case ENQUEUE_WRITE_BUFFER: return enqueue_write_buffer_decl;
+ case RELEASE_MEMORY_OBJ: return release_memory_obj_decl;
+ case RELEASE_CONTEXT: return release_context_decl;
+ case RELEASE_COMMAND_QUEUE: return release_command_queue_decl;
+ case RELEASE_PROGRAM: return release_program_decl;
+ case RELEASE_KERNEL: return release_kernel_decl;
+ case GET_PLATFORM_IDS: return get_platform_ids_decl;
+ case WAIT_FOR_EVENTS: return get_wait_for_events_decl;
+ default: gcc_unreachable ();
+ }
+}
+
+/* Add clWaitForEvent (1, EVENT_VAR); call to CODE_GEN->main_edge. */
+
+static void
+opencl_wait_for_event (opencl_main code_gen, tree event_var)
+{
+ tree function = opencl_create_function_decl (WAIT_FOR_EVENTS);
+ tree call = build_call_expr (function, 2,
+ integer_one_node,
+ event_var);
+ opencl_add_safe_call (code_gen, call, true);
+}
+
/* Add host to device memory transfer. DATA - data, which must be
transfered to device. CODE_GEN holds information related to code
generation. */
@@ -839,6 +1197,68 @@ opencl_fflush_all_device_buffers_to_host (opencl_main code_gen)
}
}
+/* Calculate correct flags for clCreateBuffer. READ means, that
+ buffer must be readable on device, WRITE - that buffer must be
+ writable on device. */
+
+static int
+opencl_get_mem_flags (bool read, bool write)
+{
+ int rw_flags;
+ int location_flags;
+ gcc_assert (read || write);
+ if (write && read)
+ rw_flags = CL_MEM_READ_WRITE;
+ else
+ {
+ if (read)
+ rw_flags = CL_MEM_READ_ONLY;
+ else
+ rw_flags = CL_MEM_WRITE_ONLY;
+ }
+ if (flag_graphite_opencl_cpu)
+ location_flags = CL_MEM_USE_HOST_PTR;
+ else
+ location_flags = CL_MEM_COPY_HOST_PTR;
+ return location_flags | rw_flags;
+}
+
+/* Create memory on device for DATA and init it by data from host.
+ ptr is pointer to host memory location. Function returns tree,
+ corresponding to memory location on device. */
+
+static tree
+opencl_create_memory_for_pointer (opencl_data data)
+{
+ tree ptr = data->object;
+ tree arr_size = data->size_variable;
+ tree function = opencl_create_function_decl (CREATE_BUFFER);
+ bool ever_read = data->ever_read_on_device;
+ bool ever_written = data->ever_written_on_device;
+ tree mem_flags = build_int_cst (NULL_TREE,
+ opencl_get_mem_flags (ever_read,
+ ever_written));
+ if (TREE_CODE (TREE_TYPE (ptr)) == ARRAY_TYPE)
+ ptr = build_addr (ptr, current_function_decl);
+
+ if (flag_graphite_opencl_debug)
+ {
+ tree result = opencl_create_tmp_var (integer_type_node,
+ "__opencl_create_buffer_result");
+
+ return build_call_expr (function, 5,
+ h_context, mem_flags,
+ arr_size, ptr,
+ build1 (ADDR_EXPR,
+ integer_ptr_type_node,
+ result));
+ }
+ else
+ return build_call_expr (function, 5,
+ h_context, mem_flags,
+ arr_size, ptr, null_pointer_node);
+}
+
/* Create memory buffers on host for all required host memory objects.
CODE_GEN holds information related to code generation. */
@@ -1010,68 +1430,6 @@ opencl_get_indirect_size (tree ptr, poly_dr_p ref)
gcc_unreachable ();
}
-/* Calculate correct flags for clCreateBuffer. READ means, that
- buffer must be readable on device, WRITE - that buffer must be
- writable on device. */
-
-static int
-opencl_get_mem_flags (bool read, bool write)
-{
- int rw_flags;
- int location_flags;
- gcc_assert (read || write);
- if (write && read)
- rw_flags = CL_MEM_READ_WRITE;
- else
- {
- if (read)
- rw_flags = CL_MEM_READ_ONLY;
- else
- rw_flags = CL_MEM_WRITE_ONLY;
- }
- if (flag_graphite_opencl_cpu)
- location_flags = CL_MEM_USE_HOST_PTR;
- else
- location_flags = CL_MEM_COPY_HOST_PTR;
- return location_flags | rw_flags;
-}
-
-/* Create memory on device for DATA and init it by data from host.
- ptr is pointer to host memory location. Function returns tree,
- corresponding to memory location on device. */
-
-static tree
-opencl_create_memory_for_pointer (opencl_data data)
-{
- tree ptr = data->object;
- tree arr_size = data->size_variable;
- tree function = opencl_create_function_decl (CREATE_BUFFER);
- bool ever_read = data->ever_read_on_device;
- bool ever_written = data->ever_written_on_device;
- tree mem_flags = build_int_cst (NULL_TREE,
- opencl_get_mem_flags (ever_read,
- ever_written));
- if (TREE_CODE (TREE_TYPE (ptr)) == ARRAY_TYPE)
- ptr = build_addr (ptr, current_function_decl);
-
- if (flag_graphite_opencl_debug)
- {
- tree result = opencl_create_tmp_var (integer_type_node,
- "__opencl_create_buffer_result");
-
- return build_call_expr (function, 5,
- h_context, mem_flags,
- arr_size, ptr,
- build1 (ADDR_EXPR,
- integer_ptr_type_node,
- result));
- }
- else
- return build_call_expr (function, 5,
- h_context, mem_flags,
- arr_size, ptr, null_pointer_node);
-}
-
/* Create variables for kernel KERNEL arguments. Each argument is
represented by new variable with it's value and it's size. If arg
is a pointer or array, it's represented by device buffer with data
@@ -1290,44 +1648,181 @@ opencl_set_data_size (opencl_main code_gen)
}
}
-/* Transform clast statement DATA from scop SCOP to OpenCL calls
- in region REGION. Place all calls to edge MAIN. PARAM_INDEX
- holds external scop params. */
+/* Find opencl_data which represents array VAR. */
-void
-opencl_transform_clast (struct clast_stmt * data, sese region,
- edge main, scop_p scop, htab_t params_index)
+static opencl_data
+opencl_get_static_data_by_tree (tree var)
{
- opencl_main code_gen;
- /* Create main data struture for code generation. */
+ map_tree_to_data tmp = map_tree_to_data_create (var, NULL);
+ map_tree_to_data * slot
+ = (map_tree_to_data *) htab_find_slot (array_data_to_tree,
+ tmp, INSERT);
+ if (*slot == NULL)
+ return NULL;
- if (dump_file && (dump_flags & TDF_DETAILS))
+ return (*slot)->value;
+
+}
+
+/* Create required OpenCL variable for given DATA. */
+
+static void
+opencl_data_init_object (opencl_data data)
+{
+ if (TREE_CODE (TREE_TYPE (data->exact_object)) == POINTER_TYPE)
{
- fprintf (dump_file, "\nGenerating OpenCL code for SCoP: \n");
- print_scop (dump_file, scop, 0);
+ data->device_object
+ = opencl_create_tmp_var (ptr_type_node, "__opencl_data");
+ data->is_static = false;
}
+ else
+ {
+ /* (TREE_CODE (TREE_TYPE (data->exact_object)) == ARRAY_TYPE) */
+ map_tree_to_data tree_ptr
+ = map_tree_to_data_create (data->exact_object, data);
- code_gen = opencl_main_create (((struct clast_root *)data)->names,
- region, main, params_index);
+ map_tree_to_data * tree_slot =
+ (map_tree_to_data *) htab_find_slot (array_data_to_tree,
+ tree_ptr, INSERT);
+ gcc_assert (*tree_slot == NULL);
+ *tree_slot = tree_ptr;
- opencl_init_basic_blocks (code_gen);
- opencl_init_data (scop, code_gen);
+ data->device_object
+ = opencl_create_static_ptr_variable ("__opencl_data");
+ data->is_static = true;
+ data->size_variable = data->size_value;
+ VEC_safe_push (opencl_data, heap, opencl_array_data, data);
+ }
+}
- code_gen->clast_meta = opencl_create_meta_from_clast (code_gen, data, 1,
- NULL);
- code_gen->curr_meta = code_gen->clast_meta;
+/* Register reference to DATA via data reference REF_KEY and
+ variable TREE_KEY in CODE_GEN structures. */
- opencl_transform_stmt_list (data, code_gen, 1);
- if (dyn_string_length (code_gen->main_program) != 0)
+static void
+opencl_register_data (opencl_main code_gen, opencl_data data,
+ tree tree_key, data_reference_p ref_key)
+{
+ htab_t ref_to_data = code_gen->ref_to_data;
+ htab_t tree_to_data = code_gen->tree_to_data;
+ map_ref_to_data ref_ptr = map_ref_to_data_create (ref_key, data);
+ map_tree_to_data tree_ptr = map_tree_to_data_create (tree_key, data);
+ map_ref_to_data * ref_slot;
+ map_tree_to_data * tree_slot;
+
+
+ ref_slot
+ = (map_ref_to_data *) htab_find_slot (ref_to_data, ref_ptr, INSERT);
+ gcc_assert (*ref_slot == NULL);
+ *ref_slot = ref_ptr;
+
+
+ tree_slot
+ = (map_tree_to_data *) htab_find_slot (tree_to_data, tree_ptr, INSERT);
+ gcc_assert (*tree_slot == NULL || (*tree_slot)->value == data);
+ *tree_slot = tree_ptr;
+}
+
+/* Analyze single data reference REF and update CODE_GEN structures.
+ If it access data, which has been accessed in data references
+ before, update it's size. Otherwise add data to array. */
+
+static void
+opencl_parse_single_data_ref (poly_dr_p ref, opencl_main code_gen)
+{
+ data_reference_p d_ref = (data_reference_p) PDR_CDR (ref);
+ tree data_ref_tree = dr_outermost_base_object (d_ref);
+ tree size = NULL_TREE;
+ opencl_data curr;
+
+ curr = opencl_get_data_by_tree (code_gen, data_ref_tree);
+ size = opencl_get_indirect_size (data_ref_tree, ref);
+ if (curr)
{
- dyn_string_append (main_program_src, code_gen->main_program);
- opencl_set_data_size (code_gen);
- opencl_init_all_device_buffers (code_gen);
- opencl_fflush_all_device_buffers_to_host (code_gen);
+ if (!curr->is_static)
+ {
+ if (!size || !curr->size_value)
+ curr->size_value = NULL;
+ else
+ curr->size_value = fold_build2 (MAX_EXPR, sizetype,
+ size, curr->size_value);
+ }
}
- recompute_all_dominators ();
- update_ssa (TODO_update_ssa);
- opencl_main_delete (code_gen);
+ else
+ {
+ curr = opencl_get_static_data_by_tree (data_ref_tree);
+ if (!curr)
+ {
+ curr = opencl_data_create (data_ref_tree, size);
+ opencl_data_init_object (curr);
+ }
+ curr->id = VEC_length (opencl_data, code_gen->opencl_function_data);
+ VEC_safe_push (opencl_data, heap, code_gen->opencl_function_data, curr);
+ }
+ opencl_register_data (code_gen, curr, data_ref_tree, d_ref);
+}
+
+/* Analyse all data reference for poly basic block PBB and update CODE_GEN
+ structures. */
+
+static void
+opencl_parse_data_refs (poly_bb_p pbb, opencl_main code_gen)
+{
+ VEC (poly_dr_p, heap) *drs = PBB_DRS (pbb);
+ int i;
+ poly_dr_p curr;
+ for (i = 0; VEC_iterate (poly_dr_p, drs, i, curr); i++)
+ opencl_parse_single_data_ref (curr, code_gen);
+}
+
+/* Analyse all data reference for scop M_SCOP and update
+ CODE_GEN structures. */
+
+static void
+opencl_init_data (scop_p m_scop, opencl_main code_gen)
+{
+ VEC (poly_bb_p, heap) * bbs = SCOP_BBS (m_scop);
+ int i;
+ poly_bb_p curr;
+ for (i = 0; VEC_iterate (poly_bb_p, bbs, i, curr); i++)
+ opencl_parse_data_refs (curr, code_gen);
+}
+
+/* Init basic block in CODE_GEN structures. */
+
+static void
+opencl_init_basic_blocks (opencl_main code_gen)
+{
+ code_gen->data_init_bb = opencl_create_bb (code_gen);
+ code_gen->kernel_edge = code_gen->main_edge;
+}
+
+/* Add function calls to create and launch kernel KERNEL to
+ CODE_GEN->main_edge. */
+
+static void
+opencl_create_gimple_for_body (opencl_body kernel, opencl_main code_gen)
+{
+ tree num_of_exec = kernel->num_of_exec;
+ tree call;
+
+ tree kernel_var
+ = opencl_insert_create_kernel_call (code_gen, (const char *)kernel->name);
+
+ tree index_type = build_index_type (build_int_cst (NULL_TREE, 2));
+ tree array_type = build_array_type (ptr_type_node, index_type);
+ tree var = opencl_create_tmp_var (array_type, "wait_event");
+ TREE_STATIC (var) = 1;
+ assemble_variable (var, 1, 0, 1);
+
+ call = build4 (ARRAY_REF, ptr_type_node, var,
+ integer_zero_node, NULL_TREE, NULL_TREE);
+ call = build_addr (call, current_function_decl);
+
+ opencl_init_local_device_memory (code_gen, kernel);
+ opencl_pass_kernel_arguments (code_gen, kernel, kernel_var);
+
+ opencl_execute_kernel (code_gen, num_of_exec, kernel_var, call);
+ opencl_wait_for_event (code_gen, call);
}
/* Prepare memory for gimple (host) statement, represented by PBB.
@@ -1406,26 +1901,19 @@ opencl_add_gimple_for_user_stmt (struct clast_user_stmt * stmt,
opencl_verify ();
}
-/* Add if statement, represented by S to current gimple.
- CODE_GEN holds information related to code generation. */
+/* Delete opencl_body DATA. */
static void
-opencl_add_gimple_for_stmt_guard (struct clast_guard * s,
- opencl_main code_gen, int depth)
+opencl_body_delete (opencl_body data)
{
- edge last_e = graphite_create_new_guard (code_gen->region,
- code_gen->main_edge, s,
- code_gen->newivs,
- code_gen->newivs_index,
- code_gen->params_index);
-
- edge true_e = get_true_edge_from_guard_bb (code_gen->main_edge->dest);
- code_gen->main_edge = true_e;
- opencl_transform_stmt_list (s->then, code_gen, depth);
- code_gen->main_edge = last_e;
-
- recompute_all_dominators ();
- opencl_verify ();
+ dyn_string_delete (data->body);
+ dyn_string_delete (data->header);
+ dyn_string_delete (data->pre_header);
+ dyn_string_delete (data->non_scalar_args);
+ VEC_free (tree, heap, data->function_args);
+ VEC_free (tree, heap, data->function_args_to_pass);
+ VEC_free (opencl_data, heap, data->data_refs);
+ free (data);
}
/* Reset data structures before processing loop, represented by META.
@@ -1493,6 +1981,8 @@ opencl_postpass_data (opencl_main code_gen, opencl_clast_meta meta)
VEC_free (opencl_data, heap, meta->post_pass_to_device);
}
+static void opencl_transform_stmt_list (struct clast_stmt *, opencl_main, int);
+
/* Add loop body, of the loop, represented by S, on host.
Loop body can contain device code.
DEPTH contains depth of given loop in current loop nest.
@@ -1578,19 +2068,26 @@ opencl_fix_meta_flags (opencl_clast_meta meta)
}
}
-/* Delete opencl_body DATA. */
+/* Add if statement, represented by S to current gimple.
+ CODE_GEN holds information related to code generation. */
static void
-opencl_body_delete (opencl_body data)
+opencl_add_gimple_for_stmt_guard (struct clast_guard * s,
+ opencl_main code_gen, int depth)
{
- dyn_string_delete (data->body);
- dyn_string_delete (data->header);
- dyn_string_delete (data->pre_header);
- dyn_string_delete (data->non_scalar_args);
- VEC_free (tree, heap, data->function_args);
- VEC_free (tree, heap, data->function_args_to_pass);
- VEC_free (opencl_data, heap, data->data_refs);
- free (data);
+ edge last_e = graphite_create_new_guard (code_gen->region,
+ code_gen->main_edge, s,
+ code_gen->newivs,
+ code_gen->newivs_index,
+ code_gen->params_index);
+
+ edge true_e = get_true_edge_from_guard_bb (code_gen->main_edge->dest);
+ code_gen->main_edge = true_e;
+ opencl_transform_stmt_list (s->then, code_gen, depth);
+ code_gen->main_edge = last_e;
+
+ recompute_all_dominators ();
+ opencl_verify ();
}
/* Parse clast statement list S, located on depth DEPTH in current loop nest.
@@ -1687,6 +2184,46 @@ opencl_transform_stmt_list (struct clast_stmt * s, opencl_main code_gen,
}
}
+/* Transform clast statement DATA from scop SCOP to OpenCL calls
+ in region REGION. Place all calls to edge MAIN. PARAM_INDEX
+ holds external scop params. */
+
+void
+opencl_transform_clast (struct clast_stmt * data, sese region,
+ edge main, scop_p scop, htab_t params_index)
+{
+ opencl_main code_gen;
+ /* Create main data struture for code generation. */
+
+ if (dump_file && (dump_flags & TDF_DETAILS))
+ {
+ fprintf (dump_file, "\nGenerating OpenCL code for SCoP: \n");
+ print_scop (dump_file, scop, 0);
+ }
+
+ code_gen = opencl_main_create (((struct clast_root *)data)->names,
+ region, main, params_index);
+
+ opencl_init_basic_blocks (code_gen);
+ opencl_init_data (scop, code_gen);
+
+ code_gen->clast_meta = opencl_create_meta_from_clast (code_gen, data, 1,
+ NULL);
+ code_gen->curr_meta = code_gen->clast_meta;
+
+ opencl_transform_stmt_list (data, code_gen, 1);
+ if (dyn_string_length (code_gen->main_program) != 0)
+ {
+ dyn_string_append (main_program_src, code_gen->main_program);
+ opencl_set_data_size (code_gen);
+ opencl_init_all_device_buffers (code_gen);
+ opencl_fflush_all_device_buffers_to_host (code_gen);
+ }
+ recompute_all_dominators ();
+ update_ssa (TODO_update_ssa);
+ opencl_main_delete (code_gen);
+}
+
/* Find opencl_data object by host object OBJ in CODE_GEN hash maps. */
opencl_data
@@ -1717,188 +2254,6 @@ opencl_get_data_by_data_ref (opencl_main code_gen, data_reference_p ref)
return (*slot)->value;
}
-/* Register reference to DATA via data reference REF_KEY and
- variable TREE_KEY in CODE_GEN structures. */
-
-static void
-opencl_register_data (opencl_main code_gen, opencl_data data,
- tree tree_key, data_reference_p ref_key)
-{
- htab_t ref_to_data = code_gen->ref_to_data;
- htab_t tree_to_data = code_gen->tree_to_data;
- map_ref_to_data ref_ptr = map_ref_to_data_create (ref_key, data);
- map_tree_to_data tree_ptr = map_tree_to_data_create (tree_key, data);
- map_ref_to_data * ref_slot;
- map_tree_to_data * tree_slot;
-
-
- ref_slot
- = (map_ref_to_data *) htab_find_slot (ref_to_data, ref_ptr, INSERT);
- gcc_assert (*ref_slot == NULL);
- *ref_slot = ref_ptr;
-
-
- tree_slot
- = (map_tree_to_data *) htab_find_slot (tree_to_data, tree_ptr, INSERT);
- gcc_assert (*tree_slot == NULL || (*tree_slot)->value == data);
- *tree_slot = tree_ptr;
-}
-
-/* Create required OpenCL variable for given DATA. */
-
-static void
-opencl_data_init_object (opencl_data data)
-{
- if (TREE_CODE (TREE_TYPE (data->exact_object)) == POINTER_TYPE)
- {
- data->device_object
- = opencl_create_tmp_var (ptr_type_node, "__opencl_data");
- data->is_static = false;
- }
- else
- {
- /* (TREE_CODE (TREE_TYPE (data->exact_object)) == ARRAY_TYPE) */
- map_tree_to_data tree_ptr
- = map_tree_to_data_create (data->exact_object, data);
-
- map_tree_to_data * tree_slot =
- (map_tree_to_data *) htab_find_slot (array_data_to_tree,
- tree_ptr, INSERT);
- gcc_assert (*tree_slot == NULL);
- *tree_slot = tree_ptr;
-
- data->device_object
- = opencl_create_static_ptr_variable ("__opencl_data");
- data->is_static = true;
- data->size_variable = data->size_value;
- VEC_safe_push (opencl_data, heap, opencl_array_data, data);
- }
-}
-
-/* Find opencl_data which represents array VAR. */
-
-static opencl_data
-opencl_get_static_data_by_tree (tree var)
-{
- map_tree_to_data tmp = map_tree_to_data_create (var, NULL);
- map_tree_to_data * slot
- = (map_tree_to_data *) htab_find_slot (array_data_to_tree,
- tmp, INSERT);
- if (*slot == NULL)
- return NULL;
-
- return (*slot)->value;
-
-}
-
-/* Analyze single data reference REF and update CODE_GEN structures.
- If it access data, which has been accessed in data references
- before, update it's size. Otherwise add data to array. */
-
-static void
-opencl_parse_single_data_ref (poly_dr_p ref, opencl_main code_gen)
-{
- data_reference_p d_ref = (data_reference_p) PDR_CDR (ref);
- tree data_ref_tree = dr_outermost_base_object (d_ref);
- tree size = NULL_TREE;
- opencl_data curr;
-
- curr = opencl_get_data_by_tree (code_gen, data_ref_tree);
- size = opencl_get_indirect_size (data_ref_tree, ref);
- if (curr)
- {
- if (!curr->is_static)
- {
- if (!size || !curr->size_value)
- curr->size_value = NULL;
- else
- curr->size_value = fold_build2 (MAX_EXPR, sizetype,
- size, curr->size_value);
- }
- }
- else
- {
- curr = opencl_get_static_data_by_tree (data_ref_tree);
- if (!curr)
- {
- curr = opencl_data_create (data_ref_tree, size);
- opencl_data_init_object (curr);
- }
- curr->id = VEC_length (opencl_data, code_gen->opencl_function_data);
- VEC_safe_push (opencl_data, heap, code_gen->opencl_function_data, curr);
- }
- opencl_register_data (code_gen, curr, data_ref_tree, d_ref);
-}
-
-/* Analyse all data reference for poly basic block PBB and update CODE_GEN
- structures. */
-
-static void
-opencl_parse_data_refs (poly_bb_p pbb, opencl_main code_gen)
-{
- VEC (poly_dr_p, heap) *drs = PBB_DRS (pbb);
- int i;
- poly_dr_p curr;
- for (i = 0; VEC_iterate (poly_dr_p, drs, i, curr); i++)
- opencl_parse_single_data_ref (curr, code_gen);
-}
-
-/* Analyse all data reference for scop M_SCOP and update
- CODE_GEN structures. */
-
-static void
-opencl_init_data (scop_p m_scop, opencl_main code_gen)
-{
- VEC (poly_bb_p, heap) * bbs = SCOP_BBS (m_scop);
- int i;
- poly_bb_p curr;
- for (i = 0; VEC_iterate (poly_bb_p, bbs, i, curr); i++)
- opencl_parse_data_refs (curr, code_gen);
-}
-
-/* Add clWaitForEvent (1, EVENT_VAR); call to CODE_GEN->main_edge. */
-
-static void
-opencl_wait_for_event (opencl_main code_gen, tree event_var)
-{
- tree function = opencl_create_function_decl (WAIT_FOR_EVENTS);
- tree call = build_call_expr (function, 2,
- integer_one_node,
- event_var);
- opencl_add_safe_call (code_gen, call, true);
-}
-
-/* This calls must be placed after outermost loop processing. */
-
-/* Add function calls to create and launch kernel KERNEL to
- CODE_GEN->main_edge. */
-
-static void
-opencl_create_gimple_for_body (opencl_body kernel, opencl_main code_gen)
-{
- tree num_of_exec = kernel->num_of_exec;
- tree call;
-
- tree kernel_var
- = opencl_insert_create_kernel_call (code_gen, (const char *)kernel->name);
-
- tree index_type = build_index_type (build_int_cst (NULL_TREE, 2));
- tree array_type = build_array_type (ptr_type_node, index_type);
- tree var = opencl_create_tmp_var (array_type, "wait_event");
- TREE_STATIC (var) = 1;
- assemble_variable (var, 1, 0, 1);
-
- call = build4 (ARRAY_REF, ptr_type_node, var,
- integer_zero_node, NULL_TREE, NULL_TREE);
- call = build_addr (call, current_function_decl);
-
- opencl_init_local_device_memory (code_gen, kernel);
- opencl_pass_kernel_arguments (code_gen, kernel, kernel_var);
-
- opencl_execute_kernel (code_gen, num_of_exec, kernel_var, call);
- opencl_wait_for_event (code_gen, call);
-}
-
/* Create global variables for opencl code. */
static void
@@ -1916,373 +2271,6 @@ opencl_create_gimple_variables (void)
h_cmd_queue = opencl_create_static_ptr_variable ("__ocl_h_cmd_queue");
}
-/* Return tree, which represents function selected by ID.
- If ID is STATIC_INIT, init all required data. */
-
-static tree
-opencl_create_function_decl (enum OPENCL_FUNCTUONS id)
-{
- static tree create_context_from_type_decl = NULL;
- static tree get_context_info_decl = NULL;
- static tree create_command_queue_decl = NULL;
- static tree create_program_with_source_decl = NULL;
- static tree build_program_decl = NULL;
- static tree create_kernel_decl = NULL;
- static tree create_buffer_decl = NULL;
- static tree set_kernel_arg_decl = NULL;
- static tree enqueue_nd_range_kernel_decl = NULL;
- static tree enqueue_read_buffer_decl = NULL;
- static tree enqueue_write_buffer_decl = NULL;
- static tree release_memory_obj_decl = NULL;
- static tree release_context_decl = NULL;
- static tree release_command_queue_decl = NULL;
- static tree release_program_decl = NULL;
- static tree release_kernel_decl = NULL;
- static tree get_platform_ids_decl = NULL;
- static tree get_wait_for_events_decl = NULL;
- switch (id)
- {
- case STATIC_INIT:
- {
- tree const_char_type = build_qualified_type (char_type_node,
- TYPE_QUAL_CONST);
- tree const_char_ptr = build_pointer_type (const_char_type);
- tree const_char_ptr_ptr = build_pointer_type (const_char_ptr);
-
- tree const_size_t = build_qualified_type (size_type_node,
- TYPE_QUAL_CONST);
- tree const_size_t_ptr = build_pointer_type (const_size_t);
-
- tree size_t_ptr = build_pointer_type (size_type_node);
-
- tree cl_device_type = integer_type_node;
- tree cl_context_info = unsigned_type_node;
- tree cl_command_queue_properties = long_unsigned_type_node;
- tree cl_mem_flags = long_unsigned_type_node;
-
- tree cl_context = ptr_type_node;
- tree cl_context_properties = ptr_type_node;
- tree cl_command_queue = ptr_type_node;
- tree cl_device_id = ptr_type_node;
- tree cl_program = ptr_type_node;
- tree cl_kernel = ptr_type_node;
- tree cl_event = ptr_type_node;
- tree cl_mem = ptr_type_node;
-
- tree const_cl_event = build_qualified_type (cl_event,
- TYPE_QUAL_CONST);
- tree cl_event_ptr = build_pointer_type (cl_event);
- tree const_cl_event_ptr = build_pointer_type (const_cl_event);
-
- tree const_cl_device_id = build_qualified_type (cl_device_id,
- TYPE_QUAL_CONST);
- tree const_cl_device_id_ptr = build_pointer_type (const_cl_device_id);
-
- tree cl_platford_id = long_integer_type_node;
- tree cl_platford_id_ptr = build_pointer_type (cl_platford_id);
-
- tree function_type;
- /* | cl_context
- | clCreateContextFromType (cl_context_properties *properties,
- | cl_device_type device_type,
- | void (*pfn_notify) (const char *errinfo,
- | const void *private_info, size_t cb,
- | void *user_data),
- | void *user_data,
- | cl_int *errcode_ret) */
- function_type
- = build_function_type_list (cl_context,
- cl_context_properties,
- cl_device_type,
- ptr_type_node,
- ptr_type_node,
- integer_ptr_type_node,
- NULL_TREE);
- create_context_from_type_decl
- = build_fn_decl (opencl_function_names[0], function_type);
-
- /* | cl_int clGetContextInfo (cl_context context,
- | cl_context_info param_name,
- | size_t param_value_size,
- | void *param_value,
- | size_t *param_value_size_ret) */
- function_type
- = build_function_type_list (integer_type_node,
- cl_context,
- cl_context_info,
- size_type_node,
- ptr_type_node,
- size_t_ptr,
- NULL_TREE);
- get_context_info_decl
- = build_fn_decl (opencl_function_names[1], function_type);
-
- /* | cl_command_queue
- | clCreateCommandQueue (cl_context context,
- | cl_device_id device,
- | cl_command_queue_properties properties,
- | cl_int *errcode_ret) */
- function_type
- = build_function_type_list (cl_command_queue,
- cl_context,
- cl_device_id,
- cl_command_queue_properties,
- integer_ptr_type_node,
- NULL_TREE);
- create_command_queue_decl
- = build_fn_decl (opencl_function_names[2], function_type);
-
- /* | cl_program clCreateProgramWithSource (cl_context context,
- | cl_uint count,
- | const char **strings,
- | const size_t *lengths,
- | cl_int *errcode_ret) */
- function_type
- = build_function_type_list (cl_program,
- cl_context,
- unsigned_type_node,
- const_char_ptr_ptr,
- const_size_t_ptr,
- integer_ptr_type_node,
- NULL_TREE);
- create_program_with_source_decl
- = build_fn_decl (opencl_function_names[3], function_type);
-
- /* | cl_int
- | clBuildProgram (cl_program program,
- | cl_uint num_devices,
- | const cl_device_id *device_list,
- | const char *options,
- | void (*pfn_notify) (cl_program, void *user_data),
- | void *user_data) */
- function_type
- = build_function_type_list (integer_type_node,
- cl_program,
- unsigned_type_node,
- const_cl_device_id_ptr,
- const_char_ptr,
- ptr_type_node,
- ptr_type_node,
- NULL_TREE);
- build_program_decl
- = build_fn_decl (opencl_function_names[4], function_type);
-
- /* | cl_kernel clCreateKernel (cl_program program,
- | const char *kernel_name,
- | cl_int *errcode_ret) */
- function_type
- = build_function_type_list (cl_kernel,
- cl_program,
- const_char_ptr,
- integer_ptr_type_node,
- NULL_TREE);
-
- create_kernel_decl
- = build_fn_decl (opencl_function_names[5], function_type);
-
- /* | cl_mem clCreateBuffer (cl_context context,
- | cl_mem_flags flags,
- | size_t size,
- | void *host_ptr,
- | cl_int *errcode_ret) */
-
- function_type
- = build_function_type_list (cl_mem,
- cl_context,
- cl_mem_flags,
- size_type_node,
- ptr_type_node,
- integer_ptr_type_node,
- NULL_TREE);
- create_buffer_decl
- = build_fn_decl (opencl_function_names[6], function_type);
-
-
- /* | cl_int clSetKernelArg (cl_kernel kernel,
- | cl_uint arg_index,
- | size_t arg_size,
- | const void *arg_value) */
-
- function_type
- = build_function_type_list (integer_type_node,
- cl_kernel,
- unsigned_type_node,
- size_type_node,
- const_ptr_type_node,
- NULL_TREE);
- set_kernel_arg_decl
- = build_fn_decl (opencl_function_names[7], function_type);
-
- /* | cl_int clEnqueueNDRangeKernel (cl_command_queue command_queue,
- | cl_kernel kernel,
- | cl_uint work_dim,
- | const size_t *global_work_offset,
- | const size_t *global_work_size,
- | const size_t *local_work_size,
- | cl_uint num_events_in_wait_list,
- | const cl_event *event_wait_list,
- | cl_event *event) */
-
- function_type
- = build_function_type_list (integer_type_node,
- cl_command_queue,
- cl_kernel,
- unsigned_type_node,
- const_size_t_ptr,
- const_size_t_ptr,
- const_size_t_ptr,
- unsigned_type_node,
- const_cl_event_ptr,
- cl_event_ptr,
- NULL_TREE);
-
- enqueue_nd_range_kernel_decl
- = build_fn_decl (opencl_function_names[8], function_type);
-
- /* | cl_int clEnqueueReadBuffer (cl_command_queue command_queue,
- | cl_mem buffer,
- | cl_bool blocking_read,
- | size_t offset,
- | size_t cb,
- | void *ptr,
- | cl_uint num_events_in_wait_list,
- | const cl_event *event_wait_list,
- | cl_event *event) */
-
- function_type
- = build_function_type_list (integer_type_node,
- cl_command_queue,
- cl_mem,
- unsigned_type_node,
- size_type_node,
- size_type_node,
- ptr_type_node,
- unsigned_type_node,
- const_cl_event_ptr,
- cl_event_ptr,
- NULL_TREE);
-
- enqueue_read_buffer_decl
- = build_fn_decl (opencl_function_names[9], function_type);
-
- /* | cl_int clEnqueueWriteBuffer (cl_command_queue command_queue,
- | cl_mem buffer,
- | cl_bool blocking_write,
- | size_t offset,
- | size_t cb,
- | const void *ptr,
- | cl_uint num_events_in_wait_list,
- | const cl_event *event_wait_list,
- | cl_event *event) */
-
- function_type
- = build_function_type_list (integer_type_node,
- cl_command_queue,
- cl_mem,
- unsigned_type_node,
- size_type_node,
- size_type_node,
- const_ptr_type_node,
- unsigned_type_node,
- const_cl_event_ptr,
- cl_event_ptr,
- NULL_TREE);
-
- enqueue_write_buffer_decl
- = build_fn_decl (opencl_function_names[10], function_type);
-
-
- /* cl_int clReleaseMemObject (cl_mem memobj) */
-
- function_type
- = build_function_type_list (integer_type_node, cl_mem, NULL_TREE);
-
- release_memory_obj_decl
- = build_fn_decl (opencl_function_names[11], function_type);
-
-
- /* cl_int clReleaseContext (cl_context context) */
- function_type
- = build_function_type_list (integer_type_node, cl_context,
- NULL_TREE);
-
- release_context_decl
- = build_fn_decl (opencl_function_names[12], function_type);
-
- /* cl_int clReleaseCommandQueue (cl_command_queue command_queue) */
- function_type
- = build_function_type_list (integer_type_node, cl_command_queue,
- NULL_TREE);
-
- release_command_queue_decl
- = build_fn_decl (opencl_function_names[13], function_type);
-
- /* cl_int clReleaseProgram (cl_program program) */
- function_type
- = build_function_type_list (integer_type_node, cl_program,
- NULL_TREE);
-
- release_program_decl
- = build_fn_decl (opencl_function_names[14], function_type);
-
- /* cl_int clReleaseKernel (cl_kernel kernel) */
- function_type
- = build_function_type_list (integer_type_node, cl_kernel, NULL_TREE);
-
- release_kernel_decl
- = build_fn_decl (opencl_function_names[15], function_type);
-
- /* | cl_int clGetPlatformIDs (cl_uint num_entries,
- | cl_platform_id *platforms,
- | cl_uint *num_platforms) */
-
-
- function_type
- = build_function_type_list (integer_type_node,
- unsigned_type_node,
- cl_platford_id_ptr,
- build_pointer_type (unsigned_type_node),
- NULL_TREE);
- get_platform_ids_decl
- = build_fn_decl (opencl_function_names [16], function_type);
-
-
- /* | cl_int clWaitForEvents (cl_uint num_events,
- | const cl_event *event_list) */
-
- function_type
- = build_function_type_list (integer_type_node,
- unsigned_type_node,
- const_cl_event_ptr,
- NULL_TREE);
-
- get_wait_for_events_decl
- = build_fn_decl (opencl_function_names [17], function_type);
-
- return NULL_TREE;
- }
- case CREATE_CONTEXT_FROM_TYPE: return create_context_from_type_decl;
- case GET_CONTEXT_INFO: return get_context_info_decl;
- case CREATE_COMMAND_QUEUE: return create_command_queue_decl;
- case CREATE_PROGRAM_WITH_SOURCE: return create_program_with_source_decl;
- case BUILD_PROGRAM: return build_program_decl;
- case CREATE_KERNEL: return create_kernel_decl;
- case CREATE_BUFFER: return create_buffer_decl;
- case SET_KERNEL_ARG: return set_kernel_arg_decl;
- case ENQUEUE_ND_RANGE_KERNEL: return enqueue_nd_range_kernel_decl;
- case ENQUEUE_READ_BUFFER: return enqueue_read_buffer_decl;
- case ENQUEUE_WRITE_BUFFER: return enqueue_write_buffer_decl;
- case RELEASE_MEMORY_OBJ: return release_memory_obj_decl;
- case RELEASE_CONTEXT: return release_context_decl;
- case RELEASE_COMMAND_QUEUE: return release_command_queue_decl;
- case RELEASE_PROGRAM: return release_program_decl;
- case RELEASE_KERNEL: return release_kernel_decl;
- case GET_PLATFORM_IDS: return get_platform_ids_decl;
- case WAIT_FOR_EVENTS: return get_wait_for_events_decl;
- default: gcc_unreachable ();
- }
-}
-
/* Create call
| clGetContextInfo (h_context, CL_CONTEXT_DEVICES, 0, 0,
| &n_context_descriptor_size);
@@ -2442,15 +2430,6 @@ opencl_create_malloc_call (tree arg)
return call;
}
-/* Init basic block in CODE_GEN structures. */
-
-static void
-opencl_init_basic_blocks (opencl_main code_gen)
-{
- code_gen->data_init_bb = opencl_create_bb (code_gen);
- code_gen->kernel_edge = code_gen->main_edge;
-}
-
/* Generate calls for opencl init functions and place them to INIT_EDGE.
Must be called only once in each function. */