Patchwork [1/7] Sort topologically static functions.

login
register
mail settings
Submitter Sebastian Pop
Date Dec. 28, 2010, 5:57 a.m.
Message ID <1293515882-16339-3-git-send-email-sebpop@gmail.com>
Download mbox | patch
Permalink /patch/76821/
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: Sort topologically static functions.
	* graphite-opencl.c: Same.
---
 gcc/ChangeLog.graphite           |    5 +
 gcc/graphite-opencl-codegen.c    | 1588 +++++++++++++++++++-------------------
 gcc/graphite-opencl-meta-clast.c |    1 -
 gcc/graphite-opencl.c            | 1399 +++++++++++++++++-----------------
 4 files changed, 1469 insertions(+), 1524 deletions(-)

Patch

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