1 /* OpencCL code generation for GRAPHITE-OpenCL.
2 Copyright (C) 2009, 2010 Free Software Foundation, Inc.
4 This file is part of GCC.
6 GCC is free software; you can redistribute it and/or modify
7 it under the terms of the GNU General Public License as published by
8 the Free Software Foundation; either version 3, or (at your option)
11 GCC is distributed in the hope that it will be useful,
12 but WITHOUT ANY WARRANTY; without even the implied warranty of
13 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
14 GNU General Public License for more details.
16 You should have received a copy of the GNU General Public License
17 along with GCC; see the file COPYING3. If not see
18 <http://www.gnu.org/licenses/>. */
20 /* OpenCL code genration for GRAPHITE-OpenCL. This file implements
21 OpenCL kernel code generation from CLAST structures. */
25 #include "coretypes.h"
30 #include "basic-block.h"
31 #include "diagnostic.h"
32 #include "tree-flow.h"
34 #include "tree-dump.h"
37 #include "tree-chrec.h"
38 #include "tree-data-ref.h"
39 #include "tree-scalar-evolution.h"
40 #include "tree-pass.h"
42 #include "value-prof.h"
43 #include "pointer-set.h"
48 #include "gimple-pretty-print.h"
52 #include "cloog/cloog.h"
54 #include "graphite-ppl.h"
56 #include "graphite-poly.h"
57 #include "graphite-scop-detection.h"
58 #include "graphite-clast-to-gimple.h"
59 #include "graphite-dependences.h"
60 #include "dyn-string.h"
61 #include "graphite-opencl.h"
63 /* Compare two clast names based on their indexes. */
66 opencl_cmp_scat (const char *scat1
, const char *scat2
)
68 int len_1
= strlen (scat1
);
69 int len_2
= strlen (scat2
);
77 return strcmp (scat1
, scat2
);
80 /* This function implements !strcmp (STR1, STR2) call. */
83 opencl_cmp_str (const void *str1
, const void *str2
)
85 const char *c_str1
= (const char *) str1
;
86 const char *c_str2
= (const char *) str2
;
88 return !strcmp (c_str1
, c_str2
);
91 /* Get body of generating kernel function from CODE_GEN. */
94 opencl_get_current_body (opencl_main code_gen
)
96 return code_gen
->current_body
->body
;
99 /* Get header of generating kernel function from CODE_GEN. */
102 opencl_get_current_header (opencl_main code_gen
)
104 return code_gen
->current_body
->header
;
107 /* Appends a string STR to the header of the generating kernel
108 function from CODE_GEN. */
111 opencl_append_string_to_header (const char *str
, opencl_main code_gen
)
113 dyn_string_t tmp
= opencl_get_current_header (code_gen
);
115 dyn_string_append_cstr (tmp
, str
);
118 /* Appends a string STR to the body of the generating kernel function
122 opencl_append_string_to_body (const char *str
, opencl_main code_gen
)
124 dyn_string_t tmp
= opencl_get_current_body (code_gen
);
126 dyn_string_append_cstr (tmp
, str
);
129 /* Appends an integer NUM to the string STR following FORMAT. */
132 opencl_append_int_to_str (dyn_string_t str
, long num
, const char *format
)
136 sprintf (tmp
, format
, num
);
137 dyn_string_append_cstr (str
, tmp
);
140 /* Appends an integer NUM to the header of the generating kernel
141 function from CODE_GEN following FORMAT. */
144 opencl_append_num_to_header (opencl_main code_gen
, long num
,
147 dyn_string_t tmp
= opencl_get_current_header (code_gen
);
149 opencl_append_int_to_str (tmp
, num
, format
);
152 /* Appends an integer NUM to the body of the generating kernel
153 function from CODE_GEN following FORMAT. */
156 opencl_append_num_to_body (opencl_main code_gen
, long num
, const char *format
)
158 dyn_string_t tmp
= opencl_get_current_body (code_gen
);
160 opencl_append_int_to_str (tmp
, num
, format
);
163 /* Get TYPE scalar (base) part. */
166 opencl_get_main_type (tree type
)
168 while (TREE_CODE (type
) == ARRAY_TYPE
169 || TREE_CODE (type
) == POINTER_TYPE
)
170 type
= TREE_TYPE (type
);
172 return build_pointer_type (type
);
175 /* Create the base part of FUNCTION declaration, similar to this:
176 "__global void __opencl_function_0". */
179 opencl_create_function_code (opencl_body function
)
181 static int opencl_function_counter
= 0;
182 dyn_string_t dest
= function
->header
;
184 dyn_string_append_cstr (dest
, "__kernel void");
185 dyn_string_append_cstr (dest
, " ");
186 dyn_string_append_cstr (dest
, "opencl_auto_function_");
187 opencl_append_int_to_str (dest
, opencl_function_counter
, "%ld");
188 dyn_string_append_cstr (dest
, " (");
189 sprintf (function
->name
, "%s%d","opencl_auto_function_",
190 opencl_function_counter
++);
193 /* Create new instance of opencl_body. */
196 opencl_body_create (void)
198 opencl_body tmp
= XNEW (struct graphite_opencl_kernel_body
);
200 tmp
->body
= dyn_string_new (OPENCL_INIT_BUFF_SIZE
);
201 tmp
->pre_header
= dyn_string_new (OPENCL_INIT_BUFF_SIZE
);
202 tmp
->header
= dyn_string_new (OPENCL_INIT_BUFF_SIZE
);
203 tmp
->non_scalar_args
= dyn_string_new (OPENCL_INIT_BUFF_SIZE
);
205 tmp
->num_of_data_writes
= 0;
206 tmp
->function_args
= VEC_alloc (tree
, heap
, OPENCL_INIT_BUFF_SIZE
);
207 tmp
->function_args_to_pass
= VEC_alloc (tree
, heap
, OPENCL_INIT_BUFF_SIZE
);
208 tmp
->data_refs
= VEC_alloc (opencl_data
, heap
, OPENCL_INIT_BUFF_SIZE
);
209 opencl_create_function_code (tmp
);
214 /* Check whether clast expression EXPT is constant in current loop nest.
215 FIRST_SCAT is the iterator of outermost loop in current loop nest. */
218 opencl_constant_expression_p (struct clast_expr
*expr
, const char *first_scat
)
222 case clast_expr_term
:
224 struct clast_term
*term
= (struct clast_term
*) expr
;
230 const char *name
= clast_name_to_str (term
->var
);
231 if (strstr (name
, "scat_") != name
)
234 return opencl_cmp_scat (first_scat
, name
) == 1;
240 struct clast_reduction
*red
= (struct clast_reduction
*) expr
;
243 for (i
= 0; i
< red
->n
; i
++)
244 if (!opencl_constant_expression_p (red
->elts
[i
], first_scat
))
252 struct clast_binary
*bin
= (struct clast_binary
*) expr
;
254 return opencl_constant_expression_p (bin
->LHS
, first_scat
);
263 /* Check whether the clast_for LOOP has constant bounds. FIRST_SCAT
264 is the iterator of outermost loop in current loop nest. */
267 opencl_constant_loop_bound_p (struct clast_for
*loop
, const char *first_scat
)
269 return opencl_constant_expression_p (loop
->UB
, first_scat
)
270 && opencl_constant_expression_p (loop
->LB
, first_scat
);
273 /* If clast loop PARENT has only one child and it's a loop too, return
274 this child. Otherwise return NULL. */
276 static struct clast_for
*
277 opencl_get_single_loop_child (struct clast_for
*parent
)
279 struct clast_stmt
*body
= parent
->body
;
282 || !CLAST_STMT_IS_A (body
, stmt_for
))
285 return (struct clast_for
*) body
;
288 /* Calculate the maximal depth of a perfect nested loop nest with LOOP
289 as outermost loop. META holds meta information for loop LOOP,
290 DEPTH is the depth of LOOP in current loop nest, FIRST_SCAT is the
291 iterator of outermost loop in current loop nest. CODE_GEN holds
292 information related to OpenCL code generation. */
295 opencl_get_perfect_nested_loop_depth (opencl_main code_gen
,
296 opencl_clast_meta meta
,
297 struct clast_for
*loop
,
298 int depth
, const char *first_scat
)
300 struct clast_for
*child
;
302 if (dependency_in_clast_loop_p (code_gen
, meta
, loop
, depth
))
305 child
= opencl_get_single_loop_child (loop
);
308 || !opencl_constant_loop_bound_p (child
, first_scat
))
311 return 1 + opencl_get_perfect_nested_loop_depth (code_gen
, meta
->body
, child
,
312 depth
+ 1, first_scat
);
315 /* Get the type of the loop iterator for loop, represented by STMT.
316 LEVEL is the depth of this loop in current loop nest. CODE_GEN
317 holds information related to OpenCL code generation. */
320 opencl_get_loop_iter_type (struct clast_for
*stmt
, opencl_main code_gen
,
323 tree lb_type
= gcc_type_for_clast_expr (stmt
->LB
, code_gen
->region
,
325 code_gen
->newivs_index
,
326 code_gen
->params_index
);
327 tree ub_type
= gcc_type_for_clast_expr (stmt
->UB
, code_gen
->region
,
329 code_gen
->newivs_index
,
330 code_gen
->params_index
);
331 tree type
= gcc_type_for_iv_of_clast_loop (stmt
, level
, lb_type
, ub_type
);
336 static const char *data_type
;
338 /* Simplified version of C-style type printing from c-aux-info.c. */
341 gen_type_1 (const char *ret_val
, tree t
)
343 switch (TREE_CODE (t
))
346 if (TYPE_READONLY (t
))
347 ret_val
= concat ("const ", ret_val
, NULL
);
349 if (TYPE_VOLATILE (t
))
350 ret_val
= concat ("volatile ", ret_val
, NULL
);
352 ret_val
= concat ("*", ret_val
, NULL
);
354 if (TREE_CODE (TREE_TYPE (t
)) == ARRAY_TYPE
355 || TREE_CODE (TREE_TYPE (t
)) == FUNCTION_TYPE
)
356 ret_val
= concat ("(", ret_val
, ")", NULL
);
358 ret_val
= gen_type_1 (ret_val
, TREE_TYPE (t
));
363 if (!COMPLETE_TYPE_P (t
) || TREE_CODE (TYPE_SIZE (t
)) != INTEGER_CST
)
364 ret_val
= gen_type_1 (concat (ret_val
, "[]", NULL
), TREE_TYPE (t
));
366 else if (int_size_in_bytes (t
) == 0)
367 ret_val
= gen_type_1 (concat (ret_val
, "[0]", NULL
), TREE_TYPE (t
));
371 int size
= int_size_in_bytes (t
) / int_size_in_bytes (TREE_TYPE (t
));
373 sprintf (buff
, "[%d]", size
);
374 ret_val
= gen_type_1 (concat (ret_val
, buff
, NULL
), TREE_TYPE (t
));
378 case IDENTIFIER_NODE
:
379 data_type
= IDENTIFIER_POINTER (t
);
383 data_type
= IDENTIFIER_POINTER (DECL_NAME (t
));
387 case FIXED_POINT_TYPE
:
388 switch (TYPE_PRECISION (t
))
390 case 8: data_type
= "char"; break;
391 case 16: data_type
= "short"; break;
392 case 32: data_type
= "int"; break;
393 case 64: data_type
= "long"; break;
394 default: gcc_unreachable ();
397 if (TYPE_UNSIGNED (t
))
398 data_type
= concat ("unsigned ", data_type
, NULL
);
403 switch (TYPE_PRECISION (t
))
405 case 32: data_type
= "float"; break;
406 case 64: data_type
= "double"; break;
407 default: gcc_unreachable ();
419 if (TYPE_READONLY (t
))
420 ret_val
= concat ("const ", ret_val
, NULL
);
422 if (TYPE_VOLATILE (t
))
423 ret_val
= concat ("volatile ", ret_val
, NULL
);
425 if (TYPE_RESTRICT (t
))
426 ret_val
= concat ("restrict ", ret_val
, NULL
);
431 /* Generate a string representation of a declaration of varable named
435 gen_type_with_name (const char *name
, tree t
)
437 const char *type_part
= gen_type_1 (name
, t
);
439 return concat (data_type
, " ", type_part
, NULL
);
442 /* Get name of the variable, represented by tree NODE. If variable is
443 temporary, generate name for it. */
446 opencl_get_var_name (tree node
)
448 bool ssa_name
= (TREE_CODE (node
) == SSA_NAME
);
454 num
= SSA_NAME_VERSION (node
);
455 node
= SSA_NAME_VAR (node
);
458 name
= DECL_NAME (node
);
463 return identifier_to_locale (IDENTIFIER_POINTER (name
));
466 const char *base
= identifier_to_locale (IDENTIFIER_POINTER (name
));
467 char *buff
= XNEWVEC (char, strlen (base
) + 5);
469 sprintf (buff
, "%s_%d", base
, num
);
475 int tmp_var_uid
= DECL_UID (node
);
476 char *tmp
= XNEWVEC (char, 30);
478 sprintf (tmp
, "opencl_var_%d_%d", tmp_var_uid
, num
);
483 /* Replace all dots to underscores in string pointed to by P. Return P. */
486 filter_dots (char *p
)
497 /* Return string with varibale definition. ARG_NAME is the name of
498 the variable and TYPE is it's type. */
501 opencl_print_function_arg_with_type (const char *arg_name
, tree type
)
503 const char *decl
= gen_type_with_name (arg_name
, type
);
504 char *ddecl
= xstrdup (decl
);
506 return filter_dots (ddecl
);
509 /* Check whether variable with name NAME has been defined as global or
510 local variable and mark it as defined. This function returns false
511 if variable has already been defined, otherwise it returns true. */
514 check_and_mark_arg (opencl_main code_gen
, const char *name
, bool local
)
518 gcc_assert (code_gen
->defined_vars
|| !local
);
520 if (code_gen
->defined_vars
)
522 slot
= (const char **) htab_find_slot (code_gen
->defined_vars
,
531 slot
= (const char **) htab_find_slot (code_gen
->global_defined_vars
,
542 /* Replace perfect nested loop nest represented by F with opencl kernel.
543 For example, loop nest like this
545 | for (scat_1 = 0; scat_1 < M; i ++)
546 | for (scat_2 = 0; scat_2 < N; j ++)
547 | for (scat_3 = 0; scat_3 < L; k ++)
550 will be replased by kernel, where scat_1, scat_2, scat_3
551 depends on thread global id. Number of threads for this kernel
552 will be M * N * L. DEPTH is the depth of F in current loop nest.
553 CODE_GEN holds information related to OpenCL code generation. BODY holds
554 information current OpenCL kernel. */
557 opencl_perfect_nested_to_kernel (opencl_main code_gen
, struct clast_for
*f
,
558 opencl_body body
, int depth
)
560 VEC (tree
, heap
) *mod
= VEC_alloc (tree
, heap
, OPENCL_INIT_BUFF_SIZE
);
562 VEC (tree
, heap
) *function_args
= body
->function_args
;
563 const int perfect_depth
564 = opencl_get_perfect_nested_loop_depth (code_gen
, code_gen
->curr_meta
, f
,
566 struct clast_for
*curr
= f
;
567 int counter
= perfect_depth
;
568 tree curr_base
= integer_one_node
;
569 basic_block calc_block
= opencl_create_bb (code_gen
);
571 opencl_append_string_to_body
572 ("size_t opencl_global_id = get_global_id (0);\n", code_gen
);
574 body
->first_iter
= curr
->iterator
;
576 /* Iterate through all loops, which form perfect loop nest. */
580 sese region
= code_gen
->region
;
581 VEC (tree
, heap
) *newivs
= code_gen
->newivs
;
582 htab_t newivs_index
= code_gen
->newivs_index
;
583 htab_t params_index
= code_gen
->params_index
;
587 tree type
= opencl_get_loop_iter_type (curr
, code_gen
, depth
);
589 const char *type_str
= gen_type_with_name (" ", type
);
591 tree low_bound
= clast_to_gcc_expression (type
, curr
->LB
, region
,
592 newivs
, newivs_index
,
595 tree up_bound
= clast_to_gcc_expression (type
, curr
->UB
, region
,
596 newivs
, newivs_index
,
601 int num
= perfect_depth
- counter
- 1;
603 gcc_assert (curr
->LB
);
604 gcc_assert (curr
->UB
);
606 body
->last_iter
= curr
->iterator
;
608 low_bound
= opencl_tree_to_var (calc_block
, low_bound
);
609 up_bound
= opencl_tree_to_var (calc_block
, up_bound
);
611 if (mpz_cmp_si (curr
->stride
, 1) > 0)
612 stride
= mpz_get_si (curr
->stride
);
614 t_stride
= build_int_cst (NULL_TREE
, stride
);
616 curr_loop_size
= build2 (MINUS_EXPR
, type
,
617 up_bound
, low_bound
);
619 curr_loop_size
= build2 (PLUS_EXPR
, type
,
621 fold_convert (type
, integer_one_node
));
625 = build2 (TRUNC_DIV_EXPR
, type
,
626 curr_loop_size
, t_stride
);
628 curr_loop_size
= opencl_tree_to_var (calc_block
, curr_loop_size
);
630 VEC_safe_push (tree
, heap
, mod
, curr_loop_size
);
632 iv
= create_tmp_var (type
, "scat_tmp_iter");
634 /* Declare loop iterator as local variable. */
635 opencl_append_string_to_body (type_str
, code_gen
);
636 opencl_append_string_to_body (curr
->iterator
, code_gen
);
638 /* Calculate the value of the iterator of current loop based of the
639 number of current thread and numbers of iterators of inner loops. */
640 opencl_append_string_to_body (" = ((opencl_global_id / ", code_gen
);
641 opencl_append_num_to_body (code_gen
, num
, "opencl_base_%d");
642 opencl_append_string_to_body (") % ", code_gen
);
643 opencl_append_num_to_body (code_gen
, num
, "opencl_mod_%d");
644 opencl_append_string_to_body (") * ", code_gen
);
645 opencl_append_num_to_body (code_gen
, stride
, "%d");
646 opencl_append_num_to_body (code_gen
, num
, " + opencl_first_%d;\n");
648 opencl_append_string_to_header (type_str
, code_gen
);
649 opencl_append_num_to_header (code_gen
, num
, "opencl_mod_%d, ");
650 opencl_append_string_to_header (type_str
, code_gen
);
651 opencl_append_num_to_header (code_gen
, num
, "opencl_first_%d, ");
653 /* Declare old loop iterator. */
654 tmp
= opencl_get_var_name (iv
);
655 check_and_mark_arg (code_gen
, tmp
, false);
656 decl
= opencl_print_function_arg_with_type (tmp
, type
);
657 opencl_append_string_to_body (decl
, code_gen
);
658 opencl_append_string_to_body (" = ", code_gen
);
659 opencl_append_string_to_body (curr
->iterator
, code_gen
);
660 opencl_append_string_to_body (";\n", code_gen
);
662 save_clast_name_index (code_gen
->newivs_index
, curr
->iterator
,
663 VEC_length (tree
, code_gen
->newivs
));
664 VEC_safe_push (tree
, heap
, code_gen
->newivs
, iv
);
666 /* Save number of iterations for loop. */
667 VEC_safe_push (tree
, heap
, function_args
, curr_loop_size
);
668 VEC_safe_push (tree
, heap
, function_args
, low_bound
);
670 body
->clast_body
= curr
->body
;
671 curr
= opencl_get_single_loop_child (curr
);
675 counter
= perfect_depth
;
677 /* Store number of iteration of inner loops for each loop in perfect
681 tree type
= TREE_TYPE (curr_base
);
682 const char *type_str
= gen_type_with_name (" ", type
);
683 tree curr
= VEC_index (tree
, mod
, counter
);
685 opencl_append_string_to_header (type_str
, code_gen
);
686 opencl_append_num_to_header (code_gen
, counter
, "opencl_base_%d, ");
688 VEC_safe_push (tree
, heap
, function_args
, curr_base
);
690 curr_base
= build2 (MULT_EXPR
, type
, curr_base
,
691 build1 (CONVERT_EXPR
, type
, curr
));
692 curr_base
= opencl_tree_to_var (calc_block
, curr_base
);
695 body
->num_of_exec
= fold_convert (integer_type_node
, curr_base
);
696 VEC_free (tree
, heap
, mod
);
699 /* Append list of names of loop iterators from CODE_GEN with same type
700 TYPE to current kernel. FIRST and LAST define outermost and
701 innermost iterators to append respectively. */
704 opencl_print_local_vars (const char *fist
, const char *last
,
705 const char *type
, opencl_main code_gen
)
707 char **names
= cloog_names_scattering (code_gen
->root_names
);
708 int len
= cloog_names_nb_scattering (code_gen
->root_names
);
711 for (i
= 0; i
< len
; i
++)
713 const char *tmp
= names
[i
];
715 if (opencl_cmp_scat (fist
, tmp
) <= 0
716 && opencl_cmp_scat (last
, tmp
) >= 0)
719 (const char **) htab_find_slot (code_gen
->global_defined_vars
,
725 if (opencl_cmp_scat (fist
, tmp
) > 0)
728 opencl_append_string_to_body (type
, code_gen
);
729 opencl_append_string_to_body (" ", code_gen
);
730 opencl_append_string_to_body (tmp
, code_gen
);
731 opencl_append_string_to_body (";\n", code_gen
);
732 *((const char **) htab_find_slot (code_gen
->global_defined_vars
,
737 /* Return tree with variable, corresponging to given clast name NAME.
738 CODE_GEN holds information related to OpenCL code generation. */
741 opencl_clast_name_to_tree (opencl_main code_gen
, clast_name_p name
)
743 return clast_name_to_gcc (name
, code_gen
->region
, code_gen
->newivs
,
744 code_gen
->newivs_index
, code_gen
->params_index
);
747 /* For a given clast name return that name, if it's local name in
748 kernel body or, otherwise, name of gimple variable created for this
749 scat_i in gimple. CODE_GEN holds information related to OpenCL
753 opencl_get_scat_real_name (opencl_main code_gen
, clast_name_p name
)
755 const char *str
= clast_name_to_str (name
);
757 /* NAME > FIRST_ITER */
758 if (opencl_cmp_scat (str
, code_gen
->current_body
->first_iter
) >= 0)
761 return opencl_get_var_name (opencl_clast_name_to_tree (code_gen
, name
));
764 /* Add variable VAR with name NAME as function argument. Append it's
765 declaration in finction header and add it as function parameter.
766 CODE_GEN holds information related to OpenCL code generation. */
769 opencl_add_function_arg (opencl_main code_gen
, tree var
, const char *name
)
771 tree type
= TREE_TYPE (var
);
772 opencl_body body
= code_gen
->current_body
;
773 const char *decl
= opencl_print_function_arg_with_type (name
, type
);
775 dyn_string_append_cstr (body
->header
, decl
);
776 dyn_string_append_cstr (body
->header
, ", ");
777 VEC_safe_push (tree
, heap
, body
->function_args
, var
);
780 /* Add clast variable (scat_i) as kernel argument. NAME is a new name
781 of loop iterator (scat_*), REAL_NAME is an old (origin) name of
782 loop iterator. CODE_GEN holds information related to OpenCL code
786 opencl_add_scat_as_arg (opencl_main code_gen
, clast_name_p name
,
787 const char *real_name
)
791 if (!check_and_mark_arg (code_gen
, real_name
, false))
794 var
= opencl_clast_name_to_tree (code_gen
, name
);
795 opencl_add_function_arg (code_gen
, var
, real_name
);
798 /* Append variable name NAME to function body. Differs from appending
799 string by replacing `.' by `_'. CODE_GEN holds information related
800 to OpenCL code generation. */
803 opencl_append_var_name (const char *name
, opencl_main code_gen
)
805 int len
= strlen (name
);
806 char *tmp
= XNEWVEC (char, len
+ 1);
809 for (i
= 0; i
<= len
; i
++)
819 opencl_append_string_to_body (tmp
, code_gen
);
823 /* Generate code for clast term T. CODE_GEN holds information
824 related to OpenCL code generation. */
827 opencl_print_term (struct clast_term
*t
, opencl_main code_gen
)
831 const char *real_name
= opencl_get_scat_real_name (code_gen
, t
->var
);
833 if (mpz_cmp_si (t
->val
, 1) == 0)
834 opencl_append_var_name (real_name
, code_gen
);
836 else if (mpz_cmp_si (t
->val
, -1) == 0)
838 opencl_append_string_to_body ("-", code_gen
);
839 opencl_append_var_name (real_name
, code_gen
);
843 opencl_append_num_to_body (code_gen
, mpz_get_si (t
->val
), "%d");
844 opencl_append_string_to_body ("*", code_gen
);
845 opencl_append_var_name (real_name
, code_gen
);
848 opencl_add_scat_as_arg (code_gen
, t
->var
, real_name
);
851 opencl_append_num_to_body (code_gen
, mpz_get_si (t
->val
), "%d");
854 /* Generate code for clast sum statement R. CODE_GEN holds information
855 related to OpenCL code generation. */
858 opencl_print_sum (struct clast_reduction
*r
, opencl_main code_gen
)
861 struct clast_term
*t
;
863 gcc_assert (r
->n
>= 1 && r
->elts
[0]->type
== clast_expr_term
);
864 t
= (struct clast_term
*) r
->elts
[0];
865 opencl_print_term (t
, code_gen
);
867 for (i
= 1; i
< r
->n
; ++i
)
869 gcc_assert (r
->elts
[i
]->type
== clast_expr_term
);
870 t
= (struct clast_term
*) r
->elts
[i
];
872 if (mpz_sgn (t
->val
) > 0)
873 opencl_append_string_to_body ("+", code_gen
);
875 opencl_print_term (t
, code_gen
);
879 static void opencl_print_expr (struct clast_expr
*, opencl_main
);
881 /* Generate code for clast min/max operation R. CODE_GEN holds
882 information related to OpenCL code generation. */
885 opencl_print_minmax_c ( struct clast_reduction
*r
, opencl_main code_gen
)
889 for (i
= 1; i
< r
->n
; ++i
)
890 opencl_append_string_to_body
891 (r
->type
== clast_red_max
? "max (" : "min (", code_gen
);
895 opencl_append_string_to_body ("(unsigned int)(", code_gen
);
896 opencl_print_expr (r
->elts
[0], code_gen
);
897 opencl_append_string_to_body (")", code_gen
);
900 for (i
= 1; i
< r
->n
; ++i
)
902 opencl_append_string_to_body (",", code_gen
);
903 opencl_append_string_to_body ("(unsigned int)(", code_gen
);
904 opencl_print_expr (r
->elts
[i
], code_gen
);
905 opencl_append_string_to_body ("))", code_gen
);
909 /* Generate code for clast reduction statement R. CODE_GEN holds
910 information related to OpenCL code generation. */
913 opencl_print_reduction (struct clast_reduction
*r
, opencl_main code_gen
)
918 opencl_print_sum (r
, code_gen
);
925 opencl_print_expr (r
->elts
[0], code_gen
);
929 opencl_print_minmax_c (r
, code_gen
);
937 /* Generate code for clast binary operation B. CODE_GEN holds
938 information related to OpenCL code generation. */
941 opencl_print_binary (struct clast_binary
*b
, opencl_main code_gen
)
943 const char *s1
= NULL
, *s2
= NULL
, *s3
= NULL
;
944 bool group
= (b
->LHS
->type
== clast_expr_red
945 && ((struct clast_reduction
*) b
->LHS
)->n
> 1);
950 s1
= "floor ((", s2
= ")/(", s3
= "))";
954 s1
= "ceil ((", s2
= ")/(", s3
= "))";
959 s1
= "(", s2
= ")/", s3
= "";
961 s1
= "", s2
= "/", s3
= "";
966 s1
= "(", s2
= ")%", s3
= "";
968 s1
= "", s2
= "%", s3
= "";
972 opencl_append_string_to_body (s1
, code_gen
);
973 opencl_print_expr (b
->LHS
, code_gen
);
974 opencl_append_string_to_body (s2
, code_gen
);
975 opencl_append_num_to_body (code_gen
, mpz_get_si (b
->RHS
), "%d");
976 opencl_append_string_to_body (s3
, code_gen
);
979 /* Generate code for clast expression E. CODE_GEN holds information
980 related to OpenCL code generation. */
983 opencl_print_expr (struct clast_expr
*e
, opencl_main code_gen
)
990 case clast_expr_term
:
991 opencl_print_term ((struct clast_term
*) e
, code_gen
);
995 opencl_print_reduction ((struct clast_reduction
*) e
, code_gen
);
999 opencl_print_binary ((struct clast_binary
*) e
, code_gen
);
1007 /* Generate OpenCL code for clast_assignment A.
1008 CODE_GEN holds information related to OpenCL code generation. */
1011 opencl_print_assignment (struct clast_assignment
*a
, opencl_main code_gen
)
1013 /* Real assignment. */
1016 opencl_append_string_to_body (a
->LHS
, code_gen
);
1017 opencl_append_string_to_body (" = ", code_gen
);
1020 /* Just expression. */
1021 opencl_print_expr (a
->RHS
, code_gen
);
1024 /* Print operation simbol (`+' `-' `*') for assignment operation GMA.
1025 CODE_GEN holds information related to OpenCL code generation. */
1028 opencl_print_gimple_assign_operation (gimple gmp
, opencl_main code_gen
)
1030 opencl_append_string_to_body
1031 (op_symbol_code (gimple_assign_rhs_code (gmp
)), code_gen
);
1034 /* Generate definition for non scalar variable VAR and place it to
1035 string DEST. Use DECL_NAME as variable name. */
1038 opencl_add_non_scalar_type_decl (tree var
, dyn_string_t dest
,
1039 const char *decl_name
)
1041 tree type
= TREE_TYPE (var
);
1042 const char *name
= opencl_get_var_name (var
);
1043 static int counter
= 0;
1044 char type_name
[30];
1045 char *tmp_name
= xstrdup (name
);
1046 const char *new_type
;
1047 tree inner_type
= TREE_TYPE (type
);
1049 filter_dots (tmp_name
);
1051 sprintf (type_name
, "oclFTmpType%d", counter
++);
1053 new_type
= opencl_print_function_arg_with_type (type_name
, inner_type
);
1055 dyn_string_append_cstr (dest
, "typedef __global ");
1056 dyn_string_append_cstr (dest
, new_type
);
1057 dyn_string_append_cstr (dest
, ";\n");
1059 dyn_string_append_cstr (dest
, type_name
);
1060 dyn_string_append_cstr (dest
, " *");
1061 dyn_string_append_cstr (dest
, tmp_name
);
1063 if (decl_name
!= NULL
)
1065 dyn_string_append_cstr (dest
, " = (");
1066 dyn_string_append_cstr (dest
, type_name
);
1067 dyn_string_append_cstr (dest
, "*)");
1068 dyn_string_append_cstr (dest
, decl_name
);
1069 dyn_string_append_cstr (dest
, ";\n");
1075 /* Append variable VAR with name VAR_NAME to current function body.
1076 If variable has been defined in current scope, but definition for
1077 it has not been generated - then generate it's definition and mark
1078 variable as defined. CODE_GEN holds information related to OpenCL
1082 opencl_add_variable (const char *var_name
, tree var
, opencl_main code_gen
)
1086 if (htab_find (code_gen
->global_defined_vars
, var_name
))
1088 opencl_append_var_name (var_name
, code_gen
);
1092 slot
= (const char **) htab_find_slot
1093 (code_gen
->defined_vars
, var_name
, INSERT
);
1095 if (!(*slot
) && defined_in_sese_p (var
, code_gen
->region
))
1098 tree type
= TREE_TYPE (var
);
1101 if (TREE_CODE (type
) == POINTER_TYPE
1102 || TREE_CODE (type
) == ARRAY_TYPE
)
1103 opencl_add_non_scalar_type_decl (var
, code_gen
->current_body
->body
,
1107 var
= SSA_NAME_VAR (var
);
1108 decl
= opencl_print_function_arg_with_type (var_name
, type
);
1109 opencl_append_string_to_body (decl
, code_gen
);
1115 opencl_append_var_name (var_name
, code_gen
);
1118 /* If variable VAR_DECL is not defined and it is not marked as a
1119 parameter, mark it as a parameter and add it to parameters list.
1120 CODE_GEN holds information related to OpenCL code generation. */
1123 opencl_try_variable (opencl_main code_gen
, tree var_decl
)
1125 const char *name
= opencl_get_var_name (var_decl
);
1127 gcc_assert (code_gen
->defined_vars
);
1129 if (check_and_mark_arg (code_gen
, name
, false))
1130 opencl_add_function_arg (code_gen
, var_decl
, name
);
1133 /* Generate operand for tree node NODE. If LSH is true, generated
1134 operand must be lvalue, otherwise it's rvalue. CODE_GEN holds
1135 information related to OpenCL code generation. Also generate
1136 definitions for variables if necessary. Variable definition is not
1137 necessary if variable has already been defined or if it has been
1138 defined in other sese. */
1141 opencl_print_operand (tree node
, bool lhs
, opencl_main code_gen
)
1143 tree scev
= scalar_evolution_in_region (code_gen
->region
,
1144 code_gen
->context_loop
,
1146 tree new_node
= chrec_apply_map (scev
, code_gen
->iv_map
);
1148 if (TREE_CODE (new_node
) != SCEV_NOT_KNOWN
)
1151 switch (TREE_CODE (node
))
1154 return opencl_print_operand (TREE_OPERAND (node
, 0), false, code_gen
);
1161 opencl_append_string_to_body ("(", code_gen
);
1162 opencl_print_operand (TREE_OPERAND (node
, 0), false, code_gen
);
1163 opencl_append_string_to_body (" + ", code_gen
);
1164 opencl_print_operand (TREE_OPERAND (node
, 1), false, code_gen
);
1165 opencl_append_string_to_body (")", code_gen
);
1174 opencl_append_string_to_body ("(", code_gen
);
1175 opencl_print_operand (TREE_OPERAND (node
, 0), false, code_gen
);
1176 opencl_append_string_to_body (" * ", code_gen
);
1177 opencl_print_operand (TREE_OPERAND (node
, 1), false, code_gen
);
1178 opencl_append_string_to_body (")", code_gen
);
1184 /* If rhs just add variable name. Otherwise
1185 it may be necessary to add variable definition. */
1186 const char *tmp
= opencl_get_var_name (node
);
1189 opencl_add_variable (tmp
, node
, code_gen
);
1191 opencl_append_var_name (tmp
, code_gen
);
1193 /* This call adds variable declaration as formal
1194 parameter in kernel header if it is necessary. */
1195 opencl_try_variable (code_gen
, node
);
1201 /* <operand>[<operand>]. */
1202 tree arr
= TREE_OPERAND (node
, 0);
1203 tree offset
= TREE_OPERAND (node
, 1);
1205 opencl_print_operand (arr
, false, code_gen
);
1206 opencl_append_string_to_body ("[", code_gen
);
1207 opencl_print_operand (offset
, false, code_gen
);
1208 opencl_append_string_to_body ("]", code_gen
);
1214 /* Just print integer constant. */
1215 unsigned HOST_WIDE_INT low
= TREE_INT_CST_LOW (node
);
1220 if (host_integerp (node
, 0))
1221 opencl_append_num_to_body (code_gen
, (long)low
, "%ld");
1224 HOST_WIDE_INT high
= TREE_INT_CST_HIGH (node
);
1229 if (tree_int_cst_sgn (node
) < 0)
1232 high
= ~high
+ !low
;
1236 sprintf (buff
+ 1, HOST_WIDE_INT_PRINT_DOUBLE_HEX
,
1237 (unsigned HOST_WIDE_INT
) high
, low
);
1238 opencl_append_string_to_body (buff
, code_gen
);
1247 REAL_VALUE_TYPE tmp
= TREE_REAL_CST (node
);
1252 real_to_decimal (buff
, &tmp
, sizeof (buff
), 0, 1);
1253 opencl_append_string_to_body (buff
, code_gen
);
1264 fixed_to_decimal (buff
, TREE_FIXED_CST_PTR (node
), sizeof (buff
));
1265 opencl_append_string_to_body (buff
, code_gen
);
1271 opencl_append_string_to_body ("\"", code_gen
);
1272 opencl_append_string_to_body (TREE_STRING_POINTER (node
), code_gen
);
1273 opencl_append_string_to_body ("\"", code_gen
);
1280 tree decl_name
= DECL_NAME (node
);
1283 gcc_assert (decl_name
);
1284 tmp
= IDENTIFIER_POINTER (decl_name
);
1286 opencl_append_var_name (tmp
, code_gen
);
1287 opencl_try_variable (code_gen
, node
);
1293 tree decl_name
= DECL_NAME (node
);
1296 gcc_assert (decl_name
);
1297 tmp
= IDENTIFIER_POINTER (decl_name
);
1298 opencl_append_var_name (tmp
, code_gen
);
1304 tree decl_name
= DECL_NAME (node
);
1308 const char *tmp
= IDENTIFIER_POINTER (decl_name
);
1309 opencl_append_var_name (tmp
, code_gen
);
1313 if (LABEL_DECL_UID (node
) != -1)
1315 opencl_append_num_to_body (code_gen
, (int) LABEL_DECL_UID (node
),
1319 opencl_append_num_to_body (code_gen
, (int) DECL_UID (node
),
1326 opencl_append_string_to_body ("(*", code_gen
);
1327 opencl_print_operand (TREE_OPERAND (node
, 0), false, code_gen
);
1328 opencl_append_string_to_body (")", code_gen
);
1334 opencl_append_string_to_body ("&", code_gen
);
1335 opencl_print_operand (TREE_OPERAND (node
, 0), false, code_gen
);
1341 tree op1
= TREE_OPERAND (node
, 0);
1342 tree op2
= TREE_OPERAND (node
, 1);
1344 opencl_print_operand (op1
, false, code_gen
);
1346 if (op1
&& TREE_CODE (op1
) == INDIRECT_REF
)
1347 opencl_append_string_to_body ("->", code_gen
);
1349 opencl_append_string_to_body (".", code_gen
);
1351 opencl_print_operand (op2
, false, code_gen
);
1363 /* Generate code for min or max gimple operand GMP. CODE_GEN holds
1364 information related to OpenCL code generation. */
1367 opencl_print_max_min_assign (gimple gmp
, opencl_main code_gen
)
1369 tree lhs
= gimple_assign_lhs (gmp
);
1370 tree rhs1
= gimple_assign_rhs1 (gmp
);
1371 tree rhs2
= gimple_assign_rhs2 (gmp
);
1372 bool max
= gimple_assign_rhs_code (gmp
) == MAX_EXPR
;
1374 opencl_print_operand (lhs
, true, code_gen
);
1375 opencl_append_string_to_body (max
?" = fmax (":"= fmin (", code_gen
);
1376 opencl_print_operand (rhs1
, false, code_gen
);
1377 opencl_append_string_to_body (",", code_gen
);
1378 opencl_print_operand (rhs2
, false, code_gen
);
1379 opencl_append_string_to_body (");\n", code_gen
);
1382 /* Print pointer expression represented by EXPR. TYPE_SIZE represents
1383 size of the base type for EXPR. CODE_GEN holds information related
1384 to OpenCL code generation. */
1387 opencl_print_addr_operand (tree expr
, tree type_size
, opencl_main code_gen
)
1389 if (TREE_CODE (TREE_TYPE (expr
)) != POINTER_TYPE
)
1391 opencl_append_string_to_body ("(", code_gen
);
1392 opencl_print_operand (expr
, false, code_gen
);
1393 opencl_append_string_to_body ("/", code_gen
);
1394 opencl_print_operand (type_size
, false, code_gen
);
1395 opencl_append_string_to_body (")", code_gen
);
1398 opencl_print_operand (expr
, false, code_gen
);
1401 /* Print unary gimple operation GMP. CODE_GEN holds information
1402 related to OpenCL code generation. */
1405 opencl_print_unary (gimple gmp
, opencl_main code_gen
)
1407 switch (gimple_assign_rhs_code (gmp
))
1410 opencl_append_string_to_body ("~", code_gen
);
1413 case TRUTH_NOT_EXPR
:
1414 opencl_append_string_to_body ("!", code_gen
);
1418 opencl_append_string_to_body ("-", code_gen
);
1427 /* Generate code for gimple assignment statement GMP. CODE_GEN holds
1428 information related to OpenCL code generation. */
1431 opencl_print_gimple_assign (gimple gmp
, opencl_main code_gen
)
1433 int num_of_ops
= gimple_num_ops (gmp
);
1439 tree result_size
= NULL
;
1441 if (gimple_assign_rhs_code (gmp
) == MAX_EXPR
1442 || gimple_assign_rhs_code (gmp
) == MIN_EXPR
)
1444 opencl_print_max_min_assign (gmp
, code_gen
);
1448 gcc_assert (num_of_ops
== 2 || num_of_ops
== 3);
1449 lhs
= gimple_assign_lhs (gmp
);
1451 addr_expr
= (TREE_CODE (TREE_TYPE (lhs
)) == POINTER_TYPE
);
1454 result_size
= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (lhs
)));
1456 rhs1
= gimple_assign_rhs1 (gmp
);
1457 rhs2
= gimple_assign_rhs2 (gmp
);
1458 result
= opencl_print_operand (lhs
, true, code_gen
);
1463 opencl_append_string_to_body (" = ", code_gen
);
1466 opencl_print_addr_operand (rhs1
, result_size
, code_gen
);
1470 opencl_print_unary (gmp
, code_gen
);
1472 opencl_print_operand (rhs1
, false, code_gen
);
1475 if (rhs2
!= NULL_TREE
)
1477 opencl_print_gimple_assign_operation (gmp
, code_gen
);
1480 opencl_print_addr_operand (rhs2
, result_size
, code_gen
);
1482 opencl_print_operand (rhs2
, false, code_gen
);
1485 opencl_append_string_to_body (";\n",code_gen
);
1488 /* Generate code for arguments for gimple call statement GMP.
1489 CODE_GEN hold information related to OpenCL code generation. */
1492 opencl_print_gimple_call_args (opencl_main code_gen
, gimple gmp
)
1494 size_t len
= gimple_call_num_args (gmp
);
1497 opencl_append_string_to_body (" (",code_gen
);
1499 for (i
= 0; i
< len
; i
++)
1501 opencl_print_operand (gimple_call_arg (gmp
, i
), false, code_gen
);
1504 opencl_append_string_to_body (", ",code_gen
);
1507 opencl_append_string_to_body (")",code_gen
);
1510 /* Replace some function names. */
1513 opencl_get_function_name (tree function
)
1515 const char *gimple_name
= IDENTIFIER_POINTER (DECL_NAME (function
));
1517 if (!strcmp (gimple_name
, "__builtin_powf"))
1523 /* Generate code for gimple call statement GMP. CODE_GEN holds information
1524 related to OpenCL code generation. */
1527 opencl_print_gimple_call (opencl_main code_gen
, gimple gmp
)
1529 tree lhs
= gimple_call_lhs (gmp
);
1530 tree function
= gimple_call_fn (gmp
);
1532 opencl_print_operand (lhs
, true, code_gen
);
1533 opencl_append_string_to_body (" = ", code_gen
);
1535 while (TREE_CODE (function
) == ADDR_EXPR
1536 || TREE_CODE (function
) == INDIRECT_REF
)
1537 function
= TREE_OPERAND (function
, 0);
1539 opencl_append_string_to_body (opencl_get_function_name (function
), code_gen
);
1540 opencl_print_gimple_call_args (code_gen
, gmp
);
1541 opencl_append_string_to_body (";\n",code_gen
);
1544 /* Generate code for gimple statment SMP. Now only assignment
1545 operation are supported, but it seems enough for clast translation.
1546 GIMPLE_COND statements are loop bound conditions and can be safely
1547 ignored. CODE_GEN holds information related to OpenCL code
1551 opencl_print_gimple (gimple gmp
, opencl_main code_gen
)
1556 switch (gimple_code (gmp
))
1559 opencl_print_gimple_assign (gmp
, code_gen
);
1569 opencl_print_gimple_call (code_gen
, gmp
);
1577 tree label
= gimple_label_label (gmp
);
1578 opencl_print_operand (label
, false, code_gen
);
1579 opencl_append_string_to_body (": ", code_gen
);
1584 debug_gimple_stmt (gmp
);
1589 static void opencl_build_defines (tree
, opencl_main
);
1591 /* For a given gimple statement STMT build definition for all names,
1592 used in this stament. If name has been defined in other sese, mark
1593 it as kernel parameter. CODE_GEN holds information related to
1594 OpenCL code generation. */
1597 opencl_expand_scalar_vars (opencl_main code_gen
, gimple stmt
)
1600 use_operand_p use_p
;
1602 FOR_EACH_SSA_USE_OPERAND (use_p
, stmt
, iter
, SSA_OP_ALL_USES
)
1604 tree use
= USE_FROM_PTR (use_p
);
1606 if (!is_gimple_reg (use
))
1609 opencl_build_defines (use
, code_gen
);
1613 /* If tree node NODE defined in current sese build and insert define
1614 statements for it, otherwise mark node as external (parameter for
1615 kernel). If tree defined in current sese, also recursively build
1616 defines for all trees in definition expression. */
1619 opencl_build_defines (tree node
, opencl_main code_gen
)
1621 switch (TREE_CODE (node
))
1625 const char *tmp
= opencl_get_var_name (node
);
1628 /* If name defined in other sese it is kernel's parameter. */
1629 if (!defined_in_sese_p (node
, code_gen
->region
))
1632 /* Bail out if this name was defined earlier either in this
1634 if (*(const char **) htab_find_slot (code_gen
->defined_vars
,
1638 /* Get definition statement. */
1639 def_stmt
= SSA_NAME_DEF_STMT (node
);
1640 opencl_expand_scalar_vars (code_gen
, def_stmt
);
1641 opencl_print_gimple (def_stmt
, code_gen
);
1647 tree arr
= TREE_OPERAND (node
, 0);
1648 tree offset
= TREE_OPERAND (node
, 1);
1649 opencl_build_defines (arr
, code_gen
);
1650 opencl_build_defines (offset
, code_gen
);
1659 /* Generate code for a single basic block BB. CODE_GEN holds
1660 information related to OpenCL code generation. */
1663 opencl_print_bb (basic_block bb
, opencl_main code_gen
)
1665 gimple_stmt_iterator gsi
;
1667 for (gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
1669 gimple stmt
= gsi_stmt (gsi
);
1670 opencl_expand_scalar_vars (code_gen
, stmt
);
1671 opencl_print_gimple (stmt
, code_gen
);
1675 /* Define non scalar variable, represented be DATA as either local
1676 variable or kernel argument. CODE_GEN holds information related to
1677 OpenCL code generation. */
1680 opencl_add_non_scalar_function_arg (opencl_main code_gen
,
1684 static int counter
= 0;
1685 opencl_body body
= code_gen
->current_body
;
1686 tree var
= data
->exact_object
;
1687 const char *name
= opencl_get_var_name (var
);
1688 tree type
= TREE_TYPE (var
);
1690 /* Check whether given variable can be privatized. */
1691 if (data
->privatized
)
1693 /* Define variable as local variable. */
1694 gcc_assert (TREE_CODE (type
) == ARRAY_TYPE
);
1695 decl
= opencl_print_function_arg_with_type (name
, type
);
1696 dyn_string_append_cstr (body
->pre_header
, decl
);
1697 dyn_string_append_cstr (body
->pre_header
, ";\n");
1702 /* Define variable as kernel argument. */
1703 char decl_name
[30];
1704 tree main_type
= opencl_get_main_type (type
);
1705 sprintf (decl_name
, "oclFTmpArg%d", counter
++);
1706 decl
= opencl_print_function_arg_with_type (decl_name
, main_type
);
1707 dyn_string_append_cstr (body
->non_scalar_args
, "__global ");
1708 opencl_add_non_scalar_type_decl (var
, body
->pre_header
, decl_name
);
1709 dyn_string_append_cstr (body
->non_scalar_args
, decl
);
1710 dyn_string_append_cstr (body
->non_scalar_args
, ", ");
1711 VEC_safe_push (opencl_data
, heap
, body
->data_refs
, data
);
1715 /* Register data reference REF to variable DATA. Do nothing, if it
1716 has already been registered. CODE_GEN holds information related to
1717 OpenCL code generation. */
1720 opencl_try_data_ref (opencl_main code_gen
, data_reference_p ref
,
1723 tree var
= dr_outermost_base_object (ref
);
1724 const char *name
= opencl_get_var_name (var
);
1727 gcc_assert (code_gen
->defined_vars
);
1729 slot
= (const char **) htab_find_slot (code_gen
->global_defined_vars
,
1735 opencl_add_non_scalar_function_arg (code_gen
, data
);
1738 /* Register data reference D_REF in current kernel. CODE_GEN hold
1739 information related to OpenCL code generation. */
1742 opencl_add_data_ref (opencl_main code_gen
, data_reference_p d_ref
)
1744 opencl_data tmp
= opencl_get_data_by_data_ref (code_gen
, d_ref
);
1748 if (!DR_IS_READ (d_ref
))
1750 bitmap_set_bit (code_gen
->curr_meta
->modified_on_device
, tmp
->id
);
1751 tmp
->written_in_current_body
= true;
1752 tmp
->ever_written_on_device
= true;
1753 code_gen
->current_body
->num_of_data_writes
++;
1757 tmp
->read_in_current_body
= true;
1758 tmp
->ever_read_on_device
= true;
1761 if (!tmp
->privatized
)
1762 tmp
->used_on_device
= true;
1764 opencl_try_data_ref (code_gen
, d_ref
, tmp
);
1767 /* Add base objects of all data references in PBB as arguments to
1768 current kernel. CODE_GEN holds information related to OpenCL code
1772 opencl_add_data_refs_pbb (poly_bb_p pbb
, opencl_main code_gen
)
1774 VEC (poly_dr_p
, heap
) *drs
= PBB_DRS (pbb
);
1778 FOR_EACH_VEC_ELT (poly_dr_p
, drs
, i
, curr
)
1780 data_reference_p d_ref
= (data_reference_p
) PDR_CDR (curr
);
1781 opencl_add_data_ref (code_gen
, d_ref
);
1785 /* Generate OpenCL code for user statement U. Code will be generated
1786 from basic block, related to U. Also induction variables mapping
1787 to old variables must be calculated to process basic block.
1788 CODE_GEN holds information related to OpenCL code generation. */
1791 opencl_print_user_stmt (struct clast_user_stmt
*u
, opencl_main code_gen
)
1798 int nb_loops
= number_of_loops ();
1800 code_gen
->iv_map
= VEC_alloc (tree
, heap
, nb_loops
);
1802 for (i
= 0; i
< nb_loops
; i
++)
1803 VEC_safe_push (tree
, heap
, code_gen
->iv_map
, NULL_TREE
);
1805 build_iv_mapping (code_gen
->iv_map
, code_gen
->region
,
1807 code_gen
->newivs_index
, u
,
1808 code_gen
->params_index
);
1810 code_gen
->defined_vars
1811 = htab_create (10, htab_hash_string
, opencl_cmp_str
, NULL
);
1812 opencl_append_string_to_body ("{\n", code_gen
);
1815 pbb
= (poly_bb_p
) cloog_statement_usr (cs
);
1816 gbbp
= PBB_BLACK_BOX (pbb
);
1818 code_gen
->context_loop
= bb
->loop_father
;
1820 opencl_add_data_refs_pbb (pbb
, code_gen
);
1821 opencl_print_bb (bb
, code_gen
);
1822 opencl_append_string_to_body ("}\n", code_gen
);
1823 htab_delete (code_gen
->defined_vars
);
1824 code_gen
->defined_vars
= NULL
;
1825 VEC_free (tree
, heap
, code_gen
->iv_map
);
1828 static void opencl_print_stmt_list (struct clast_stmt
*, opencl_main
, int);
1830 /* Generate code for clast for statement F, locate on depth LEVEL.
1831 CODE_GEN holds information related to OpenCL code generation. */
1834 opencl_print_for (struct clast_for
*f
, opencl_main code_gen
, int level
)
1841 opencl_append_string_to_body ("for (", code_gen
);
1845 opencl_append_string_to_body (f
->iterator
, code_gen
);
1846 opencl_append_string_to_body ("=", code_gen
);
1847 opencl_print_expr (f
->LB
, code_gen
);
1850 opencl_append_string_to_body (";", code_gen
);
1854 opencl_append_string_to_body (f
->iterator
, code_gen
);
1855 opencl_append_string_to_body ("<=", code_gen
);
1856 opencl_print_expr (f
->UB
, code_gen
);
1859 opencl_append_string_to_body (";", code_gen
);
1861 if (mpz_cmp_si (f
->stride
, 1) > 0)
1863 opencl_append_string_to_body (f
->iterator
, code_gen
);
1864 opencl_append_string_to_body ("+=", code_gen
);
1865 opencl_append_num_to_body (code_gen
, mpz_get_si (f
->stride
), "%d)\n{\n");
1869 opencl_append_string_to_body (f
->iterator
, code_gen
);
1870 opencl_append_string_to_body ("++", code_gen
);
1871 opencl_append_string_to_body (")\n{\n", code_gen
);
1874 iv_type
= opencl_get_loop_iter_type (f
, code_gen
, level
);
1875 iv
= create_tmp_var (iv_type
, "scat_tmp_iter");
1877 tmp
= opencl_get_var_name (iv
);
1878 check_and_mark_arg (code_gen
, tmp
, false);
1879 decl
= opencl_print_function_arg_with_type (tmp
, iv_type
);
1880 opencl_append_string_to_body (decl
, code_gen
);
1882 opencl_append_string_to_body (" = ", code_gen
);
1883 opencl_append_string_to_body (f
->iterator
, code_gen
);
1884 opencl_append_string_to_body (";\n", code_gen
);
1886 save_clast_name_index (code_gen
->newivs_index
, f
->iterator
,
1887 VEC_length (tree
, code_gen
->newivs
));
1888 VEC_safe_push (tree
, heap
, code_gen
->newivs
, iv
);
1890 opencl_print_stmt_list (f
->body
, code_gen
, level
+ 1);
1891 opencl_append_string_to_body ("}\n", code_gen
);
1894 /* Generate code for clast equation EQ. CODE_GEN holds information
1895 related to OpenCL code generation. */
1898 opencl_print_equation (struct clast_equation
*eq
, opencl_main code_gen
)
1900 opencl_print_expr (eq
->LHS
, code_gen
);
1903 opencl_append_string_to_body (" == ", code_gen
);
1905 else if (eq
->sign
> 0)
1906 opencl_append_string_to_body (" >= ", code_gen
);
1909 opencl_append_string_to_body (" <= ", code_gen
);
1911 opencl_print_expr (eq
->RHS
, code_gen
);
1914 /* Generate code for clast conditional statement G, locate on depth DEPTH.
1915 CODE_GEN holds information related to OpenCL code generation. */
1918 opencl_print_guard (struct clast_guard
*g
, opencl_main code_gen
, int depth
)
1922 opencl_append_string_to_body ("if ", code_gen
);
1925 opencl_append_string_to_body ("(", code_gen
);
1927 for (k
= 0; k
< g
->n
; ++k
)
1930 opencl_append_string_to_body (" && ", code_gen
);
1932 opencl_append_string_to_body ("(", code_gen
);
1933 opencl_print_equation (&g
->eq
[k
], code_gen
);
1934 opencl_append_string_to_body (")", code_gen
);
1938 opencl_append_string_to_body (")", code_gen
);
1940 opencl_append_string_to_body (" {\n", code_gen
);
1941 opencl_print_stmt_list (g
->then
, code_gen
, depth
);
1942 opencl_append_string_to_body ("}\n", code_gen
);
1945 /* Generate code for clast statement S, located on depth DEPTH.
1946 CODE_GEN holds information related to OpenCL code generation. */
1949 opencl_print_stmt_list (struct clast_stmt
*s
, opencl_main code_gen
, int depth
)
1951 for ( ; s
; s
= s
->next
)
1953 gcc_assert (!CLAST_STMT_IS_A (s
, stmt_root
));
1955 if (CLAST_STMT_IS_A (s
, stmt_ass
))
1957 opencl_print_assignment ((struct clast_assignment
*) s
, code_gen
);
1958 opencl_append_string_to_body (";\n", code_gen
);
1961 else if (CLAST_STMT_IS_A (s
, stmt_user
))
1962 opencl_print_user_stmt ((struct clast_user_stmt
*) s
, code_gen
);
1964 else if (CLAST_STMT_IS_A (s
, stmt_for
))
1965 opencl_print_for ((struct clast_for
*) s
, code_gen
, depth
);
1967 else if (CLAST_STMT_IS_A (s
, stmt_guard
))
1968 opencl_print_guard ((struct clast_guard
*) s
, code_gen
, depth
);
1970 else if (CLAST_STMT_IS_A (s
, stmt_block
))
1972 opencl_append_string_to_body ("{\n", code_gen
);
1973 opencl_print_stmt_list (((struct clast_block
*) s
)->body
, code_gen
,
1975 opencl_append_string_to_body ("}\n", code_gen
);
1983 /* Generate code for loop statement F. DEPTH is the depth of F in
1984 current loop nest. CODE_GEN holds information related to OpenCL
1988 opencl_print_loop (struct clast_for
*f
, opencl_main code_gen
, int depth
)
1990 opencl_body current_body
= code_gen
->current_body
;
1992 code_gen
->global_defined_vars
1993 = htab_create (10, htab_hash_string
, opencl_cmp_str
, NULL
);
1995 opencl_perfect_nested_to_kernel (code_gen
, f
, current_body
, depth
);
1997 /* Define local loop iterators. */
1998 opencl_print_local_vars (current_body
->first_iter
,
1999 current_body
->last_iter
,
2000 "unsigned int", code_gen
);
2002 /* Generate code for kernel body. */
2003 opencl_print_stmt_list (current_body
->clast_body
, code_gen
, depth
+ 1);
2004 opencl_append_string_to_body ("}\n", code_gen
);
2006 if (current_body
->num_of_data_writes
)
2008 dyn_string_t header
= current_body
->header
;
2011 dyn_string_append (header
, current_body
->non_scalar_args
);
2012 offset
= dyn_string_length (header
) - 2;
2014 if (*(dyn_string_buf (header
) + offset
) == ',')
2015 *(dyn_string_buf (header
) + offset
) = ' ';
2017 opencl_append_string_to_header (")\n{\n", code_gen
);
2020 return current_body
;
2023 /* Generate kernel function code for clast for statement F, located on
2024 depth DEPTH. CODE_GEN holds information related to OpenCL code
2028 opencl_clast_to_kernel (struct clast_for
*f
, opencl_main code_gen
,
2031 opencl_body tmp
= opencl_body_create ();
2033 code_gen
->current_body
= tmp
;
2034 return opencl_print_loop (f
, code_gen
, depth
);