示例#1
0
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;
}
示例#2
0
static __isl_give isl_printer *allocate_device_arrays(
	__isl_take isl_printer *p, struct gpu_prog *prog)
{
	int i;

	for (i = 0; i < prog->n_array; ++i) {
		struct gpu_array_info *array = &prog->array[i];

		if (!gpu_array_requires_device_allocation(&prog->array[i]))
			continue;
		p = ppcg_ast_expr_print_macros(array->bound_expr, p);
	
		if(print_device_arrays_or_not(&prog->array[i]))
		{
			p = isl_printer_start_line(p);
			p = isl_printer_print_str(p,
				"cudaCheckReturn(cudaMalloc((void **) &dev_");
			p = isl_printer_print_str(p, prog->array[i].name);
			p = isl_printer_print_str(p, ", ");
			p = gpu_array_info_print_size(p, &prog->array[i]);
			p = isl_printer_print_str(p, "));");
			p = isl_printer_end_line(p);
		}
	}
	p = isl_printer_start_line(p);
	p = isl_printer_end_line(p);
	return p;
}
示例#3
0
int main(int argc, char **argv)
{
	struct isl_ctx *ctx;
	struct isl_map *map;
	struct isl_options *options;
	isl_printer *p;
	int exact;

	options = isl_options_new_with_defaults();
	assert(options);
	argc = isl_options_parse(options, argc, argv, ISL_ARG_ALL);

	ctx = isl_ctx_alloc_with_options(&isl_options_args, options);

	p = isl_printer_to_file(ctx, stdout);

	map = isl_map_read_from_file(ctx, stdin);
	map = isl_map_transitive_closure(map, &exact);
	if (!exact)
		p = isl_printer_print_str(p, "# NOT exact\n");
	p = isl_printer_print_map(p, map);
	p = isl_printer_end_line(p);
	map = isl_map_compute_divs(map);
	map = isl_map_coalesce(map);
	p = isl_printer_print_str(p, "# coalesced\n");
	p = isl_printer_print_map(p, map);
	p = isl_printer_end_line(p);
	isl_map_free(map);

	isl_printer_free(p);

	isl_ctx_free(ctx);

	return 0;
}
示例#4
0
/* Allocate a device array for "array'.
 *
 * Emit a max-expression to ensure the device array can contain at least one
 * element if the array's positive size guard expression is not trivial.
 */
static __isl_give isl_printer *allocate_device_array(__isl_take isl_printer *p,
	struct gpu_array_info *array)
{
	int need_lower_bound;

	p = ppcg_start_block(p);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "dev_");
	p = isl_printer_print_str(p, array->name);
	p = isl_printer_print_str(p, " = clCreateBuffer(context, ");
	p = isl_printer_print_str(p, "CL_MEM_READ_WRITE, ");

	need_lower_bound = !is_array_positive_size_guard_trivial(array);
	if (need_lower_bound) {
		p = isl_printer_print_str(p, "max(sizeof(");
		p = isl_printer_print_str(p, array->type);
		p = isl_printer_print_str(p, "), ");
	}
	p = gpu_array_info_print_size(p, array);
	if (need_lower_bound)
		p = isl_printer_print_str(p, ")");

	p = isl_printer_print_str(p, ", NULL, &err);");
	p = isl_printer_end_line(p);
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn(err);");
	p = isl_printer_end_line(p);

	p = ppcg_end_block(p);

	return p;
}
示例#5
0
/* Print the header of the given kernel to both gen->cuda.kernel_h
 * and gen->cuda.kernel_c.
 */
static void print_kernel_headers(struct gpu_prog *prog,
	struct ppcg_kernel *kernel, struct cuda_info *cuda)
{
	isl_printer *p;

	p = isl_printer_to_file(prog->ctx, cuda->kernel_h);
	p = isl_printer_set_output_format(p, ISL_FORMAT_C);

	p = isl_printer_start_line(p);
	p = isl_printer_end_line(p);

	p = print_kernel_header(p, prog, kernel);
	p = isl_printer_print_str(p, ";");
	p = isl_printer_end_line(p);
	isl_printer_free(p);

	p = isl_printer_to_file(prog->ctx, cuda->kernel_c);
	p = isl_printer_set_output_format(p, ISL_FORMAT_C);

	p = print_kernel_header(p, prog, kernel);
	p = isl_printer_end_line(p);
	isl_printer_free(p);

	//isl_printer *p;

}
示例#6
0
/* Prints a #pragma to enable support for double floating-point
 * precision.  OpenCL 1.0 adds support for double precision floating-point as
 * an optional extension. An application that wants to use double will need to
 * include the #pragma OPENCL EXTENSION cl_khr_fp64 : enable directive before
 * any double precision data type is declared in the kernel code.
 */
static __isl_give isl_printer *opencl_enable_double_support(
	__isl_take isl_printer *p)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "#pragma OPENCL EXTENSION cl_khr_fp64 :"
		" enable");
	p = isl_printer_end_line(p);
	p = isl_printer_start_line(p);
	p = isl_printer_end_line(p);

	return p;
}
int main(int argc, char **argv)
{
	struct isl_ctx *ctx = isl_ctx_alloc();
	struct isl_basic_set *bset;
	struct isl_vec *obj;
	struct isl_vec *sol;
	isl_int opt;
	unsigned dim;
	enum isl_lp_result res;
	isl_printer *p;

	isl_int_init(opt);
	bset = isl_basic_set_read_from_file(ctx, stdin);
	assert(bset);
	obj = isl_vec_read_from_file(ctx, stdin);
	assert(obj);
	dim = isl_basic_set_total_dim(bset);
	assert(obj->size >= dim && obj->size <= dim + 1);
	if (obj->size != dim + 1)
		obj = isl_vec_lin_to_aff(obj);
	else
		obj = vec_ror(obj);
	res = isl_basic_set_solve_ilp(bset, 0, obj->el, &opt, &sol);
	switch (res) {
	case isl_lp_error:
		fprintf(stderr, "error\n");
		return -1;
	case isl_lp_empty:
		fprintf(stdout, "empty\n");
		break;
	case isl_lp_unbounded:
		fprintf(stdout, "unbounded\n");
		break;
	case isl_lp_ok:
		p = isl_printer_to_file(ctx, stdout);
		p = isl_printer_print_vec(p, sol);
		p = isl_printer_end_line(p);
		p = isl_printer_print_isl_int(p, opt);
		p = isl_printer_end_line(p);
		isl_printer_free(p);
	}
	isl_basic_set_free(bset);
	isl_vec_free(obj);
	isl_vec_free(sol);
	isl_ctx_free(ctx);
	isl_int_clear(opt);

	return 0;
}
示例#8
0
/* Print a declaration for the device array corresponding to "array" on "p".
 */
static __isl_give isl_printer *declare_device_array(__isl_take isl_printer *p,
	struct gpu_array_info *array)
{
	int i;
	if(print_device_arrays_or_not(array))
	{
		p = isl_printer_start_line(p);
		p = isl_printer_print_str(p, array->type);
		p = isl_printer_print_str(p, " ");
		if (!array->linearize && array->n_index > 1)
			p = isl_printer_print_str(p, "(");
		p = isl_printer_print_str(p, "*dev_");
		p = isl_printer_print_str(p, array->name);
		if (!array->linearize && array->n_index > 1) {
			p = isl_printer_print_str(p, ")");
			for (i = 1; i < array->n_index; i++) {
				isl_ast_expr *bound;
				bound = isl_ast_expr_get_op_arg(array->bound_expr,
								1 + i);
				p = isl_printer_print_str(p, "[");
				p = isl_printer_print_ast_expr(p, bound);
				p = isl_printer_print_str(p, "]");
				isl_ast_expr_free(bound);
			}
		}
		p = isl_printer_print_str(p, ";");
		p = isl_printer_end_line(p);
	}
	return p;
}
示例#9
0
/* Print a list of iterators of type "type" with names "ids" to "p".
 * Each iterator is assigned the corresponding opencl identifier returned
 * by the function "opencl_id".
 * Unlike the equivalent function in the CUDA backend which prints iterators
 * in reverse order to promote coalescing, this function does not print
 * iterators in reverse order.  The OpenCL backend currently does not take
 * into account any coalescing considerations.
 */
static __isl_give isl_printer *print_iterators(__isl_take isl_printer *p,
	const char *type, __isl_keep isl_id_list *ids, const char *opencl_id)
{
	int i, n;

	n = isl_id_list_n_id(ids);
	if (n <= 0)
		return p;
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, type);
	p = isl_printer_print_str(p, " ");
	for (i = 0; i < n; ++i) {
		isl_id *id;

		if (i)
			p = isl_printer_print_str(p, ", ");
		id = isl_id_list_get_id(ids, i);
		p = isl_printer_print_id(p, id);
		isl_id_free(id);
		p = isl_printer_print_str(p, " = ");
		p = isl_printer_print_str(p, opencl_id);
		p = isl_printer_print_str(p, "(");
		p = isl_printer_print_int(p, i);
		p = isl_printer_print_str(p, ")");
	}
	p = isl_printer_print_str(p, ";");
	p = isl_printer_end_line(p);

	return p;
}
示例#10
0
static __isl_give isl_printer *print_kernel_var(__isl_take isl_printer *p,
	struct ppcg_kernel_var *var)
{
	int j;

	p = isl_printer_start_line(p);
	if (var->type == ppcg_access_shared)
		p = isl_printer_print_str(p, "__shared__ ");
	p = isl_printer_print_str(p, var->array->type);
	p = isl_printer_print_str(p, " ");
	p = isl_printer_print_str(p,  var->name);
	for (j = 0; j < var->array->n_index; ++j) {
		isl_val *v;

		p = isl_printer_print_str(p, "[");
		v = isl_vec_get_element_val(var->size, j);
		p = isl_printer_print_val(p, v);
		isl_val_free(v);
		p = isl_printer_print_str(p, "]");
	}
	p = isl_printer_print_str(p, ";");
	p = isl_printer_end_line(p);

	return p;
}
示例#11
0
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");
}
示例#12
0
/* Print a sync statement.
 */
static __isl_give isl_printer *print_sync(__isl_take isl_printer *p,
	struct ppcg_kernel_stmt *stmt)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "__syncthreads();");
	p = isl_printer_end_line(p);

	return p;
}
示例#13
0
static __isl_give isl_printer *opencl_release_cl_objects(
	__isl_take isl_printer *p, struct opencl_info *info)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn(clReleaseCommandQueue"
					"(queue));");
	p = isl_printer_end_line(p);
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn(clReleaseProgram"
					"(program));");
	p = isl_printer_end_line(p);
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn(clReleaseContext"
					"(context));");
	p = isl_printer_end_line(p);

	return p;
}
示例#14
0
static __isl_give isl_printer *opencl_declare_device_arrays(
	__isl_take isl_printer *p, struct gpu_prog *prog)
{
	int i;

	for (i = 0; i < prog->n_array; ++i) {
		if (!gpu_array_requires_device_allocation(&prog->array[i]))
			continue;
		p = isl_printer_start_line(p);
		p = isl_printer_print_str(p, "cl_mem dev_");
		p = isl_printer_print_str(p, prog->array[i].name);
		p = isl_printer_print_str(p, ";");
		p = isl_printer_end_line(p);
	}
	p = isl_printer_start_line(p);
	p = isl_printer_end_line(p);
	return p;
}
示例#15
0
/* Print a call to barrier() which is a sync statement.
 * All work-items in a work-group executing the kernel on a processor must
 * execute the barrier() function before any are allowed to continue execution
 * beyond the barrier.
 * The flag CLK_LOCAL_MEM_FENCE makes the barrier function either flush any
 * variables stored in local memory or queue a memory fence to ensure correct
 * ordering of memory operations to local memory.
 * The flag CLK_GLOBAL_MEM_FENCE makes the barrier function queue a memory
 * fence to ensure correct ordering of memory operations to global memory.
 */
static __isl_give isl_printer *opencl_print_sync(__isl_take isl_printer *p,
	struct ppcg_kernel_stmt *stmt)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p,
		"barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);");
	p = isl_printer_end_line(p);

	return p;
}
示例#16
0
/* Free the device array corresponding to "array"
 */
static __isl_give isl_printer *release_device_array(__isl_take isl_printer *p,
	struct gpu_array_info *array)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn("
					"clReleaseMemObject(dev_");
	p = isl_printer_print_str(p, array->name);
	p = isl_printer_print_str(p, "));");
	p = isl_printer_end_line(p);

	return p;
}
示例#17
0
void isl_id_dump(__isl_keep isl_id *id)
{
	isl_printer *printer;

	if (!id)
		return;

	printer = isl_printer_to_file(isl_id_get_ctx(id), stderr);
	printer = isl_printer_print_id(printer, id);
	printer = isl_printer_end_line(printer);

	isl_printer_free(printer);
}
示例#18
0
/* Print the grid definition.
 */
static __isl_give isl_printer *print_grid(__isl_take isl_printer *p,
	struct ppcg_kernel *kernel)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "dim3 k");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, "_dimGrid");
	p = print_grid_size(p, kernel);
	p = isl_printer_print_str(p, ";");
	p = isl_printer_end_line(p);

	return p;
}
示例#19
0
/* Print the header of the given kernel.
 */
static __isl_give isl_printer *opencl_print_kernel_header(
	__isl_take isl_printer *p, struct gpu_prog *prog,
	struct ppcg_kernel *kernel)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "__kernel void kernel");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, "(");
	p = opencl_print_kernel_arguments(p, prog, kernel, 1);
	p = isl_printer_print_str(p, ")");
	p = isl_printer_end_line(p);

	return p;
}
示例#20
0
static __isl_give isl_printer *declare_device_arrays(__isl_take isl_printer *p,
	struct gpu_prog *prog)
{
	int i;

	for (i = 0; i < prog->n_array; ++i) {
		if (!gpu_array_requires_device_allocation(&prog->array[i]))
			continue;

		p = declare_device_array(p, &prog->array[i]);
	}
	p = isl_printer_start_line(p);
	p = isl_printer_end_line(p);
	return p;
}
示例#21
0
/* Print code to "p" for copying "array" back from the device to the host
 * in its entirety.  The bounds on the extent of "array" have
 * been precomputed in extract_array_info and are used in
 * gpu_array_info_print_size.
 */
static __isl_give isl_printer *copy_array_from_device_global_memory(
	__isl_take isl_printer *p, struct gpu_array_info *array)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "cudaCheckReturn(cudaMemcpy(");
	if (gpu_array_is_scalar(array))
		p = isl_printer_print_str(p, "&");
	p = isl_printer_print_str(p, array->name);
	p = isl_printer_print_str(p, ", dev_");
	p = isl_printer_print_str(p, array->name);
	p = isl_printer_print_str(p, ", ");
	p = gpu_array_info_print_size(p, array);
	p = isl_printer_print_str(p, ", cudaMemcpyDeviceToHost));");
	p = isl_printer_end_line(p);
	return p;
}
示例#22
0
/* Allocate accessed device arrays.
 */
static __isl_give isl_printer *opencl_allocate_device_arrays(
	__isl_take isl_printer *p, struct gpu_prog *prog)
{
	int i;

	for (i = 0; i < prog->n_array; ++i) {
		struct gpu_array_info *array = &prog->array[i];

		if (!gpu_array_requires_device_allocation(array))
			continue;

		p = allocate_device_array(p, array);
	}
	p = isl_printer_start_line(p);
	p = isl_printer_end_line(p);
	return p;
}
示例#23
0
static __isl_give isl_printer *free_device_arrays(__isl_take isl_printer *p,
	struct gpu_prog *prog)
{
	int i;

	for (i = 0; i < prog->n_array; ++i) {
		if (!gpu_array_requires_device_allocation(&prog->array[i]))
			continue;
		if(print_device_arrays_or_not(&prog->array[i]))
		{
			p = isl_printer_start_line(p);
			p = isl_printer_print_str(p, "cudaCheckReturn(cudaFree(dev_");
			p = isl_printer_print_str(p, prog->array[i].name);
			p = isl_printer_print_str(p, "));");
			p = isl_printer_end_line(p);
		}
	}
	return p;
}
示例#24
0
static __isl_give isl_printer *opencl_print_host_macros(
	__isl_take isl_printer *p)
{
	const char *macros =
		"#define openclCheckReturn(ret) \\\n"
		"  if (ret != CL_SUCCESS) {\\\n"
		"    fprintf(stderr, \"OpenCL error: %s\\n\", "
		"opencl_error_string(ret)); \\\n"
		"    fflush(stderr); \\\n"
		"    assert(ret == CL_SUCCESS);\\\n  }\n";

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, macros);
	p = isl_printer_end_line(p);

	p = isl_ast_op_type_print_macro(isl_ast_op_max, p);

	return p;
}
示例#25
0
/* Print code to "p" for copying "array" from the host to the device
 * in its entirety.  The bounds on the extent of "array" have
 * been precomputed in extract_array_info and are used in
 * gpu_array_info_print_size.
 */
static __isl_give isl_printer *copy_array_to_device(__isl_take isl_printer *p,
	struct gpu_array_info *array)
{

	if(print_device_arrays_or_not(array))
	{
		p = isl_printer_start_line(p);
		p = isl_printer_print_str(p, "cudaCheckReturn(cudaMemcpy(dev_");
		p = isl_printer_print_str(p, array->name);
		p = isl_printer_print_str(p, ", ");

		if (gpu_array_is_scalar(array))
			p = isl_printer_print_str(p, "&");
		p = isl_printer_print_str(p, array->name);
		p = isl_printer_print_str(p, ", ");

		p = gpu_array_info_print_size(p, array);
		p = isl_printer_print_str(p, ", cudaMemcpyHostToDevice));");
		p = isl_printer_end_line(p);
	}
	return p;
}
示例#26
0
int main(int argc, char **argv)
{
    struct isl_ctx *ctx = isl_ctx_alloc();
    struct isl_basic_set *bset;
    struct isl_vec *sample;
    isl_printer *p;

    bset = isl_basic_set_read_from_file(ctx, stdin);
    sample = isl_basic_set_sample_vec(isl_basic_set_copy(bset));
    p = isl_printer_to_file(ctx, stdout);
    p = isl_printer_print_vec(p, sample);
    p = isl_printer_end_line(p);
    isl_printer_free(p);
    assert(sample);
    if (sample->size > 0)
        assert(isl_basic_set_contains(bset, sample));
    isl_basic_set_free(bset);
    isl_vec_free(sample);
    isl_ctx_free(ctx);

    return 0;
}
示例#27
0
/* Copy "array" from the host to the device (to_host = 0) or
 * back from the device to the host (to_host = 1).
 */
static __isl_give isl_printer *copy_array(__isl_take isl_printer *p,
	struct gpu_array_info *array, int to_host)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn(");
	if (to_host)
		p = isl_printer_print_str(p, "clEnqueueReadBuffer");
	else
		p = isl_printer_print_str(p, "clEnqueueWriteBuffer");
	p = isl_printer_print_str(p, "(queue, dev_");
	p = isl_printer_print_str(p, array->name);
	p = isl_printer_print_str(p, ", CL_TRUE, 0, ");
	p = gpu_array_info_print_size(p, array);

	if (gpu_array_is_scalar(array))
		p = isl_printer_print_str(p, ", &");
	else
		p = isl_printer_print_str(p, ", ");
	p = isl_printer_print_str(p, array->name);
	p = isl_printer_print_str(p, ", 0, NULL, NULL));");
	p = isl_printer_end_line(p);

	return p;
}
示例#28
0
/* Print a call to the OpenCL clSetKernelArg() function which sets
 * the arguments of the kernel.  arg_name and arg_index are the name and the
 * index of the kernel argument.  The index of the leftmost argument of
 * the kernel is 0 whereas the index of the rightmost argument of the kernel
 * is n - 1, where n is the total number of the kernel arguments.
 * read_only_scalar is a boolean that indicates whether the argument is a read
 * only scalar.
 */
static __isl_give isl_printer *opencl_set_kernel_argument(
	__isl_take isl_printer *p, int kernel_id,
	const char *arg_name, int arg_index, int read_only_scalar)
{
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p,
		"openclCheckReturn(clSetKernelArg(kernel");
	p = isl_printer_print_int(p, kernel_id);
	p = isl_printer_print_str(p, ", ");
	p = isl_printer_print_int(p, arg_index);
	p = isl_printer_print_str(p, ", sizeof(");

	if (read_only_scalar) {
		p = isl_printer_print_str(p, arg_name);
		p = isl_printer_print_str(p, "), &");
	} else
		p = isl_printer_print_str(p, "cl_mem), (void *) &dev_");

	p = isl_printer_print_str(p, arg_name);
	p = isl_printer_print_str(p, "));");
	p = isl_printer_end_line(p);

	return p;
}
示例#29
0
/* Print the user statement of the host code to "p".
 *
 * The host code may contain original user statements, kernel launches,
 * statements that copy data to/from the device and statements
 * the initialize or clear the device.
 * The original user statements and the kernel launches have
 * an associated annotation, while the other statements do not.
 * The latter are handled by print_device_node.
 * The annotation on the user statements is called "user".
 *
 * In case of a kernel launch, print a block of statements that
 * defines the grid and the block and then launches the kernel.
 */
static __isl_give isl_printer *print_host_user(__isl_take isl_printer *p,
	__isl_take isl_ast_print_options *print_options,
	__isl_keep isl_ast_node *node, void *user)
{
	isl_id *id;
	int is_user;
	struct ppcg_kernel *kernel;
	struct ppcg_kernel_stmt *stmt;
	struct print_host_user_data *data;

	isl_ast_print_options_free(print_options);

	data = (struct print_host_user_data *) user;

	id = isl_ast_node_get_annotation(node);
	if (!id)
	{
		//p = isl_printer_print_str(p,"marker_NO_ID_CASE");
		return print_device_node(p, node, data->prog);
	}	
	is_user = !strcmp(isl_id_get_name(id), "user");
	kernel = is_user ? NULL : isl_id_get_user(id);
	stmt = is_user ? isl_id_get_user(id) : NULL;
	isl_id_free(id);

	if (is_user)
		return ppcg_kernel_print_domain(p, stmt);

	p = ppcg_start_block(p);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "dim3 k");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, "_dimBlock");
	print_reverse_list(isl_printer_get_file(p),
				kernel->n_block, kernel->block_dim);
	p = isl_printer_print_str(p, ";");
	p = isl_printer_end_line(p);

	p = print_grid(p, kernel);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "kernel");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, " <<<k");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, "_dimGrid, k");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, "_dimBlock>>> (");
	p = print_kernel_arguments(p, data->prog, kernel, 0);
	p = isl_printer_print_str(p, ");");
	p = isl_printer_end_line(p);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "cudaCheckKernel();");
	p = isl_printer_end_line(p);

	p = ppcg_end_block(p);

	p = isl_printer_start_line(p);
	p = isl_printer_end_line(p);

	p = copy_data_from_device_to_device(p,kernel); 

	printf("printing kernel");

	print_kernel(data->prog, kernel, data->cuda);

	printf("printing kernel done");

	return p;
}
示例#30
0
/* Print the user statement of the host code to "p".
 *
 * The host code may contain original user statements, kernel launches and
 * statements that copy data to/from the device.
 * The original user statements and the kernel launches have
 * an associated annotation, while the data copy statements do not.
 * The latter are handled by print_to_from_device.
 * The annotation on the user statements is called "user".
 *
 * In case of a kernel launch, print a block of statements that
 * defines the grid and the work group and then launches the kernel.
 *
 * A grid is composed of many work groups (blocks), each work group holds
 * many work-items (threads).
 *
 * global_work_size[kernel->n_block] represents the total number of work
 * items.  It points to an array of kernel->n_block unsigned
 * values that describe the total number of work-items that will execute
 * the kernel.  The total number of work-items is computed as:
 * global_work_size[0] *...* global_work_size[kernel->n_block - 1].
 *
 * The size of each work group (i.e. the number of work-items in each work
 * group) is described using block_size[kernel->n_block].  The total
 * number of work-items in a block (work-group) is computed as:
 * block_size[0] *... * block_size[kernel->n_block - 1].
 *
 * For more information check:
 * http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html
 */
static __isl_give isl_printer *opencl_print_host_user(
	__isl_take isl_printer *p,
	__isl_take isl_ast_print_options *print_options,
	__isl_keep isl_ast_node *node, void *user)
{
	isl_id *id;
	int is_user;
	struct ppcg_kernel *kernel;
	struct ppcg_kernel_stmt *stmt;
	struct print_host_user_data_opencl *data;

	isl_ast_print_options_free(print_options);

	data = (struct print_host_user_data_opencl *) user;

	id = isl_ast_node_get_annotation(node);
	if (!id)
		return print_to_from_device(p, node, data->prog);

	is_user = !strcmp(isl_id_get_name(id), "user");
	kernel = is_user ? NULL : isl_id_get_user(id);
	stmt = is_user ? isl_id_get_user(id) : NULL;
	isl_id_free(id);

	if (is_user)
		return ppcg_kernel_print_domain(p, stmt);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "{");
	p = isl_printer_end_line(p);
	p = isl_printer_indent(p, 2);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "size_t global_work_size[");

	if (kernel->n_block > 0)
		p = isl_printer_print_int(p, kernel->n_block);
	else
		p = isl_printer_print_int(p, 1);

	p = isl_printer_print_str(p, "] = {");
	p = opencl_print_total_number_of_work_items_as_list(p, kernel);
	p = isl_printer_print_str(p, "};");
	p = isl_printer_end_line(p);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "size_t block_size[");

	if (kernel->n_block > 0)
		p = isl_printer_print_int(p, kernel->n_block);
	else
		p = isl_printer_print_int(p, 1);

	p = isl_printer_print_str(p, "] = {");
	p = opencl_print_block_sizes(p, kernel);
	p = isl_printer_print_str(p, "};");
	p = isl_printer_end_line(p);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "cl_kernel kernel");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, " = clCreateKernel(program, \"kernel");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, "\", &err);");
	p = isl_printer_end_line(p);
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn(err);");
	p = isl_printer_end_line(p);

	opencl_set_kernel_arguments(p, data->prog, kernel);

	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn(clEnqueueNDRangeKernel"
		"(queue, kernel");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, ", ");
	if (kernel->n_block > 0)
		p = isl_printer_print_int(p, kernel->n_block);
	else
		p = isl_printer_print_int(p, 1);

	p = isl_printer_print_str(p, ", NULL, global_work_size, "
					"block_size, "
					"0, NULL, NULL));");
	p = isl_printer_end_line(p);
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "openclCheckReturn("
					"clReleaseKernel(kernel");
	p = isl_printer_print_int(p, kernel->id);
	p = isl_printer_print_str(p, "));");
	p = isl_printer_end_line(p);
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "clFinish(queue);");
	p = isl_printer_end_line(p);
	p = isl_printer_indent(p, -2);
	p = isl_printer_start_line(p);
	p = isl_printer_print_str(p, "}");
	p = isl_printer_end_line(p);

	p = isl_printer_start_line(p);
	p = isl_printer_end_line(p);

	data->opencl->kprinter = opencl_print_kernel(data->prog, kernel,
						data->opencl->kprinter);

	return p;
}