123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470 |
- /**
- * 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 <OpenCL/cl.h>
- #else
- #include <CL/cl.h>
- #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;
- }
|