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;
}
