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