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"
25 #define min(a, b) (((a) < (b)) ? (a) : (b))
26 #define max(a, b) (((a) > (b)) ? (a) : (b))
28 /* options are the global options passed to generate_opencl.
29 * input is the name of the input file.
30 * output is the user-specified output file name and may be NULL
31 * if not specified by the user.
32 * kernel_c_name is the name of the kernel_c file.
33 * kprinter is an isl_printer for the kernel file.
34 * host_c is the generated source file for the host code. kernel_c is
35 * the generated source file for the kernel.
38 struct ppcg_options
*options
;
41 char kernel_c_name
[PATH_MAX
];
43 isl_printer
*kprinter
;
49 /* Open the file called "name" for writing or print an error message.
51 static FILE *open_or_croak(const char *name
)
55 file
= fopen(name
, "w");
57 fprintf(stderr
, "Failed to open \"%s\" for writing\n", name
);
61 /* Open the host .c file and the kernel .h and .cl files for writing.
62 * Their names are derived from info->output (or info->input if
63 * the user did not specify an output file name).
64 * Add the necessary includes to these files, including those specified
67 * Return 0 on success and -1 on failure.
69 static int opencl_open_files(struct opencl_info
*info
)
78 ext
= strrchr(info
->output
, '.');
79 len
= ext
? ext
- info
->output
: strlen(info
->output
);
80 memcpy(name
, info
->output
, len
);
82 info
->host_c
= open_or_croak(info
->output
);
84 len
= ppcg_extract_base_name(name
, info
->input
);
86 strcpy(name
+ len
, "_host.c");
87 info
->host_c
= open_or_croak(name
);
90 memcpy(info
->kernel_c_name
, name
, len
);
91 strcpy(info
->kernel_c_name
+ len
, "_kernel.cl");
92 info
->kernel_c
= open_or_croak(info
->kernel_c_name
);
94 if (!info
->host_c
|| !info
->kernel_c
)
97 fprintf(info
->host_c
, "#include <assert.h>\n");
98 fprintf(info
->host_c
, "#include <stdio.h>\n");
99 fprintf(info
->host_c
, "#include \"ocl_utilities.h\"\n");
100 if (info
->options
->opencl_embed_kernel_code
) {
101 fprintf(info
->host_c
, "#include \"%s\"\n\n",
102 info
->kernel_c_name
);
105 for (i
= 0; i
< info
->options
->opencl_n_include_file
; ++i
) {
106 info
->kprinter
= isl_printer_print_str(info
->kprinter
,
108 info
->kprinter
= isl_printer_print_str(info
->kprinter
,
109 info
->options
->opencl_include_files
[i
]);
110 info
->kprinter
= isl_printer_print_str(info
->kprinter
, ">\n");
116 /* Write text to a file and escape some special characters that would break a
119 static void opencl_print_escaped(const char *str
, const char *end
, FILE *file
)
121 const char *prev
= str
;
123 while ((str
= strpbrk(prev
, "\"\\")) && str
< end
) {
124 fwrite(prev
, 1, str
- prev
, file
);
125 fprintf(file
, "\\%c", *str
);
130 fwrite(prev
, 1, end
- prev
, file
);
133 /* Write text to a file as a C string literal.
135 * This function also prints any characters after the last newline, although
136 * normally the input string should end with a newline.
138 static void opencl_print_as_c_string(const char *str
, FILE *file
)
140 const char *prev
= str
;
142 while ((str
= strchr(prev
, '\n'))) {
143 fprintf(file
, "\n\"");
144 opencl_print_escaped(prev
, str
, file
);
145 fprintf(file
, "\\n\"");
151 fprintf(file
, "\n\"");
152 opencl_print_escaped(prev
, prev
+ strlen(prev
), file
);
157 /* Write the code that we have accumulated in the kernel isl_printer to the
158 * kernel.cl file. If the opencl_embed_kernel_code option has been set, print
159 * the code as a C string literal. Start that string literal with an empty
160 * line, such that line numbers reported by the OpenCL C compiler match those
161 * of the kernel file.
163 * Return 0 on success and -1 on failure.
165 static int opencl_write_kernel_file(struct opencl_info
*opencl
)
167 char *raw
= isl_printer_get_str(opencl
->kprinter
);
172 if (opencl
->options
->opencl_embed_kernel_code
) {
173 fprintf(opencl
->kernel_c
,
174 "static const char kernel_code[] = \"\\n\"");
175 opencl_print_as_c_string(raw
, opencl
->kernel_c
);
176 fprintf(opencl
->kernel_c
, ";\n");
178 fprintf(opencl
->kernel_c
, "%s", raw
);
185 /* Close all output files. Write the kernel contents to the kernel file before
188 * Return 0 on success and -1 on failure.
190 static int opencl_close_files(struct opencl_info
*info
)
194 if (info
->kernel_c
) {
195 r
= opencl_write_kernel_file(info
);
196 fclose(info
->kernel_c
);
199 fclose(info
->host_c
);
204 static __isl_give isl_printer
*opencl_print_host_macros(
205 __isl_take isl_printer
*p
)
208 "#define openclCheckReturn(ret) \\\n"
209 " if (ret != CL_SUCCESS) {\\\n"
210 " fprintf(stderr, \"OpenCL error: %s\\n\", "
211 "opencl_error_string(ret)); \\\n"
212 " fflush(stderr); \\\n"
213 " assert(ret == CL_SUCCESS);\\\n }\n";
215 p
= isl_printer_start_line(p
);
216 p
= isl_printer_print_str(p
, macros
);
217 p
= isl_printer_end_line(p
);
222 static __isl_give isl_printer
*opencl_declare_device_arrays(
223 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
227 for (i
= 0; i
< prog
->n_array
; ++i
) {
228 if (!gpu_array_requires_device_allocation(&prog
->array
[i
]))
230 p
= isl_printer_start_line(p
);
231 p
= isl_printer_print_str(p
, "cl_mem dev_");
232 p
= isl_printer_print_str(p
, prog
->array
[i
].name
);
233 p
= isl_printer_print_str(p
, ";");
234 p
= isl_printer_end_line(p
);
236 p
= isl_printer_start_line(p
);
237 p
= isl_printer_end_line(p
);
241 /* Given an array, check whether its positive size guard expression is
244 static int is_array_positive_size_guard_trivial(struct gpu_array_info
*array
)
249 guard
= gpu_array_positive_size_guard(array
);
250 is_trivial
= isl_set_plain_is_universe(guard
);
255 /* Allocate a device array for "array'.
257 * Emit a max-expression to ensure the device array can contain at least one
258 * element if the array's positive size guard expression is not trivial.
260 static __isl_give isl_printer
*allocate_device_array(__isl_take isl_printer
*p
,
261 struct gpu_array_info
*array
)
263 int need_lower_bound
;
265 need_lower_bound
= !is_array_positive_size_guard_trivial(array
);
266 if (need_lower_bound
)
267 p
= isl_ast_op_type_print_macro(isl_ast_op_max
, p
);
269 p
= ppcg_start_block(p
);
271 p
= isl_printer_start_line(p
);
272 p
= isl_printer_print_str(p
, "dev_");
273 p
= isl_printer_print_str(p
, array
->name
);
274 p
= isl_printer_print_str(p
, " = clCreateBuffer(context, ");
275 p
= isl_printer_print_str(p
, "CL_MEM_READ_WRITE, ");
277 if (need_lower_bound
) {
278 p
= isl_printer_print_str(p
, ppcg_max
);
279 p
= isl_printer_print_str(p
, "(sizeof(");
280 p
= isl_printer_print_str(p
, array
->type
);
281 p
= isl_printer_print_str(p
, "), ");
283 p
= gpu_array_info_print_size(p
, array
);
284 if (need_lower_bound
)
285 p
= isl_printer_print_str(p
, ")");
287 p
= isl_printer_print_str(p
, ", NULL, &err);");
288 p
= isl_printer_end_line(p
);
289 p
= isl_printer_start_line(p
);
290 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
291 p
= isl_printer_end_line(p
);
293 p
= ppcg_end_block(p
);
298 /* Allocate accessed device arrays.
300 static __isl_give isl_printer
*opencl_allocate_device_arrays(
301 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
305 for (i
= 0; i
< prog
->n_array
; ++i
) {
306 struct gpu_array_info
*array
= &prog
->array
[i
];
308 if (!gpu_array_requires_device_allocation(array
))
311 p
= allocate_device_array(p
, array
);
313 p
= isl_printer_start_line(p
);
314 p
= isl_printer_end_line(p
);
318 /* Free the device array corresponding to "array"
320 static __isl_give isl_printer
*release_device_array(__isl_take isl_printer
*p
,
321 struct gpu_array_info
*array
)
323 p
= isl_printer_start_line(p
);
324 p
= isl_printer_print_str(p
, "openclCheckReturn("
325 "clReleaseMemObject(dev_");
326 p
= isl_printer_print_str(p
, array
->name
);
327 p
= isl_printer_print_str(p
, "));");
328 p
= isl_printer_end_line(p
);
333 /* Free the accessed device arrays.
335 static __isl_give isl_printer
*opencl_release_device_arrays(
336 __isl_take isl_printer
*p
, struct gpu_prog
*prog
)
340 for (i
= 0; i
< prog
->n_array
; ++i
) {
341 struct gpu_array_info
*array
= &prog
->array
[i
];
342 if (!gpu_array_requires_device_allocation(array
))
345 p
= release_device_array(p
, array
);
350 /* Create an OpenCL device, context, command queue and build the kernel.
351 * input is the name of the input file provided to ppcg.
353 static __isl_give isl_printer
*opencl_setup(__isl_take isl_printer
*p
,
354 const char *input
, struct opencl_info
*info
)
356 p
= isl_printer_start_line(p
);
357 p
= isl_printer_print_str(p
, "cl_device_id device;");
358 p
= isl_printer_end_line(p
);
359 p
= isl_printer_start_line(p
);
360 p
= isl_printer_print_str(p
, "cl_context context;");
361 p
= isl_printer_end_line(p
);
362 p
= isl_printer_start_line(p
);
363 p
= isl_printer_print_str(p
, "cl_program program;");
364 p
= isl_printer_end_line(p
);
365 p
= isl_printer_start_line(p
);
366 p
= isl_printer_print_str(p
, "cl_command_queue queue;");
367 p
= isl_printer_end_line(p
);
368 p
= isl_printer_start_line(p
);
369 p
= isl_printer_print_str(p
, "cl_int err;");
370 p
= isl_printer_end_line(p
);
371 p
= isl_printer_start_line(p
);
372 p
= isl_printer_print_str(p
, "device = opencl_create_device(");
373 p
= isl_printer_print_int(p
, info
->options
->opencl_use_gpu
);
374 p
= isl_printer_print_str(p
, ");");
375 p
= isl_printer_end_line(p
);
376 p
= isl_printer_start_line(p
);
377 p
= isl_printer_print_str(p
, "context = clCreateContext(NULL, 1, "
378 "&device, NULL, NULL, &err);");
379 p
= isl_printer_end_line(p
);
380 p
= isl_printer_start_line(p
);
381 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
382 p
= isl_printer_end_line(p
);
383 p
= isl_printer_start_line(p
);
384 p
= isl_printer_print_str(p
, "queue = clCreateCommandQueue"
385 "(context, device, 0, &err);");
386 p
= isl_printer_end_line(p
);
387 p
= isl_printer_start_line(p
);
388 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
389 p
= isl_printer_end_line(p
);
391 p
= isl_printer_start_line(p
);
392 p
= isl_printer_print_str(p
, "program = ");
394 if (info
->options
->opencl_embed_kernel_code
) {
395 p
= isl_printer_print_str(p
, "opencl_build_program_from_string("
396 "context, device, kernel_code, "
397 "sizeof(kernel_code), \"");
399 p
= isl_printer_print_str(p
, "opencl_build_program_from_file("
400 "context, device, \"");
401 p
= isl_printer_print_str(p
, info
->kernel_c_name
);
402 p
= isl_printer_print_str(p
, "\", \"");
405 if (info
->options
->opencl_compiler_options
)
406 p
= isl_printer_print_str(p
,
407 info
->options
->opencl_compiler_options
);
409 p
= isl_printer_print_str(p
, "\");");
410 p
= isl_printer_end_line(p
);
411 p
= isl_printer_start_line(p
);
412 p
= isl_printer_end_line(p
);
417 static __isl_give isl_printer
*opencl_release_cl_objects(
418 __isl_take isl_printer
*p
, struct opencl_info
*info
)
420 p
= isl_printer_start_line(p
);
421 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseCommandQueue"
423 p
= isl_printer_end_line(p
);
424 p
= isl_printer_start_line(p
);
425 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseProgram"
427 p
= isl_printer_end_line(p
);
428 p
= isl_printer_start_line(p
);
429 p
= isl_printer_print_str(p
, "openclCheckReturn(clReleaseContext"
431 p
= isl_printer_end_line(p
);
436 /* Print a call to the OpenCL clSetKernelArg() function which sets
437 * the arguments of the kernel. arg_name and arg_index are the name and the
438 * index of the kernel argument. The index of the leftmost argument of
439 * the kernel is 0 whereas the index of the rightmost argument of the kernel
440 * is n - 1, where n is the total number of the kernel arguments.
441 * read_only_scalar is a boolean that indicates whether the argument is a read
444 static __isl_give isl_printer
*opencl_set_kernel_argument(
445 __isl_take isl_printer
*p
, int kernel_id
,
446 const char *arg_name
, int arg_index
, int read_only_scalar
)
448 p
= isl_printer_start_line(p
);
449 p
= isl_printer_print_str(p
,
450 "openclCheckReturn(clSetKernelArg(kernel");
451 p
= isl_printer_print_int(p
, kernel_id
);
452 p
= isl_printer_print_str(p
, ", ");
453 p
= isl_printer_print_int(p
, arg_index
);
454 p
= isl_printer_print_str(p
, ", sizeof(");
456 if (read_only_scalar
) {
457 p
= isl_printer_print_str(p
, arg_name
);
458 p
= isl_printer_print_str(p
, "), &");
460 p
= isl_printer_print_str(p
, "cl_mem), (void *) &dev_");
462 p
= isl_printer_print_str(p
, arg_name
);
463 p
= isl_printer_print_str(p
, "));");
464 p
= isl_printer_end_line(p
);
469 /* Print the block sizes as a list of the sizes in each
472 static __isl_give isl_printer
*opencl_print_block_sizes(
473 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
477 if (kernel
->n_block
> 0)
478 for (i
= 0; i
< kernel
->n_block
; ++i
) {
480 p
= isl_printer_print_str(p
, ", ");
481 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
484 p
= isl_printer_print_str(p
, "1");
489 /* Set the arguments of the OpenCL kernel by printing a call to the OpenCL
490 * clSetKernelArg() function for each kernel argument.
492 static __isl_give isl_printer
*opencl_set_kernel_arguments(
493 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
494 struct ppcg_kernel
*kernel
)
501 for (i
= 0; i
< prog
->n_array
; ++i
) {
504 required
= ppcg_kernel_requires_array_argument(kernel
, i
);
506 return isl_printer_free(p
);
509 ro
= gpu_array_is_read_only_scalar(&prog
->array
[i
]);
510 opencl_set_kernel_argument(p
, kernel
->id
, prog
->array
[i
].name
,
515 space
= isl_union_set_get_space(kernel
->arrays
);
516 nparam
= isl_space_dim(space
, isl_dim_param
);
517 for (i
= 0; i
< nparam
; ++i
) {
520 name
= isl_space_get_dim_name(space
, isl_dim_param
, i
);
521 opencl_set_kernel_argument(p
, kernel
->id
, name
, arg_index
, 1);
524 isl_space_free(space
);
526 n
= isl_space_dim(kernel
->space
, isl_dim_set
);
527 for (i
= 0; i
< n
; ++i
) {
530 name
= isl_space_get_dim_name(kernel
->space
, isl_dim_set
, i
);
531 opencl_set_kernel_argument(p
, kernel
->id
, name
, arg_index
, 1);
538 /* Print the arguments to a kernel declaration or call. If "types" is set,
539 * then print a declaration (including the types of the arguments).
541 * The arguments are printed in the following order
542 * - the arrays accessed by the kernel
544 * - the host loop iterators
546 static __isl_give isl_printer
*opencl_print_kernel_arguments(
547 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
548 struct ppcg_kernel
*kernel
, int types
)
556 for (i
= 0; i
< prog
->n_array
; ++i
) {
559 required
= ppcg_kernel_requires_array_argument(kernel
, i
);
561 return isl_printer_free(p
);
566 p
= isl_printer_print_str(p
, ", ");
569 p
= gpu_array_info_print_declaration_argument(p
,
570 &prog
->array
[i
], "__global");
572 p
= gpu_array_info_print_call_argument(p
,
578 space
= isl_union_set_get_space(kernel
->arrays
);
579 nparam
= isl_space_dim(space
, isl_dim_param
);
580 for (i
= 0; i
< nparam
; ++i
) {
583 name
= isl_space_get_dim_name(space
, isl_dim_param
, i
);
586 p
= isl_printer_print_str(p
, ", ");
588 p
= isl_printer_print_str(p
, "int ");
589 p
= isl_printer_print_str(p
, name
);
593 isl_space_free(space
);
595 n
= isl_space_dim(kernel
->space
, isl_dim_set
);
596 type
= isl_options_get_ast_iterator_type(prog
->ctx
);
597 for (i
= 0; i
< n
; ++i
) {
601 p
= isl_printer_print_str(p
, ", ");
602 name
= isl_space_get_dim_name(kernel
->space
, isl_dim_set
, i
);
604 p
= isl_printer_print_str(p
, type
);
605 p
= isl_printer_print_str(p
, " ");
607 p
= isl_printer_print_str(p
, name
);
615 /* Print the header of the given kernel.
617 static __isl_give isl_printer
*opencl_print_kernel_header(
618 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
619 struct ppcg_kernel
*kernel
)
621 p
= isl_printer_start_line(p
);
622 p
= isl_printer_print_str(p
, "__kernel void kernel");
623 p
= isl_printer_print_int(p
, kernel
->id
);
624 p
= isl_printer_print_str(p
, "(");
625 p
= opencl_print_kernel_arguments(p
, prog
, kernel
, 1);
626 p
= isl_printer_print_str(p
, ")");
627 p
= isl_printer_end_line(p
);
632 /* Print a list of iterators of type "type" with names "ids" to "p".
633 * Each iterator is assigned the corresponding opencl identifier returned
634 * by the function "opencl_id".
635 * Unlike the equivalent function in the CUDA backend which prints iterators
636 * in reverse order to promote coalescing, this function does not print
637 * iterators in reverse order. The OpenCL backend currently does not take
638 * into account any coalescing considerations.
640 static __isl_give isl_printer
*print_iterators(__isl_take isl_printer
*p
,
641 const char *type
, __isl_keep isl_id_list
*ids
, const char *opencl_id
)
645 n
= isl_id_list_n_id(ids
);
648 p
= isl_printer_start_line(p
);
649 p
= isl_printer_print_str(p
, type
);
650 p
= isl_printer_print_str(p
, " ");
651 for (i
= 0; i
< n
; ++i
) {
655 p
= isl_printer_print_str(p
, ", ");
656 id
= isl_id_list_get_id(ids
, i
);
657 p
= isl_printer_print_id(p
, id
);
659 p
= isl_printer_print_str(p
, " = ");
660 p
= isl_printer_print_str(p
, opencl_id
);
661 p
= isl_printer_print_str(p
, "(");
662 p
= isl_printer_print_int(p
, i
);
663 p
= isl_printer_print_str(p
, ")");
665 p
= isl_printer_print_str(p
, ";");
666 p
= isl_printer_end_line(p
);
671 static __isl_give isl_printer
*opencl_print_kernel_iterators(
672 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
674 isl_ctx
*ctx
= isl_ast_node_get_ctx(kernel
->tree
);
677 type
= isl_options_get_ast_iterator_type(ctx
);
679 p
= print_iterators(p
, type
, kernel
->block_ids
, "get_group_id");
680 p
= print_iterators(p
, type
, kernel
->thread_ids
, "get_local_id");
685 static __isl_give isl_printer
*opencl_print_kernel_var(
686 __isl_take isl_printer
*p
, struct ppcg_kernel_var
*var
)
691 p
= isl_printer_start_line(p
);
692 if (var
->type
== ppcg_access_shared
)
693 p
= isl_printer_print_str(p
, "__local ");
694 p
= isl_printer_print_str(p
, var
->array
->type
);
695 p
= isl_printer_print_str(p
, " ");
696 p
= isl_printer_print_str(p
, var
->name
);
697 for (j
= 0; j
< var
->array
->n_index
; ++j
) {
698 p
= isl_printer_print_str(p
, "[");
699 v
= isl_vec_get_element_val(var
->size
, j
);
700 p
= isl_printer_print_val(p
, v
);
701 p
= isl_printer_print_str(p
, "]");
704 p
= isl_printer_print_str(p
, ";");
705 p
= isl_printer_end_line(p
);
710 static __isl_give isl_printer
*opencl_print_kernel_vars(
711 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
715 for (i
= 0; i
< kernel
->n_var
; ++i
)
716 p
= opencl_print_kernel_var(p
, &kernel
->var
[i
]);
721 /* Print a call to barrier() which is a sync statement.
722 * All work-items in a work-group executing the kernel on a processor must
723 * execute the barrier() function before any are allowed to continue execution
724 * beyond the barrier.
725 * The flag CLK_LOCAL_MEM_FENCE makes the barrier function either flush any
726 * variables stored in local memory or queue a memory fence to ensure correct
727 * ordering of memory operations to local memory.
728 * The flag CLK_GLOBAL_MEM_FENCE makes the barrier function queue a memory
729 * fence to ensure correct ordering of memory operations to global memory.
731 static __isl_give isl_printer
*opencl_print_sync(__isl_take isl_printer
*p
,
732 struct ppcg_kernel_stmt
*stmt
)
734 p
= isl_printer_start_line(p
);
735 p
= isl_printer_print_str(p
,
736 "barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);");
737 p
= isl_printer_end_line(p
);
742 /* Data structure containing function names for which the calls
743 * should be changed from
749 * opencl_name((type) (arg))
751 static struct ppcg_opencl_fn
{
753 const char *opencl_name
;
756 { "expf", "exp", "float" },
757 { "powf", "pow", "float" },
758 { "sqrtf", "sqrt", "float" },
761 #define ARRAY_SIZE(array) (sizeof(array)/sizeof(*array))
763 /* If the name of function called by "expr" matches any of those
764 * in ppcg_opencl_fn, then replace the call by a cast to the corresponding
765 * type in ppcg_opencl_fn and a call to corresponding OpenCL function.
767 static __isl_give pet_expr
*map_opencl_call(__isl_take pet_expr
*expr
,
773 name
= pet_expr_call_get_name(expr
);
774 for (i
= 0; i
< ARRAY_SIZE(opencl_fn
); ++i
) {
777 if (strcmp(name
, opencl_fn
[i
].name
))
779 expr
= pet_expr_call_set_name(expr
, opencl_fn
[i
].opencl_name
);
780 arg
= pet_expr_get_arg(expr
, 0);
781 arg
= pet_expr_new_cast(opencl_fn
[i
].type
, arg
);
782 expr
= pet_expr_set_arg(expr
, 0, arg
);
787 /* Print the body of a statement from the input program,
788 * for use in OpenCL code.
790 * Before calling ppcg_kernel_print_domain to print the actual statement body,
791 * we first modify this body to take into account that the output code
792 * is OpenCL code. In particular, if the statement calls any function
793 * with a "f" suffix, then it needs to be replaced by a call to
794 * the corresponding function without suffix after casting the argument
797 static __isl_give isl_printer
*print_opencl_kernel_domain(
798 __isl_take isl_printer
*p
, struct ppcg_kernel_stmt
*stmt
)
803 ps
= stmt
->u
.d
.stmt
->stmt
;
804 tree
= pet_tree_copy(ps
->body
);
805 ps
->body
= pet_tree_map_call_expr(ps
->body
, &map_opencl_call
, NULL
);
806 p
= ppcg_kernel_print_domain(p
, stmt
);
807 pet_tree_free(ps
->body
);
813 /* This function is called for each user statement in the AST,
814 * i.e., for each kernel body statement, copy statement or sync statement.
816 static __isl_give isl_printer
*opencl_print_kernel_stmt(
817 __isl_take isl_printer
*p
,
818 __isl_take isl_ast_print_options
*print_options
,
819 __isl_keep isl_ast_node
*node
, void *user
)
822 struct ppcg_kernel_stmt
*stmt
;
824 id
= isl_ast_node_get_annotation(node
);
825 stmt
= isl_id_get_user(id
);
828 isl_ast_print_options_free(print_options
);
830 switch (stmt
->type
) {
831 case ppcg_kernel_copy
:
832 return ppcg_kernel_print_copy(p
, stmt
);
833 case ppcg_kernel_sync
:
834 return opencl_print_sync(p
, stmt
);
835 case ppcg_kernel_domain
:
836 return print_opencl_kernel_domain(p
, stmt
);
842 /* Return true if there is a double array in prog->array or
843 * if any of the types in prog->scop involve any doubles.
844 * To check the latter condition, we simply search for the string "double"
845 * in the type definitions, which may result in false positives.
847 static __isl_give
int any_double_elements(struct gpu_prog
*prog
)
851 for (i
= 0; i
< prog
->n_array
; ++i
)
852 if (strcmp(prog
->array
[i
].type
, "double") == 0)
855 for (i
= 0; i
< prog
->scop
->pet
->n_type
; ++i
) {
856 struct pet_type
*type
= prog
->scop
->pet
->types
[i
];
858 if (strstr(type
->definition
, "double"))
865 /* Prints a #pragma to enable support for double floating-point
866 * precision. OpenCL 1.0 adds support for double precision floating-point as
867 * an optional extension. An application that wants to use double will need to
868 * include the #pragma OPENCL EXTENSION cl_khr_fp64 : enable directive before
869 * any double precision data type is declared in the kernel code.
871 static __isl_give isl_printer
*opencl_enable_double_support(
872 __isl_take isl_printer
*p
)
874 p
= isl_printer_start_line(p
);
875 p
= isl_printer_print_str(p
, "#pragma OPENCL EXTENSION cl_khr_fp64 :"
877 p
= isl_printer_end_line(p
);
878 p
= isl_printer_start_line(p
);
879 p
= isl_printer_end_line(p
);
884 static __isl_give isl_printer
*opencl_print_kernel(struct gpu_prog
*prog
,
885 struct ppcg_kernel
*kernel
, __isl_take isl_printer
*p
)
887 isl_ctx
*ctx
= isl_ast_node_get_ctx(kernel
->tree
);
888 isl_ast_print_options
*print_options
;
890 print_options
= isl_ast_print_options_alloc(ctx
);
891 print_options
= isl_ast_print_options_set_print_user(print_options
,
892 &opencl_print_kernel_stmt
, NULL
);
894 p
= isl_printer_set_output_format(p
, ISL_FORMAT_C
);
895 p
= opencl_print_kernel_header(p
, prog
, kernel
);
896 p
= isl_printer_print_str(p
, "{");
897 p
= isl_printer_end_line(p
);
898 p
= isl_printer_indent(p
, 4);
899 p
= opencl_print_kernel_iterators(p
, kernel
);
900 p
= opencl_print_kernel_vars(p
, kernel
);
901 p
= isl_printer_end_line(p
);
902 p
= ppcg_set_macro_names(p
);
903 p
= isl_ast_op_type_print_macro(isl_ast_op_fdiv_q
, p
);
904 p
= ppcg_print_macros(p
, kernel
->tree
);
905 p
= isl_ast_node_print(kernel
->tree
, p
, print_options
);
906 p
= isl_printer_indent(p
, -4);
907 p
= isl_printer_start_line(p
);
908 p
= isl_printer_print_str(p
, "}");
909 p
= isl_printer_end_line(p
);
914 struct print_host_user_data_opencl
{
915 struct opencl_info
*opencl
;
916 struct gpu_prog
*prog
;
919 /* This function prints the i'th block size multiplied by the i'th grid size,
920 * where i (a parameter to this function) is one of the possible dimensions of
921 * grid sizes and block sizes.
922 * If the dimension of block sizes is not equal to the dimension of grid sizes
923 * the output is calculated as follows:
926 * block_sizes[dim1] is the list of blocks sizes and it contains dim1 elements.
927 * grid_sizes[dim2] is the list of grid sizes and it contains dim2 elements.
930 * If (i > dim2) then the output is block_sizes[i]
931 * If (i > dim1) then the output is grid_sizes[i]
933 static __isl_give isl_printer
*opencl_print_total_number_of_work_items_for_dim(
934 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
, int i
)
936 int grid_dim
, block_dim
;
937 isl_ast_expr
*grid_size_expr
;
938 isl_ast_expr
*bound_grid
;
940 grid_dim
= isl_multi_pw_aff_dim(kernel
->grid_size
, isl_dim_set
);
941 block_dim
= kernel
->n_block
;
943 if (i
< min(grid_dim
, block_dim
)) {
944 grid_size_expr
= kernel
->grid_size_expr
;
945 bound_grid
= isl_ast_expr_get_op_arg(grid_size_expr
, 1 + i
);
946 p
= isl_printer_print_str(p
, "(");
947 p
= isl_printer_print_ast_expr(p
, bound_grid
);
948 p
= isl_printer_print_str(p
, ") * ");
949 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
950 isl_ast_expr_free(bound_grid
);
951 } else if (i
>= grid_dim
) {
952 p
= isl_printer_print_int(p
, kernel
->block_dim
[i
]);
954 grid_size_expr
= kernel
->grid_size_expr
;
955 bound_grid
= isl_ast_expr_get_op_arg(grid_size_expr
, 1 + i
);
956 p
= isl_printer_print_ast_expr(p
, bound_grid
);
957 isl_ast_expr_free(bound_grid
);
963 /* Print a list that represents the total number of work items. The list is
964 * constructed by performing an element-wise multiplication of the block sizes
965 * and the grid sizes. To explain how the list is constructed, suppose that:
966 * block_sizes[dim1] is the list of blocks sizes and it contains dim1 elements.
967 * grid_sizes[dim2] is the list of grid sizes and it contains dim2 elements.
969 * The output of this function is constructed as follows:
970 * If (dim1 > dim2) then the output is the following list:
971 * grid_sizes[0]*block_sizes[0], ..., grid_sizes[dim2-1]*block_sizes[dim2-1],
972 * block_sizes[dim2], ..., block_sizes[dim1-2], block_sizes[dim1-1].
974 * If (dim2 > dim1) then the output is the following list:
975 * grid_sizes[0]*block_sizes[0], ..., grid_sizes[dim1-1] * block_sizes[dim1-1],
976 * grid_sizes[dim1], grid_sizes[dim2-2], grid_sizes[dim2-1].
978 * To calculate the total number of work items out of the list constructed by
979 * this function, the user should multiply the elements of the list.
981 static __isl_give isl_printer
*opencl_print_total_number_of_work_items_as_list(
982 __isl_take isl_printer
*p
, struct ppcg_kernel
*kernel
)
985 int grid_dim
, block_dim
;
987 grid_dim
= isl_multi_pw_aff_dim(kernel
->grid_size
, isl_dim_set
);
988 block_dim
= kernel
->n_block
;
990 if ((grid_dim
<= 0) || (block_dim
<= 0)) {
991 p
= isl_printer_print_str(p
, "1");
995 for (i
= 0; i
<= max(grid_dim
, block_dim
) - 1; i
++) {
997 p
= isl_printer_print_str(p
, ", ");
999 p
= opencl_print_total_number_of_work_items_for_dim(p
,
1006 /* Copy "array" from the host to the device (to_host = 0) or
1007 * back from the device to the host (to_host = 1).
1009 static __isl_give isl_printer
*copy_array(__isl_take isl_printer
*p
,
1010 struct gpu_array_info
*array
, int to_host
)
1012 p
= isl_printer_start_line(p
);
1013 p
= isl_printer_print_str(p
, "openclCheckReturn(");
1015 p
= isl_printer_print_str(p
, "clEnqueueReadBuffer");
1017 p
= isl_printer_print_str(p
, "clEnqueueWriteBuffer");
1018 p
= isl_printer_print_str(p
, "(queue, dev_");
1019 p
= isl_printer_print_str(p
, array
->name
);
1020 p
= isl_printer_print_str(p
, ", CL_TRUE, 0, ");
1021 p
= gpu_array_info_print_size(p
, array
);
1023 if (gpu_array_is_scalar(array
))
1024 p
= isl_printer_print_str(p
, ", &");
1026 p
= isl_printer_print_str(p
, ", ");
1027 p
= isl_printer_print_str(p
, array
->name
);
1028 p
= isl_printer_print_str(p
, ", 0, NULL, NULL));");
1029 p
= isl_printer_end_line(p
);
1034 /* Print code for initializing the device for execution of the transformed
1035 * code. This includes declaring locally defined variables as well as
1036 * declaring and allocating the required copies of arrays on the device.
1038 static __isl_give isl_printer
*init_device(__isl_take isl_printer
*p
,
1039 struct gpu_prog
*prog
, struct opencl_info
*opencl
)
1041 p
= opencl_print_host_macros(p
);
1043 p
= gpu_print_local_declarations(p
, prog
);
1044 p
= opencl_declare_device_arrays(p
, prog
);
1045 p
= opencl_setup(p
, opencl
->input
, opencl
);
1046 p
= opencl_allocate_device_arrays(p
, prog
);
1051 /* Print code for clearing the device after execution of the transformed code.
1052 * In particular, free the memory that was allocated on the device.
1054 static __isl_give isl_printer
*clear_device(__isl_take isl_printer
*p
,
1055 struct gpu_prog
*prog
, struct opencl_info
*opencl
)
1057 p
= opencl_release_device_arrays(p
, prog
);
1058 p
= opencl_release_cl_objects(p
, opencl
);
1063 /* Print a statement for copying an array to or from the device,
1064 * or for initializing or clearing the device.
1065 * The statement identifier of a copying node is called
1066 * "to_device_<array name>" or "from_device_<array name>" and
1067 * its user pointer points to the gpu_array_info of the array
1068 * that needs to be copied.
1069 * The node for initializing the device is called "init_device".
1070 * The node for clearing the device is called "clear_device".
1072 * Extract the array (if any) from the identifier and call
1073 * init_device, clear_device, copy_array_to_device or copy_array_from_device.
1075 static __isl_give isl_printer
*print_device_node(__isl_take isl_printer
*p
,
1076 __isl_keep isl_ast_node
*node
, struct gpu_prog
*prog
,
1077 struct opencl_info
*opencl
)
1079 isl_ast_expr
*expr
, *arg
;
1082 struct gpu_array_info
*array
;
1084 expr
= isl_ast_node_user_get_expr(node
);
1085 arg
= isl_ast_expr_get_op_arg(expr
, 0);
1086 id
= isl_ast_expr_get_id(arg
);
1087 name
= isl_id_get_name(id
);
1088 array
= isl_id_get_user(id
);
1090 isl_ast_expr_free(arg
);
1091 isl_ast_expr_free(expr
);
1094 return isl_printer_free(p
);
1095 if (!strcmp(name
, "init_device"))
1096 return init_device(p
, prog
, opencl
);
1097 if (!strcmp(name
, "clear_device"))
1098 return clear_device(p
, prog
, opencl
);
1100 return isl_printer_free(p
);
1102 if (!prefixcmp(name
, "to_device"))
1103 return copy_array(p
, array
, 0);
1105 return copy_array(p
, array
, 1);
1108 /* Print the user statement of the host code to "p".
1110 * The host code may contain original user statements, kernel launches,
1111 * statements that copy data to/from the device and statements
1112 * the initialize or clear the device.
1113 * The original user statements and the kernel launches have
1114 * an associated annotation, while the other statements do not.
1115 * The latter are handled by print_device_node.
1116 * The annotation on the user statements is called "user".
1118 * In case of a kernel launch, print a block of statements that
1119 * defines the grid and the work group and then launches the kernel.
1121 * A grid is composed of many work groups (blocks), each work group holds
1122 * many work-items (threads).
1124 * global_work_size[kernel->n_block] represents the total number of work
1125 * items. It points to an array of kernel->n_block unsigned
1126 * values that describe the total number of work-items that will execute
1127 * the kernel. The total number of work-items is computed as:
1128 * global_work_size[0] *...* global_work_size[kernel->n_block - 1].
1130 * The size of each work group (i.e. the number of work-items in each work
1131 * group) is described using block_size[kernel->n_block]. The total
1132 * number of work-items in a block (work-group) is computed as:
1133 * block_size[0] *... * block_size[kernel->n_block - 1].
1135 * For more information check:
1136 * http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html
1138 static __isl_give isl_printer
*opencl_print_host_user(
1139 __isl_take isl_printer
*p
,
1140 __isl_take isl_ast_print_options
*print_options
,
1141 __isl_keep isl_ast_node
*node
, void *user
)
1145 struct ppcg_kernel
*kernel
;
1146 struct ppcg_kernel_stmt
*stmt
;
1147 struct print_host_user_data_opencl
*data
;
1149 isl_ast_print_options_free(print_options
);
1151 data
= (struct print_host_user_data_opencl
*) user
;
1153 id
= isl_ast_node_get_annotation(node
);
1155 return print_device_node(p
, node
, data
->prog
, data
->opencl
);
1157 is_user
= !strcmp(isl_id_get_name(id
), "user");
1158 kernel
= is_user
? NULL
: isl_id_get_user(id
);
1159 stmt
= is_user
? isl_id_get_user(id
) : NULL
;
1163 return ppcg_kernel_print_domain(p
, stmt
);
1165 p
= isl_printer_start_line(p
);
1166 p
= isl_printer_print_str(p
, "{");
1167 p
= isl_printer_end_line(p
);
1168 p
= isl_printer_indent(p
, 2);
1170 p
= isl_printer_start_line(p
);
1171 p
= isl_printer_print_str(p
, "size_t global_work_size[");
1173 if (kernel
->n_block
> 0)
1174 p
= isl_printer_print_int(p
, kernel
->n_block
);
1176 p
= isl_printer_print_int(p
, 1);
1178 p
= isl_printer_print_str(p
, "] = {");
1179 p
= opencl_print_total_number_of_work_items_as_list(p
, kernel
);
1180 p
= isl_printer_print_str(p
, "};");
1181 p
= isl_printer_end_line(p
);
1183 p
= isl_printer_start_line(p
);
1184 p
= isl_printer_print_str(p
, "size_t block_size[");
1186 if (kernel
->n_block
> 0)
1187 p
= isl_printer_print_int(p
, kernel
->n_block
);
1189 p
= isl_printer_print_int(p
, 1);
1191 p
= isl_printer_print_str(p
, "] = {");
1192 p
= opencl_print_block_sizes(p
, kernel
);
1193 p
= isl_printer_print_str(p
, "};");
1194 p
= isl_printer_end_line(p
);
1196 p
= isl_printer_start_line(p
);
1197 p
= isl_printer_print_str(p
, "cl_kernel kernel");
1198 p
= isl_printer_print_int(p
, kernel
->id
);
1199 p
= isl_printer_print_str(p
, " = clCreateKernel(program, \"kernel");
1200 p
= isl_printer_print_int(p
, kernel
->id
);
1201 p
= isl_printer_print_str(p
, "\", &err);");
1202 p
= isl_printer_end_line(p
);
1203 p
= isl_printer_start_line(p
);
1204 p
= isl_printer_print_str(p
, "openclCheckReturn(err);");
1205 p
= isl_printer_end_line(p
);
1207 opencl_set_kernel_arguments(p
, data
->prog
, kernel
);
1209 p
= isl_printer_start_line(p
);
1210 p
= isl_printer_print_str(p
, "openclCheckReturn(clEnqueueNDRangeKernel"
1212 p
= isl_printer_print_int(p
, kernel
->id
);
1213 p
= isl_printer_print_str(p
, ", ");
1214 if (kernel
->n_block
> 0)
1215 p
= isl_printer_print_int(p
, kernel
->n_block
);
1217 p
= isl_printer_print_int(p
, 1);
1219 p
= isl_printer_print_str(p
, ", NULL, global_work_size, "
1221 "0, NULL, NULL));");
1222 p
= isl_printer_end_line(p
);
1223 p
= isl_printer_start_line(p
);
1224 p
= isl_printer_print_str(p
, "openclCheckReturn("
1225 "clReleaseKernel(kernel");
1226 p
= isl_printer_print_int(p
, kernel
->id
);
1227 p
= isl_printer_print_str(p
, "));");
1228 p
= isl_printer_end_line(p
);
1229 p
= isl_printer_start_line(p
);
1230 p
= isl_printer_print_str(p
, "clFinish(queue);");
1231 p
= isl_printer_end_line(p
);
1232 p
= isl_printer_indent(p
, -2);
1233 p
= isl_printer_start_line(p
);
1234 p
= isl_printer_print_str(p
, "}");
1235 p
= isl_printer_end_line(p
);
1237 p
= isl_printer_start_line(p
);
1238 p
= isl_printer_end_line(p
);
1240 data
->opencl
->kprinter
= opencl_print_kernel(data
->prog
, kernel
,
1241 data
->opencl
->kprinter
);
1246 static __isl_give isl_printer
*opencl_print_host_code(
1247 __isl_take isl_printer
*p
, struct gpu_prog
*prog
,
1248 __isl_keep isl_ast_node
*tree
, struct opencl_info
*opencl
)
1250 isl_ast_print_options
*print_options
;
1251 isl_ctx
*ctx
= isl_ast_node_get_ctx(tree
);
1252 struct print_host_user_data_opencl data
= { opencl
, prog
};
1254 print_options
= isl_ast_print_options_alloc(ctx
);
1255 print_options
= isl_ast_print_options_set_print_user(print_options
,
1256 &opencl_print_host_user
, &data
);
1258 p
= ppcg_print_macros(p
, tree
);
1259 p
= isl_ast_node_print(tree
, p
, print_options
);
1264 /* Given a gpu_prog "prog" and the corresponding transformed AST
1265 * "tree", print the entire OpenCL code to "p".
1267 static __isl_give isl_printer
*print_opencl(__isl_take isl_printer
*p
,
1268 struct gpu_prog
*prog
, __isl_keep isl_ast_node
*tree
,
1269 struct gpu_types
*types
, void *user
)
1271 struct opencl_info
*opencl
= user
;
1273 opencl
->kprinter
= isl_printer_set_output_format(opencl
->kprinter
,
1275 if (any_double_elements(prog
))
1276 opencl
->kprinter
= opencl_enable_double_support(
1278 if (opencl
->options
->opencl_print_kernel_types
)
1279 opencl
->kprinter
= gpu_print_types(opencl
->kprinter
, types
,
1282 if (!opencl
->kprinter
)
1283 return isl_printer_free(p
);
1285 p
= opencl_print_host_code(p
, prog
, tree
, opencl
);
1290 /* Transform the code in the file called "input" by replacing
1291 * all scops by corresponding OpenCL code.
1292 * The host code is written to "output" or a name derived from
1293 * "input" if "output" is NULL.
1294 * The kernel code is placed in separate files with names
1295 * derived from "output" or "input".
1297 * We let generate_gpu do all the hard work and then let it call
1298 * us back for printing the AST in print_opencl.
1300 * To prepare for this printing, we first open the output files
1301 * and we close them after generate_gpu has finished.
1303 int generate_opencl(isl_ctx
*ctx
, struct ppcg_options
*options
,
1304 const char *input
, const char *output
)
1306 struct opencl_info opencl
= { options
, input
, output
};
1309 opencl
.kprinter
= isl_printer_to_str(ctx
);
1310 r
= opencl_open_files(&opencl
);
1313 r
= generate_gpu(ctx
, input
, opencl
.host_c
, options
,
1314 &print_opencl
, &opencl
);
1316 if (opencl_close_files(&opencl
) < 0)
1318 isl_printer_free(opencl
.kprinter
);