@@ -1,5 +1,11 @@
2010-12-27 Sebastian Pop <sebastian.pop@amd.com>
+ * graphite-opencl-codegen.c: Fix formatting.
+ * graphite-opencl-meta-clast.c: Same.
+ * graphite-opencl.c: Same.
+
+2010-12-27 Sebastian Pop <sebastian.pop@amd.com>
+
* graphite-opencl-codegen.c: Fix formating problems.
* graphite-opencl-meta-clast.c: Same.
* graphite-opencl.c: Same.
@@ -85,6 +85,7 @@ opencl_cmp_str (const void *str1, const void *str2)
{
const char *c_str1 = (const char *) str1;
const char *c_str2 = (const char *) str2;
+
return !strcmp (c_str1, c_str2);
}
@@ -111,6 +112,7 @@ static void
opencl_append_string_to_header (const char *str, opencl_main code_gen)
{
dyn_string_t tmp = opencl_get_current_header (code_gen);
+
dyn_string_append_cstr (tmp, str);
}
@@ -121,6 +123,7 @@ static void
opencl_append_string_to_body (const char *str, opencl_main code_gen)
{
dyn_string_t tmp = opencl_get_current_body (code_gen);
+
dyn_string_append_cstr (tmp, str);
}
@@ -130,6 +133,7 @@ static void
opencl_append_int_to_str (dyn_string_t str, long num, const char *format)
{
char tmp[100];
+
sprintf (tmp, format, num);
dyn_string_append_cstr (str, tmp);
}
@@ -142,6 +146,7 @@ opencl_append_num_to_header (opencl_main code_gen, long num,
const char *format)
{
dyn_string_t tmp = opencl_get_current_header (code_gen);
+
opencl_append_int_to_str (tmp, num, format);
}
@@ -152,6 +157,7 @@ static void
opencl_append_num_to_body (opencl_main code_gen, long num, const char *format)
{
dyn_string_t tmp = opencl_get_current_body (code_gen);
+
opencl_append_int_to_str (tmp, num, format);
}
@@ -163,6 +169,7 @@ opencl_get_main_type (tree type)
while (TREE_CODE (type) == ARRAY_TYPE
|| TREE_CODE (type) == POINTER_TYPE)
type = TREE_TYPE (type);
+
return build_pointer_type (type);
}
@@ -174,6 +181,7 @@ opencl_create_function_code (opencl_body function)
{
static int opencl_function_counter = 0;
dyn_string_t dest = function->header;
+
dyn_string_append_cstr (dest, "__kernel void");
dyn_string_append_cstr (dest, " ");
dyn_string_append_cstr (dest, "opencl_auto_function_");
@@ -227,21 +235,26 @@ opencl_constant_expression_p (struct clast_expr *expr, const char *first_scat)
return opencl_cmp_scat (first_scat, name) == 1;
}
}
+
case clast_expr_red:
{
struct clast_reduction *red = (struct clast_reduction *) expr;
int i;
+
for (i = 0; i < red->n; i++)
if (!opencl_constant_expression_p (red->elts [i], first_scat))
return false;
return true;
}
+
case clast_expr_bin:
{
struct clast_binary *bin = (struct clast_binary *) expr;
+
return opencl_constant_expression_p (bin->LHS, first_scat);
}
+
default:
gcc_unreachable ();
return false;
@@ -286,6 +299,7 @@ opencl_get_perfect_nested_loop_depth (opencl_main code_gen,
int depth, const char *first_scat)
{
struct clast_for *child;
+
if (dependency_in_clast_loop_p (code_gen, meta, loop, depth))
return 0;
@@ -332,6 +346,7 @@ gen_type_1 (const char *ret_val, tree t)
case POINTER_TYPE:
if (TYPE_READONLY (t))
ret_val = concat ("const ", ret_val, NULL);
+
if (TYPE_VOLATILE (t))
ret_val = concat ("volatile ", ret_val, NULL);
@@ -348,8 +363,10 @@ gen_type_1 (const char *ret_val, tree t)
case ARRAY_TYPE:
if (!COMPLETE_TYPE_P (t) || TREE_CODE (TYPE_SIZE (t)) != INTEGER_CST)
ret_val = gen_type_1 (concat (ret_val, "[]", NULL), TREE_TYPE (t));
+
else if (int_size_in_bytes (t) == 0)
ret_val = gen_type_1 (concat (ret_val, "[0]", NULL), TREE_TYPE (t));
+
else
{
int size = int_size_in_bytes (t) / int_size_in_bytes (TREE_TYPE (t));
@@ -377,8 +394,10 @@ gen_type_1 (const char *ret_val, tree t)
case 64: data_type = "long"; break;
default: gcc_unreachable ();
}
+
if (TYPE_UNSIGNED (t))
data_type = concat ("unsigned ", data_type, NULL);
+
break;
case REAL_TYPE:
@@ -400,10 +419,13 @@ gen_type_1 (const char *ret_val, tree t)
if (TYPE_READONLY (t))
ret_val = concat ("const ", ret_val, NULL);
+
if (TYPE_VOLATILE (t))
ret_val = concat ("volatile ", ret_val, NULL);
+
if (TYPE_RESTRICT (t))
ret_val = concat ("restrict ", ret_val, NULL);
+
return ret_val;
}
@@ -414,6 +436,7 @@ static const char *
gen_type_with_name (const char *name, tree t)
{
const char *type_part = gen_type_1 (name, t);
+
return concat (data_type, " ", type_part, NULL);
}
@@ -423,15 +446,18 @@ gen_type_with_name (const char *name, tree t)
static const char *
opencl_get_var_name (tree node)
{
- bool ssa_name = TREE_CODE (node) == SSA_NAME;
+ 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)
@@ -440,6 +466,7 @@ opencl_get_var_name (tree node)
{
const char *base = identifier_to_locale (IDENTIFIER_POINTER (name));
char *buff = XNEWVEC (char, strlen (base) + 5);
+
sprintf (buff, "%s_%d", base, num);
return buff;
}
@@ -448,6 +475,7 @@ opencl_get_var_name (tree node)
{
int tmp_var_uid = DECL_UID (node);
char *tmp = XNEWVEC (char, 30);
+
sprintf (tmp, "opencl_var_%d_%d", tmp_var_uid, num);
return tmp;
}
@@ -459,9 +487,11 @@ static char *
filter_dots (char *p)
{
char *s;
+
for (s = p; *s; s++)
if (*s == '.')
*s = '_';
+
return p;
}
@@ -472,8 +502,8 @@ 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);
+ char *ddecl = xstrdup (decl);
+
return filter_dots (ddecl);
}
@@ -485,13 +515,16 @@ 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;
}
@@ -500,8 +533,10 @@ check_and_mark_arg (opencl_main code_gen, const char *name, bool local)
name, INSERT);
if (*slot)
return false;
+
if (!local)
*slot = name;
+
return true;
}
@@ -533,6 +568,7 @@ opencl_perfect_nested_to_kernel (opencl_main code_gen, struct clast_for *f,
int counter = perfect_depth;
tree curr_base = integer_one_node;
basic_block calc_block = opencl_create_bb (code_gen);
+
opencl_append_string_to_body
("size_t opencl_global_id = get_global_id (0);\n", code_gen);
@@ -672,9 +708,11 @@ opencl_print_local_vars (const char *fist, const char *last,
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++)
{
const char *tmp = names[i];
+
if (opencl_cmp_scat (fist, tmp) <= 0
&& opencl_cmp_scat (last, tmp) >= 0)
{
@@ -731,12 +769,10 @@ opencl_get_scat_real_name (opencl_main code_gen, clast_name_p name)
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);
+ tree type = TREE_TYPE (var);
+ opencl_body body = code_gen->current_body;
+ const char *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);
@@ -770,13 +806,17 @@ opencl_append_var_name (const char *name, opencl_main code_gen)
int len = strlen (name);
char *tmp = XNEWVEC (char, len + 1);
int i;
+
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);
}
@@ -793,6 +833,7 @@ opencl_print_term (struct clast_term *t, opencl_main code_gen)
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);
@@ -804,6 +845,7 @@ opencl_print_term (struct clast_term *t, opencl_main code_gen)
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
@@ -827,8 +869,10 @@ opencl_print_sum (struct clast_reduction *r, opencl_main 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);
}
}
@@ -842,15 +886,18 @@ static void
opencl_print_minmax_c ( struct clast_reduction *r, opencl_main code_gen)
{
int i;
+
for (i = 1; i < r->n; ++i)
- opencl_append_string_to_body (r->type == clast_red_max ? "max (" : "min (",
- code_gen);
+ 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);
@@ -871,6 +918,7 @@ opencl_print_reduction (struct clast_reduction *r, opencl_main code_gen)
case clast_red_sum:
opencl_print_sum (r, code_gen);
break;
+
case clast_red_min:
case clast_red_max:
if (r->n == 1)
@@ -878,8 +926,10 @@ opencl_print_reduction (struct clast_reduction *r, opencl_main code_gen)
opencl_print_expr (r->elts[0], code_gen);
break;
}
+
opencl_print_minmax_c (r, code_gen);
break;
+
default:
gcc_unreachable ();
}
@@ -900,15 +950,18 @@ opencl_print_binary (struct clast_binary *b, opencl_main 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
s1 = "", s2 = "/", s3 = "";
break;
+
case clast_bin_mod:
if (group)
s1 = "(", s2 = ")%", s3 = "";
@@ -932,17 +985,21 @@ opencl_print_expr (struct clast_expr *e, opencl_main code_gen)
{
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:
gcc_unreachable ();
}
@@ -1003,6 +1060,7 @@ opencl_add_non_scalar_type_decl (tree var, dyn_string_t dest,
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, " = (");
@@ -1011,6 +1069,7 @@ opencl_add_non_scalar_type_decl (tree var, dyn_string_t dest,
dyn_string_append_cstr (dest, decl_name);
dyn_string_append_cstr (dest, ";\n");
}
+
free (tmp_name);
}
@@ -1024,6 +1083,7 @@ static void
opencl_add_variable (const char *var_name, tree var, opencl_main code_gen)
{
const char **slot;
+
if (htab_find (code_gen->global_defined_vars, var_name))
{
opencl_append_var_name (var_name, code_gen);
@@ -1033,11 +1093,12 @@ opencl_add_variable (const char *var_name, tree var, opencl_main code_gen)
slot = (const char **) htab_find_slot
(code_gen->defined_vars, var_name, INSERT);
- if (! (*slot) && defined_in_sese_p (var, code_gen->region))
+ if (!(*slot) && defined_in_sese_p (var, code_gen->region))
{
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,
@@ -1048,8 +1109,10 @@ opencl_add_variable (const char *var_name, tree var, opencl_main code_gen)
decl = opencl_print_function_arg_with_type (var_name, type);
opencl_append_string_to_body (decl, code_gen);
}
+
return;
}
+
opencl_append_var_name (var_name, code_gen);
}
@@ -1061,6 +1124,7 @@ static void
opencl_try_variable (opencl_main code_gen, tree var_decl)
{
const char *name = opencl_get_var_name (var_decl);
+
gcc_assert (code_gen->defined_vars);
if (check_and_mark_arg (code_gen, name, false))
@@ -1089,10 +1153,12 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
{
case NOP_EXPR:
return opencl_print_operand (TREE_OPERAND (node, 0), false, code_gen);
+
case PLUS_EXPR:
{
if (lhs)
return -1;
+
opencl_append_string_to_body ("(", code_gen);
opencl_print_operand (TREE_OPERAND (node, 0), false, code_gen);
opencl_append_string_to_body (" + ", code_gen);
@@ -1100,10 +1166,12 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
opencl_append_string_to_body (")", code_gen);
return 0;
}
+
case MULT_EXPR:
{
if (lhs)
return -1;
+
opencl_append_string_to_body ("(", code_gen);
opencl_print_operand (TREE_OPERAND (node, 0), false, code_gen);
opencl_append_string_to_body (" * ", code_gen);
@@ -1117,6 +1185,7 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
/* If rhs just add variable name. Otherwise
it may be necessary to add variable definition. */
const char *tmp = opencl_get_var_name (node);
+
if (lhs)
opencl_add_variable (tmp, node, code_gen);
else
@@ -1127,62 +1196,77 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
opencl_try_variable (code_gen, node);
return 0;
}
+
case ARRAY_REF:
{
/* <operand>[<operand>]. */
tree arr = TREE_OPERAND (node, 0);
tree offset = TREE_OPERAND (node, 1);
- opencl_print_operand (arr, false, code_gen);
+ opencl_print_operand (arr, false, code_gen);
opencl_append_string_to_body ("[", code_gen);
opencl_print_operand (offset, false, code_gen);
opencl_append_string_to_body ("]", code_gen);
return 0;
}
+
case INTEGER_CST:
{
/* Just print integer constant. */
unsigned HOST_WIDE_INT low = TREE_INT_CST_LOW (node);
+
if (lhs)
return -1;
+
if (host_integerp (node, 0))
opencl_append_num_to_body (code_gen, (long)low, "%ld");
else
{
HOST_WIDE_INT high = TREE_INT_CST_HIGH (node);
char buff[100];
+
buff[0] = ' ';
+
if (tree_int_cst_sgn (node) < 0)
{
buff[0] = '-';
high = ~high + !low;
low = -low;
}
+
sprintf (buff + 1, HOST_WIDE_INT_PRINT_DOUBLE_HEX,
(unsigned HOST_WIDE_INT) high, low);
opencl_append_string_to_body (buff, code_gen);
}
+
return 0;
}
+
case REAL_CST:
{
char buff[100];
REAL_VALUE_TYPE tmp = TREE_REAL_CST (node);
+
if (lhs)
return -1;
+
real_to_decimal (buff, &tmp, sizeof (buff), 0, 1);
opencl_append_string_to_body (buff, code_gen);
return 0;
}
+
case FIXED_CST:
{
char buff[100];
+
if (lhs)
return -1;
+
fixed_to_decimal (buff, TREE_FIXED_CST_PTR (node), sizeof (buff));
opencl_append_string_to_body (buff, code_gen);
return 0;
}
+
case STRING_CST:
{
opencl_append_string_to_body ("\"", code_gen);
@@ -1190,11 +1274,13 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
opencl_append_string_to_body ("\"", code_gen);
return 0;
}
+
case VAR_DECL:
case PARM_DECL:
{
tree decl_name = DECL_NAME (node);
const char *tmp;
+
gcc_assert (decl_name);
tmp = IDENTIFIER_POINTER (decl_name);
@@ -1202,18 +1288,22 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
opencl_try_variable (code_gen, node);
return 0;
}
+
case FIELD_DECL:
{
tree decl_name = DECL_NAME (node);
const char *tmp;
+
gcc_assert (decl_name);
tmp = IDENTIFIER_POINTER (decl_name);
opencl_append_var_name (tmp, code_gen);
return 0;
}
+
case LABEL_DECL:
{
tree decl_name = DECL_NAME (node);
+
if (decl_name)
{
const char *tmp = IDENTIFIER_POINTER (decl_name);
@@ -1231,6 +1321,7 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
"D_%u");
return 0;
}
+
case INDIRECT_REF:
{
opencl_append_string_to_body ("(*", code_gen);
@@ -1238,16 +1329,19 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
opencl_append_string_to_body (")", code_gen);
return 0;
}
+
case ADDR_EXPR:
{
opencl_append_string_to_body ("&", code_gen);
opencl_print_operand (TREE_OPERAND (node, 0), false, code_gen);
return 0;
}
+
case COMPONENT_REF:
{
tree op1 = TREE_OPERAND (node, 0);
tree op2 = TREE_OPERAND (node, 1);
+
opencl_print_operand (op1, false, code_gen);
if (op1 && TREE_CODE (op1) == INDIRECT_REF)
@@ -1258,6 +1352,7 @@ opencl_print_operand (tree node, bool lhs, opencl_main code_gen)
opencl_print_operand (op2, false, code_gen);
return 0;
}
+
default:
debug_tree (node);
gcc_unreachable ();
@@ -1315,12 +1410,15 @@ opencl_print_unary (gimple gmp, opencl_main 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;
@@ -1347,18 +1445,22 @@ opencl_print_gimple_assign (gimple gmp, opencl_main code_gen)
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);
if (addr_expr)
@@ -1367,16 +1469,20 @@ opencl_print_gimple_assign (gimple gmp, opencl_main code_gen)
{
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);
}
@@ -1388,13 +1494,17 @@ 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);
}
+
opencl_append_string_to_body (")",code_gen);
}
@@ -1404,8 +1514,10 @@ 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;
}
@@ -1417,12 +1529,14 @@ 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);
@@ -1445,15 +1559,20 @@ opencl_print_gimple (gimple gmp, opencl_main code_gen)
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);
@@ -1461,6 +1580,7 @@ opencl_print_gimple (gimple gmp, opencl_main code_gen)
opencl_append_string_to_body (": ", code_gen);
}
break;
+
default:
debug_gimple_stmt (gmp);
gcc_unreachable ();
@@ -1479,11 +1599,14 @@ 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)
{
tree use = USE_FROM_PTR (use_p);
+
if (!is_gimple_reg (use))
continue;
+
opencl_build_defines (use, code_gen);
}
}
@@ -1496,8 +1619,7 @@ opencl_expand_scalar_vars (opencl_main code_gen, gimple stmt)
static void
opencl_build_defines (tree node, opencl_main code_gen)
{
- enum tree_code code = TREE_CODE (node);
- switch (code)
+ switch (TREE_CODE (node))
{
case SSA_NAME:
{
@@ -1520,6 +1642,7 @@ opencl_build_defines (tree node, opencl_main code_gen)
opencl_print_gimple (def_stmt, code_gen);
return;
}
+
case ARRAY_REF:
{
tree arr = TREE_OPERAND (node, 0);
@@ -1528,6 +1651,7 @@ opencl_build_defines (tree node, opencl_main code_gen)
opencl_build_defines (offset, code_gen);
return;
}
+
default:
gcc_unreachable ();
}
@@ -1540,6 +1664,7 @@ 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);
@@ -1599,12 +1724,14 @@ opencl_try_data_ref (opencl_main code_gen, data_reference_p ref,
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;
+
*slot = name;
opencl_add_non_scalar_function_arg (code_gen, data);
}
@@ -1618,6 +1745,7 @@ opencl_add_data_ref (opencl_main code_gen, data_reference_p d_ref)
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);
@@ -1630,6 +1758,7 @@ opencl_add_data_ref (opencl_main code_gen, data_reference_p d_ref)
tmp->read_in_current_body = true;
tmp->ever_read_on_device = true;
}
+
if (!tmp->privatized)
tmp->used_on_device = true;
@@ -1668,10 +1797,12 @@ opencl_print_user_stmt (struct clast_user_stmt *u, opencl_main code_gen)
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,
@@ -1707,13 +1838,16 @@ opencl_print_for (struct clast_for *f, opencl_main code_gen, int level)
tree iv_type;
const char *tmp;
const char *decl;
+
opencl_append_string_to_body ("for (", code_gen);
+
if (f->LB)
{
opencl_append_string_to_body (f->iterator, code_gen);
opencl_append_string_to_body ("=", code_gen);
opencl_print_expr (f->LB, code_gen);
}
+
opencl_append_string_to_body (";", code_gen);
if (f->UB)
@@ -1722,6 +1856,7 @@ opencl_print_for (struct clast_for *f, opencl_main code_gen, int level)
opencl_append_string_to_body ("<=", code_gen);
opencl_print_expr (f->UB, code_gen);
}
+
opencl_append_string_to_body (";", code_gen);
if (mpz_cmp_si (f->stride, 1) > 0)
@@ -1736,6 +1871,7 @@ opencl_print_for (struct clast_for *f, opencl_main code_gen, int level)
opencl_append_string_to_body ("++", code_gen);
opencl_append_string_to_body (")\n{\n", code_gen);
}
+
iv_type = opencl_get_loop_iter_type (f, code_gen, level);
iv = create_tmp_var (iv_type, "scat_tmp_iter");
@@ -1763,12 +1899,16 @@ 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);
}
@@ -1779,19 +1919,25 @@ static void
opencl_print_guard (struct clast_guard *g, opencl_main code_gen, int depth)
{
int k;
+
opencl_append_string_to_body ("if ", code_gen);
+
if (g->n > 1)
opencl_append_string_to_body ("(", code_gen);
+
for (k = 0; k < g->n; ++k)
{
if (k > 0)
opencl_append_string_to_body (" && ", code_gen);
+
opencl_append_string_to_body ("(", code_gen);
opencl_print_equation (&g->eq[k], code_gen);
opencl_append_string_to_body (")", code_gen);
}
+
if (g->n > 1)
opencl_append_string_to_body (")", code_gen);
+
opencl_append_string_to_body (" {\n", code_gen);
opencl_print_stmt_list (g->then, code_gen, depth);
opencl_append_string_to_body ("}\n", code_gen);
@@ -1803,29 +1949,36 @@ opencl_print_guard (struct clast_guard *g, opencl_main code_gen, int depth)
static void
opencl_print_stmt_list (struct clast_stmt *s, opencl_main code_gen, int depth)
{
- 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 ();
- }
+ 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 loop statement F. DEPTH is the depth of F in
@@ -1877,6 +2030,7 @@ opencl_clast_to_kernel (struct clast_for *f, opencl_main code_gen,
int depth)
{
opencl_body tmp = opencl_body_create ();
+
code_gen->current_body = tmp;
return opencl_print_loop (f, code_gen, depth);
}
@@ -76,6 +76,7 @@ static hashval_t
opencl_pair_to_hash (const void *data)
{
const struct opencl_pair_def *obj = (const struct opencl_pair_def *) data;
+
return (hashval_t) (obj->id);
}
@@ -96,6 +97,7 @@ static opencl_pair
opencl_pair_create (int new_id, int new_val)
{
opencl_pair tmp = XNEW (struct opencl_pair_def);
+
tmp->id = new_id;
tmp->val = new_val;
return tmp;
@@ -117,6 +119,7 @@ opencl_clast_meta_create (int depth, opencl_clast_meta parent,
bool access_init)
{
opencl_clast_meta tmp = XNEW (struct opencl_clast_meta_def);
+
tmp->out_depth = depth;
tmp->in_depth = 0;
tmp->next = NULL;
@@ -126,6 +129,7 @@ opencl_clast_meta_create (int depth, opencl_clast_meta parent,
tmp->modified_on_host = BITMAP_ALLOC (NULL);
tmp->modified_on_device = BITMAP_ALLOC (NULL);
tmp->access_unsupported = false;
+
if (access_init)
{
tmp->can_be_private = BITMAP_ALLOC (NULL);
@@ -136,6 +140,7 @@ opencl_clast_meta_create (int depth, opencl_clast_meta parent,
tmp->access = NULL;
tmp->can_be_private = NULL;
}
+
return tmp;
}
@@ -156,12 +161,14 @@ opencl_supported_type_p (tree type, bool ptr, bool array)
return false;
return opencl_supported_type_p (TREE_TYPE (type), true, false);
}
+
case ARRAY_TYPE:
{
if (ptr)
return false;
return opencl_supported_type_p (TREE_TYPE (type), false, true);
}
+
case FUNCTION_DECL:
case FUNCTION_TYPE:
case COMPLEX_TYPE:
@@ -176,9 +183,9 @@ opencl_supported_type_p (tree type, bool ptr, bool array)
case BOOLEAN_TYPE:
case INTEGER_TYPE:
case REAL_TYPE:
- return true;
case VOID_TYPE:
return true;
+
case OFFSET_TYPE:
case FIXED_POINT_TYPE:
case VECTOR_TYPE:
@@ -209,9 +216,10 @@ opencl_supported_arg_p (opencl_main code_gen, tree arg)
case PARM_DECL:
{
tree type = TREE_TYPE (arg);
- if (TREE_CODE (type) == POINTER_TYPE)
- if (!opencl_get_data_by_tree (code_gen, arg))
- return false;
+
+ if (TREE_CODE (type) == POINTER_TYPE
+ && !opencl_get_data_by_tree (code_gen, arg))
+ return false;
return opencl_supported_type_p (type, false, false);
}
@@ -246,23 +254,28 @@ opencl_gimple_assign_with_supported_types_p (opencl_main code_gen, gimple gmp)
{
tree curr_tree;
int num_of_ops = gimple_num_ops (gmp);
- gcc_assert (gimple_code (gmp) == GIMPLE_ASSIGN);
- gcc_assert (num_of_ops == 2 || num_of_ops == 3);
+
+ gcc_assert (gimple_code (gmp) == GIMPLE_ASSIGN
+ && (num_of_ops == 2 || num_of_ops == 3));
curr_tree = gimple_assign_lhs (gmp);
+
if (!opencl_supported_arg_p (code_gen, curr_tree))
return false;
curr_tree = gimple_assign_rhs1 (gmp);
+
if (!opencl_supported_arg_p (code_gen, curr_tree))
return false;
if (num_of_ops == 3)
{
curr_tree = gimple_assign_rhs2 (gmp);
+
if (!opencl_supported_arg_p (code_gen, curr_tree))
return false;
}
+
return true;
}
@@ -274,11 +287,14 @@ static bool
opencl_supported_type_access_p (opencl_main code_gen, basic_block bb)
{
gimple_stmt_iterator gsi;
+
for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
{
gimple stmt = gsi_stmt (gsi);
+
if (!stmt)
continue;
+
switch (gimple_code (stmt))
{
case GIMPLE_DEBUG:
@@ -286,6 +302,7 @@ opencl_supported_type_access_p (opencl_main code_gen, basic_block bb)
case GIMPLE_PHI:
case GIMPLE_LABEL:
continue;
+
case GIMPLE_ASSIGN:
if (!opencl_gimple_assign_with_supported_types_p (code_gen, stmt))
{
@@ -298,13 +315,16 @@ opencl_supported_type_access_p (opencl_main code_gen, basic_block bb)
return false;
}
continue;
+
case GIMPLE_CALL:
return false;
+
default:
debug_gimple_stmt (stmt);
gcc_unreachable ();
}
}
+
return true;
}
@@ -318,8 +338,10 @@ opencl_def_use_data (opencl_main code_gen, tree obj, bitmap visited,
opencl_clast_meta meta, bool def)
{
opencl_data data;
+
if (obj == NULL)
return;
+
data = opencl_get_data_by_tree (code_gen,
opencl_get_base_object_by_tree (obj));
if (data == NULL)
@@ -347,11 +369,14 @@ opencl_calc_bb_privatization (opencl_main code_gen, basic_block bb,
{
gimple_stmt_iterator gsi;
bitmap visited = BITMAP_ALLOC (NULL);
+
for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
{
gimple stmt = gsi_stmt (gsi);
+
if (gimple_code (stmt) != GIMPLE_ASSIGN)
continue;
+
opencl_def_use_data (code_gen, gimple_assign_lhs (stmt),
visited, meta, true);
@@ -361,6 +386,7 @@ opencl_calc_bb_privatization (opencl_main code_gen, basic_block bb,
opencl_def_use_data (code_gen, gimple_assign_rhs2 (stmt),
visited, meta, false);
}
+
BITMAP_FREE (visited);
}
@@ -382,6 +408,7 @@ opencl_set_meta_rw_flags (opencl_clast_meta meta,
basic_block bb = GBB_BB (gbb);
int i;
poly_dr_p curr;
+
if (!opencl_supported_type_access_p (code_gen, bb))
{
if (dump_file && (dump_flags & TDF_DETAILS))
@@ -394,7 +421,9 @@ opencl_set_meta_rw_flags (opencl_clast_meta meta,
meta->access_unsupported = true;
}
+
opencl_calc_bb_privatization (code_gen, bb, meta);
+
for (i = 0; VEC_iterate (poly_dr_p, drs, i, curr); i++)
{
data_reference_p d_ref = (data_reference_p) PDR_CDR (curr);
@@ -412,6 +441,7 @@ opencl_set_meta_rw_flags (opencl_clast_meta meta,
dump_data_reference (dump_file, d_ref);
}
+
continue;
}
@@ -424,6 +454,7 @@ opencl_set_meta_rw_flags (opencl_clast_meta meta,
if (!graphite_outer_subscript_bound (curr, false))
{
meta->access_unsupported = true;
+
if (dump_file && (dump_flags & TDF_DETAILS))
{
fprintf (dump_file, "Can not determine subscript bound "
@@ -438,6 +469,7 @@ opencl_set_meta_rw_flags (opencl_clast_meta meta,
if (data->size_value == NULL)
{
meta->access_unsupported = true;
+
if (dump_file && (dump_flags & TDF_DETAILS))
{
fprintf (dump_file,
@@ -445,6 +477,7 @@ opencl_set_meta_rw_flags (opencl_clast_meta meta,
dump_data_reference (dump_file, d_ref);
}
}
+
bitmap_set_bit (meta->access, data->id);
}
}
@@ -456,9 +489,11 @@ opencl_collect_definitions_info (opencl_clast_meta meta)
{
opencl_clast_meta curr = meta->body->next;
bitmap tmp_access = BITMAP_ALLOC (NULL);
+
bitmap_copy (tmp_access, meta->body->access);
meta->can_be_private = BITMAP_ALLOC (NULL);
bitmap_copy (meta->can_be_private, meta->body->can_be_private);
+
while (curr)
{
bitmap new_defs = BITMAP_ALLOC (NULL);
@@ -468,6 +503,7 @@ opencl_collect_definitions_info (opencl_clast_meta meta)
curr = curr->next;
BITMAP_FREE (new_defs);
}
+
meta->access = tmp_access;
}
@@ -486,12 +522,15 @@ opencl_create_meta_from_clast (opencl_main code_gen,
int max_depth = 0;
opencl_clast_meta result = NULL;
opencl_clast_meta curr = NULL;
- struct clast_stmt *curr_stmt = body;
- for ( ; curr_stmt; curr_stmt = curr_stmt->next)
+ struct clast_stmt *curr_stmt;
+
+ for (curr_stmt = body; curr_stmt; curr_stmt = curr_stmt->next)
{
opencl_clast_meta tmp_result = NULL;
+
if (CLAST_STMT_IS_A (curr_stmt, stmt_root))
continue;
+
if (CLAST_STMT_IS_A (curr_stmt, stmt_user))
{
tmp_result = opencl_clast_meta_create (depth, parent, true);
@@ -499,6 +538,7 @@ opencl_create_meta_from_clast (opencl_main code_gen,
(struct clast_user_stmt*) curr_stmt,
code_gen);
}
+
if (CLAST_STMT_IS_A (curr_stmt, stmt_guard))
{
struct clast_guard *if_stmt = (struct clast_guard *) curr_stmt;
@@ -507,15 +547,19 @@ opencl_create_meta_from_clast (opencl_main code_gen,
tmp_result = opencl_create_meta_from_clast (code_gen, if_stmt->then,
depth, parent);
}
+
if (CLAST_STMT_IS_A (curr_stmt, stmt_block))
{
struct clast_block *bl_stmt = (struct clast_block *) curr_stmt;
+
tmp_result = opencl_create_meta_from_clast (code_gen, bl_stmt->body,
depth, parent);
}
+
if (CLAST_STMT_IS_A (curr_stmt, stmt_for))
{
struct clast_for *for_stmt = (struct clast_for *) curr_stmt;
+
tmp_result = opencl_clast_meta_create (depth, parent, false);
tmp_result->body
= opencl_create_meta_from_clast (code_gen, for_stmt->body,
@@ -525,12 +569,16 @@ opencl_create_meta_from_clast (opencl_main code_gen,
? max_depth : tmp_result->in_depth + 1;
opencl_collect_definitions_info (tmp_result);
}
+
if (!result)
curr = result = tmp_result;
else
curr->next = tmp_result;
- while (curr->next != NULL) curr = curr->next;
+
+ while (curr->next)
+ curr = curr->next;
}
+
if (parent)
parent->in_depth = max_depth;
@@ -578,19 +626,23 @@ opencl_calc_max_depth_tab (opencl_clast_meta meta, htab_t data, int depth)
bitmap stmt_access = meta->access;
unsigned i;
bitmap_iterator bi;
+
if (meta->access_unsupported)
return false;
+
EXECUTE_IF_SET_IN_BITMAP (stmt_access, 0, i, bi)
{
opencl_pair curr_pair = opencl_pair_create (i, depth);
struct opencl_pair_def **slot
= (struct opencl_pair_def **) htab_find_slot (data, curr_pair,
INSERT);
+
if (*slot == NULL)
*slot = curr_pair;
else
{
opencl_pair old_pair = *slot;
+
if (old_pair->val > curr_pair->val)
opencl_pair_delete (curr_pair);
else
@@ -601,8 +653,10 @@ opencl_calc_max_depth_tab (opencl_clast_meta meta, htab_t data, int depth)
}
}
}
+
meta = meta->next;
}
+
return true;
}
@@ -653,20 +707,23 @@ opencl_evaluate_data_access_p (opencl_data obj, opencl_clast_meta meta)
int depth = obj->depth;
int data_id = obj->id;
opencl_clast_meta parent = meta->parent;
- if (obj->privatized)
- return false;
- if (depth < obj->data_dim)
+
+ if (obj->privatized
+ || depth < obj->data_dim)
return false;
+
if (parent)
{
/* We have outer loop. */
bitmap curr_bitmap = parent->modified_on_host;
+
/* Memory transfer for this statement has been placed outside
outer loop, so for one memory transfer will be executing more
then one kernel (first case). */
if (!bitmap_bit_p (curr_bitmap, data_id))
return true;
}
+
/* Check max depth of memory access (second case). */
return (depth > obj->data_dim);
}
@@ -678,6 +735,7 @@ opencl_get_data_by_id (opencl_main code_gen, int id)
{
VEC (opencl_data, heap) *main_data = code_gen->opencl_function_data;
opencl_data res = VEC_index (opencl_data, main_data, id);
+
gcc_assert (res->id == id);
return res;
}
@@ -705,9 +763,12 @@ opencl_analyse_data_access_p (opencl_main code_gen,
{
int id = curr->id;
opencl_data obj = opencl_get_data_by_id (code_gen, id);
+
VEC_safe_push (opencl_data, heap, data_objs, obj);
+
if (max_dim < obj->data_dim)
max_dim = obj->data_dim;
+
obj->depth = curr->val;
}
@@ -715,9 +776,11 @@ opencl_analyse_data_access_p (opencl_main code_gen,
{
if (curr_data->data_dim != max_dim)
continue;
+
if (opencl_evaluate_data_access_p (curr_data, meta))
return true;
}
+
return false;
}
@@ -738,28 +801,30 @@ opencl_should_be_parallel_p (opencl_main code_gen,
fprintf (dump_file, "opencl_should_be_parallel_p: ");
/* Avoid launching a lot of small kernels in a deep loop. */
- if (!flag_graphite_opencl_no_depth_check)
- if (depth > i_depth + opencl_base_depth_const)
- {
- if (dump_p)
- fprintf (dump_file, "avoiding small kernel in deep loop\n");
- return false;
- }
+ if (!flag_graphite_opencl_no_depth_check
+ && depth > i_depth + opencl_base_depth_const)
+ {
+ if (dump_p)
+ fprintf (dump_file, "avoiding small kernel in deep loop\n");
+
+ return false;
+ }
max_access_depth = htab_create (OPENCL_INIT_BUFF_SIZE,
opencl_pair_to_hash,
opencl_pair_cmp, free);
/* Can't parallelize if statements in loop contain unsupported types. */
- if (!flag_graphite_opencl_no_types_check)
- if (!opencl_calc_max_depth_tab (meta, max_access_depth, 0))
- {
- htab_delete (max_access_depth);
- if (dump_p)
- fprintf (dump_file, "unsupported types\n");
+ if (!flag_graphite_opencl_no_types_check
+ && !opencl_calc_max_depth_tab (meta, max_access_depth, 0))
+ {
+ htab_delete (max_access_depth);
- return false;
- }
+ if (dump_p)
+ fprintf (dump_file, "unsupported types\n");
+
+ return false;
+ }
/* Can't parallelize if memory transfer is not reasonable. */
if (!flag_graphite_opencl_no_memory_transfer_check
@@ -767,8 +832,10 @@ opencl_should_be_parallel_p (opencl_main code_gen,
&& !opencl_analyse_data_access_p (code_gen, max_access_depth, meta))
{
htab_delete (max_access_depth);
+
if (dump_p)
fprintf (dump_file, "avoiding large memory transfer\n");
+
return false;
}
@@ -128,6 +128,7 @@ map_ref_to_data_to_hash (const void *data)
{
const struct map_ref_to_data_def *obj
= (const struct map_ref_to_data_def *) data;
+
return htab_hash_pointer (obj->key);
}
@@ -151,6 +152,7 @@ map_ref_to_data_create (data_reference_p new_key,
opencl_data new_value)
{
map_ref_to_data tmp = XNEW (struct map_ref_to_data_def);
+
tmp->key = new_key;
tmp->value = new_value;
return tmp;
@@ -173,6 +175,7 @@ map_tree_to_data_to_hash (const void *data)
{
const struct map_tree_to_data_def *obj
= (const struct map_tree_to_data_def *) data;
+
return htab_hash_pointer (obj->key);
}
@@ -196,6 +199,7 @@ map_tree_to_data_create (tree new_key,
opencl_data new_value)
{
map_tree_to_data tmp = XNEW (struct map_tree_to_data_def);
+
tmp->key = new_key;
tmp->value = new_value;
return tmp;
@@ -208,6 +212,7 @@ static tree
opencl_create_tmp_var (tree type, const char *name)
{
tree tmp = create_tmp_var (type, name);
+
TREE_ADDRESSABLE (tmp) = 1;
return tmp;
}
@@ -237,6 +242,7 @@ opencl_fflush_rw_flags (opencl_main code_gen)
VEC (opencl_data, heap) *datas = code_gen->opencl_function_data;
int i;
opencl_data curr;
+
for (i = 0; VEC_iterate (opencl_data, datas, i, curr); i ++)
{
curr->written_in_current_body = false;
@@ -251,6 +257,7 @@ basic_block
opencl_create_bb (opencl_main code_gen)
{
basic_block tmp = split_edge (code_gen->main_edge);
+
code_gen->main_edge = single_succ_edge (tmp);
return tmp;
}
@@ -268,14 +275,18 @@ opencl_clast_meta_delete (opencl_clast_meta data)
{
if (!data)
return;
+
opencl_clast_meta_delete (data->body);
opencl_clast_meta_delete (data->next);
BITMAP_FREE (data->modified_on_device);
BITMAP_FREE (data->modified_on_host);
+
if (data->access != NULL)
BITMAP_FREE (data->access);
+
if (data->can_be_private)
BITMAP_FREE (data->can_be_private);
+
free (data);
}
@@ -326,10 +337,11 @@ static const char *opencl_function_names[] =
};
#endif
+
/* Variable, which holds OpenCL context. */
static GTY(()) tree h_context;
-/* Variable, which holds OpenCL comman queue. */
+/* Variable, which holds OpenCL command queue. */
static GTY(()) tree h_cmd_queue;
/* Variable, which holds OpenCL program for current function. */
@@ -355,16 +367,16 @@ zero_dim_array_p (tree var)
tree domain;
tree up_bound;
- if (TREE_CODE (type) != ARRAY_TYPE)
- return false;
- if (TREE_CODE (TREE_TYPE (type)) == ARRAY_TYPE)
- return false;
- domain = TYPE_DOMAIN (type);
- if (domain == NULL)
+ if (TREE_CODE (type) != ARRAY_TYPE
+ || TREE_CODE (TREE_TYPE (type)) == ARRAY_TYPE
+ || (domain = TYPE_DOMAIN (type)) == NULL)
return false;
- up_bound = TYPE_MAX_VALUE (domain);
+
+ up_bound = TYPE_MAX_VALUE (domain);
+
if (TREE_CODE (up_bound) != INTEGER_CST)
return false;
+
return TREE_INT_CST_LOW (up_bound) == 0;
}
@@ -378,8 +390,10 @@ opencl_private_var_name_p (const char *name)
static const char *close_phi = "Close_Phi";
static const char *cross_bb = "Cross_BB_scalar_dependence";
static const char *commutative = "Commutative_Associative_Reduction";
+
if (!name)
return false;
+
return
((strstr (name, general_reduction) == name)
|| (strstr (name, close_phi) == name)
@@ -393,13 +407,12 @@ static bool
graphite_artificial_array_p (tree var)
{
tree name;
- if (TREE_CODE (var) != VAR_DECL)
- return false;
- if (!zero_dim_array_p (var))
- return false;
- name = DECL_NAME (var);
- if (!name)
+
+ if (TREE_CODE (var) != VAR_DECL
+ || !zero_dim_array_p (var)
+ || !(name = DECL_NAME (var)))
return false;
+
return opencl_private_var_name_p (IDENTIFIER_POINTER (name));
}
@@ -409,12 +422,14 @@ 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;
}
@@ -462,6 +477,7 @@ opencl_main_create (CloogNames *names, sese region, edge main_edge,
htab_t params_index)
{
opencl_main tmp = XNEW (struct graphite_opencl_creator);
+
tmp->root_names = names;
tmp->defined_vars = NULL;
tmp->global_defined_vars = NULL;
@@ -490,15 +506,18 @@ opencl_main_delete (opencl_main data)
{
int i;
opencl_data curr;
+
dyn_string_delete (data->main_program);
htab_delete (data->newivs_index);
htab_delete (data->ref_to_data);
htab_delete (data->tree_to_data);
opencl_clast_meta_delete (data->clast_meta);
+
for (i = 0; VEC_iterate (opencl_data, data->opencl_function_data, i, curr);
i++)
if (!curr->is_static)
opencl_data_delete (curr);
+
VEC_free (tree, heap, data->newivs);
VEC_free (opencl_data, heap, data->opencl_function_data);
free (data);
@@ -520,6 +539,7 @@ opencl_add_safe_call_on_edge (tree call, bool zero_return, edge src)
{
basic_block bb = split_edge (src);
gimple_stmt_iterator g_iter = gsi_last_bb (bb);
+
force_gimple_operand_gsi (&g_iter, call, true, NULL, false,
GSI_CONTINUE_LINKING);
return single_succ_edge (bb);
@@ -532,6 +552,7 @@ opencl_add_safe_call_on_edge (tree call, bool zero_return, edge src)
tree abort_funtion;
tree abort_call;
gimple_stmt_iterator g_iter;
+
if (zero_return)
{
tree correct_result = build1 (CONVERT_EXPR, TREE_TYPE (call),
@@ -547,6 +568,7 @@ opencl_add_safe_call_on_edge (tree call, bool zero_return, edge src)
cmp = build2 (NE_EXPR, boolean_type_node,
call, incorrect_result);
}
+
result = create_empty_if_region_on_edge (src, cmp);
abort_bb = get_false_edge_from_guard_bb (src->dest)->dest;
abort_funtion = build_fn_decl ("abort", build_function_type_list
@@ -579,6 +601,7 @@ opencl_get_base_object_by_tree (tree obj)
while (TREE_CODE (obj) == INDIRECT_REF
|| TREE_CODE (obj) == ARRAY_REF)
obj = TREE_OPERAND (obj, 0);
+
return obj;
}
@@ -588,6 +611,7 @@ tree
dr_outermost_base_object (data_reference_p dr)
{
tree addr = DR_BASE_ADDRESS (dr);
+
if (!addr)
{
/* In case, we don't know base object. For example:
@@ -600,11 +624,13 @@ dr_outermost_base_object (data_reference_p dr)
Just return the innermost object when the base address is unknown. */
tree ref = DR_REF (dr);
+
return opencl_get_base_object_by_tree (ref);
}
if (TREE_CODE (addr) == ADDR_EXPR)
addr = TREE_OPERAND (addr, 0);
+
return addr;
}
@@ -617,6 +643,7 @@ static edge
opencl_get_edge_for_init (opencl_main code_gen, int data_id, bool device)
{
opencl_clast_meta curr = code_gen->curr_meta;
+
if (!curr)
return NULL;
@@ -625,10 +652,13 @@ opencl_get_edge_for_init (opencl_main code_gen, int data_id, bool device)
opencl_clast_meta parent = curr->parent;
bitmap curr_bitmap
= device ? parent->modified_on_host : parent->modified_on_device;
+
if (bitmap_bit_p (curr_bitmap, data_id))
break;
+
curr = curr->parent;
}
+
return curr->init_edge;
}
@@ -656,6 +686,7 @@ opencl_create_function_decl (enum OPENCL_FUNCTIONS id)
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:
@@ -977,24 +1008,61 @@ opencl_create_function_decl (enum OPENCL_FUNCTIONS id)
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;
+
+ 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 ();
}
}
@@ -1008,6 +1076,7 @@ opencl_wait_for_event (opencl_main code_gen, tree event_var)
tree call = build_call_expr (function, 2,
integer_one_node,
event_var);
+
opencl_add_safe_call (code_gen, call, true);
}
@@ -1087,10 +1156,12 @@ opencl_pass_to_device (opencl_main code_gen, opencl_data data)
integer_zero_node,
null_pointer_node,
event_call);
+
if (init_edge)
opencl_add_safe_call_on_edge (call, true, init_edge);
else
opencl_add_safe_call (code_gen, call, true);
+
data->up_to_date_on_device = true;
opencl_wait_for_event (code_gen, event_call);
return data->device_object;
@@ -1108,7 +1179,6 @@ opencl_pass_to_host (opencl_main code_gen, opencl_data data)
tree curr_type;
tree curr;
tree call;
-
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");
@@ -1116,8 +1186,10 @@ opencl_pass_to_host (opencl_main code_gen, opencl_data data)
TREE_STATIC (var) = 1;
assemble_variable (var, 1, 0, 1);
+
if (data->written_in_current_body)
data->up_to_date_on_device = false;
+
if (data->up_to_date_on_host)
{
if (!data->inited_in_current_loop_on_host
@@ -1132,10 +1204,8 @@ opencl_pass_to_host (opencl_main code_gen, opencl_data data)
data->inited_in_current_loop_on_host = true;
- if (flag_graphite_opencl_cpu)
- return;
-
- if (data->privatized)
+ if (flag_graphite_opencl_cpu
+ || data->privatized)
return;
init_edge = opencl_get_edge_for_init (code_gen, data->id, false);
@@ -1151,7 +1221,6 @@ opencl_pass_to_host (opencl_main code_gen, opencl_data data)
integer_zero_node, NULL_TREE, NULL_TREE);
event_call = build_addr (event_call, current_function_decl);
-
call = build_call_expr (function, 9,
h_cmd_queue,
data->device_object,
@@ -1166,6 +1235,7 @@ opencl_pass_to_host (opencl_main code_gen, opencl_data data)
opencl_add_safe_call_on_edge (call, true, init_edge);
else
opencl_add_safe_call (code_gen, call, true);
+
opencl_wait_for_event (code_gen, event_call);
data->up_to_date_on_host = true;
}
@@ -1181,20 +1251,21 @@ opencl_fflush_all_device_buffers_to_host (opencl_main code_gen)
int i;
opencl_data curr;
tree function = opencl_create_function_decl (RELEASE_MEMORY_OBJ);
+
for (i = 0; VEC_iterate (opencl_data, datas, i, curr); i ++)
{
curr->written_in_current_body = true;
opencl_pass_to_host (code_gen, curr);
}
+
for (i = 0; VEC_iterate (opencl_data, datas, i, curr); i ++)
- {
- if (curr->used_on_device && !curr->is_static)
- {
- tree var = curr->device_object;
- tree call = build_call_expr (function, 1, var);
- opencl_add_safe_call (code_gen, call, true);
- }
- }
+ if (curr->used_on_device && !curr->is_static)
+ {
+ tree var = curr->device_object;
+ tree call = build_call_expr (function, 1, var);
+
+ opencl_add_safe_call (code_gen, call, true);
+ }
}
/* Calculate correct flags for clCreateBuffer. READ means, that
@@ -1206,7 +1277,9 @@ 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
@@ -1216,10 +1289,12 @@ opencl_get_mem_flags (bool read, bool write)
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;
}
@@ -1269,9 +1344,11 @@ opencl_init_all_device_buffers (opencl_main code_gen)
int i;
opencl_data curr;
edge data_init_edge = single_succ_edge (code_gen->data_init_bb);
+
for (i = 0; VEC_iterate (opencl_data, datas, i, curr); i ++)
{
tree tmp;
+
if (!curr->used_on_device || curr->is_static)
continue;
@@ -1402,6 +1479,7 @@ static tree
opencl_get_indirect_size (tree ptr, poly_dr_p ref)
{
ptr = TREE_TYPE (ptr);
+
switch (TREE_CODE (ptr))
{
case ARRAY_TYPE:
@@ -1412,6 +1490,7 @@ opencl_get_indirect_size (tree ptr, poly_dr_p ref)
tree inner_type = TREE_TYPE (ptr);
tree t = graphite_outer_subscript_bound (ref, false);
tree inner_type_size = TYPE_SIZE_UNIT (inner_type);
+
if (inner_type_size == NULL)
return NULL;
@@ -1424,9 +1503,11 @@ opencl_get_indirect_size (tree ptr, poly_dr_p ref)
t = fold_build2 (MULT_EXPR, sizetype, t, inner_type_size);
return t;
}
+
default:
return NULL_TREE;
}
+
gcc_unreachable ();
}
@@ -1447,7 +1528,9 @@ opencl_init_local_device_memory (opencl_main code_gen, opencl_body kernel)
int i;
basic_block bb = opencl_create_bb (code_gen);
basic_block kernel_bb = split_edge (code_gen->kernel_edge);
+
code_gen->kernel_edge = single_succ_edge (kernel_bb);
+
for (i = 0; VEC_iterate (tree, *args, i, curr); i ++)
{
gimple_stmt_iterator g_iter = gsi_last_bb (bb);
@@ -1457,6 +1540,7 @@ opencl_init_local_device_memory (opencl_main code_gen, opencl_body kernel)
tree tmp_var;
tree mov;
tree curr_var = opencl_create_tmp_var (curr_type, "__ocl_iv");
+
if (TREE_CODE (curr) != PARM_DECL
&& TREE_CODE (curr) != VAR_DECL)
{
@@ -1480,13 +1564,13 @@ opencl_init_local_device_memory (opencl_main code_gen, opencl_body kernel)
GSI_CONTINUE_LINKING);
VEC_safe_push (tree, heap, *args_to_pass, tmp_var);
}
+
for (i = 0; VEC_iterate (opencl_data, *refs, i, curr_data); i++)
{
gimple_stmt_iterator kernel_g_iter = gsi_last_bb (kernel_bb);
tree new_type;
tree tmp_var;
tree mov;
-
tree curr = opencl_pass_to_device (code_gen, curr_data);
tree curr_type = ptr_type_node;
@@ -1521,6 +1605,7 @@ opencl_pass_kernel_arguments (opencl_main code_gen, opencl_body kernel,
tree arg;
int i;
tree function = opencl_create_function_decl (SET_KERNEL_ARG);
+
for (i = 0; VEC_iterate (tree, args_to_pass, i, arg); i++)
{
tree call
@@ -1554,7 +1639,6 @@ opencl_execute_kernel (opencl_main code_gen, tree num_of_exec,
call = build2 (MODIFY_EXPR, integer_type_node, num_of_threads, num_of_exec);
-
force_gimple_operand_gsi (&g_iter, call, true, NULL, false,
GSI_CONTINUE_LINKING);
@@ -1593,11 +1677,13 @@ opencl_create_function_call (edge base)
dyn_string_prepend_cstr
(main_program_src, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n ");
src = dyn_string_buf (main_program_src);
+
if (dump_file && (dump_flags & TDF_DETAILS))
{
fprintf (dump_file, "\nGenerated OpenCL code: \n");
fprintf (dump_file, "%s", src);
}
+
new_edge = opencl_insert_create_program_with_source_call (src, base);
return opencl_insert_build_program_call (new_edge);
@@ -1614,6 +1700,7 @@ opencl_mark_privatized_data (opencl_main code_gen)
int i;
opencl_data curr;
bitmap can_be_private = code_gen->curr_meta->can_be_private;
+
for (i = 0; VEC_iterate (opencl_data, datas, i, curr); i ++)
curr->privatized = bitmap_bit_p (can_be_private, curr->id);
@@ -1629,15 +1716,14 @@ opencl_set_data_size (opencl_main code_gen)
int i;
opencl_data curr;
gimple_stmt_iterator g_iter = gsi_last_bb (code_gen->data_init_bb);
+
for (i = 0; VEC_iterate (opencl_data, datas, i, curr); i ++)
{
tree call;
- if (curr->is_static)
- continue;
- if (!curr->used_on_device)
- continue;
- if (curr->size_value == NULL)
+ if (curr->is_static
+ || !curr->used_on_device
+ || curr->size_value == NULL)
continue;
call = build2 (MODIFY_EXPR, size_type_node,
@@ -1680,10 +1766,10 @@ opencl_data_init_object (opencl_data data)
/* (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);
- 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;
@@ -1709,13 +1795,11 @@ opencl_register_data (opencl_main code_gen, opencl_data 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);
@@ -1731,11 +1815,9 @@ 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;
+ opencl_data curr = opencl_get_data_by_tree (code_gen, data_ref_tree);
+ tree size = opencl_get_indirect_size (data_ref_tree, ref);
- 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)
@@ -1750,14 +1832,17 @@ opencl_parse_single_data_ref (poly_dr_p ref, opencl_main 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);
}
@@ -1770,6 +1855,7 @@ 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);
}
@@ -1783,6 +1869,7 @@ 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);
}
@@ -1811,6 +1898,7 @@ opencl_create_gimple_for_body (opencl_body kernel, opencl_main code_gen)
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);
@@ -1835,6 +1923,7 @@ opencl_prepare_memory_for_gimple_stmt (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++)
{
data_reference_p d_ref = (data_reference_p) PDR_CDR (curr);
@@ -1843,9 +1932,11 @@ opencl_prepare_memory_for_gimple_stmt (poly_bb_p pbb, opencl_main code_gen)
/* Scalar variables can be passed directly. */
data = opencl_get_data_by_data_ref (code_gen, d_ref);
+
/* Private variables should not be passed from device to host. */
if (data->privatized)
continue;
+
is_read = DR_IS_READ (d_ref);
gcc_assert (data);
@@ -1874,6 +1965,7 @@ opencl_add_gimple_for_user_stmt (struct clast_user_stmt *stmt,
VEC (tree, heap) *iv_map = VEC_alloc (tree, heap, nb_loops);
htab_t newivs_index = code_gen->newivs_index;
VEC (tree, heap) *newivs = code_gen->newivs;
+
/* Get basic block to add. */
gbb = PBB_BLACK_BOX (pbb);
@@ -1924,6 +2016,7 @@ opencl_init_new_loop (opencl_clast_meta meta, opencl_main code_gen)
{
opencl_data curr;
unsigned i;
+
meta->post_pass_to_host
= VEC_alloc (opencl_data, heap, OPENCL_INIT_BUFF_SIZE);
meta->post_pass_to_device
@@ -1948,6 +2041,7 @@ opencl_postpass_data (opencl_main code_gen, opencl_clast_meta meta)
{
opencl_data curr;
unsigned i;
+
for (i = 0;
VEC_iterate (opencl_data, meta->post_pass_to_host, i, curr); i++)
{
@@ -1962,6 +2056,7 @@ opencl_postpass_data (opencl_main code_gen, opencl_clast_meta meta)
curr->written_in_current_body = false;
opencl_pass_to_device (code_gen, curr);
}
+
if (meta->parent)
{
VEC (opencl_data, heap) *parent_vec_host
@@ -1977,6 +2072,7 @@ opencl_postpass_data (opencl_main code_gen, opencl_clast_meta meta)
VEC_iterate (opencl_data, meta->post_pass_to_device, i, curr); i++)
VEC_safe_push (opencl_data, heap, parent_vec_device, curr);
}
+
VEC_free (opencl_data, heap, meta->post_pass_to_host);
VEC_free (opencl_data, heap, meta->post_pass_to_device);
}
@@ -1994,19 +2090,14 @@ opencl_add_gimple_for_loop (struct clast_for *s, opencl_main code_gen,
int depth, bool dependency)
{
loop_p old_parent = code_gen->context_loop;
- loop_p new_loop
- = graphite_create_new_loop (code_gen->region,
- code_gen->main_edge,
- s, code_gen->context_loop,
- & code_gen->newivs,
- code_gen->newivs_index,
- code_gen->params_index,
- depth);
-
+ loop_p new_loop = graphite_create_new_loop
+ (code_gen->region, code_gen->main_edge, s, code_gen->context_loop,
+ &code_gen->newivs, code_gen->newivs_index, code_gen->params_index, depth);
edge last_e = single_exit (new_loop);
edge to_body = single_succ_edge (new_loop->header);
basic_block after = to_body->dest;
opencl_clast_meta parent = code_gen->curr_meta->parent;
+
last_e = single_succ_edge (split_edge (last_e));
code_gen->context_loop = new_loop;
@@ -2026,6 +2117,7 @@ opencl_add_gimple_for_loop (struct clast_for *s, opencl_main code_gen,
if (flag_loop_parallelize_all && !dependency)
new_loop->can_be_parallel = true;
+
opencl_verify ();
}
@@ -2043,10 +2135,9 @@ opencl_add_gimple_for_stmt_for (struct clast_for *s, opencl_main code_gen,
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;
+ code_gen->main_edge = true_e;
opencl_add_gimple_for_loop (s, code_gen, depth, dependency);
code_gen->main_edge = last_e;
}
@@ -2060,6 +2151,7 @@ static void
opencl_fix_meta_flags (opencl_clast_meta meta)
{
opencl_clast_meta curr = meta->body;
+
while (curr)
{
bitmap_ior_into (meta->modified_on_host, curr->modified_on_host);
@@ -2080,8 +2172,8 @@ opencl_add_gimple_for_stmt_guard (struct clast_guard *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;
@@ -2100,11 +2192,14 @@ opencl_transform_stmt_list (struct clast_stmt *s, opencl_main code_gen,
int depth)
{
bool dump_p = dump_file && (dump_flags & TDF_DETAILS);
+
for ( ; s; s = s->next)
{
opencl_clast_meta tmp = code_gen->curr_meta;
+
if (CLAST_STMT_IS_A (s, stmt_root))
continue;
+
else if (CLAST_STMT_IS_A (s, stmt_user))
{
code_gen->curr_meta->init_edge = code_gen->main_edge;
@@ -2118,6 +2213,7 @@ opencl_transform_stmt_list (struct clast_stmt *s, opencl_main code_gen,
struct clast_for *for_stmt = (struct clast_for *) s;
bool dependency = false;
bool parallel = false;
+
/* If there are dependencies in loop, it can't be parallelized. */
if (!flag_graphite_opencl_no_dep_check &&
dependency_in_clast_loop_p (code_gen, current_clast,
@@ -2125,8 +2221,10 @@ opencl_transform_stmt_list (struct clast_stmt *s, opencl_main code_gen,
{
if (dump_p)
fprintf (dump_file, "dependency in loop\n");
+
dependency = true;
}
+
if (!dependency)
parallel = opencl_should_be_parallel_p (code_gen, current_clast,
depth);
@@ -2137,6 +2235,7 @@ opencl_transform_stmt_list (struct clast_stmt *s, opencl_main code_gen,
if (parallel && !dependency)
{
opencl_body current_body;
+
opencl_fflush_rw_flags (code_gen);
opencl_mark_privatized_data (code_gen);
current_clast->on_device = true;
@@ -2179,6 +2278,7 @@ opencl_transform_stmt_list (struct clast_stmt *s, opencl_main code_gen,
code_gen, depth);
else
gcc_unreachable ();
+
if (tmp->parent)
opencl_fix_meta_flags (tmp->parent);
}
@@ -2212,6 +2312,7 @@ opencl_transform_clast (struct clast_stmt *data, sese region,
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);
@@ -2219,6 +2320,7 @@ opencl_transform_clast (struct clast_stmt *data, sese region,
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);
@@ -2286,13 +2388,13 @@ opencl_create_clGetContextInfo_1 (tree pointer_to_size)
tree zero_pointer = null_pointer_node;
tree cl_contex_devices = build_int_cst (NULL_TREE, CL_CONTEXT_DEVICES);
tree context_var = h_context;
- tree call = build_call_expr (function, 5,
- context_var,
- cl_contex_devices,
- integer_zero_node,
- zero_pointer,
- pointer_to_size);
- return call;
+
+ return build_call_expr (function, 5,
+ context_var,
+ cl_contex_devices,
+ integer_zero_node,
+ zero_pointer,
+ pointer_to_size);
}
/* Create call
@@ -2308,13 +2410,13 @@ opencl_create_clGetContextInfo_2 (tree size, tree a_devices)
tree zero_pointer = null_pointer_node;
tree cl_contex_devices = build_int_cst (NULL_TREE, CL_CONTEXT_DEVICES);
tree context_var = h_context;
- tree call = build_call_expr (function, 5,
- context_var,
- cl_contex_devices,
- size,
- a_devices,
- zero_pointer);
- return call;
+
+ return build_call_expr (function, 5,
+ context_var,
+ cl_contex_devices,
+ size,
+ a_devices,
+ zero_pointer);
}
/* Create context_properties array variable. */
@@ -2327,6 +2429,7 @@ opencl_create_context_properties (void)
tree index_type = build_index_type (build_int_cst (NULL_TREE, 3));
tree array_type = build_array_type (cl_context_properties_type,
index_type);
+
return opencl_create_tmp_var (array_type, "context_properties");
}
@@ -2389,14 +2492,13 @@ opencl_create_clCreateContextFromType (tree properties)
tree device
= build_int_cst (NULL_TREE, flag_graphite_opencl_cpu
? CL_DEVICE_TYPE_CPU : CL_DEVICE_TYPE_GPU);
- tree call;
- call = build_call_expr (function, 5,
+
+ return build_call_expr (function, 5,
build_addr (properties, current_function_decl),
device,
zero_pointer,
zero_pointer,
zero_pointer);
- return call;
}
/* Create call
@@ -2409,12 +2511,12 @@ opencl_create_clCreateCommandQueue (tree dev_id)
tree function = opencl_create_function_decl (CREATE_COMMAND_QUEUE);
tree zero_pointer = null_pointer_node;
tree context = h_context;
- tree call = build_call_expr (function, 4,
- context,
- dev_id,
- zero_pointer,
- zero_pointer);
- return call;
+
+ return build_call_expr (function, 4,
+ context,
+ dev_id,
+ zero_pointer,
+ zero_pointer);
}
/* Create call malloc (ARG). */
@@ -2422,14 +2524,12 @@ opencl_create_clCreateCommandQueue (tree dev_id)
static tree
opencl_create_malloc_call (tree arg)
{
- tree function_type =
- build_function_type_list (ptr_type_node,
- integer_type_node,
- NULL_TREE);
+ tree function_type = build_function_type_list (ptr_type_node,
+ integer_type_node,
+ NULL_TREE);
tree function = build_fn_decl ("malloc", function_type);
- tree call = build_call_expr (function, 1, arg);
- return call;
+ return build_call_expr (function, 1, arg);
}
/* Generate calls for opencl init functions and place them to INIT_EDGE.
@@ -2445,6 +2545,7 @@ opencl_create_init_context (edge init_edge)
tree a_devices = opencl_create_tmp_var (build_pointer_type (ptr_type_node),
"__ocl_a_devices");
tree properties = opencl_create_context_properties ();
+
init_edge = opencl_set_context_properties (init_edge, properties);
call = opencl_create_clCreateContextFromType (properties);
call = build2 (MODIFY_EXPR, TREE_TYPE (h_context),
@@ -2485,14 +2586,19 @@ build_poly_bb_vec (struct clast_stmt *root,
poly_bb_p tmp
= (poly_bb_p) cloog_statement_usr
(((struct clast_user_stmt *) root)->statement);
+
VEC_safe_push (poly_bb_p, heap, *vec, tmp);
}
+
else if (CLAST_STMT_IS_A (root, stmt_for))
build_poly_bb_vec (((struct clast_for *) root)->body, vec);
+
else if (CLAST_STMT_IS_A (root, stmt_guard))
build_poly_bb_vec (((struct clast_guard *) root)->then, vec);
+
else if (CLAST_STMT_IS_A (root, stmt_block))
build_poly_bb_vec (((struct clast_block *) root)->body, vec);
+
root = root->next;
}
}
@@ -2508,11 +2614,14 @@ opencl_dependency_between_pbbs_p (opencl_main code_gen, poly_bb_p pbb1,
{
int i, j;
poly_dr_p pdr1, pdr2;
+
timevar_push (TV_GRAPHITE_DATA_DEPS);
+
for (i = 0; VEC_iterate (poly_dr_p, PBB_DRS (pbb1), i, pdr1); i++)
{
data_reference_p ref1 = (data_reference_p)PDR_CDR (pdr1);
opencl_data data_1 = opencl_get_data_by_data_ref (code_gen, ref1);
+
if (bitmap_bit_p (can_be_private, data_1->id))
continue;
@@ -2549,14 +2658,15 @@ dependency_in_clast_loop_p (opencl_main code_gen, opencl_clast_meta meta,
int i;
poly_bb_p pbb1;
bitmap can_be_private;
- build_poly_bb_vec (stmt->body, &pbbs);
+ build_poly_bb_vec (stmt->body, &pbbs);
can_be_private = meta->can_be_private;
for (i = 0; VEC_iterate (poly_bb_p, pbbs, i, pbb1); i++)
{
int j;
poly_bb_p pbb2;
+
for (j = 0; VEC_iterate (poly_bb_p, pbbs, j, pbb2); j++)
if (opencl_dependency_between_pbbs_p (code_gen, pbb1, pbb1,
level, can_be_private))
@@ -2565,6 +2675,7 @@ dependency_in_clast_loop_p (opencl_main code_gen, opencl_clast_meta meta,
return true;
}
}
+
VEC_free (poly_bb_p, heap, pbbs);
return false;
}
@@ -2599,9 +2710,11 @@ opencl_init_static_data (edge init_edge)
{
int i;
opencl_data curr;
+
for (i = 0; VEC_iterate (opencl_data, opencl_array_data, i, curr); i ++)
{
tree tmp;
+
if (!curr->used_on_device)
continue;
@@ -2609,6 +2722,7 @@ opencl_init_static_data (edge init_edge)
tmp = build2 (MODIFY_EXPR, ptr_type_node, curr->device_object, tmp);
init_edge = opencl_add_safe_call_on_edge (tmp, false, init_edge);
}
+
return init_edge;
}
@@ -2621,6 +2735,7 @@ graphite_opencl_finalize (edge static_init_edge)
{
int i;
opencl_data curr;
+
if (dyn_string_length (main_program_src) != 0)
{
tree call = build2 (EQ_EXPR, boolean_type_node,
@@ -2644,7 +2759,9 @@ graphite_opencl_finalize (edge static_init_edge)
init_edge = get_true_edge_from_guard_bb (before_init->dest);
init_edge = opencl_create_init_context (init_edge);
}
+
dyn_string_delete (main_program_src);
+
for (i = 0; VEC_iterate (opencl_data, opencl_array_data, i, curr); i++)
opencl_data_delete (curr);
@@ -2810,6 +2927,7 @@ print_char_p_htab (void **h, void *v)
{
char **ptr = (char **) h;
FILE *file = (FILE *) v;
+
fprintf (file, " %s\n", *ptr);
return 1;
}
@@ -2821,6 +2939,7 @@ print_tree_to_data_htab (void **h, void *v)
FILE *file = (FILE *) v;
tree key = (*map)->key;
opencl_data data = (*map)->value;
+
print_node_brief (file, "key = ", key, 2);
fprintf (file, " data_id = %d\n", data->id);
return 1;
@@ -2833,6 +2952,7 @@ print_ref_to_data_htab (void **h, void *v)
FILE *file = (FILE *) v;
data_reference_p key = (*map)->key;
opencl_data data = (*map)->value;
+
fprintf (file, "key::\n");
dump_data_reference (file, key);
fprintf (file, "data_id = %d\n\n", data->id);
@@ -2845,11 +2965,13 @@ dump_opencl_main (opencl_main code_gen, FILE *file, bool verbose)
fprintf (file, "Current meta::\n");
dump_opencl_clast_meta (code_gen->curr_meta, file, false, 2);
fprintf (file, "\n");
+
if (code_gen->current_body)
{
fprintf (file, "Current body::\n");
dump_opencl_body (code_gen->current_body, file, verbose);
}
+
fprintf (file, "\n\nData init basic block::\n");
dump_bb (code_gen->data_init_bb, stderr, 0);
@@ -2866,6 +2988,7 @@ dump_opencl_main (opencl_main code_gen, FILE *file, bool verbose)
htab_traverse_noresize (code_gen->global_defined_vars,
print_char_p_htab, file);
}
+
fprintf (file, "Refs to data::\n");
htab_traverse_noresize (code_gen->ref_to_data,
print_ref_to_data_htab, file);