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; }
/* Print the arguments to a kernel declaration or call. If "types" is set, * then print a declaration (including the types of the arguments). * * The arguments are printed in the following order * - the arrays accessed by the kernel * - the parameters * - the host loop iterators */ static __isl_give isl_printer *print_kernel_arguments(__isl_take isl_printer *p, struct gpu_prog *prog, struct ppcg_kernel *kernel, int types) { int i, n; int first = 1; unsigned nparam; isl_space *space; const char *type; for (i = 0; i < prog->n_array; ++i) { int required; required = ppcg_kernel_requires_array_argument(kernel, i); if (required < 0) return isl_printer_free(p); if (!required) continue; if(!print_device_arrays_or_not(&prog->array[i])) continue; if (!first) p = isl_printer_print_str(p, ", "); if (types) p = gpu_array_info_print_declaration_argument(p, &prog->array[i], NULL); else p = gpu_array_info_print_call_argument(p, &prog->array[i]); first = 0; } space = isl_union_set_get_space(kernel->arrays); nparam = isl_space_dim(space, isl_dim_param); for (i = 0; i < nparam; ++i) { const char *name; name = isl_space_get_dim_name(space, isl_dim_param, i); if (!first) p = isl_printer_print_str(p, ", "); if (types) p = isl_printer_print_str(p, "int "); p = isl_printer_print_str(p, name); first = 0; } isl_space_free(space); n = isl_space_dim(kernel->space, isl_dim_set); type = isl_options_get_ast_iterator_type(prog->ctx); for (i = 0; i < n; ++i) { const char *name; if (!first) p = isl_printer_print_str(p, ", "); name = isl_space_get_dim_name(kernel->space, isl_dim_set, i); if (types) { p = isl_printer_print_str(p, type); p = isl_printer_print_str(p, " "); } p = isl_printer_print_str(p, name); first = 0; } return p; }