Patchwork [3/7] Fix formatting.

login
register
mail settings
Submitter Sebastian Pop
Date Dec. 28, 2010, 5:57 a.m.
Message ID <1293515882-16339-5-git-send-email-sebpop@gmail.com>
Download mbox | patch
Permalink /patch/76822/
State New
Headers show

Comments

Sebastian Pop - Dec. 28, 2010, 5:57 a.m.
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.
---
 gcc/ChangeLog.graphite           |    6 +
 gcc/graphite-opencl-codegen.c    |  230 ++++++++++++++++++++++-----
 gcc/graphite-opencl-meta-clast.c |  121 +++++++++++---
 gcc/graphite-opencl.c            |  331 ++++++++++++++++++++++++++------------
 4 files changed, 519 insertions(+), 169 deletions(-)

Patch

diff --git a/gcc/ChangeLog.graphite b/gcc/ChangeLog.graphite
index 256ae25..aea2188 100644
--- a/gcc/ChangeLog.graphite
+++ b/gcc/ChangeLog.graphite
@@ -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.
diff --git a/gcc/graphite-opencl-codegen.c b/gcc/graphite-opencl-codegen.c
index fdaaa7a..db1a59a 100644
--- a/gcc/graphite-opencl-codegen.c
+++ b/gcc/graphite-opencl-codegen.c
@@ -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);
 }
diff --git a/gcc/graphite-opencl-meta-clast.c b/gcc/graphite-opencl-meta-clast.c
index d562afe..ed81e77 100644
--- a/gcc/graphite-opencl-meta-clast.c
+++ b/gcc/graphite-opencl-meta-clast.c
@@ -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;
     }
 
diff --git a/gcc/graphite-opencl.c b/gcc/graphite-opencl.c
index 2aa659c..b0ef21d 100644
--- a/gcc/graphite-opencl.c
+++ b/gcc/graphite-opencl.c
@@ -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);