From patchwork Tue Dec 28 05:57:56 2010 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sebastian Pop X-Patchwork-Id: 76821 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) by ozlabs.org (Postfix) with SMTP id 0FE0DB70A3 for ; Tue, 28 Dec 2010 16:59:45 +1100 (EST) Received: (qmail 14804 invoked by alias); 28 Dec 2010 05:59:15 -0000 Received: (qmail 12909 invoked by uid 22791); 28 Dec 2010 05:58:41 -0000 X-SWARE-Spam-Status: No, hits=-0.5 required=5.0 tests=AWL, BAYES_20, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, FREEMAIL_FROM, RCVD_IN_DNSWL_LOW, RFC_ABUSE_POST, TW_GB, TW_TM, T_TO_NO_BRKTS_FREEMAIL X-Spam-Check-By: sourceware.org Received: from mail-gw0-f47.google.com (HELO mail-gw0-f47.google.com) (74.125.83.47) by sourceware.org (qpsmtpd/0.43rc1) with ESMTP; Tue, 28 Dec 2010 05:58:21 +0000 Received: by mail-gw0-f47.google.com with SMTP id a12so1542864gwa.20 for ; Mon, 27 Dec 2010 21:58:20 -0800 (PST) Received: by 10.150.158.16 with SMTP id g16mr17038836ybe.157.1293515900361; Mon, 27 Dec 2010 21:58:20 -0800 (PST) Received: from napoca (adsl-75-54-87-199.dsl.austtx.sbcglobal.net [75.54.87.199]) by mx.google.com with ESMTPS id p32sm9287007ybk.20.2010.12.27.21.58.15 (version=TLSv1/SSLv3 cipher=RC4-MD5); Mon, 27 Dec 2010 21:58:19 -0800 (PST) Received: by napoca (sSMTP sendmail emulation); Mon, 27 Dec 2010 23:58:14 -0600 From: Sebastian Pop To: gcc-patches@gcc.gnu.org Cc: gcc-graphite@googlegroups.com, amonakov@ispras.ru, kayrick@ispras.ru, Sebastian Pop Subject: [PATCH 1/7] Sort topologically static functions. Date: Mon, 27 Dec 2010 23:57:56 -0600 Message-Id: <1293515882-16339-3-git-send-email-sebpop@gmail.com> In-Reply-To: <1293515882-16339-1-git-send-email-sebpop@gmail.com> References: <1293515882-16339-1-git-send-email-sebpop@gmail.com> X-IsSubscribed: yes Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org 2010-12-27 Sebastian Pop * 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(-) 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 + + * graphite-opencl-codegen.c: Sort topologically static functions. + * graphite-opencl.c: Same. + 2010-12-25 Sebastian Pop * 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. _create - creates a new object of such type and returns it. _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. */