blob: c61382f2066b61912ceccbedd6d42724d742cb49 [file] [log] [blame]
/*
* Copyright 2013 Ecole Normale Superieure
*
* Use of this software is governed by the MIT license
*
* Written by Sven Verdoolaege and Riyadh Baghdadi,
* Ecole Normale Superieure, 45 rue d’Ulm, 75230 Paris, France
*/
#include <ctype.h>
#include <limits.h>
#include <string.h>
#include <isl/aff.h>
#include <isl/ast.h>
#include "opencl.h"
#include "gpu_print.h"
#include "gpu.h"
#include "ppcg.h"
#include "print.h"
#include "schedule.h"
#include "util.h"
#define min(a, b) (((a) < (b)) ? (a) : (b))
#define max(a, b) (((a) > (b)) ? (a) : (b))
/* options are the global options passed to generate_opencl.
* input is the name of the input file.
* output is the user-specified output file name and may be NULL
* if not specified by the user.
* kernel_c_name is the name of the kernel_c file.
* kprinter is an isl_printer for the kernel file.
* host_c is the generated source file for the host code. kernel_c is
* the generated source file for the kernel.
*/
struct opencl_info {
struct ppcg_options *options;
const char *input;
const char *output;
char kernel_c_name[PATH_MAX];
isl_printer *kprinter;
FILE *host_c;
FILE *kernel_c;
};
/* Open the file called "name" for writing or print an error message.
*/
static FILE *open_or_croak(const char *name)
{
FILE *file;
file = fopen(name, "w");
if (!file)
fprintf(stderr, "Failed to open \"%s\" for writing\n", name);
return file;
}
/* 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;
}
/* Write text to a file and escape some special characters that would break a
* C string.
*/
static void opencl_print_escaped(const char *str, const char *end, FILE *file)
{
const char *prev = str;
while ((str = strpbrk(prev, "\"\\")) && str < end) {
fwrite(prev, 1, str - prev, file);
fprintf(file, "\\%c", *str);
prev = str + 1;
}
if (*prev)
fwrite(prev, 1, end - prev, file);
}
/* Write text to a file as a C string literal.
*
* This function also prints any characters after the last newline, although
* normally the input string should end with a newline.
*/
static void opencl_print_as_c_string(const char *str, FILE *file)
{
const char *prev = str;
while ((str = strchr(prev, '\n'))) {
fprintf(file, "\n\"");
opencl_print_escaped(prev, str, file);
fprintf(file, "\\n\"");
prev = str + 1;
}
if (*prev) {
fprintf(file, "\n\"");
opencl_print_escaped(prev, prev + strlen(prev), file);
fprintf(file, "\"");
}
}
/* Write the code that we have accumulated in the kernel isl_printer to the
* kernel.cl file. If the opencl_embed_kernel_code option has been set, print
* the code as a C string literal. Start that string literal with an empty
* line, such that line numbers reported by the OpenCL C compiler match those
* of the kernel file.
*
* Return 0 on success and -1 on failure.
*/
static int opencl_write_kernel_file(struct opencl_info *opencl)
{
char *raw = isl_printer_get_str(opencl->kprinter);
if (!raw)
return -1;
if (opencl->options->opencl_embed_kernel_code) {
fprintf(opencl->kernel_c,
"static const char kernel_code[] = \"\\n\"");
opencl_print_as_c_string(raw, opencl->kernel_c);
fprintf(opencl->kernel_c, ";\n");
} else
fprintf(opencl->kernel_c, "%s", raw);
free(raw);
return 0;
}
/* Close all output files. Write the kernel contents to the kernel file before
* closing it.
*
* Return 0 on success and -1 on failure.
*/
static int opencl_close_files(struct opencl_info *info)
{
int r = 0;
if (info->kernel_c) {
r = opencl_write_kernel_file(info);
fclose(info->kernel_c);
}
if (info->host_c)
fclose(info->host_c);
return r;
}
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 *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;
}
/* Given an array, check whether its positive size guard expression is
* trivial.
*/
static int is_array_positive_size_guard_trivial(struct gpu_array_info *array)
{
isl_set *guard;
int is_trivial;
guard = gpu_array_positive_size_guard(array);
is_trivial = isl_set_plain_is_universe(guard);
isl_set_free(guard);
return is_trivial;
}
/* 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;
}
/* 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;
}
/* 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 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;
}
/* Set the arguments of the OpenCL kernel by printing a call to the OpenCL
* clSetKernelArg() function for each kernel argument.
*/
static __isl_give isl_printer *opencl_set_kernel_arguments(
__isl_take isl_printer *p, struct gpu_prog *prog,
struct ppcg_kernel *kernel)
{
int i, n, ro;
unsigned nparam;
isl_space *space;
int arg_index = 0;
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;
ro = gpu_array_is_read_only_scalar(&prog->array[i]);
opencl_set_kernel_argument(p, kernel->id, prog->array[i].name,
arg_index, ro);
arg_index++;
}
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);
opencl_set_kernel_argument(p, kernel->id, name, arg_index, 1);
arg_index++;
}
isl_space_free(space);
n = isl_space_dim(kernel->space, isl_dim_set);
for (i = 0; i < n; ++i) {
const char *name;
name = isl_space_get_dim_name(kernel->space, isl_dim_set, i);
opencl_set_kernel_argument(p, kernel->id, name, arg_index, 1);
arg_index++;
}
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 *opencl_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 (!first)
p = isl_printer_print_str(p, ", ");
if (types)
p = gpu_array_info_print_declaration_argument(p,
&prog->array[i], "__global");
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;
}
/* 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;
}
/* 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;
}
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;
}
static __isl_give isl_printer *opencl_print_kernel_var(
__isl_take isl_printer *p, struct ppcg_kernel_var *var)
{
int j;
isl_val *v;
p = isl_printer_start_line(p);
if (var->type == ppcg_access_shared)
p = isl_printer_print_str(p, "__local ");
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) {
p = isl_printer_print_str(p, "[");
v = isl_vec_get_element_val(var->size, j);
p = isl_printer_print_val(p, v);
p = isl_printer_print_str(p, "]");
isl_val_free(v);
}
p = isl_printer_print_str(p, ";");
p = isl_printer_end_line(p);
return p;
}
static __isl_give isl_printer *opencl_print_kernel_vars(
__isl_take isl_printer *p, struct ppcg_kernel *kernel)
{
int i;
for (i = 0; i < kernel->n_var; ++i)
p = opencl_print_kernel_var(p, &kernel->var[i]);
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;
}
/* Data structure containing function names for which the calls
* should be changed from
*
* name(arg)
*
* to
*
* opencl_name((type) (arg))
*/
static struct ppcg_opencl_fn {
const char *name;
const char *opencl_name;
const char *type;
} opencl_fn[] = {
{ "expf", "exp", "float" },
{ "powf", "pow", "float" },
{ "sqrtf", "sqrt", "float" },
};
#define ARRAY_SIZE(array) (sizeof(array)/sizeof(*array))
/* If the name of function called by "expr" matches any of those
* in ppcg_opencl_fn, then replace the call by a cast to the corresponding
* type in ppcg_opencl_fn and a call to corresponding OpenCL function.
*/
static __isl_give pet_expr *map_opencl_call(__isl_take pet_expr *expr,
void *user)
{
const char *name;
int i;
name = pet_expr_call_get_name(expr);
for (i = 0; i < ARRAY_SIZE(opencl_fn); ++i) {
pet_expr *arg;
if (strcmp(name, opencl_fn[i].name))
continue;
expr = pet_expr_call_set_name(expr, opencl_fn[i].opencl_name);
arg = pet_expr_get_arg(expr, 0);
arg = pet_expr_new_cast(opencl_fn[i].type, arg);
expr = pet_expr_set_arg(expr, 0, arg);
}
return expr;
}
/* Print the body of a statement from the input program,
* for use in OpenCL code.
*
* Before calling ppcg_kernel_print_domain to print the actual statement body,
* we first modify this body to take into account that the output code
* is OpenCL code. In particular, if the statement calls any function
* with a "f" suffix, then it needs to be replaced by a call to
* the corresponding function without suffix after casting the argument
* to a float.
*/
static __isl_give isl_printer *print_opencl_kernel_domain(
__isl_take isl_printer *p, struct ppcg_kernel_stmt *stmt)
{
struct pet_stmt *ps;
pet_tree *tree;
ps = stmt->u.d.stmt->stmt;
tree = pet_tree_copy(ps->body);
ps->body = pet_tree_map_call_expr(ps->body, &map_opencl_call, NULL);
p = ppcg_kernel_print_domain(p, stmt);
pet_tree_free(ps->body);
ps->body = tree;
return p;
}
/* This function is called for each user statement in the AST,
* i.e., for each kernel body statement, copy statement or sync statement.
*/
static __isl_give isl_printer *opencl_print_kernel_stmt(
__isl_take isl_printer *p,
__isl_take isl_ast_print_options *print_options,
__isl_keep isl_ast_node *node, void *user)
{
isl_id *id;
struct ppcg_kernel_stmt *stmt;
id = isl_ast_node_get_annotation(node);
stmt = isl_id_get_user(id);
isl_id_free(id);
isl_ast_print_options_free(print_options);
switch (stmt->type) {
case ppcg_kernel_copy:
return ppcg_kernel_print_copy(p, stmt);
case ppcg_kernel_sync:
return opencl_print_sync(p, stmt);
case ppcg_kernel_domain:
return print_opencl_kernel_domain(p, stmt);
}
return p;
}
/* Return true if there is a double array in prog->array or
* if any of the types in prog->scop involve any doubles.
* To check the latter condition, we simply search for the string "double"
* in the type definitions, which may result in false positives.
*/
static __isl_give int any_double_elements(struct gpu_prog *prog)
{
int i;
for (i = 0; i < prog->n_array; ++i)
if (strcmp(prog->array[i].type, "double") == 0)
return 1;
for (i = 0; i < prog->scop->pet->n_type; ++i) {
struct pet_type *type = prog->scop->pet->types[i];
if (strstr(type->definition, "double"))
return 1;
}
return 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;
}
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;
}
struct print_host_user_data_opencl {
struct opencl_info *opencl;
struct gpu_prog *prog;
};
/* 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;
}
/* 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 a statement for copying an array to or from the device.
* The statement identifier is called "to_device_<array name>" or
* "from_device_<array name>" and its user pointer points
* to the gpu_array_info of the array that needs to be copied.
*
* Extract the array from the identifier and call
* copy_array_to_device or copy_array_from_device.
*/
static __isl_give isl_printer *print_to_from_device(__isl_take isl_printer *p,
__isl_keep isl_ast_node *node, struct gpu_prog *prog)
{
isl_ast_expr *expr, *arg;
isl_id *id;
const char *name;
struct gpu_array_info *array;
expr = isl_ast_node_user_get_expr(node);
arg = isl_ast_expr_get_op_arg(expr, 0);
id = isl_ast_expr_get_id(arg);
name = isl_id_get_name(id);
array = isl_id_get_user(id);
isl_id_free(id);
isl_ast_expr_free(arg);
isl_ast_expr_free(expr);
if (!name)
array = NULL;
if (!array)
return isl_printer_free(p);
if (!prefixcmp(name, "to_device"))
return copy_array(p, array, 0);
else
return copy_array(p, array, 1);
}
/* 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;
}
static __isl_give isl_printer *opencl_print_host_code(
__isl_take isl_printer *p, struct gpu_prog *prog,
__isl_keep isl_ast_node *tree, struct opencl_info *opencl)
{
isl_ast_print_options *print_options;
isl_ctx *ctx = isl_ast_node_get_ctx(tree);
struct print_host_user_data_opencl data = { opencl, prog };
print_options = isl_ast_print_options_alloc(ctx);
print_options = isl_ast_print_options_set_print_user(print_options,
&opencl_print_host_user, &data);
p = ppcg_print_macros(p, tree);
p = isl_ast_node_print(tree, p, print_options);
return p;
}
/* Create an OpenCL device, context, command queue and build the kernel.
* input is the name of the input file provided to ppcg.
*/
static __isl_give isl_printer *opencl_setup(__isl_take isl_printer *p,
const char *input, struct opencl_info *info)
{
p = isl_printer_start_line(p);
p = isl_printer_print_str(p, "cl_device_id device;");
p = isl_printer_end_line(p);
p = isl_printer_start_line(p);
p = isl_printer_print_str(p, "cl_context context;");
p = isl_printer_end_line(p);
p = isl_printer_start_line(p);
p = isl_printer_print_str(p, "cl_program program;");
p = isl_printer_end_line(p);
p = isl_printer_start_line(p);
p = isl_printer_print_str(p, "cl_command_queue queue;");
p = isl_printer_end_line(p);
p = isl_printer_start_line(p);
p = isl_printer_print_str(p, "cl_int err;");
p = isl_printer_end_line(p);
p = isl_printer_start_line(p);
p = isl_printer_print_str(p, "device = opencl_create_device(");
p = isl_printer_print_int(p, info->options->opencl_use_gpu);
p = isl_printer_print_str(p, ");");
p = isl_printer_end_line(p);
p = isl_printer_start_line(p);
p = isl_printer_print_str(p, "context = clCreateContext(NULL, 1, "
"&device, NULL, 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 = isl_printer_start_line(p);
p = isl_printer_print_str(p, "queue = clCreateCommandQueue"
"(context, device, 0, &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 = isl_printer_start_line(p);
p = isl_printer_print_str(p, "program = ");
if (info->options->opencl_embed_kernel_code) {
p = isl_printer_print_str(p, "opencl_build_program_from_string("
"context, device, kernel_code, "
"sizeof(kernel_code), \"");
} else {
p = isl_printer_print_str(p, "opencl_build_program_from_file("
"context, device, \"");
p = isl_printer_print_str(p, info->kernel_c_name);
p = isl_printer_print_str(p, "\", \"");
}
if (info->options->opencl_compiler_options)
p = isl_printer_print_str(p,
info->options->opencl_compiler_options);
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;
}
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;
}
/* 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;
}
/* Free the accessed device arrays.
*/
static __isl_give isl_printer *opencl_release_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 = release_device_array(p, array);
}
return p;
}
/* Given a gpu_prog "prog" and the corresponding transformed AST
* "tree", print the entire OpenCL code to "p".
*/
static __isl_give isl_printer *print_opencl(__isl_take isl_printer *p,
struct gpu_prog *prog, __isl_keep isl_ast_node *tree,
struct gpu_types *types, void *user)
{
struct opencl_info *opencl = user;
opencl->kprinter = isl_printer_set_output_format(opencl->kprinter,
ISL_FORMAT_C);
if (any_double_elements(prog))
opencl->kprinter = opencl_enable_double_support(
opencl->kprinter);
if (opencl->options->opencl_print_kernel_types)
opencl->kprinter = gpu_print_types(opencl->kprinter, types,
prog);
if (!opencl->kprinter)
return isl_printer_free(p);
p = ppcg_start_block(p);
p = opencl_print_host_macros(p);
p = gpu_print_local_declarations(p, prog);
p = opencl_declare_device_arrays(p, prog);
p = opencl_setup(p, opencl->input, opencl);
p = opencl_allocate_device_arrays(p, prog);
p = opencl_print_host_code(p, prog, tree, opencl);
p = opencl_release_device_arrays(p, prog);
p = opencl_release_cl_objects(p, opencl);
p = ppcg_end_block(p);
return p;
}
/* Transform the code in the file called "input" by replacing
* all scops by corresponding OpenCL code.
* The host code is written to "output" or a name derived from
* "input" if "output" is NULL.
* The kernel code is placed in separate files with names
* derived from "output" or "input".
*
* We let generate_gpu do all the hard work and then let it call
* us back for printing the AST in print_opencl.
*
* To prepare for this printing, we first open the output files
* and we close them after generate_gpu has finished.
*/
int generate_opencl(isl_ctx *ctx, struct ppcg_options *options,
const char *input, const char *output)
{
struct opencl_info opencl = { options, input, output };
int r;
opencl.kprinter = isl_printer_to_str(ctx);
r = opencl_open_files(&opencl);
if (r >= 0)
r = generate_gpu(ctx, input, opencl.host_c, options,
&print_opencl, &opencl);
if (opencl_close_files(&opencl) < 0)
r = -1;
isl_printer_free(opencl.kprinter);
return r;
}