/** * SECTION:ufo-filter-task * @Short_description: Process arbitrary Filter kernels * @Title: filter * * This module is used to load an arbitrary #UfoPadding2DTask:kernel from * #UfoPadding2DTask:filename and execute it on each input. The kernel must have * only two global float array parameters, the first represents the input, the * second one the output. #UfoPadding2DTask:num-dims must be changed, if the kernel * accesses either one or three dimensional index spaces. */ #ifdef __APPLE__ #include #else #include #endif #include "ufo-padding-2d-task.h" typedef enum { PADDING_ZERO = 0, PADDING_CONST, PADDING_GAVG, PADDING_BREP } PaddingMode; struct _UfoPadding2DTaskPrivate { guint in_width; guint in_height; guint out_width; guint out_height; // extent adds guint xl; guint xr; guint yt; guint yb; size_t global_work_size_small[2]; size_t global_work_size_large[2]; PaddingMode mode; // padding constant float pconst; cl_kernel kernel_iconst; cl_kernel kernel_cpyimg; cl_kernel kernel_brep; }; static void ufo_task_interface_init (UfoTaskIface *iface); G_DEFINE_TYPE_WITH_CODE (UfoPadding2DTask, ufo_padding_2d_task, UFO_TYPE_TASK_NODE, G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK, ufo_task_interface_init)) #define UFO_PADDING_2D_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_PADDING_2D_TASK, UfoPadding2DTaskPrivate)) enum { PROP_0, PROP_XL, PROP_XR, PROP_YT, PROP_YB, PROP_MODE, PROP_PCONST, N_PROPERTIES }; static GParamSpec *properties[N_PROPERTIES] = { NULL, }; UfoNode * ufo_padding_2d_task_new (void) { return UFO_NODE (g_object_new (UFO_TYPE_PADDING_2D_TASK, NULL)); } static gboolean ufo_padding_2d_task_process (UfoTask *task, UfoBuffer **inputs, UfoBuffer *output, UfoRequisition *requisition) { UfoPadding2DTaskPrivate *priv; UfoGpuNode *node; cl_command_queue cmd_queue; cl_mem in_mem; cl_mem out_mem; priv = UFO_PADDING_2D_TASK (task)->priv; const PaddingMode mode = priv->mode; const guint pxl = priv->xl; const guint pyt = priv->yt; float pval = priv->pconst; const guint ixs = priv->in_width; const guint iys = priv->in_height; const guint oxs = priv->out_width; if (mode == PADDING_GAVG) { gfloat *indata = ufo_buffer_get_host_array (inputs[0], NULL); gfloat sum = 0; guint psz = ixs * iys; for (guint i =0; i < psz; i++) sum += indata[i]; pval = sum / (gfloat) psz; } node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task))); cmd_queue = ufo_gpu_node_get_cmd_queue (node); in_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue); out_mem = ufo_buffer_get_device_array (output, cmd_queue); if ((mode == PADDING_ZERO) || (mode == PADDING_CONST) || (mode == PADDING_GAVG)) { cl_kernel k_iconst = priv->kernel_iconst; cl_kernel k_cpyimg = priv->kernel_cpyimg; /// fill with constant UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_iconst, 0, sizeof(cl_mem), (void *) &out_mem)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_iconst, 1, sizeof(int), &oxs)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_iconst, 2, sizeof(float), &pval)); UFO_RESOURCES_CHECK_CLERR(clEnqueueNDRangeKernel(cmd_queue, k_iconst, 2, NULL, priv->global_work_size_large, NULL, 0, NULL, NULL)); /// copy old image UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 0, sizeof(cl_mem), (void *) &in_mem)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 1, sizeof(cl_mem), (void *) &out_mem)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 2, sizeof(int), &ixs)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 3, sizeof(int), &oxs)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 4, sizeof(int), &pxl)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 5, sizeof(int), &pyt)); UFO_RESOURCES_CHECK_CLERR(clEnqueueNDRangeKernel(cmd_queue, k_cpyimg, 2, NULL, priv->global_work_size_small, NULL, 0, NULL, NULL)); } if (mode == PADDING_BREP) { cl_kernel k_brep = priv->kernel_brep; UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 0, sizeof(cl_mem), (void *) &in_mem)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 1, sizeof(cl_mem), (void *) &out_mem)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 2, sizeof(int), &ixs)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 3, sizeof(int), &iys)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 4, sizeof(int), &oxs)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 5, sizeof(int), &pxl)); UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 6, sizeof(int), &pyt)); UFO_RESOURCES_CHECK_CLERR(clEnqueueNDRangeKernel(cmd_queue, k_brep, 2, NULL, priv->global_work_size_large, NULL, 0, NULL, NULL)); } return TRUE; } static void ufo_padding_2d_task_setup (UfoTask *task, UfoResources *resources, GError **error) { UfoPadding2DTaskPrivate *priv; priv = UFO_PADDING_2D_TASK_GET_PRIVATE (task); priv->kernel_iconst = ufo_resources_get_kernel (resources, "padding_2d.cl", "padding_2d_init_const", error); priv->kernel_cpyimg = ufo_resources_get_kernel (resources, "padding_2d.cl", "padding_2d_copy_in", error); priv->kernel_brep = ufo_resources_get_kernel (resources, "padding_2d.cl", "padding_2d_brep", error); if (priv->kernel_iconst) { UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel_iconst)); } if (priv->kernel_cpyimg) { UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel_cpyimg)); } if (priv->kernel_brep) { UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel_brep)); } } static void ufo_padding_2d_task_get_requisition (UfoTask *task, UfoBuffer **inputs, UfoRequisition *requisition) { UfoPadding2DTaskPrivate *priv; UfoRequisition in_req; priv = UFO_PADDING_2D_TASK_GET_PRIVATE (task); ufo_buffer_get_requisition (inputs[0], &in_req); priv->in_width = in_req.dims[0]; priv->in_height = in_req.dims[1]; requisition->n_dims = 2; requisition->dims[0] = priv->out_width = priv->xl + priv->in_width + priv->xr; requisition->dims[1] = priv->out_height = priv->yt + priv->in_height + priv->yb; priv->global_work_size_small[0] = (size_t) priv->in_width; priv->global_work_size_small[1] = (size_t) priv->in_height; priv->global_work_size_large[0] = requisition->dims[0]; priv->global_work_size_large[1] = requisition->dims[1]; } static guint ufo_padding_2d_task_get_num_inputs (UfoTask *task) { return 1; } static guint ufo_padding_2d_task_get_num_dimensions (UfoTask *task, guint input) { g_return_val_if_fail (input == 0, 0); return 2; } static UfoTaskMode ufo_padding_2d_task_get_mode (UfoTask *task) { return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU; } static UfoNode * ufo_padding_2d_task_copy_real (UfoNode *node, GError **error) { UfoPadding2DTask *orig; UfoPadding2DTask *copy; orig = UFO_PADDING_2D_TASK (node); copy = UFO_PADDING_2D_TASK (ufo_padding_2d_task_new ()); copy->priv->xl = orig->priv->xl; copy->priv->xr = orig->priv->xr; copy->priv->yb = orig->priv->yb; copy->priv->yt = orig->priv->yt; copy->priv->mode = orig->priv->mode; copy->priv->pconst = orig->priv->pconst; return UFO_NODE (copy); } static gboolean ufo_padding_2d_task_equal_real (UfoNode *n1, UfoNode *n2) { g_return_val_if_fail (UFO_IS_PADDING_2D_TASK (n1) && UFO_IS_PADDING_2D_TASK (n2), FALSE); return TRUE; } static void ufo_padding_2d_task_finalize (GObject *object) { UfoPadding2DTaskPrivate *priv; priv = UFO_PADDING_2D_TASK_GET_PRIVATE (object); if (priv->kernel_brep) { UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->kernel_brep)); priv->kernel_brep = NULL; } if (priv->kernel_cpyimg) { UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->kernel_cpyimg)); priv->kernel_cpyimg = NULL; } if (priv->kernel_iconst) { UFO_RESOURCES_CHECK_CLERR (clReleaseKernel (priv->kernel_iconst)); priv->kernel_iconst = NULL; } G_OBJECT_CLASS (ufo_padding_2d_task_parent_class)->finalize (object); } static void ufo_task_interface_init (UfoTaskIface *iface) { iface->setup = ufo_padding_2d_task_setup; iface->get_requisition = ufo_padding_2d_task_get_requisition; iface->get_num_inputs = ufo_padding_2d_task_get_num_inputs; iface->get_num_dimensions = ufo_padding_2d_task_get_num_dimensions; iface->get_mode = ufo_padding_2d_task_get_mode; iface->process = ufo_padding_2d_task_process; } static void ufo_padding_2d_task_set_property (GObject *object, guint property_id, const GValue *value, GParamSpec *pspec) { UfoPadding2DTaskPrivate *priv = UFO_PADDING_2D_TASK_GET_PRIVATE (object); switch (property_id) { case PROP_XL: priv->xl = g_value_get_uint(value); break; case PROP_XR: priv->xr = g_value_get_uint(value); break; case PROP_YT: priv->yt = g_value_get_uint(value); break; case PROP_YB: priv->yb = g_value_get_uint(value); break; case PROP_MODE: if (!g_strcmp0(g_value_get_string(value), "zero")) priv->mode = PADDING_ZERO; else if (!g_strcmp0(g_value_get_string(value), "const")) priv->mode = PADDING_CONST; else if (!g_strcmp0(g_value_get_string(value), "gavg")) priv->mode = PADDING_GAVG; else if (!g_strcmp0(g_value_get_string(value), "brep")) priv->mode = PADDING_BREP; else G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec); break; case PROP_PCONST: priv->pconst = (float) g_value_get_double(value); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec); break; } } static void ufo_padding_2d_task_get_property (GObject *object, guint property_id, GValue *value, GParamSpec *pspec) { UfoPadding2DTaskPrivate *priv = UFO_PADDING_2D_TASK_GET_PRIVATE (object); switch (property_id) { case PROP_XL: g_value_set_uint(value, priv->xl); break; case PROP_XR: g_value_set_uint(value, priv->xr); break; case PROP_YT: g_value_set_uint(value, priv->yt); break; case PROP_YB: g_value_set_uint(value, priv->yb); break; case PROP_MODE: switch (priv->mode) { case PADDING_ZERO: g_value_set_string(value, "zero"); break; case PADDING_CONST: g_value_set_string(value, "const"); break; case PADDING_GAVG: g_value_set_string(value, "gavg"); break; case PADDING_BREP: g_value_set_string(value, "brep"); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec); break; } break; case PROP_PCONST: g_value_set_double(value, priv->pconst); break; default: G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec); break; } } static void ufo_padding_2d_task_class_init (UfoPadding2DTaskClass *klass) { GObjectClass *oclass; UfoNodeClass *node_class; oclass = G_OBJECT_CLASS (klass); node_class = UFO_NODE_CLASS (klass); oclass->finalize = ufo_padding_2d_task_finalize; oclass->set_property = ufo_padding_2d_task_set_property; oclass->get_property = ufo_padding_2d_task_get_property; properties[PROP_XL] = g_param_spec_uint("xl", "Number of additional pixel on the left hand image side", "Number of additional pixel on the left hand image side", 0, 16384, 1, G_PARAM_READWRITE); properties[PROP_XR] = g_param_spec_uint("xr", "Number of additional pixel on the right hand image side", "Number of additional pixel on the right hand image side", 0, 16384, 1, G_PARAM_READWRITE); properties[PROP_YT] = g_param_spec_uint("yt", "Number of additional pixel on the top image side", "Number of additional pixel on the top image side", 0, 16384, 1, G_PARAM_READWRITE); properties[PROP_YB] = g_param_spec_uint("yb", "Number of additional pixel on the bottom image side", "Number of additional pixel on the bottom image side", 0, 16384, 1, G_PARAM_READWRITE); properties[PROP_MODE] = g_param_spec_string("mode", "Padding mode can be 'zero', 'const', 'gavg' or 'brep' ", "Padding mode can be 'zero', 'const', 'gavg' or 'brep' ", "zero", G_PARAM_READWRITE); properties[PROP_PCONST] = g_param_spec_double("pconst", "Padding constant", "Padding constant", -320000.0, 320000.0, 0.0, G_PARAM_READWRITE); for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++) g_object_class_install_property (oclass, i, properties[i]); node_class->copy = ufo_padding_2d_task_copy_real; node_class->equal = ufo_padding_2d_task_equal_real; g_type_class_add_private(klass, sizeof(UfoPadding2DTaskPrivate)); } static void ufo_padding_2d_task_init (UfoPadding2DTask *self) { UfoPadding2DTaskPrivate *priv; self->priv = priv = UFO_PADDING_2D_TASK_GET_PRIVATE (self); priv->xl = 1; priv->xr = 1; priv->yt = 1; priv->yb = 1; priv->mode = PADDING_ZERO; priv->pconst = 0.0; priv->kernel_iconst = NULL; priv->kernel_cpyimg = NULL; priv->kernel_brep = NULL; }