20 static __isl_give isl_printer *print_cuda_macros(__isl_take isl_printer *p)  in print_cuda_macros()  argument
38 	p = isl_printer_print_str(p, macros);  in print_cuda_macros()
39 	return p;  in print_cuda_macros()
42 /* Print a declaration for the device array corresponding to "array" on "p".
44 static __isl_give isl_printer *declare_device_array(__isl_take isl_printer *p,  in declare_device_array()  argument
49 	p = isl_printer_start_line(p);  in declare_device_array()
50 	p = isl_printer_print_str(p, array->type);  in declare_device_array()
51 	p = isl_printer_print_str(p, " ");  in declare_device_array()
53 		p = isl_printer_print_str(p, "(");  in declare_device_array()
54 	p = isl_printer_print_str(p, "*dev_");  in declare_device_array()
55 	p = isl_printer_print_str(p, array->name);  in declare_device_array()
57 		p = isl_printer_print_str(p, ")");  in declare_device_array()
62 			p = isl_printer_print_str(p, "[");  in declare_device_array()
63 			p = isl_printer_print_ast_expr(p, bound);  in declare_device_array()
64 			p = isl_printer_print_str(p, "]");  in declare_device_array()
68 	p = isl_printer_print_str(p, ";");  in declare_device_array()
69 	p = isl_printer_end_line(p);  in declare_device_array()
71 	return p;  in declare_device_array()
74 static __isl_give isl_printer *declare_device_arrays(__isl_take isl_printer *p,  in declare_device_arrays()  argument
83 		p = declare_device_array(p, &prog->array[i]);  in declare_device_arrays()
85 	p = isl_printer_start_line(p);  in declare_device_arrays()
86 	p = isl_printer_end_line(p);  in declare_device_arrays()
87 	return p;  in declare_device_arrays()
91 	__isl_take isl_printer *p, struct gpu_prog *prog)  in allocate_device_arrays()  argument
100 		p = ppcg_ast_expr_print_macros(array->bound_expr, p);  in allocate_device_arrays()
101 		p = isl_printer_start_line(p);  in allocate_device_arrays()
102 		p = isl_printer_print_str(p,  in allocate_device_arrays()
104 		p = isl_printer_print_str(p, prog->array[i].name);  in allocate_device_arrays()
105 		p = isl_printer_print_str(p, ", ");  in allocate_device_arrays()
106 		p = gpu_array_info_print_size(p, &prog->array[i]);  in allocate_device_arrays()
107 		p = isl_printer_print_str(p, "));");  in allocate_device_arrays()
108 		p = isl_printer_end_line(p);  in allocate_device_arrays()
110 	p = isl_printer_start_line(p);  in allocate_device_arrays()
111 	p = isl_printer_end_line(p);  in allocate_device_arrays()
112 	return p;  in allocate_device_arrays()
115 static __isl_give isl_printer *free_device_arrays(__isl_take isl_printer *p,  in free_device_arrays()  argument
123 		p = isl_printer_start_line(p);  in free_device_arrays()
124 		p = isl_printer_print_str(p, "cudaCheckReturn(cudaFree(dev_");  in free_device_arrays()
125 		p = isl_printer_print_str(p, prog->array[i].name);  in free_device_arrays()
126 		p = isl_printer_print_str(p, "));");  in free_device_arrays()
127 		p = isl_printer_end_line(p);  in free_device_arrays()
130 	return p;  in free_device_arrays()
133 /* Print code to "p" for copying "array" from the host to the device
138 static __isl_give isl_printer *copy_array_to_device(__isl_take isl_printer *p,  in copy_array_to_device()  argument
141 	p = isl_printer_start_line(p);  in copy_array_to_device()
142 	p = isl_printer_print_str(p, "cudaCheckReturn(cudaMemcpy(dev_");  in copy_array_to_device()
143 	p = isl_printer_print_str(p, array->name);  in copy_array_to_device()
144 	p = isl_printer_print_str(p, ", ");  in copy_array_to_device()
147 		p = isl_printer_print_str(p, "&");  in copy_array_to_device()
148 	p = isl_printer_print_str(p, array->name);  in copy_array_to_device()
149 	p = isl_printer_print_str(p, ", ");  in copy_array_to_device()
151 	p = gpu_array_info_print_size(p, array);  in copy_array_to_device()
152 	p = isl_printer_print_str(p, ", cudaMemcpyHostToDevice));");  in copy_array_to_device()
153 	p = isl_printer_end_line(p);  in copy_array_to_device()
155 	return p;  in copy_array_to_device()
158 /* Print code to "p" for copying "array" back from the device to the host
164 	__isl_take isl_printer *p, struct gpu_array_info *array)  in copy_array_from_device()  argument
166 	p = isl_printer_start_line(p);  in copy_array_from_device()
167 	p = isl_printer_print_str(p, "cudaCheckReturn(cudaMemcpy(");  in copy_array_from_device()
169 		p = isl_printer_print_str(p, "&");  in copy_array_from_device()
170 	p = isl_printer_print_str(p, array->name);  in copy_array_from_device()
171 	p = isl_printer_print_str(p, ", dev_");  in copy_array_from_device()
172 	p = isl_printer_print_str(p, array->name);  in copy_array_from_device()
173 	p = isl_printer_print_str(p, ", ");  in copy_array_from_device()
174 	p = gpu_array_info_print_size(p, array);  in copy_array_from_device()
175 	p = isl_printer_print_str(p, ", cudaMemcpyDeviceToHost));");  in copy_array_from_device()
176 	p = isl_printer_end_line(p);  in copy_array_from_device()
178 	return p;  in copy_array_from_device()
181 static __isl_give isl_printer* print_reverse_list(__isl_take isl_printer *p, int len, int *list)  in print_reverse_list()  argument
186 		return p;  in print_reverse_list()
188 	p = isl_printer_print_str(p, "(");  in print_reverse_list()
191 			p = isl_printer_print_str(p, ", ");  in print_reverse_list()
192 		p = isl_printer_print_int(p, list[len - 1 - i]);  in print_reverse_list()
194 	return isl_printer_print_str(p, ")");  in print_reverse_list()
200 static __isl_give isl_printer *print_grid_size(__isl_take isl_printer *p,  in print_grid_size()  argument
208 		return p;  in print_grid_size()
210 	p = isl_printer_print_str(p, "(");  in print_grid_size()
215 		p = isl_printer_print_ast_expr(p, bound);  in print_grid_size()
219 			p = isl_printer_print_str(p, ", ");  in print_grid_size()
222 	p = isl_printer_print_str(p, ")");  in print_grid_size()
224 	return p;  in print_grid_size()
229 static __isl_give isl_printer *print_grid(__isl_take isl_printer *p,  in print_grid()  argument
232 	p = isl_printer_start_line(p);  in print_grid()
233 	p = isl_printer_print_str(p, "dim3 k");  in print_grid()
234 	p = isl_printer_print_int(p, kernel->id);  in print_grid()
235 	p = isl_printer_print_str(p, "_dimGrid");  in print_grid()
236 	p = print_grid_size(p, kernel);  in print_grid()
237 	p = isl_printer_print_str(p, ";");  in print_grid()
238 	p = isl_printer_end_line(p);  in print_grid()
240 	return p;  in print_grid()
251 static __isl_give isl_printer *print_kernel_arguments(__isl_take isl_printer *p,  in print_kernel_arguments()  argument
265 			return isl_printer_free(p);  in print_kernel_arguments()
270 			p = isl_printer_print_str(p, ", ");  in print_kernel_arguments()
273 			p = gpu_array_info_print_declaration_argument(p,  in print_kernel_arguments()
276 			p = gpu_array_info_print_call_argument(p,  in print_kernel_arguments()
290 			p = isl_printer_print_str(p, ", ");  in print_kernel_arguments()
292 			p = isl_printer_print_str(p, "int ");  in print_kernel_arguments()
293 		p = isl_printer_print_str(p, name);  in print_kernel_arguments()
305 			p = isl_printer_print_str(p, ", ");  in print_kernel_arguments()
308 			p = isl_printer_print_str(p, type);  in print_kernel_arguments()
309 			p = isl_printer_print_str(p, " ");  in print_kernel_arguments()
311 		p = isl_printer_print_str(p, name);  in print_kernel_arguments()
316 	return p;  in print_kernel_arguments()
321 static __isl_give isl_printer *print_kernel_header(__isl_take isl_printer *p,  in print_kernel_header()  argument
324 	p = isl_printer_start_line(p);  in print_kernel_header()
325 	p = isl_printer_print_str(p, "__global__ void kernel");  in print_kernel_header()
326 	p = isl_printer_print_int(p, kernel->id);  in print_kernel_header()
327 	p = isl_printer_print_str(p, "(");  in print_kernel_header()
328 	p = print_kernel_arguments(p, prog, kernel, 1);  in print_kernel_header()
329 	p = isl_printer_print_str(p, ")");  in print_kernel_header()
331 	return p;  in print_kernel_header()
340 	isl_printer *p;  in print_kernel_headers()  local
342 	p = isl_printer_to_file(prog->ctx, cuda->kernel_h);  in print_kernel_headers()
343 	p = isl_printer_set_output_format(p, ISL_FORMAT_C);  in print_kernel_headers()
344 	p = print_kernel_header(p, prog, kernel);  in print_kernel_headers()
345 	p = isl_printer_print_str(p, ";");  in print_kernel_headers()
346 	p = isl_printer_end_line(p);  in print_kernel_headers()
347 	isl_printer_free(p);  in print_kernel_headers()
349 	p = isl_printer_to_file(prog->ctx, cuda->kernel_c);  in print_kernel_headers()
350 	p = isl_printer_set_output_format(p, ISL_FORMAT_C);  in print_kernel_headers()
351 	p = print_kernel_header(p, prog, kernel);  in print_kernel_headers()
352 	p = isl_printer_end_line(p);  in print_kernel_headers()
353 	isl_printer_free(p);  in print_kernel_headers()
403 static __isl_give isl_printer *print_kernel_var(__isl_take isl_printer *p,  in print_kernel_var()  argument
408 	p = isl_printer_start_line(p);  in print_kernel_var()
410 		p = isl_printer_print_str(p, "__shared__ ");  in print_kernel_var()
411 	p = isl_printer_print_str(p, var->array->type);  in print_kernel_var()
412 	p = isl_printer_print_str(p, " ");  in print_kernel_var()
413 	p = isl_printer_print_str(p,  var->name);  in print_kernel_var()
417 		p = isl_printer_print_str(p, "[");  in print_kernel_var()
419 		p = isl_printer_print_val(p, v);  in print_kernel_var()
421 		p = isl_printer_print_str(p, "]");  in print_kernel_var()
423 	p = isl_printer_print_str(p, ";");  in print_kernel_var()
424 	p = isl_printer_end_line(p);  in print_kernel_var()
426 	return p;  in print_kernel_var()
429 static __isl_give isl_printer *print_kernel_vars(__isl_take isl_printer *p,  in print_kernel_vars()  argument
435 		p = print_kernel_var(p, &kernel->var[i]);  in print_kernel_vars()
437 	return p;  in print_kernel_vars()
442 static __isl_give isl_printer *print_sync(__isl_take isl_printer *p,  in print_sync()  argument
445 	p = isl_printer_start_line(p);  in print_sync()
446 	p = isl_printer_print_str(p, "__syncthreads();");  in print_sync()
447 	p = isl_printer_end_line(p);  in print_sync()
449 	return p;  in print_sync()
455 static __isl_give isl_printer *print_kernel_stmt(__isl_take isl_printer *p,  in print_kernel_stmt()  argument
470 		return ppcg_kernel_print_copy(p, stmt);  in print_kernel_stmt()
472 		return print_sync(p, stmt);  in print_kernel_stmt()
474 		return ppcg_kernel_print_domain(p, stmt);  in print_kernel_stmt()
477 	return p;  in print_kernel_stmt()
485 	isl_printer *p;  in print_kernel()  local
491 	p = isl_printer_to_file(ctx, cuda->kernel_c);  in print_kernel()
492 	p = isl_printer_set_output_format(p, ISL_FORMAT_C);  in print_kernel()
493 	p = isl_printer_indent(p, 4);  in print_kernel()
495 	p = print_kernel_vars(p, kernel);  in print_kernel()
496 	p = isl_printer_end_line(p);  in print_kernel()
497 	p = ppcg_set_macro_names(p);  in print_kernel()
498 	p = gpu_print_macros(p, kernel->tree);  in print_kernel()
503 	p = isl_ast_node_print(kernel->tree, p, print_options);  in print_kernel()
504 	isl_printer_free(p);  in print_kernel()
513 static __isl_give isl_printer *init_device(__isl_take isl_printer *p,  in init_device()  argument
516 	p = print_cuda_macros(p);  in init_device()
518 	p = gpu_print_local_declarations(p, prog);  in init_device()
519 	p = declare_device_arrays(p, prog);  in init_device()
520 	p = allocate_device_arrays(p, prog);  in init_device()
522 	return p;  in init_device()
528 static __isl_give isl_printer *clear_device(__isl_take isl_printer *p,  in clear_device()  argument
531 	p = free_device_arrays(p, prog);  in clear_device()
533 	return p;  in clear_device()
548 static __isl_give isl_printer *print_device_node(__isl_take isl_printer *p,  in print_device_node()  argument
566 		return isl_printer_free(p);  in print_device_node()
568 		return init_device(p, prog);  in print_device_node()
570 		return clear_device(p, prog);  in print_device_node()
572 		return isl_printer_free(p);  in print_device_node()
575 		return copy_array_to_device(p, array);  in print_device_node()
577 		return copy_array_from_device(p, array);  in print_device_node()
585 /* Print the user statement of the host code to "p".
598 __isl_give isl_printer *print_host_user(__isl_take isl_printer *p,  in print_host_user()  argument
614 		return print_device_node(p, node, data->prog);  in print_host_user()
622 		return ppcg_kernel_print_domain(p, stmt);  in print_host_user()
624 	p = ppcg_start_block(p);  in print_host_user()
626 	p = isl_printer_start_line(p);  in print_host_user()
627 	p = isl_printer_print_str(p, "dim3 k");  in print_host_user()
628 	p = isl_printer_print_int(p, kernel->id);  in print_host_user()
629 	p = isl_printer_print_str(p, "_dimBlock");  in print_host_user()
630 	p = print_reverse_list(p, kernel->n_block, kernel->block_dim);  in print_host_user()
631 	p = isl_printer_print_str(p, ";");  in print_host_user()
632 	p = isl_printer_end_line(p);  in print_host_user()
634 	p = print_grid(p, kernel);  in print_host_user()
636 	p = isl_printer_start_line(p);  in print_host_user()
637 	p = isl_printer_print_str(p, "kernel");  in print_host_user()
638 	p = isl_printer_print_int(p, kernel->id);  in print_host_user()
639 	p = isl_printer_print_str(p, " <<<k");  in print_host_user()
640 	p = isl_printer_print_int(p, kernel->id);  in print_host_user()
641 	p = isl_printer_print_str(p, "_dimGrid, k");  in print_host_user()
642 	p = isl_printer_print_int(p, kernel->id);  in print_host_user()
643 	p = isl_printer_print_str(p, "_dimBlock>>> (");  in print_host_user()
644 	p = print_kernel_arguments(p, data->prog, kernel, 0);  in print_host_user()
645 	p = isl_printer_print_str(p, ");");  in print_host_user()
646 	p = isl_printer_end_line(p);  in print_host_user()
648 	p = isl_printer_start_line(p);  in print_host_user()
649 	p = isl_printer_print_str(p, "cudaCheckKernel();");  in print_host_user()
650 	p = isl_printer_end_line(p);  in print_host_user()
652 	p = ppcg_end_block(p);  in print_host_user()
654 	p = isl_printer_start_line(p);  in print_host_user()
655 	p = isl_printer_end_line(p);  in print_host_user()
661 	return p;  in print_host_user()
664 static __isl_give isl_printer *print_host_code(__isl_take isl_printer *p,  in print_host_code()  argument
676 	p = gpu_print_macros(p, tree);  in print_host_code()
677 	p = isl_ast_node_print(tree, p, print_options);  in print_host_code()
679 	return p;  in print_host_code()
683  * "tree", print the entire CUDA code to "p".
687 static __isl_give isl_printer *print_cuda(__isl_take isl_printer *p,  in print_cuda()  argument
694 	kernel = isl_printer_to_file(isl_printer_get_ctx(p), cuda->kernel_c);  in print_cuda()
700 		return isl_printer_free(p);  in print_cuda()
702 	p = print_host_code(p, prog, tree, cuda);  in print_cuda()
704 	return p;  in print_cuda()