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;
}
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;
}
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;
}
/* This function prints the i'th block size multiplied by the i'th grid size,
* where i (a parameter to this function) is one of the possible dimensions of
* grid sizes and block sizes.
* If the dimension of block sizes is not equal to the dimension of grid sizes
* the output is calculated as follows:
*
* Suppose that:
* block_sizes[dim1] is the list of blocks sizes and it contains dim1 elements.
* grid_sizes[dim2] is the list of grid sizes and it contains dim2 elements.
*
* The output is:
* If (i > dim2) then the output is block_sizes[i]
* If (i > dim1) then the output is grid_sizes[i]
*/
static __isl_give isl_printer *opencl_print_total_number_of_work_items_for_dim(
__isl_take isl_printer *p, struct ppcg_kernel *kernel, int i)
{
int grid_dim, block_dim;
isl_pw_aff *bound_grid;
grid_dim = isl_multi_pw_aff_dim(kernel->grid_size, isl_dim_set);
block_dim = kernel->n_block;
if (i < min(grid_dim, block_dim)) {
bound_grid = isl_multi_pw_aff_get_pw_aff(kernel->grid_size, i);
p = isl_printer_print_str(p, "(");
p = isl_printer_print_pw_aff(p, bound_grid);
p = isl_printer_print_str(p, ") * ");
p = isl_printer_print_int(p, kernel->block_dim[i]);
isl_pw_aff_free(bound_grid);
} else if (i >= grid_dim)
p = isl_printer_print_int(p, kernel->block_dim[i]);
else {
bound_grid = isl_multi_pw_aff_get_pw_aff(kernel->grid_size, i);
p = isl_printer_print_pw_aff(p, bound_grid);
isl_pw_aff_free(bound_grid);
}
return p;
}
/* 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;
}
/* Print the effective grid size as a list of the sizes in each
* dimension, from innermost to outermost.
*/
static __isl_give isl_printer *print_grid_size(__isl_take isl_printer *p,
struct ppcg_kernel *kernel)
{
int i;
int dim;
dim = isl_multi_pw_aff_dim(kernel->grid_size, isl_dim_set);
if (dim == 0)
return p;
p = isl_printer_print_str(p, "(");
for (i = dim - 1; i >= 0; --i) {
isl_ast_expr *bound;
bound = isl_ast_expr_get_op_arg(kernel->grid_size_expr, 1 + i);
p = isl_printer_print_ast_expr(p, bound);
isl_ast_expr_free(bound);
if (i > 0)
p = isl_printer_print_str(p, ", ");
}
p = isl_printer_print_str(p, ")");
return p;
}
/* 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;
}
/* 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;
}
/* Print the header of the given kernel.
*/
static __isl_give isl_printer *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, "__global__ void kernel");
p = isl_printer_print_int(p, kernel->id);
p = isl_printer_print_str(p, "(");
p = print_kernel_arguments(p, prog, kernel, 1);
p = isl_printer_print_str(p, ")");
return p;
}
/* 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;
}
/* 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;
}
/* 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;
}
/* Print the block sizes as a list of the sizes in each
* dimension.
*/
static __isl_give isl_printer *opencl_print_block_sizes(
__isl_take isl_printer *p, struct ppcg_kernel *kernel)
{
int i;
if (kernel->n_block > 0)
for (i = 0; i < kernel->n_block; ++i) {
if (i)
p = isl_printer_print_str(p, ", ");
p = isl_printer_print_int(p, kernel->block_dim[i]);
}
else
p = isl_printer_print_str(p, "1");
return p;
}
/* Open the host .c file and the kernel .h and .cl files for writing.
* Their names are derived from info->output (or info->input if
* the user did not specify an output file name).
* Add the necessary includes to these files, including those specified
* by the user.
*
* Return 0 on success and -1 on failure.
*/
static int opencl_open_files(struct opencl_info *info)
{
char name[PATH_MAX];
int i;
int len;
if (info->output) {
const char *ext;
ext = strrchr(info->output, '.');
len = ext ? ext - info->output : strlen(info->output);
memcpy(name, info->output, len);
info->host_c = open_or_croak(info->output);
} else {
len = ppcg_extract_base_name(name, info->input);
strcpy(name + len, "_host.c");
info->host_c = open_or_croak(name);
}
memcpy(info->kernel_c_name, name, len);
strcpy(info->kernel_c_name + len, "_kernel.cl");
info->kernel_c = open_or_croak(info->kernel_c_name);
if (!info->host_c || !info->kernel_c)
return -1;
fprintf(info->host_c, "#include <assert.h>\n");
fprintf(info->host_c, "#include <stdio.h>\n");
fprintf(info->host_c, "#include \"ocl_utilities.h\"\n");
if (info->options->opencl_embed_kernel_code) {
fprintf(info->host_c, "#include \"%s\"\n\n",
info->kernel_c_name);
}
for (i = 0; i < info->options->opencl_n_include_file; ++i) {
info->kprinter = isl_printer_print_str(info->kprinter,
"#include <");
info->kprinter = isl_printer_print_str(info->kprinter,
info->options->opencl_include_files[i]);
info->kprinter = isl_printer_print_str(info->kprinter, ">\n");
}
return 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;
}
__isl_give isl_printer *isl_printer_print_id(__isl_take isl_printer *p,
__isl_keep isl_id *id)
{
if (!id)
goto error;
if (id->name)
p = isl_printer_print_str(p, id->name);
if (id->user) {
char buffer[50];
snprintf(buffer, sizeof(buffer), "@%p", id->user);
p = isl_printer_print_str(p, buffer);
}
return p;
error:
isl_printer_free(p);
return NULL;
}
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;
}
/* 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;
}
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;
}
/* 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;
}
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;
}
/* 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;
}
/* Print a list that represents the total number of work items. The list is
* constructed by performing an element-wise multiplication of the block sizes
* and the grid sizes. To explain how the list is constructed, suppose that:
* block_sizes[dim1] is the list of blocks sizes and it contains dim1 elements.
* grid_sizes[dim2] is the list of grid sizes and it contains dim2 elements.
*
* The output of this function is constructed as follows:
* If (dim1 > dim2) then the output is the following list:
* grid_sizes[0]*block_sizes[0], ..., grid_sizes[dim2-1]*block_sizes[dim2-1],
* block_sizes[dim2], ..., block_sizes[dim1-2], block_sizes[dim1-1].
*
* If (dim2 > dim1) then the output is the following list:
* grid_sizes[0]*block_sizes[0], ..., grid_sizes[dim1-1] * block_sizes[dim1-1],
* grid_sizes[dim1], grid_sizes[dim2-2], grid_sizes[dim2-1].
*
* To calculate the total number of work items out of the list constructed by
* this function, the user should multiply the elements of the list.
*/
static __isl_give isl_printer *opencl_print_total_number_of_work_items_as_list(
__isl_take isl_printer *p, struct ppcg_kernel *kernel)
{
int i;
int grid_dim, block_dim;
grid_dim = isl_multi_pw_aff_dim(kernel->grid_size, isl_dim_set);
block_dim = kernel->n_block;
if ((grid_dim <= 0) || (block_dim <= 0)) {
p = isl_printer_print_str(p, "1");
return p;
}
for (i = 0; i <= max(grid_dim, block_dim) - 1; i++) {
if (i > 0)
p = isl_printer_print_str(p, ", ");
p = opencl_print_total_number_of_work_items_for_dim(p,
kernel, i);
}
return p;
}
/* 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;
}
/* 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;
}
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;
}
static __isl_give isl_printer *print_cuda_macros(__isl_take isl_printer *p)
{
const char *macros =
"#define cudaCheckReturn(ret) \\\n"
" do { \\\n"
" cudaError_t cudaCheckReturn_e = (ret); \\\n"
" if (cudaCheckReturn_e != cudaSuccess) { \\\n"
" fprintf(stderr, \"CUDA error: %s\\n\", "
"cudaGetErrorString(cudaCheckReturn_e)); \\\n"
" fflush(stderr); \\\n"
" } \\\n"
" assert(cudaCheckReturn_e == cudaSuccess); \\\n"
" } while(0)\n"
"#define cudaCheckKernel() \\\n"
" do { \\\n"
" cudaCheckReturn(cudaGetLastError()); \\\n"
" } while(0)\n\n";
p = isl_printer_print_str(p, macros);
return p;
}
/* 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;
}
/* 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;
}
/* 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;
}
