2 * Copyright 2013 Ecole Normale Superieure
4 * Use of this software is governed by the GNU LGPLv2.1 license
6 * Written by Sven Verdoolaege and Riyadh Baghdadi,
7 * Ecole Normale Superieure, 45 rue d’Ulm, 75230 Paris, France
18 #include "gpu_print.h"
24 #define min(a, b) (((a) < (b)) ? (a) : (b))
25 #define max(a, b) (((a) > (b)) ? (a) : (b))
27 /* options are the global options passed to generate_opencl.
28 * input is the name of the input file.
29 * output is the user-specified output file name and may be NULL
30 * if not specified by the user.
31 * kernel_c_name is the name of the kernel_c file.
32 * host_c is the generated source file for the host code. kernel_c is
33 * the generated source file for the kernel. kernel_h is the generated
34 * header file for the kernel.
37 struct ppcg_options
*options
;
40 char kernel_c_name
[PATH_MAX
];
47 /* Open the file called "name" for writing or print an error message.
49 static FILE *open_or_croak(const char *name
)
53 file
= fopen(name
, "w");
55 fprintf(stderr
, "Failed to open \"%s\" for writing\n", name
);
59 /* Open the host .c file and the kernel .h and .cl files for writing.
60 * Their names are derived from info->output (or info->input if
61 * the user did not specify an output file name).
62 * Add the necessary includes to these files, including those specified
65 * Return 0 on success and -1 on failure.
67 static int opencl_open_files(struct opencl_info
*info
)
76 ext
= strrchr(info
->output
, '.');
77 len
= ext
? ext
- info
->output
: strlen(info
->output
);
78 memcpy(name
, info
->output
, len
);
80 info
->host_c
= open_or_croak(info
->output
);
82 len
= ppcg_extract_base_name(name
, info
->input
);
84 strcpy(name
+ len
, "_host.c");
85 info
->host_c
= open_or_croak(name
);
88 memcpy(info
->kernel_c_name
, name
, len
);
89 strcpy(info
->kernel_c_name
+ len
, "_kernel.cl");
90 info
->kernel_c
= open_or_croak(info
->kernel_c_name
);
92 strcpy(name
+ len
, "_kernel.h");
93 info
->kernel_h
= open_or_croak(name
);
95 if (!info
->host_c
|| !info
->kernel_c
|| !info
->host_c
)
98 fprintf(info
->host_c
, "#include <assert.h>\n");
99 fprintf(info
->host_c
, "#include <stdio.h>\n");
100 fprintf(info
->host_c
, "#include \"%s\"\n\n", ppcg_base_name(name
));
101 fprintf(info
->kernel_h
, "#if defined(__APPLE__)\n");
102 fprintf(info
->kernel_h
, "#include <OpenCL/opencl.h>\n");
103 fprintf(info
->kernel_h
, "#else\n");
104 fprintf(info
->kernel_h
, "#include <CL/opencl.h>\n");
105 fprintf(info
->kernel_h
, "#endif\n\n");
106 fprintf(info
->kernel_h
, "cl_device_id opencl_create_device("
108 fprintf(info
->kernel_h
, "cl_program opencl_build_program("
110 "cl_device_id dev, const char *filename, "
111 "const char *opencl_options);\n");
112 fprintf(info
->kernel_h
,
113 "const char *opencl_error_string(cl_int error);\n");
114 for (i
= 0; i
< info
->options
->opencl_n_include_file
; ++i
)
115 fprintf(info
->kernel_c
, "#include <%s>\n",
116 info
->options
->opencl_include_files
[i
]);
121 /* Close all output files.
123 static void opencl_close_files(struct opencl_info
*info
)
126 fclose(info
->kernel_c
);
128 fclose(info
->kernel_h
);
130 fclose(info
->host_c
);
133 static __isl_give isl_printer
*opencl_print_host_macros(__isl_take isl_printer
*p
)
136 "#define openclCheckReturn(ret) \\\n"
137 " if (ret != CL_SUCCESS) {\\\n"
138 " fprintf(stderr, \"OpenCL error: %s\\n\", "
139 "opencl_error_string(ret)); \\\n"
140 " fflush(stderr); \\\n"
141 " assert(ret == CL_SUCCESS);\\\n }\n";
143 p
= isl_printer_start_line(p
);
144 p
= isl_printer_print_str(p
, macros
);
145 p
= isl_printer_end_line(p
);
147 p
= isl_ast_op_type_print_macro(isl_ast_op_max
, p
);
152 static __isl_give isl_printer
*opencl_declare_device_arrays(
153 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
157 for (i
= 0; i
< prog
->n_array
; ++i
) {
158 if (gpu_array_is_read_only_scalar(&prog
->array
[i
]))
160 p
= isl_printer_start_line(p
);
161 p
= isl_printer_print_str(p
, "cl_mem dev_");
162 p
= isl_printer_print_str(p
, prog
->array
[i
].name
);
163 p
= isl_printer_print_str(p
, ";");
164 p
= isl_printer_end_line(p
);
166 p
= isl_printer_start_line(p
);
167 p
= isl_printer_end_line(p
);
171 /* Given an array, check whether its positive size guard expression is
174 static int is_array_positive_size_guard_trivial(struct gpu_array_info
*array
)
179 guard
= gpu_array_positive_size_guard(array
);
180 is_trivial
= isl_set_plain_is_universe(guard
);
185 /* Allocate a device array for array and copy the contents to the device
188 * Emit a max-expression to ensure the device array can contain at least one
189 * element if the array's positive size guard expression is not trivial.
191 static __isl_give isl_printer
*allocate_device_array(__isl_take isl_printer
*p
,
192 struct gpu_array_info
*array
, int copy
)
194 int need_lower_bound
;
196 p
= ppcg_start_block(p
);
198 p
= isl_printer_start_line(p
);
199 p
= isl_printer_print_str(p
, "dev_");
200 p
= isl_printer_print_str(p
, array
->name
);
201 p
= isl_printer_print_str(p
, " = clCreateBuffer(context, ");
202 p
= isl_printer_print_str(p
, "CL_MEM_READ_WRITE");
205 p
= isl_printer_print_str(p
, ", ");
207 p
= isl_printer_print_str(p
, " | CL_MEM_COPY_HOST_PTR, ");
209 need_lower_bound
= !is_array_positive_size_guard_trivial(array
);
210 if (need_lower_bound
) {
211 p
= isl_printer_print_str(p
, "max(sizeof(");
212 p
= isl_printer_print_str(p
, array
->type
);
213 p
= isl_printer_print_str(p
, "), ");
215 p
= gpu_array_info_print_size(p
, array
);
216 if (need_lower_bound
)
217 p
= isl_printer_print_str(p
, ")");
220 p
= isl_printer_print_str(p
, ", NULL");
221 else if (gpu_array_is_scalar(array
)) {
222 p
= isl_printer_print_str(p
, ", &");
223 p
= isl_printer_print_str(p
, array
->name
);
225 p
= isl_printer_print_str(p
, ", ");
226 p
= isl_printer_print_str(p
, array
->name
);
229 p
= isl_printer_print_str(p
, ", &err);");
230 p
= isl_printer_end_line(p
);
231 p
= isl_printer_start_line(p
);
232 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
233 p
= isl_printer_end_line(p
);
235 p
= ppcg_end_block(p
);
240 /* Allocate device arrays and copy the contents of copy_in arrays into device.
242 static __isl_give isl_printer
*opencl_allocate_device_arrays(
243 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
247 for (i
= 0; i
< prog
->n_array
; ++i
) {
248 struct gpu_array_info
*array
= &prog
->array
[i
];
253 if (gpu_array_is_read_only_scalar(array
))
256 space
= isl_space_copy(array
->space
);
257 read_i
= isl_union_set_extract_set(prog
->copy_in
, space
);
258 empty
= isl_set_plain_is_empty(read_i
);
259 isl_set_free(read_i
);
261 p
= allocate_device_array(p
, array
, !empty
);
263 p
= isl_printer_start_line(p
);
264 p
= isl_printer_end_line(p
);
268 /* Print a call to the OpenCL clSetKernelArg() function which sets
269 * the arguments of the kernel. arg_name and arg_index are the name and the
270 * index of the kernel argument. The index of the leftmost argument of
271 * the kernel is 0 whereas the index of the rightmost argument of the kernel
272 * is n - 1, where n is the total number of the kernel arguments.
273 * read_only_scalar is a boolean that indicates whether the argument is a read
276 static __isl_give isl_printer
*opencl_set_kernel_argument(
277 __isl_take isl_printer
*p
, int kernel_id
,
278 const char *arg_name
, int arg_index
, int read_only_scalar
)
280 p
= isl_printer_start_line(p
);
281 p
= isl_printer_print_str(p
,
282 "openclCheckReturn(clSetKernelArg(kernel");
283 p
= isl_printer_print_int(p
, kernel_id
);
284 p
= isl_printer_print_str(p
, ", ");
285 p
= isl_printer_print_int(p
, arg_index
);
286 p
= isl_printer_print_str(p
, ", sizeof(");
288 if (read_only_scalar
) {
289 p
= isl_printer_print_str(p
, arg_name
);
290 p
= isl_printer_print_str(p
, "), &");
292 p
= isl_printer_print_str(p
, "cl_mem), (void *) &dev_");
294 p
= isl_printer_print_str(p
, arg_name
);
295 p
= isl_printer_print_str(p
, "));");
296 p
= isl_printer_end_line(p
);
301 /* Print the block sizes as a list of the sizes in each
304 static __isl_give isl_printer
*opencl_print_block_sizes(
305 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
309 if (kernel
->n_block
> 0)
310 for (i
= 0; i
< kernel
->n_block
; ++i
) {
312 p
= isl_printer_print_str(p
, ", ");
313 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
316 p
= isl_printer_print_str(p
, "1");
321 /* Set the arguments of the OpenCL kernel by printing a call to the OpenCL
322 * clSetKernelArg() function for each kernel argument.
324 static __isl_give isl_printer
*opencl_set_kernel_arguments(
325 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
326 struct ppcg_kernel
*kernel
)
334 for (i
= 0; i
< prog
->n_array
; ++i
) {
338 space
= isl_space_copy(prog
->array
[i
].space
);
339 arr
= isl_union_set_extract_set(kernel
->arrays
, space
);
340 empty
= isl_set_plain_is_empty(arr
);
344 ro
= gpu_array_is_read_only_scalar(&prog
->array
[i
]);
345 opencl_set_kernel_argument(p
, kernel
->id
, prog
->array
[i
].name
,
350 space
= isl_union_set_get_space(kernel
->arrays
);
351 nparam
= isl_space_dim(space
, isl_dim_param
);
352 for (i
= 0; i
< nparam
; ++i
) {
355 name
= isl_space_get_dim_name(space
, isl_dim_param
, i
);
356 opencl_set_kernel_argument(p
, kernel
->id
, name
, arg_index
, 1);
359 isl_space_free(space
);
361 n
= isl_space_dim(kernel
->space
, isl_dim_set
);
362 for (i
= 0; i
< n
; ++i
) {
366 name
= isl_space_get_dim_name(kernel
->space
, isl_dim_set
, i
);
367 opencl_set_kernel_argument(p
, kernel
->id
, name
, arg_index
, 1);
374 /* Print the arguments to a kernel declaration or call. If "types" is set,
375 * then print a declaration (including the types of the arguments).
377 * The arguments are printed in the following order
378 * - the arrays accessed by the kernel
380 * - the host loop iterators
382 static __isl_give isl_printer
*opencl_print_kernel_arguments(
383 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
384 struct ppcg_kernel
*kernel
, int types
)
392 for (i
= 0; i
< prog
->n_array
; ++i
) {
396 space
= isl_space_copy(prog
->array
[i
].space
);
397 arr
= isl_union_set_extract_set(kernel
->arrays
, space
);
398 empty
= isl_set_plain_is_empty(arr
);
404 p
= isl_printer_print_str(p
, ", ");
407 p
= gpu_array_info_print_declaration_argument(p
,
408 &prog
->array
[i
], "__global");
410 p
= gpu_array_info_print_call_argument(p
,
416 space
= isl_union_set_get_space(kernel
->arrays
);
417 nparam
= isl_space_dim(space
, isl_dim_param
);
418 for (i
= 0; i
< nparam
; ++i
) {
421 name
= isl_space_get_dim_name(space
, isl_dim_param
, i
);
424 p
= isl_printer_print_str(p
, ", ");
426 p
= isl_printer_print_str(p
, "int ");
427 p
= isl_printer_print_str(p
, name
);
431 isl_space_free(space
);
433 n
= isl_space_dim(kernel
->space
, isl_dim_set
);
434 type
= isl_options_get_ast_iterator_type(prog
->ctx
);
435 for (i
= 0; i
< n
; ++i
) {
440 p
= isl_printer_print_str(p
, ", ");
441 name
= isl_space_get_dim_name(kernel
->space
, isl_dim_set
, i
);
443 p
= isl_printer_print_str(p
, type
);
444 p
= isl_printer_print_str(p
, " ");
446 p
= isl_printer_print_str(p
, name
);
454 /* Print the header of the given kernel.
456 static __isl_give isl_printer
*opencl_print_kernel_header(
457 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
458 struct ppcg_kernel
*kernel
)
460 p
= isl_printer_start_line(p
);
461 p
= isl_printer_print_str(p
, "__kernel void kernel");
462 p
= isl_printer_print_int(p
, kernel
->id
);
463 p
= isl_printer_print_str(p
, "(");
464 p
= opencl_print_kernel_arguments(p
, prog
, kernel
, 1);
465 p
= isl_printer_print_str(p
, ")");
466 p
= isl_printer_end_line(p
);
471 /* Unlike the equivalent function in the CUDA backend which prints iterators
472 * in reverse order to promote coalescing, this function does not print
473 * iterators in reverse order. The OpenCL backend currently does not take
474 * into account any coalescing considerations.
476 static __isl_give isl_printer
*opencl_print_kernel_iterators(
477 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
480 isl_ctx
*ctx
= isl_ast_node_get_ctx(kernel
->tree
);
483 type
= isl_options_get_ast_iterator_type(ctx
);
485 n_grid
= isl_multi_pw_aff_dim(kernel
->grid_size
, isl_dim_set
);
487 p
= isl_printer_start_line(p
);
488 p
= isl_printer_print_str(p
, type
);
489 p
= isl_printer_print_str(p
, " ");
490 for (i
= 0; i
< n_grid
; ++i
) {
492 p
= isl_printer_print_str(p
, ", ");
493 p
= isl_printer_print_str(p
, "b");
494 p
= isl_printer_print_int(p
, i
);
495 p
= isl_printer_print_str(p
, " = get_group_id(");
496 p
= isl_printer_print_int(p
, i
);
497 p
= isl_printer_print_str(p
, ")");
499 p
= isl_printer_print_str(p
, ";");
500 p
= isl_printer_end_line(p
);
503 if (kernel
->n_block
> 0) {
504 p
= isl_printer_start_line(p
);
505 p
= isl_printer_print_str(p
, type
);
506 p
= isl_printer_print_str(p
, " ");
507 for (i
= 0; i
< kernel
->n_block
; ++i
) {
509 p
= isl_printer_print_str(p
, ", ");
510 p
= isl_printer_print_str(p
, "t");
511 p
= isl_printer_print_int(p
, i
);
512 p
= isl_printer_print_str(p
, " = get_local_id(");
513 p
= isl_printer_print_int(p
, i
);
514 p
= isl_printer_print_str(p
, ")");
516 p
= isl_printer_print_str(p
, ";");
517 p
= isl_printer_end_line(p
);
523 static __isl_give isl_printer
*opencl_print_kernel_var(
524 __isl_take isl_printer
*p
, struct ppcg_kernel_var
*var
)
529 p
= isl_printer_start_line(p
);
530 if (var
->type
== ppcg_access_shared
)
531 p
= isl_printer_print_str(p
, "__local ");
532 p
= isl_printer_print_str(p
, var
->array
->type
);
533 p
= isl_printer_print_str(p
, " ");
534 p
= isl_printer_print_str(p
, var
->name
);
535 for (j
= 0; j
< var
->array
->n_index
; ++j
) {
536 p
= isl_printer_print_str(p
, "[");
537 v
= isl_vec_get_element_val(var
->size
, j
);
538 p
= isl_printer_print_val(p
, v
);
539 p
= isl_printer_print_str(p
, "]");
542 p
= isl_printer_print_str(p
, ";");
543 p
= isl_printer_end_line(p
);
548 static __isl_give isl_printer
*opencl_print_kernel_vars(
549 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
553 for (i
= 0; i
< kernel
->n_var
; ++i
)
554 p
= opencl_print_kernel_var(p
, &kernel
->var
[i
]);
559 /* Print a call to barrier() which is a sync statement.
560 * All work-items in a work-group executing the kernel on a processor must
561 * execute the barrier() function before any are allowed to continue execution
562 * beyond the barrier.
563 * The flag CLK_LOCAL_MEM_FENCE makes the barrier function either flush any
564 * variables stored in local memory or queue a memory fence to ensure correct
565 * ordering of memory operations to local memory.
566 * The flag CLK_GLOBAL_MEM_FENCE makes the barrier function queue a memory
567 * fence to ensure correct ordering of memory operations to global memory.
569 static __isl_give isl_printer
*opencl_print_sync(__isl_take isl_printer
*p
,
570 struct ppcg_kernel_stmt
*stmt
)
572 p
= isl_printer_start_line(p
);
573 p
= isl_printer_print_str(p
,
574 "barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);");
575 p
= isl_printer_end_line(p
);
580 /* This function is called for each user statement in the AST,
581 * i.e., for each kernel body statement, copy statement or sync statement.
583 static __isl_give isl_printer
*opencl_print_kernel_stmt(
584 __isl_take isl_printer
*p
,
585 __isl_take isl_ast_print_options
*print_options
,
586 __isl_keep isl_ast_node
*node
, void *user
)
589 struct ppcg_kernel_stmt
*stmt
;
591 id
= isl_ast_node_get_annotation(node
);
592 stmt
= isl_id_get_user(id
);
595 isl_ast_print_options_free(print_options
);
597 switch (stmt
->type
) {
598 case ppcg_kernel_copy
:
599 return ppcg_kernel_print_copy(p
, stmt
);
600 case ppcg_kernel_sync
:
601 return opencl_print_sync(p
, stmt
);
602 case ppcg_kernel_domain
:
603 return ppcg_kernel_print_domain(p
, stmt
);
609 /* Return true if there is a double array in prog->array or
610 * if any of the types in prog->scop involve any doubles.
611 * To check the latter condition, we simply search for the string "double"
612 * in the type definitions, which may result in false positives.
614 static __isl_give
int any_double_elements(struct gpu_prog
*prog
)
618 for (i
= 0; i
< prog
->n_array
; ++i
)
619 if (strcmp(prog
->array
[i
].type
, "double") == 0)
622 for (i
= 0; i
< prog
->scop
->pet
->n_type
; ++i
) {
623 struct pet_type
*type
= prog
->scop
->pet
->types
[i
];
625 if (strstr(type
->definition
, "double"))
632 /* Prints a #pragma to enable support for double floating-point
633 * precision. OpenCL 1.0 adds support for double precision floating-point as
634 * an optional extension. An application that wants to use double will need to
635 * include the #pragma OPENCL EXTENSION cl_khr_fp64 : enable directive before
636 * any double precision data type is declared in the kernel code.
638 static __isl_give isl_printer
*opencl_enable_double_support(
639 __isl_take isl_printer
*p
)
643 p
= isl_printer_start_line(p
);
644 p
= isl_printer_print_str(p
, "#pragma OPENCL EXTENSION cl_khr_fp64 :"
646 p
= isl_printer_end_line(p
);
647 p
= isl_printer_start_line(p
);
648 p
= isl_printer_end_line(p
);
653 static void opencl_print_kernel(struct gpu_prog
*prog
,
654 struct ppcg_kernel
*kernel
, struct opencl_info
*opencl
)
656 isl_ctx
*ctx
= isl_ast_node_get_ctx(kernel
->tree
);
657 isl_ast_print_options
*print_options
;
660 p
= isl_printer_to_file(ctx
, opencl
->kernel_c
);
661 print_options
= isl_ast_print_options_alloc(ctx
);
662 print_options
= isl_ast_print_options_set_print_user(print_options
,
663 &opencl_print_kernel_stmt
, NULL
);
665 p
= isl_printer_set_output_format(p
, ISL_FORMAT_C
);
666 p
= opencl_print_kernel_header(p
, prog
, kernel
);
667 p
= isl_printer_print_str(p
, "{");
668 p
= isl_printer_end_line(p
);
669 p
= isl_printer_indent(p
, 4);
670 p
= opencl_print_kernel_iterators(p
, kernel
);
671 p
= opencl_print_kernel_vars(p
, kernel
);
672 p
= isl_printer_end_line(p
);
673 p
= gpu_print_macros(p
, kernel
->tree
);
674 p
= isl_ast_node_print(kernel
->tree
, p
, print_options
);
675 p
= isl_printer_print_str(p
, "}");
676 p
= isl_printer_end_line(p
);
680 struct print_host_user_data_opencl
{
681 struct opencl_info
*opencl
;
682 struct gpu_prog
*prog
;
685 /* This function prints the i'th block size multiplied by the i'th grid size,
686 * where i (a parameter to this function) is one of the possible dimensions of
687 * grid sizes and block sizes.
688 * If the dimension of block sizes is not equal to the dimension of grid sizes
689 * the output is calculated as follows:
692 * block_sizes[dim1] is the list of blocks sizes and it contains dim1 elements.
693 * grid_sizes[dim2] is the list of grid sizes and it contains dim2 elements.
696 * If (i > dim2) then the output is block_sizes[i]
697 * If (i > dim1) then the output is grid_sizes[i]
699 static __isl_give isl_printer
*opencl_print_total_number_of_work_items_for_dim(
700 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
, int i
)
702 int grid_dim
, block_dim
;
704 grid_dim
= isl_multi_pw_aff_dim(kernel
->grid_size
, isl_dim_set
);
705 block_dim
= kernel
->n_block
;
707 isl_pw_aff
*bound_grid
;
709 if (i
< min(grid_dim
, block_dim
)) {
710 bound_grid
= isl_multi_pw_aff_get_pw_aff(kernel
->grid_size
, i
);
711 p
= isl_printer_print_str(p
, "(");
712 p
= isl_printer_print_pw_aff(p
, bound_grid
);
713 p
= isl_printer_print_str(p
, ") * ");
714 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
715 isl_pw_aff_free(bound_grid
);
716 } else if (i
>= grid_dim
)
717 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
719 bound_grid
= isl_multi_pw_aff_get_pw_aff(kernel
->grid_size
, i
);
720 p
= isl_printer_print_pw_aff(p
, bound_grid
);
721 isl_pw_aff_free(bound_grid
);
727 /* Print a list that represents the total number of work items. The list is
728 * constructed by performing an element-wise multiplication of the block sizes
729 * and the grid sizes. To explain how the list is constructed, suppose that:
730 * block_sizes[dim1] is the list of blocks sizes and it contains dim1 elements.
731 * grid_sizes[dim2] is the list of grid sizes and it contains dim2 elements.
733 * The output of this function is constructed as follows:
734 * If (dim1 > dim2) then the output is the following list:
735 * grid_sizes[0]*block_sizes[0], ..., grid_sizes[dim2-1]*block_sizes[dim2-1],
736 * block_sizes[dim2], ..., block_sizes[dim1-2], block_sizes[dim1-1].
738 * If (dim2 > dim1) then the output is the following list:
739 * grid_sizes[0]*block_sizes[0], ..., grid_sizes[dim1-1] * block_sizes[dim1-1],
740 * grid_sizes[dim1], grid_sizes[dim2-2], grid_sizes[dim2-1].
742 * To calculate the total number of work items out of the list constructed by
743 * this function, the user should multiply the elements of the list.
745 static __isl_give isl_printer
*opencl_print_total_number_of_work_items_as_list(
746 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
749 int grid_dim
, block_dim
;
751 grid_dim
= isl_multi_pw_aff_dim(kernel
->grid_size
, isl_dim_set
);
752 block_dim
= kernel
->n_block
;
754 if ((grid_dim
<= 0) || (block_dim
<= 0)) {
755 p
= isl_printer_print_str(p
, "1");
759 for (i
= 0; i
<= max(grid_dim
, block_dim
) - 1; i
++) {
761 p
= isl_printer_print_str(p
, ", ");
763 p
= opencl_print_total_number_of_work_items_for_dim(p
,
770 /* Print the user statement of the host code to "p".
772 * In particular, print a block of statements that defines the grid
773 * and the work group and then launches the kernel.
775 * A grid is composed of many work groups (blocks), each work group holds
776 * many work-items (threads).
778 * global_work_size[kernel->n_block] represents the total number of work
779 * items. It points to an array of kernel->n_block unsigned
780 * values that describe the total number of work-items that will execute
781 * the kernel. The total number of work-items is computed as:
782 * global_work_size[0] *...* global_work_size[kernel->n_block - 1].
784 * The size of each work group (i.e. the number of work-items in each work
785 * group) is described using block_size[kernel->n_block]. The total
786 * number of work-items in a block (work-group) is computed as:
787 * block_size[0] *... * block_size[kernel->n_block - 1].
789 * For more information check:
790 * http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html
792 static __isl_give isl_printer
*opencl_print_host_user(
793 __isl_take isl_printer
*p
,
794 __isl_take isl_ast_print_options
*print_options
,
795 __isl_keep isl_ast_node
*node
, void *user
)
798 struct ppcg_kernel
*kernel
;
799 struct print_host_user_data_opencl
*data
;
802 id
= isl_ast_node_get_annotation(node
);
803 kernel
= isl_id_get_user(id
);
806 data
= (struct print_host_user_data_opencl
*) user
;
808 p
= isl_printer_start_line(p
);
809 p
= isl_printer_print_str(p
, "{");
810 p
= isl_printer_end_line(p
);
811 p
= isl_printer_indent(p
, 2);
813 p
= isl_printer_start_line(p
);
814 p
= isl_printer_print_str(p
, "size_t global_work_size[");
816 if (kernel
->n_block
> 0)
817 p
= isl_printer_print_int(p
, kernel
->n_block
);
819 p
= isl_printer_print_int(p
, 1);
821 p
= isl_printer_print_str(p
, "] = {");
822 p
= opencl_print_total_number_of_work_items_as_list(p
, kernel
);
823 p
= isl_printer_print_str(p
, "};");
824 p
= isl_printer_end_line(p
);
826 p
= isl_printer_start_line(p
);
827 p
= isl_printer_print_str(p
, "size_t block_size[");
829 if (kernel
->n_block
> 0)
830 p
= isl_printer_print_int(p
, kernel
->n_block
);
832 p
= isl_printer_print_int(p
, 1);
834 p
= isl_printer_print_str(p
, "] = {");
835 p
= opencl_print_block_sizes(p
, kernel
);
836 p
= isl_printer_print_str(p
, "};");
837 p
= isl_printer_end_line(p
);
839 p
= isl_printer_start_line(p
);
840 p
= isl_printer_print_str(p
, "cl_kernel kernel");
841 p
= isl_printer_print_int(p
, kernel
->id
);
842 p
= isl_printer_print_str(p
, " = clCreateKernel(program, \"kernel");
843 p
= isl_printer_print_int(p
, kernel
->id
);
844 p
= isl_printer_print_str(p
, "\", &err);");
845 p
= isl_printer_end_line(p
);
846 p
= isl_printer_start_line(p
);
847 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
848 p
= isl_printer_end_line(p
);
850 opencl_set_kernel_arguments(p
, data
->prog
, kernel
);
852 p
= isl_printer_start_line(p
);
853 p
= isl_printer_print_str(p
, "openclCheckReturn(clEnqueueNDRangeKernel"
855 p
= isl_printer_print_int(p
, kernel
->id
);
856 p
= isl_printer_print_str(p
, ", ");
857 if (kernel
->n_block
> 0)
858 p
= isl_printer_print_int(p
, kernel
->n_block
);
860 p
= isl_printer_print_int(p
, 1);
862 p
= isl_printer_print_str(p
, ", NULL, global_work_size, "
865 p
= isl_printer_end_line(p
);
866 p
= isl_printer_start_line(p
);
867 p
= isl_printer_print_str(p
, "openclCheckReturn("
868 "clReleaseKernel(kernel");
869 p
= isl_printer_print_int(p
, kernel
->id
);
870 p
= isl_printer_print_str(p
, "));");
871 p
= isl_printer_end_line(p
);
872 p
= isl_printer_start_line(p
);
873 p
= isl_printer_print_str(p
, "clFinish(queue);");
874 p
= isl_printer_end_line(p
);
875 p
= isl_printer_indent(p
, -2);
876 p
= isl_printer_start_line(p
);
877 p
= isl_printer_print_str(p
, "}");
878 p
= isl_printer_end_line(p
);
880 p
= isl_printer_start_line(p
);
881 p
= isl_printer_end_line(p
);
883 opencl_print_kernel(data
->prog
, kernel
, data
->opencl
);
885 isl_ast_print_options_free(print_options
);
890 static __isl_give isl_printer
*opencl_print_host_code(
891 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
892 __isl_keep isl_ast_node
*tree
, struct opencl_info
*opencl
)
894 isl_ast_print_options
*print_options
;
895 isl_ctx
*ctx
= isl_ast_node_get_ctx(tree
);
896 struct print_host_user_data_opencl data
= { opencl
, prog
};
898 print_options
= isl_ast_print_options_alloc(ctx
);
899 print_options
= isl_ast_print_options_set_print_user(print_options
,
900 &opencl_print_host_user
, &data
);
902 p
= gpu_print_macros(p
, tree
);
903 p
= isl_ast_node_print(tree
, p
, print_options
);
908 /* Copy "array" back from the GPU to the host.
910 static __isl_give isl_printer
*copy_array_from_device(__isl_take isl_printer
*p
,
913 struct gpu_array_info
*array
= user
;
915 p
= isl_printer_start_line(p
);
916 p
= isl_printer_print_str(p
, "openclCheckReturn("
917 "clEnqueueReadBuffer(queue,"
919 p
= isl_printer_print_str(p
, array
->name
);
920 p
= isl_printer_print_str(p
, ", CL_TRUE, 0, ");
921 p
= gpu_array_info_print_size(p
, array
);
923 if (gpu_array_is_scalar(array
))
924 p
= isl_printer_print_str(p
, ", &");
926 p
= isl_printer_print_str(p
, ", ");
927 p
= isl_printer_print_str(p
, array
->name
);
928 p
= isl_printer_print_str(p
, ", 0, NULL, NULL));");
929 p
= isl_printer_end_line(p
);
934 /* Copy copy_out arrays back from the GPU to the host.
936 * Only perform the copying for arrays with strictly positive size.
938 static __isl_give isl_printer
*opencl_copy_arrays_from_device(
939 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
942 isl_union_set
*copy_out
;
943 copy_out
= isl_union_set_copy(prog
->copy_out
);
945 for (i
= 0; i
< prog
->n_array
; ++i
) {
946 struct gpu_array_info
*array
= &prog
->array
[i
];
952 space
= isl_space_copy(array
->space
);
953 copy_out_i
= isl_union_set_extract_set(copy_out
, space
);
954 empty
= isl_set_plain_is_empty(copy_out_i
);
955 isl_set_free(copy_out_i
);
959 guard
= gpu_array_positive_size_guard(array
);
960 p
= ppcg_print_guarded(p
, guard
, isl_set_copy(prog
->context
),
961 ©_array_from_device
, array
);
964 isl_union_set_free(copy_out
);
965 p
= isl_printer_start_line(p
);
966 p
= isl_printer_end_line(p
);
970 /* Create an OpenCL device, context, command queue and build the kernel.
971 * input is the name of the input file provided to ppcg.
973 static __isl_give isl_printer
*opencl_setup(__isl_take isl_printer
*p
,
974 const char *input
, struct opencl_info
*info
)
978 p
= isl_printer_start_line(p
);
979 p
= isl_printer_print_str(p
, "cl_device_id device;");
980 p
= isl_printer_end_line(p
);
981 p
= isl_printer_start_line(p
);
982 p
= isl_printer_print_str(p
, "cl_context context;");
983 p
= isl_printer_end_line(p
);
984 p
= isl_printer_start_line(p
);
985 p
= isl_printer_print_str(p
, "cl_program program;");
986 p
= isl_printer_end_line(p
);
987 p
= isl_printer_start_line(p
);
988 p
= isl_printer_print_str(p
, "cl_command_queue queue;");
989 p
= isl_printer_end_line(p
);
990 p
= isl_printer_start_line(p
);
991 p
= isl_printer_print_str(p
, "cl_int err;");
992 p
= isl_printer_end_line(p
);
993 p
= isl_printer_start_line(p
);
994 p
= isl_printer_print_str(p
, "device = opencl_create_device(");
995 p
= isl_printer_print_int(p
, info
->options
->opencl_use_gpu
);
996 p
= isl_printer_print_str(p
, ");");
997 p
= isl_printer_end_line(p
);
998 p
= isl_printer_start_line(p
);
999 p
= isl_printer_print_str(p
, "context = clCreateContext(NULL, 1, "
1000 "&device, NULL, NULL, &err);");
1001 p
= isl_printer_end_line(p
);
1002 p
= isl_printer_start_line(p
);
1003 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
1004 p
= isl_printer_end_line(p
);
1005 p
= isl_printer_start_line(p
);
1006 p
= isl_printer_print_str(p
, "queue = clCreateCommandQueue"
1007 "(context, device, 0, &err);");
1008 p
= isl_printer_end_line(p
);
1009 p
= isl_printer_start_line(p
);
1010 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
1011 p
= isl_printer_end_line(p
);
1013 p
= isl_printer_start_line(p
);
1014 p
= isl_printer_print_str(p
, "program = opencl_build_program("
1015 "context, device, \"");
1016 p
= isl_printer_print_str(p
, info
->kernel_c_name
);
1017 p
= isl_printer_print_str(p
, "\", \"");
1019 if (info
->options
->opencl_compiler_options
)
1020 p
= isl_printer_print_str(p
,
1021 info
->options
->opencl_compiler_options
);
1023 p
= isl_printer_print_str(p
, "\");");
1024 p
= isl_printer_end_line(p
);
1025 p
= isl_printer_start_line(p
);
1026 p
= isl_printer_end_line(p
);
1031 static __isl_give isl_printer
*opencl_release_cl_objects(
1032 __isl_take isl_printer
*p
, struct opencl_info
*info
)
1034 p
= isl_printer_start_line(p
);
1035 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseCommandQueue"
1037 p
= isl_printer_end_line(p
);
1038 p
= isl_printer_start_line(p
);
1039 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseProgram"
1041 p
= isl_printer_end_line(p
);
1042 p
= isl_printer_start_line(p
);
1043 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseContext"
1045 p
= isl_printer_end_line(p
);
1050 /* Free the device array corresponding to "array"
1052 static __isl_give isl_printer
*release_device_array(__isl_take isl_printer
*p
,
1053 struct gpu_array_info
*array
)
1055 p
= isl_printer_start_line(p
);
1056 p
= isl_printer_print_str(p
, "openclCheckReturn("
1057 "clReleaseMemObject(dev_");
1058 p
= isl_printer_print_str(p
, array
->name
);
1059 p
= isl_printer_print_str(p
, "));");
1060 p
= isl_printer_end_line(p
);
1065 /* Free the device arrays.
1067 static __isl_give isl_printer
*opencl_release_device_arrays(
1068 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
1072 for (i
= 0; i
< prog
->n_array
; ++i
) {
1073 struct gpu_array_info
*array
= &prog
->array
[i
];
1074 if (gpu_array_is_read_only_scalar(array
))
1077 p
= release_device_array(p
, array
);
1082 /* Given a gpu_prog "prog" and the corresponding transformed AST
1083 * "tree", print the entire OpenCL code to "p".
1085 static __isl_give isl_printer
*print_opencl(__isl_take isl_printer
*p
,
1086 struct gpu_prog
*prog
, __isl_keep isl_ast_node
*tree
,
1087 struct gpu_types
*types
, void *user
)
1089 struct opencl_info
*opencl
= user
;
1090 isl_printer
*kernel
;
1092 kernel
= isl_printer_to_file(isl_printer_get_ctx(p
), opencl
->kernel_c
);
1093 kernel
= isl_printer_set_output_format(kernel
, ISL_FORMAT_C
);
1094 if (any_double_elements(prog
))
1095 kernel
= opencl_enable_double_support(kernel
);
1096 kernel
= gpu_print_types(kernel
, types
, prog
);
1097 isl_printer_free(kernel
);
1100 return isl_printer_free(p
);
1102 p
= ppcg_start_block(p
);
1104 p
= opencl_print_host_macros(p
);
1106 p
= opencl_declare_device_arrays(p
, prog
);
1107 p
= opencl_setup(p
, opencl
->input
, opencl
);
1108 p
= opencl_allocate_device_arrays(p
, prog
);
1110 p
= opencl_print_host_code(p
, prog
, tree
, opencl
);
1112 p
= opencl_copy_arrays_from_device(p
, prog
);
1113 p
= opencl_release_device_arrays(p
, prog
);
1114 p
= opencl_release_cl_objects(p
, opencl
);
1116 p
= ppcg_end_block(p
);
1121 /* Transform the code in the file called "input" by replacing
1122 * all scops by corresponding OpenCL code.
1123 * The host code is written to "output" or a name derived from
1124 * "input" if "output" is NULL.
1125 * The kernel code is placed in separate files with names
1126 * derived from "output" or "input".
1128 * We let generate_gpu do all the hard work and then let it call
1129 * us back for printing the AST in print_cuda.
1131 * To prepare for this printing, we first open the output files
1132 * and we close them after generate_gpu has finished.
1134 int generate_opencl(isl_ctx
*ctx
, struct ppcg_options
*options
,
1135 const char *input
, const char *output
)
1137 struct opencl_info opencl
= { options
, input
, output
};
1140 r
= opencl_open_files(&opencl
);
1143 r
= generate_gpu(ctx
, input
, opencl
.host_c
, options
,
1144 &print_opencl
, &opencl
);
1146 opencl_close_files(&opencl
);