static void print_kernel(struct gpu_prog *prog, struct ppcg_kernel *kernel,
struct cuda_info *cuda)
{
isl_ctx *ctx = isl_ast_node_get_ctx(kernel->tree);
isl_ast_print_options *print_options;
isl_printer *p;
print_kernel_headers(prog, kernel, cuda);
fprintf(cuda->kernel_c, "{\n");
print_kernel_iterators(cuda->kernel_c, kernel);
p = isl_printer_to_file(ctx, cuda->kernel_c);
p = isl_printer_set_output_format(p, ISL_FORMAT_C);
p = isl_printer_indent(p, 4);
p = print_kernel_vars(p, kernel);
p = isl_printer_end_line(p);
p = ppcg_set_macro_names(p);
p = gpu_print_macros(p, kernel->tree);
//p = print_surface_read_to_temp(p, kernel);
print_options = isl_ast_print_options_alloc(ctx);
print_options = isl_ast_print_options_set_print_user(print_options,
&print_kernel_stmt, NULL);
p = isl_ast_node_print(kernel->tree, p, print_options);
isl_printer_free(p);
fprintf(cuda->kernel_c, "}\n");
}
static __isl_give isl_printer *opencl_print_kernel(struct gpu_prog *prog,
struct ppcg_kernel *kernel, __isl_take isl_printer *p)
{
isl_ctx *ctx = isl_ast_node_get_ctx(kernel->tree);
isl_ast_print_options *print_options;
print_options = isl_ast_print_options_alloc(ctx);
print_options = isl_ast_print_options_set_print_user(print_options,
&opencl_print_kernel_stmt, NULL);
p = isl_printer_set_output_format(p, ISL_FORMAT_C);
p = opencl_print_kernel_header(p, prog, kernel);
p = isl_printer_print_str(p, "{");
p = isl_printer_end_line(p);
p = isl_printer_indent(p, 4);
p = opencl_print_kernel_iterators(p, kernel);
p = opencl_print_kernel_vars(p, kernel);
p = isl_printer_end_line(p);
p = isl_ast_op_type_print_macro(isl_ast_op_fdiv_q, p);
p = ppcg_print_macros(p, kernel->tree);
p = isl_ast_node_print(kernel->tree, p, print_options);
p = isl_printer_indent(p, -4);
p = isl_printer_start_line(p);
p = isl_printer_print_str(p, "}");
p = isl_printer_end_line(p);
return p;
}
/* Create a graft for "node" with no guards and no enforced conditions.
*/
__isl_give isl_ast_graft *isl_ast_graft_alloc(
__isl_take isl_ast_node *node, __isl_keep isl_ast_build *build)
{
isl_ctx *ctx;
isl_space *space;
isl_ast_graft *graft;
if (!node)
return NULL;
ctx = isl_ast_node_get_ctx(node);
graft = isl_calloc_type(ctx, isl_ast_graft);
if (!graft)
goto error;
space = isl_ast_build_get_space(build, 1);
graft->ref = 1;
graft->node = node;
graft->guard = isl_set_universe(isl_space_copy(space));
graft->enforced = isl_basic_set_universe(space);
if (!graft->guard || !graft->enforced)
return isl_ast_graft_free(graft);
return graft;
error:
isl_ast_node_free(node);
return NULL;
}
static void print_kernel_iterators(FILE *out, struct ppcg_kernel *kernel)
{
isl_ctx *ctx = isl_ast_node_get_ctx(kernel->tree);
const char *type;
const char *block_dims[] = { "blockIdx.x", "blockIdx.y" };
const char *thread_dims[] = { "threadIdx.x", "threadIdx.y",
"threadIdx.z" };
type = isl_options_get_ast_iterator_type(ctx);
print_iterators(out, type, kernel->block_ids, block_dims);
print_iterators(out, type, kernel->thread_ids, thread_dims);
}
static __isl_give isl_printer *opencl_print_kernel_iterators(
__isl_take isl_printer *p, struct ppcg_kernel *kernel)
{
isl_ctx *ctx = isl_ast_node_get_ctx(kernel->tree);
const char *type;
type = isl_options_get_ast_iterator_type(ctx);
p = print_iterators(p, type, kernel->block_ids, "get_group_id");
p = print_iterators(p, type, kernel->thread_ids, "get_local_id");
return p;
}
static __isl_give isl_printer *opencl_print_host_code(
__isl_take isl_printer *p, struct gpu_prog *prog,
__isl_keep isl_ast_node *tree, struct opencl_info *opencl)
{
isl_ast_print_options *print_options;
isl_ctx *ctx = isl_ast_node_get_ctx(tree);
struct print_host_user_data_opencl data = { opencl, prog };
print_options = isl_ast_print_options_alloc(ctx);
print_options = isl_ast_print_options_set_print_user(print_options,
&opencl_print_host_user, &data);
p = ppcg_print_macros(p, tree);
p = isl_ast_node_print(tree, p, print_options);
return p;
}
static __isl_give isl_printer *print_host_code(__isl_take isl_printer *p,
struct gpu_prog *prog, __isl_keep isl_ast_node *tree,
struct cuda_info *cuda)
{
isl_ast_print_options *print_options;
isl_ctx *ctx = isl_ast_node_get_ctx(tree);
struct print_host_user_data data = { cuda, prog };
print_options = isl_ast_print_options_alloc(ctx);
print_options = isl_ast_print_options_set_print_user(print_options,
&print_host_user, &data);
p = gpu_print_macros(p, tree);
p = isl_ast_node_print(tree, p, print_options);
printf("\n freeing done 1");
fflush(stdout);
return p;
}
