2 * Copyright 2013 Ecole Normale Superieure
4 * Use of this software is governed by the MIT 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 * kprinter is an isl_printer for the kernel file.
33 * host_c is the generated source file for the host code. kernel_c is
34 * the generated source file for the kernel. kernel_h is the generated
35 * header file for the kernel.
38 struct ppcg_options
*options
;
41 char kernel_c_name
[PATH_MAX
];
43 isl_printer
*kprinter
;
50 /* Open the file called "name" for writing or print an error message.
52 static FILE *open_or_croak(const char *name
)
56 file
= fopen(name
, "w");
58 fprintf(stderr
, "Failed to open \"%s\" for writing\n", name
);
62 /* Open the host .c file and the kernel .h and .cl files for writing.
63 * Their names are derived from info->output (or info->input if
64 * the user did not specify an output file name).
65 * Add the necessary includes to these files, including those specified
68 * Return 0 on success and -1 on failure.
70 static int opencl_open_files(struct opencl_info
*info
)
79 ext
= strrchr(info
->output
, '.');
80 len
= ext
? ext
- info
->output
: strlen(info
->output
);
81 memcpy(name
, info
->output
, len
);
83 info
->host_c
= open_or_croak(info
->output
);
85 len
= ppcg_extract_base_name(name
, info
->input
);
87 strcpy(name
+ len
, "_host.c");
88 info
->host_c
= open_or_croak(name
);
91 memcpy(info
->kernel_c_name
, name
, len
);
92 strcpy(info
->kernel_c_name
+ len
, "_kernel.cl");
93 info
->kernel_c
= open_or_croak(info
->kernel_c_name
);
95 strcpy(name
+ len
, "_kernel.h");
96 info
->kernel_h
= open_or_croak(name
);
98 if (!info
->host_c
|| !info
->kernel_c
|| !info
->kernel_h
)
101 fprintf(info
->host_c
, "#include <assert.h>\n");
102 fprintf(info
->host_c
, "#include <stdio.h>\n");
103 fprintf(info
->host_c
, "#include \"%s\"\n\n", ppcg_base_name(name
));
104 if (info
->options
->opencl_embed_kernel_code
) {
105 fprintf(info
->host_c
, "#include \"%s\"\n\n",
106 info
->kernel_c_name
);
109 fprintf(info
->kernel_h
, "#if defined(__APPLE__)\n");
110 fprintf(info
->kernel_h
, "#include <OpenCL/opencl.h>\n");
111 fprintf(info
->kernel_h
, "#else\n");
112 fprintf(info
->kernel_h
, "#include <CL/opencl.h>\n");
113 fprintf(info
->kernel_h
, "#endif\n\n");
114 fprintf(info
->kernel_h
, "cl_device_id opencl_create_device("
116 fprintf(info
->kernel_h
, "cl_program opencl_build_program_from_string("
118 "cl_device_id dev, const char *program_source, "
119 "size_t program_size, "
120 "const char *opencl_options);\n");
121 fprintf(info
->kernel_h
, "cl_program opencl_build_program_from_file("
123 "cl_device_id dev, const char *filename, "
124 "const char *opencl_options);\n");
125 fprintf(info
->kernel_h
,
126 "const char *opencl_error_string(cl_int error);\n");
127 for (i
= 0; i
< info
->options
->opencl_n_include_file
; ++i
) {
128 info
->kprinter
= isl_printer_print_str(info
->kprinter
,
130 info
->kprinter
= isl_printer_print_str(info
->kprinter
,
131 info
->options
->opencl_include_files
[i
]);
132 info
->kprinter
= isl_printer_print_str(info
->kprinter
, ">\n");
138 /* Write text to a file and escape some special characters that would break a
141 static void opencl_print_escaped(const char *str
, const char *end
, FILE *file
)
143 const char *prev
= str
;
145 while ((str
= strpbrk(prev
, "\"\\")) && str
< end
) {
146 fwrite(prev
, 1, str
- prev
, file
);
147 fprintf(file
, "\\%c", *str
);
152 fwrite(prev
, 1, end
- prev
, file
);
155 /* Write text to a file as a C string literal.
157 * This function also prints any characters after the last newline, although
158 * normally the input string should end with a newline.
160 static void opencl_print_as_c_string(const char *str
, FILE *file
)
162 const char *prev
= str
;
164 while ((str
= strchr(prev
, '\n'))) {
165 fprintf(file
, "\n\"");
166 opencl_print_escaped(prev
, str
, file
);
167 fprintf(file
, "\\n\"");
173 fprintf(file
, "\n\"");
174 opencl_print_escaped(prev
, prev
+ strlen(prev
), file
);
179 /* Write the code that we have accumulated in the kernel isl_printer to the
180 * kernel.cl file. If the opencl_embed_kernel_code option has been set, print
181 * the code as a C string literal. Start that string literal with an empty
182 * line, such that line numbers reported by the OpenCL C compiler match those
183 * of the kernel file.
185 * Return 0 on success and -1 on failure.
187 static int opencl_write_kernel_file(struct opencl_info
*opencl
)
189 char *raw
= isl_printer_get_str(opencl
->kprinter
);
194 if (opencl
->options
->opencl_embed_kernel_code
) {
195 fprintf(opencl
->kernel_c
,
196 "static const char kernel_code[] = \"\\n\"");
197 opencl_print_as_c_string(raw
, opencl
->kernel_c
);
198 fprintf(opencl
->kernel_c
, ";\n");
200 fprintf(opencl
->kernel_c
, "%s", raw
);
207 /* Close all output files. Write the kernel contents to the kernel file before
210 * Return 0 on success and -1 on failure.
212 static int opencl_close_files(struct opencl_info
*info
)
216 if (info
->kernel_c
) {
217 r
= opencl_write_kernel_file(info
);
218 fclose(info
->kernel_c
);
221 fclose(info
->kernel_h
);
223 fclose(info
->host_c
);
228 static __isl_give isl_printer
*opencl_print_host_macros(
229 __isl_take isl_printer
*p
)
232 "#define openclCheckReturn(ret) \\\n"
233 " if (ret != CL_SUCCESS) {\\\n"
234 " fprintf(stderr, \"OpenCL error: %s\\n\", "
235 "opencl_error_string(ret)); \\\n"
236 " fflush(stderr); \\\n"
237 " assert(ret == CL_SUCCESS);\\\n }\n";
239 p
= isl_printer_start_line(p
);
240 p
= isl_printer_print_str(p
, macros
);
241 p
= isl_printer_end_line(p
);
243 p
= isl_ast_op_type_print_macro(isl_ast_op_max
, p
);
248 static __isl_give isl_printer
*opencl_declare_device_arrays(
249 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
253 for (i
= 0; i
< prog
->n_array
; ++i
) {
254 if (gpu_array_is_read_only_scalar(&prog
->array
[i
]))
256 p
= isl_printer_start_line(p
);
257 p
= isl_printer_print_str(p
, "cl_mem dev_");
258 p
= isl_printer_print_str(p
, prog
->array
[i
].name
);
259 p
= isl_printer_print_str(p
, ";");
260 p
= isl_printer_end_line(p
);
262 p
= isl_printer_start_line(p
);
263 p
= isl_printer_end_line(p
);
267 /* Given an array, check whether its positive size guard expression is
270 static int is_array_positive_size_guard_trivial(struct gpu_array_info
*array
)
275 guard
= gpu_array_positive_size_guard(array
);
276 is_trivial
= isl_set_plain_is_universe(guard
);
281 /* Allocate a device array for "array'.
283 * Emit a max-expression to ensure the device array can contain at least one
284 * element if the array's positive size guard expression is not trivial.
286 static __isl_give isl_printer
*allocate_device_array(__isl_take isl_printer
*p
,
287 struct gpu_array_info
*array
)
289 int need_lower_bound
;
291 p
= ppcg_start_block(p
);
293 p
= isl_printer_start_line(p
);
294 p
= isl_printer_print_str(p
, "dev_");
295 p
= isl_printer_print_str(p
, array
->name
);
296 p
= isl_printer_print_str(p
, " = clCreateBuffer(context, ");
297 p
= isl_printer_print_str(p
, "CL_MEM_READ_WRITE, ");
299 need_lower_bound
= !is_array_positive_size_guard_trivial(array
);
300 if (need_lower_bound
) {
301 p
= isl_printer_print_str(p
, "max(sizeof(");
302 p
= isl_printer_print_str(p
, array
->type
);
303 p
= isl_printer_print_str(p
, "), ");
305 p
= gpu_array_info_print_size(p
, array
);
306 if (need_lower_bound
)
307 p
= isl_printer_print_str(p
, ")");
309 p
= isl_printer_print_str(p
, ", NULL, &err);");
310 p
= isl_printer_end_line(p
);
311 p
= isl_printer_start_line(p
);
312 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
313 p
= isl_printer_end_line(p
);
315 p
= ppcg_end_block(p
);
320 /* Allocate device arrays.
322 static __isl_give isl_printer
*opencl_allocate_device_arrays(
323 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
327 for (i
= 0; i
< prog
->n_array
; ++i
) {
328 struct gpu_array_info
*array
= &prog
->array
[i
];
330 if (gpu_array_is_read_only_scalar(array
))
333 p
= allocate_device_array(p
, array
);
335 p
= isl_printer_start_line(p
);
336 p
= isl_printer_end_line(p
);
340 /* Print a call to the OpenCL clSetKernelArg() function which sets
341 * the arguments of the kernel. arg_name and arg_index are the name and the
342 * index of the kernel argument. The index of the leftmost argument of
343 * the kernel is 0 whereas the index of the rightmost argument of the kernel
344 * is n - 1, where n is the total number of the kernel arguments.
345 * read_only_scalar is a boolean that indicates whether the argument is a read
348 static __isl_give isl_printer
*opencl_set_kernel_argument(
349 __isl_take isl_printer
*p
, int kernel_id
,
350 const char *arg_name
, int arg_index
, int read_only_scalar
)
352 p
= isl_printer_start_line(p
);
353 p
= isl_printer_print_str(p
,
354 "openclCheckReturn(clSetKernelArg(kernel");
355 p
= isl_printer_print_int(p
, kernel_id
);
356 p
= isl_printer_print_str(p
, ", ");
357 p
= isl_printer_print_int(p
, arg_index
);
358 p
= isl_printer_print_str(p
, ", sizeof(");
360 if (read_only_scalar
) {
361 p
= isl_printer_print_str(p
, arg_name
);
362 p
= isl_printer_print_str(p
, "), &");
364 p
= isl_printer_print_str(p
, "cl_mem), (void *) &dev_");
366 p
= isl_printer_print_str(p
, arg_name
);
367 p
= isl_printer_print_str(p
, "));");
368 p
= isl_printer_end_line(p
);
373 /* Print the block sizes as a list of the sizes in each
376 static __isl_give isl_printer
*opencl_print_block_sizes(
377 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
381 if (kernel
->n_block
> 0)
382 for (i
= 0; i
< kernel
->n_block
; ++i
) {
384 p
= isl_printer_print_str(p
, ", ");
385 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
388 p
= isl_printer_print_str(p
, "1");
393 /* Set the arguments of the OpenCL kernel by printing a call to the OpenCL
394 * clSetKernelArg() function for each kernel argument.
396 static __isl_give isl_printer
*opencl_set_kernel_arguments(
397 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
398 struct ppcg_kernel
*kernel
)
405 for (i
= 0; i
< prog
->n_array
; ++i
) {
409 space
= isl_space_copy(prog
->array
[i
].space
);
410 arr
= isl_union_set_extract_set(kernel
->arrays
, space
);
411 empty
= isl_set_plain_is_empty(arr
);
415 ro
= gpu_array_is_read_only_scalar(&prog
->array
[i
]);
416 opencl_set_kernel_argument(p
, kernel
->id
, prog
->array
[i
].name
,
421 space
= isl_union_set_get_space(kernel
->arrays
);
422 nparam
= isl_space_dim(space
, isl_dim_param
);
423 for (i
= 0; i
< nparam
; ++i
) {
426 name
= isl_space_get_dim_name(space
, isl_dim_param
, i
);
427 opencl_set_kernel_argument(p
, kernel
->id
, name
, arg_index
, 1);
430 isl_space_free(space
);
432 n
= isl_space_dim(kernel
->space
, isl_dim_set
);
433 for (i
= 0; i
< n
; ++i
) {
436 name
= isl_space_get_dim_name(kernel
->space
, isl_dim_set
, i
);
437 opencl_set_kernel_argument(p
, kernel
->id
, name
, arg_index
, 1);
444 /* Print the arguments to a kernel declaration or call. If "types" is set,
445 * then print a declaration (including the types of the arguments).
447 * The arguments are printed in the following order
448 * - the arrays accessed by the kernel
450 * - the host loop iterators
452 static __isl_give isl_printer
*opencl_print_kernel_arguments(
453 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
454 struct ppcg_kernel
*kernel
, int types
)
462 for (i
= 0; i
< prog
->n_array
; ++i
) {
466 space
= isl_space_copy(prog
->array
[i
].space
);
467 arr
= isl_union_set_extract_set(kernel
->arrays
, space
);
468 empty
= isl_set_plain_is_empty(arr
);
474 p
= isl_printer_print_str(p
, ", ");
477 p
= gpu_array_info_print_declaration_argument(p
,
478 &prog
->array
[i
], "__global");
480 p
= gpu_array_info_print_call_argument(p
,
486 space
= isl_union_set_get_space(kernel
->arrays
);
487 nparam
= isl_space_dim(space
, isl_dim_param
);
488 for (i
= 0; i
< nparam
; ++i
) {
491 name
= isl_space_get_dim_name(space
, isl_dim_param
, i
);
494 p
= isl_printer_print_str(p
, ", ");
496 p
= isl_printer_print_str(p
, "int ");
497 p
= isl_printer_print_str(p
, name
);
501 isl_space_free(space
);
503 n
= isl_space_dim(kernel
->space
, isl_dim_set
);
504 type
= isl_options_get_ast_iterator_type(prog
->ctx
);
505 for (i
= 0; i
< n
; ++i
) {
509 p
= isl_printer_print_str(p
, ", ");
510 name
= isl_space_get_dim_name(kernel
->space
, isl_dim_set
, i
);
512 p
= isl_printer_print_str(p
, type
);
513 p
= isl_printer_print_str(p
, " ");
515 p
= isl_printer_print_str(p
, name
);
523 /* Print the header of the given kernel.
525 static __isl_give isl_printer
*opencl_print_kernel_header(
526 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
527 struct ppcg_kernel
*kernel
)
529 p
= isl_printer_start_line(p
);
530 p
= isl_printer_print_str(p
, "__kernel void kernel");
531 p
= isl_printer_print_int(p
, kernel
->id
);
532 p
= isl_printer_print_str(p
, "(");
533 p
= opencl_print_kernel_arguments(p
, prog
, kernel
, 1);
534 p
= isl_printer_print_str(p
, ")");
535 p
= isl_printer_end_line(p
);
540 /* Print a list of "n" iterators of type "type" called "prefix%d" to "p".
541 * Each iterator is assigned the corresponding opencl identifier returned
542 * by the function "opencl_id".
543 * Unlike the equivalent function in the CUDA backend which prints iterators
544 * in reverse order to promote coalescing, this function does not print
545 * iterators in reverse order. The OpenCL backend currently does not take
546 * into account any coalescing considerations.
548 static __isl_give isl_printer
*print_iterators(__isl_take isl_printer
*p
,
549 const char *type
, int n
, const char *prefix
, const char *opencl_id
)
555 p
= isl_printer_start_line(p
);
556 p
= isl_printer_print_str(p
, type
);
557 p
= isl_printer_print_str(p
, " ");
558 for (i
= 0; i
< n
; ++i
) {
560 p
= isl_printer_print_str(p
, ", ");
561 p
= isl_printer_print_str(p
, prefix
);
562 p
= isl_printer_print_int(p
, i
);
563 p
= isl_printer_print_str(p
, " = ");
564 p
= isl_printer_print_str(p
, opencl_id
);
565 p
= isl_printer_print_str(p
, "(");
566 p
= isl_printer_print_int(p
, i
);
567 p
= isl_printer_print_str(p
, ")");
569 p
= isl_printer_print_str(p
, ";");
570 p
= isl_printer_end_line(p
);
575 static __isl_give isl_printer
*opencl_print_kernel_iterators(
576 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
579 isl_ctx
*ctx
= isl_ast_node_get_ctx(kernel
->tree
);
582 type
= isl_options_get_ast_iterator_type(ctx
);
584 n_grid
= isl_multi_pw_aff_dim(kernel
->grid_size
, isl_dim_set
);
585 p
= print_iterators(p
, type
, n_grid
, "b", "get_group_id");
586 p
= print_iterators(p
, type
, kernel
->n_block
, "t", "get_local_id");
591 static __isl_give isl_printer
*opencl_print_kernel_var(
592 __isl_take isl_printer
*p
, struct ppcg_kernel_var
*var
)
597 p
= isl_printer_start_line(p
);
598 if (var
->type
== ppcg_access_shared
)
599 p
= isl_printer_print_str(p
, "__local ");
600 p
= isl_printer_print_str(p
, var
->array
->type
);
601 p
= isl_printer_print_str(p
, " ");
602 p
= isl_printer_print_str(p
, var
->name
);
603 for (j
= 0; j
< var
->array
->n_index
; ++j
) {
604 p
= isl_printer_print_str(p
, "[");
605 v
= isl_vec_get_element_val(var
->size
, j
);
606 p
= isl_printer_print_val(p
, v
);
607 p
= isl_printer_print_str(p
, "]");
610 p
= isl_printer_print_str(p
, ";");
611 p
= isl_printer_end_line(p
);
616 static __isl_give isl_printer
*opencl_print_kernel_vars(
617 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
621 for (i
= 0; i
< kernel
->n_var
; ++i
)
622 p
= opencl_print_kernel_var(p
, &kernel
->var
[i
]);
627 /* Print a call to barrier() which is a sync statement.
628 * All work-items in a work-group executing the kernel on a processor must
629 * execute the barrier() function before any are allowed to continue execution
630 * beyond the barrier.
631 * The flag CLK_LOCAL_MEM_FENCE makes the barrier function either flush any
632 * variables stored in local memory or queue a memory fence to ensure correct
633 * ordering of memory operations to local memory.
634 * The flag CLK_GLOBAL_MEM_FENCE makes the barrier function queue a memory
635 * fence to ensure correct ordering of memory operations to global memory.
637 static __isl_give isl_printer
*opencl_print_sync(__isl_take isl_printer
*p
,
638 struct ppcg_kernel_stmt
*stmt
)
640 p
= isl_printer_start_line(p
);
641 p
= isl_printer_print_str(p
,
642 "barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);");
643 p
= isl_printer_end_line(p
);
648 /* This function is called for each user statement in the AST,
649 * i.e., for each kernel body statement, copy statement or sync statement.
651 static __isl_give isl_printer
*opencl_print_kernel_stmt(
652 __isl_take isl_printer
*p
,
653 __isl_take isl_ast_print_options
*print_options
,
654 __isl_keep isl_ast_node
*node
, void *user
)
657 struct ppcg_kernel_stmt
*stmt
;
659 id
= isl_ast_node_get_annotation(node
);
660 stmt
= isl_id_get_user(id
);
663 isl_ast_print_options_free(print_options
);
665 switch (stmt
->type
) {
666 case ppcg_kernel_copy
:
667 return ppcg_kernel_print_copy(p
, stmt
);
668 case ppcg_kernel_sync
:
669 return opencl_print_sync(p
, stmt
);
670 case ppcg_kernel_domain
:
671 return ppcg_kernel_print_domain(p
, stmt
);
677 /* Return true if there is a double array in prog->array or
678 * if any of the types in prog->scop involve any doubles.
679 * To check the latter condition, we simply search for the string "double"
680 * in the type definitions, which may result in false positives.
682 static __isl_give
int any_double_elements(struct gpu_prog
*prog
)
686 for (i
= 0; i
< prog
->n_array
; ++i
)
687 if (strcmp(prog
->array
[i
].type
, "double") == 0)
690 for (i
= 0; i
< prog
->scop
->pet
->n_type
; ++i
) {
691 struct pet_type
*type
= prog
->scop
->pet
->types
[i
];
693 if (strstr(type
->definition
, "double"))
700 /* Prints a #pragma to enable support for double floating-point
701 * precision. OpenCL 1.0 adds support for double precision floating-point as
702 * an optional extension. An application that wants to use double will need to
703 * include the #pragma OPENCL EXTENSION cl_khr_fp64 : enable directive before
704 * any double precision data type is declared in the kernel code.
706 static __isl_give isl_printer
*opencl_enable_double_support(
707 __isl_take isl_printer
*p
)
709 p
= isl_printer_start_line(p
);
710 p
= isl_printer_print_str(p
, "#pragma OPENCL EXTENSION cl_khr_fp64 :"
712 p
= isl_printer_end_line(p
);
713 p
= isl_printer_start_line(p
);
714 p
= isl_printer_end_line(p
);
719 static __isl_give isl_printer
*opencl_print_kernel(struct gpu_prog
*prog
,
720 struct ppcg_kernel
*kernel
, __isl_take isl_printer
*p
)
722 isl_ctx
*ctx
= isl_ast_node_get_ctx(kernel
->tree
);
723 isl_ast_print_options
*print_options
;
725 print_options
= isl_ast_print_options_alloc(ctx
);
726 print_options
= isl_ast_print_options_set_print_user(print_options
,
727 &opencl_print_kernel_stmt
, NULL
);
729 p
= isl_printer_set_output_format(p
, ISL_FORMAT_C
);
730 p
= opencl_print_kernel_header(p
, prog
, kernel
);
731 p
= isl_printer_print_str(p
, "{");
732 p
= isl_printer_end_line(p
);
733 p
= isl_printer_indent(p
, 4);
734 p
= opencl_print_kernel_iterators(p
, kernel
);
735 p
= opencl_print_kernel_vars(p
, kernel
);
736 p
= isl_printer_end_line(p
);
737 p
= gpu_print_macros(p
, kernel
->tree
);
738 p
= isl_ast_node_print(kernel
->tree
, p
, print_options
);
739 p
= isl_printer_indent(p
, -4);
740 p
= isl_printer_start_line(p
);
741 p
= isl_printer_print_str(p
, "}");
742 p
= isl_printer_end_line(p
);
747 struct print_host_user_data_opencl
{
748 struct opencl_info
*opencl
;
749 struct gpu_prog
*prog
;
752 /* This function prints the i'th block size multiplied by the i'th grid size,
753 * where i (a parameter to this function) is one of the possible dimensions of
754 * grid sizes and block sizes.
755 * If the dimension of block sizes is not equal to the dimension of grid sizes
756 * the output is calculated as follows:
759 * block_sizes[dim1] is the list of blocks sizes and it contains dim1 elements.
760 * grid_sizes[dim2] is the list of grid sizes and it contains dim2 elements.
763 * If (i > dim2) then the output is block_sizes[i]
764 * If (i > dim1) then the output is grid_sizes[i]
766 static __isl_give isl_printer
*opencl_print_total_number_of_work_items_for_dim(
767 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
, int i
)
769 int grid_dim
, block_dim
;
770 isl_pw_aff
*bound_grid
;
772 grid_dim
= isl_multi_pw_aff_dim(kernel
->grid_size
, isl_dim_set
);
773 block_dim
= kernel
->n_block
;
775 if (i
< min(grid_dim
, block_dim
)) {
776 bound_grid
= isl_multi_pw_aff_get_pw_aff(kernel
->grid_size
, i
);
777 p
= isl_printer_print_str(p
, "(");
778 p
= isl_printer_print_pw_aff(p
, bound_grid
);
779 p
= isl_printer_print_str(p
, ") * ");
780 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
781 isl_pw_aff_free(bound_grid
);
782 } else if (i
>= grid_dim
)
783 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
785 bound_grid
= isl_multi_pw_aff_get_pw_aff(kernel
->grid_size
, i
);
786 p
= isl_printer_print_pw_aff(p
, bound_grid
);
787 isl_pw_aff_free(bound_grid
);
793 /* Print a list that represents the total number of work items. The list is
794 * constructed by performing an element-wise multiplication of the block sizes
795 * and the grid sizes. To explain how the list is constructed, suppose that:
796 * block_sizes[dim1] is the list of blocks sizes and it contains dim1 elements.
797 * grid_sizes[dim2] is the list of grid sizes and it contains dim2 elements.
799 * The output of this function is constructed as follows:
800 * If (dim1 > dim2) then the output is the following list:
801 * grid_sizes[0]*block_sizes[0], ..., grid_sizes[dim2-1]*block_sizes[dim2-1],
802 * block_sizes[dim2], ..., block_sizes[dim1-2], block_sizes[dim1-1].
804 * If (dim2 > dim1) then the output is the following list:
805 * grid_sizes[0]*block_sizes[0], ..., grid_sizes[dim1-1] * block_sizes[dim1-1],
806 * grid_sizes[dim1], grid_sizes[dim2-2], grid_sizes[dim2-1].
808 * To calculate the total number of work items out of the list constructed by
809 * this function, the user should multiply the elements of the list.
811 static __isl_give isl_printer
*opencl_print_total_number_of_work_items_as_list(
812 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
815 int grid_dim
, block_dim
;
817 grid_dim
= isl_multi_pw_aff_dim(kernel
->grid_size
, isl_dim_set
);
818 block_dim
= kernel
->n_block
;
820 if ((grid_dim
<= 0) || (block_dim
<= 0)) {
821 p
= isl_printer_print_str(p
, "1");
825 for (i
= 0; i
<= max(grid_dim
, block_dim
) - 1; i
++) {
827 p
= isl_printer_print_str(p
, ", ");
829 p
= opencl_print_total_number_of_work_items_for_dim(p
,
836 /* Copy "array" from the host to the device (to_host = 0) or
837 * back from the device to the host (to_host = 1).
839 static __isl_give isl_printer
*copy_array(__isl_take isl_printer
*p
,
840 struct gpu_array_info
*array
, int to_host
)
842 p
= isl_printer_start_line(p
);
843 p
= isl_printer_print_str(p
, "openclCheckReturn(");
845 p
= isl_printer_print_str(p
, "clEnqueueReadBuffer");
847 p
= isl_printer_print_str(p
, "clEnqueueWriteBuffer");
848 p
= isl_printer_print_str(p
, "(queue, dev_");
849 p
= isl_printer_print_str(p
, array
->name
);
850 p
= isl_printer_print_str(p
, ", CL_TRUE, 0, ");
851 p
= gpu_array_info_print_size(p
, array
);
853 if (gpu_array_is_scalar(array
))
854 p
= isl_printer_print_str(p
, ", &");
856 p
= isl_printer_print_str(p
, ", ");
857 p
= isl_printer_print_str(p
, array
->name
);
858 p
= isl_printer_print_str(p
, ", 0, NULL, NULL));");
859 p
= isl_printer_end_line(p
);
864 /* Copy "array" from the host to the device.
866 static __isl_give isl_printer
*copy_array_to_device(__isl_take isl_printer
*p
,
869 struct gpu_array_info
*array
= user
;
871 return copy_array(p
, array
, 0);
874 /* Copy "array" back from the device to the host.
876 static __isl_give isl_printer
*copy_array_from_device(__isl_take isl_printer
*p
,
879 struct gpu_array_info
*array
= user
;
881 return copy_array(p
, array
, 1);
884 /* Copy the "copy" arrays from the host to the device (to_host = 0) or
885 * back from the device to the host (to_host = 1).
887 * Only perform the copying for arrays with strictly positive size.
889 static __isl_give isl_printer
*opencl_copy_arrays(__isl_take isl_printer
*p
,
890 struct gpu_prog
*prog
, __isl_keep isl_union_set
*copy
, int to_host
)
894 for (i
= 0; i
< prog
->n_array
; ++i
) {
895 struct gpu_array_info
*array
= &prog
->array
[i
];
901 if (gpu_array_is_read_only_scalar(array
))
904 space
= isl_space_copy(array
->space
);
905 copy_i
= isl_union_set_extract_set(copy
, space
);
906 empty
= isl_set_plain_is_empty(copy_i
);
907 isl_set_free(copy_i
);
911 guard
= gpu_array_positive_size_guard(array
);
912 p
= ppcg_print_guarded(p
, guard
, isl_set_copy(prog
->context
),
913 to_host
? ©_array_from_device
:
914 ©_array_to_device
, array
);
917 p
= isl_printer_start_line(p
);
918 p
= isl_printer_end_line(p
);
922 /* Copy the prog->copy_in arrays from the host to the device.
924 static __isl_give isl_printer
*opencl_copy_arrays_to_device(
925 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
927 return opencl_copy_arrays(p
, prog
, prog
->copy_in
, 0);
930 /* Copy the prog->copy_out arrays back from the device to the host.
932 static __isl_give isl_printer
*opencl_copy_arrays_from_device(
933 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
935 return opencl_copy_arrays(p
, prog
, prog
->copy_out
, 1);
938 /* Print the user statement of the host code to "p".
940 * In particular, print a block of statements that defines the grid
941 * and the work group and then launches the kernel.
943 * A grid is composed of many work groups (blocks), each work group holds
944 * many work-items (threads).
946 * global_work_size[kernel->n_block] represents the total number of work
947 * items. It points to an array of kernel->n_block unsigned
948 * values that describe the total number of work-items that will execute
949 * the kernel. The total number of work-items is computed as:
950 * global_work_size[0] *...* global_work_size[kernel->n_block - 1].
952 * The size of each work group (i.e. the number of work-items in each work
953 * group) is described using block_size[kernel->n_block]. The total
954 * number of work-items in a block (work-group) is computed as:
955 * block_size[0] *... * block_size[kernel->n_block - 1].
957 * For more information check:
958 * http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html
960 static __isl_give isl_printer
*opencl_print_host_user(
961 __isl_take isl_printer
*p
,
962 __isl_take isl_ast_print_options
*print_options
,
963 __isl_keep isl_ast_node
*node
, void *user
)
966 struct ppcg_kernel
*kernel
;
967 struct print_host_user_data_opencl
*data
;
969 id
= isl_ast_node_get_annotation(node
);
970 kernel
= isl_id_get_user(id
);
973 data
= (struct print_host_user_data_opencl
*) user
;
975 p
= isl_printer_start_line(p
);
976 p
= isl_printer_print_str(p
, "{");
977 p
= isl_printer_end_line(p
);
978 p
= isl_printer_indent(p
, 2);
980 p
= isl_printer_start_line(p
);
981 p
= isl_printer_print_str(p
, "size_t global_work_size[");
983 if (kernel
->n_block
> 0)
984 p
= isl_printer_print_int(p
, kernel
->n_block
);
986 p
= isl_printer_print_int(p
, 1);
988 p
= isl_printer_print_str(p
, "] = {");
989 p
= opencl_print_total_number_of_work_items_as_list(p
, kernel
);
990 p
= isl_printer_print_str(p
, "};");
991 p
= isl_printer_end_line(p
);
993 p
= isl_printer_start_line(p
);
994 p
= isl_printer_print_str(p
, "size_t block_size[");
996 if (kernel
->n_block
> 0)
997 p
= isl_printer_print_int(p
, kernel
->n_block
);
999 p
= isl_printer_print_int(p
, 1);
1001 p
= isl_printer_print_str(p
, "] = {");
1002 p
= opencl_print_block_sizes(p
, kernel
);
1003 p
= isl_printer_print_str(p
, "};");
1004 p
= isl_printer_end_line(p
);
1006 p
= isl_printer_start_line(p
);
1007 p
= isl_printer_print_str(p
, "cl_kernel kernel");
1008 p
= isl_printer_print_int(p
, kernel
->id
);
1009 p
= isl_printer_print_str(p
, " = clCreateKernel(program, \"kernel");
1010 p
= isl_printer_print_int(p
, kernel
->id
);
1011 p
= isl_printer_print_str(p
, "\", &err);");
1012 p
= isl_printer_end_line(p
);
1013 p
= isl_printer_start_line(p
);
1014 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
1015 p
= isl_printer_end_line(p
);
1017 opencl_set_kernel_arguments(p
, data
->prog
, kernel
);
1019 p
= isl_printer_start_line(p
);
1020 p
= isl_printer_print_str(p
, "openclCheckReturn(clEnqueueNDRangeKernel"
1022 p
= isl_printer_print_int(p
, kernel
->id
);
1023 p
= isl_printer_print_str(p
, ", ");
1024 if (kernel
->n_block
> 0)
1025 p
= isl_printer_print_int(p
, kernel
->n_block
);
1027 p
= isl_printer_print_int(p
, 1);
1029 p
= isl_printer_print_str(p
, ", NULL, global_work_size, "
1031 "0, NULL, NULL));");
1032 p
= isl_printer_end_line(p
);
1033 p
= isl_printer_start_line(p
);
1034 p
= isl_printer_print_str(p
, "openclCheckReturn("
1035 "clReleaseKernel(kernel");
1036 p
= isl_printer_print_int(p
, kernel
->id
);
1037 p
= isl_printer_print_str(p
, "));");
1038 p
= isl_printer_end_line(p
);
1039 p
= isl_printer_start_line(p
);
1040 p
= isl_printer_print_str(p
, "clFinish(queue);");
1041 p
= isl_printer_end_line(p
);
1042 p
= isl_printer_indent(p
, -2);
1043 p
= isl_printer_start_line(p
);
1044 p
= isl_printer_print_str(p
, "}");
1045 p
= isl_printer_end_line(p
);
1047 p
= isl_printer_start_line(p
);
1048 p
= isl_printer_end_line(p
);
1050 data
->opencl
->kprinter
= opencl_print_kernel(data
->prog
, kernel
,
1051 data
->opencl
->kprinter
);
1053 isl_ast_print_options_free(print_options
);
1058 static __isl_give isl_printer
*opencl_print_host_code(
1059 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
1060 __isl_keep isl_ast_node
*tree
, struct opencl_info
*opencl
)
1062 isl_ast_print_options
*print_options
;
1063 isl_ctx
*ctx
= isl_ast_node_get_ctx(tree
);
1064 struct print_host_user_data_opencl data
= { opencl
, prog
};
1066 print_options
= isl_ast_print_options_alloc(ctx
);
1067 print_options
= isl_ast_print_options_set_print_user(print_options
,
1068 &opencl_print_host_user
, &data
);
1070 p
= gpu_print_macros(p
, tree
);
1071 p
= isl_ast_node_print(tree
, p
, print_options
);
1076 /* Create an OpenCL device, context, command queue and build the kernel.
1077 * input is the name of the input file provided to ppcg.
1079 static __isl_give isl_printer
*opencl_setup(__isl_take isl_printer
*p
,
1080 const char *input
, struct opencl_info
*info
)
1082 p
= isl_printer_start_line(p
);
1083 p
= isl_printer_print_str(p
, "cl_device_id device;");
1084 p
= isl_printer_end_line(p
);
1085 p
= isl_printer_start_line(p
);
1086 p
= isl_printer_print_str(p
, "cl_context context;");
1087 p
= isl_printer_end_line(p
);
1088 p
= isl_printer_start_line(p
);
1089 p
= isl_printer_print_str(p
, "cl_program program;");
1090 p
= isl_printer_end_line(p
);
1091 p
= isl_printer_start_line(p
);
1092 p
= isl_printer_print_str(p
, "cl_command_queue queue;");
1093 p
= isl_printer_end_line(p
);
1094 p
= isl_printer_start_line(p
);
1095 p
= isl_printer_print_str(p
, "cl_int err;");
1096 p
= isl_printer_end_line(p
);
1097 p
= isl_printer_start_line(p
);
1098 p
= isl_printer_print_str(p
, "device = opencl_create_device(");
1099 p
= isl_printer_print_int(p
, info
->options
->opencl_use_gpu
);
1100 p
= isl_printer_print_str(p
, ");");
1101 p
= isl_printer_end_line(p
);
1102 p
= isl_printer_start_line(p
);
1103 p
= isl_printer_print_str(p
, "context = clCreateContext(NULL, 1, "
1104 "&device, NULL, NULL, &err);");
1105 p
= isl_printer_end_line(p
);
1106 p
= isl_printer_start_line(p
);
1107 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
1108 p
= isl_printer_end_line(p
);
1109 p
= isl_printer_start_line(p
);
1110 p
= isl_printer_print_str(p
, "queue = clCreateCommandQueue"
1111 "(context, device, 0, &err);");
1112 p
= isl_printer_end_line(p
);
1113 p
= isl_printer_start_line(p
);
1114 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
1115 p
= isl_printer_end_line(p
);
1117 p
= isl_printer_start_line(p
);
1118 p
= isl_printer_print_str(p
, "program = ");
1120 if (info
->options
->opencl_embed_kernel_code
) {
1121 p
= isl_printer_print_str(p
, "opencl_build_program_from_string("
1122 "context, device, kernel_code, "
1123 "sizeof(kernel_code), \"");
1125 p
= isl_printer_print_str(p
, "opencl_build_program_from_file("
1126 "context, device, \"");
1127 p
= isl_printer_print_str(p
, info
->kernel_c_name
);
1128 p
= isl_printer_print_str(p
, "\", \"");
1131 if (info
->options
->opencl_compiler_options
)
1132 p
= isl_printer_print_str(p
,
1133 info
->options
->opencl_compiler_options
);
1135 p
= isl_printer_print_str(p
, "\");");
1136 p
= isl_printer_end_line(p
);
1137 p
= isl_printer_start_line(p
);
1138 p
= isl_printer_end_line(p
);
1143 static __isl_give isl_printer
*opencl_release_cl_objects(
1144 __isl_take isl_printer
*p
, struct opencl_info
*info
)
1146 p
= isl_printer_start_line(p
);
1147 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseCommandQueue"
1149 p
= isl_printer_end_line(p
);
1150 p
= isl_printer_start_line(p
);
1151 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseProgram"
1153 p
= isl_printer_end_line(p
);
1154 p
= isl_printer_start_line(p
);
1155 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseContext"
1157 p
= isl_printer_end_line(p
);
1162 /* Free the device array corresponding to "array"
1164 static __isl_give isl_printer
*release_device_array(__isl_take isl_printer
*p
,
1165 struct gpu_array_info
*array
)
1167 p
= isl_printer_start_line(p
);
1168 p
= isl_printer_print_str(p
, "openclCheckReturn("
1169 "clReleaseMemObject(dev_");
1170 p
= isl_printer_print_str(p
, array
->name
);
1171 p
= isl_printer_print_str(p
, "));");
1172 p
= isl_printer_end_line(p
);
1177 /* Free the device arrays.
1179 static __isl_give isl_printer
*opencl_release_device_arrays(
1180 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
1184 for (i
= 0; i
< prog
->n_array
; ++i
) {
1185 struct gpu_array_info
*array
= &prog
->array
[i
];
1186 if (gpu_array_is_read_only_scalar(array
))
1189 p
= release_device_array(p
, array
);
1194 /* Given a gpu_prog "prog" and the corresponding transformed AST
1195 * "tree", print the entire OpenCL code to "p".
1197 static __isl_give isl_printer
*print_opencl(__isl_take isl_printer
*p
,
1198 struct gpu_prog
*prog
, __isl_keep isl_ast_node
*tree
,
1199 struct gpu_types
*types
, void *user
)
1201 struct opencl_info
*opencl
= user
;
1203 opencl
->kprinter
= isl_printer_set_output_format(opencl
->kprinter
,
1205 if (any_double_elements(prog
))
1206 opencl
->kprinter
= opencl_enable_double_support(
1208 if (opencl
->options
->opencl_print_kernel_types
)
1209 opencl
->kprinter
= gpu_print_types(opencl
->kprinter
, types
,
1212 if (!opencl
->kprinter
)
1213 return isl_printer_free(p
);
1215 p
= ppcg_start_block(p
);
1217 p
= opencl_print_host_macros(p
);
1219 p
= opencl_declare_device_arrays(p
, prog
);
1220 p
= opencl_setup(p
, opencl
->input
, opencl
);
1221 p
= opencl_allocate_device_arrays(p
, prog
);
1222 p
= opencl_copy_arrays_to_device(p
, prog
);
1224 p
= opencl_print_host_code(p
, prog
, tree
, opencl
);
1226 p
= opencl_copy_arrays_from_device(p
, prog
);
1227 p
= opencl_release_device_arrays(p
, prog
);
1228 p
= opencl_release_cl_objects(p
, opencl
);
1230 p
= ppcg_end_block(p
);
1235 /* Transform the code in the file called "input" by replacing
1236 * all scops by corresponding OpenCL code.
1237 * The host code is written to "output" or a name derived from
1238 * "input" if "output" is NULL.
1239 * The kernel code is placed in separate files with names
1240 * derived from "output" or "input".
1242 * We let generate_gpu do all the hard work and then let it call
1243 * us back for printing the AST in print_opencl.
1245 * To prepare for this printing, we first open the output files
1246 * and we close them after generate_gpu has finished.
1248 int generate_opencl(isl_ctx
*ctx
, struct ppcg_options
*options
,
1249 const char *input
, const char *output
)
1251 struct opencl_info opencl
= { options
, input
, output
};
1254 opencl
.kprinter
= isl_printer_to_str(ctx
);
1255 r
= opencl_open_files(&opencl
);
1258 r
= generate_gpu(ctx
, input
, opencl
.host_c
, options
,
1259 &print_opencl
, &opencl
);
1261 if (opencl_close_files(&opencl
) < 0)
1263 isl_printer_free(opencl
.kprinter
);