|
@@ -19,164 +19,114 @@
|
|
|
*/
|
|
|
|
|
|
struct _UfoFilterLaminoFTConvPrivate {
|
|
|
- // float example;
|
|
|
- cl_kernel kernel;
|
|
|
+ cl_kernel kernel;
|
|
|
+ cl_mem filter_mem;
|
|
|
+ size_t global_work_size[2];
|
|
|
+ guint img_width;
|
|
|
+ guint img_height;
|
|
|
};
|
|
|
|
|
|
-G_DEFINE_TYPE(UfoFilterLaminoFTConv, ufo_filter_lamino_ft_conv, UFO_TYPE_FILTER)
|
|
|
-
|
|
|
-#define UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_LAMINO_FT_CONV, UfoFilterLaminoFTConvPrivate))
|
|
|
-
|
|
|
enum {
|
|
|
- PROP_0,
|
|
|
-// PROP_EXAMPLE,
|
|
|
- N_PROPERTIES
|
|
|
+ INPUT_FILTER,
|
|
|
+ INPUT_IMAGE
|
|
|
};
|
|
|
|
|
|
-// static GParamSpec *lamino_ft_conv_properties[N_PROPERTIES] = { NULL, };
|
|
|
+G_DEFINE_TYPE(UfoFilterLaminoFTConv, ufo_filter_lamino_ft_conv, UFO_TYPE_FILTER)
|
|
|
|
|
|
+#define UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_LAMINO_FT_CONV, UfoFilterLaminoFTConvPrivate))
|
|
|
|
|
|
-static void ufo_filter_lamino_ft_conv_initialize(UfoFilter *filter)
|
|
|
+static void
|
|
|
+ufo_filter_lamino_ft_conv_initialize (UfoFilter *filter, UfoBuffer *params[], guint **dim_sizes, GError **error)
|
|
|
{
|
|
|
-
|
|
|
- UfoFilterLaminoFTConv *self = UFO_FILTER_LAMINO_FT_CONV(filter);
|
|
|
- UfoResourceManager *manager = ufo_resource_manager();
|
|
|
- GError *error = NULL;
|
|
|
- self->priv->kernel = ufo_resource_manager_get_kernel(manager, "lamino_ft_conv.cl", "lamino_c", &error);
|
|
|
-
|
|
|
-
|
|
|
- if (error != NULL) {
|
|
|
- g_warning("%s", error->message);
|
|
|
- g_error_free(error);
|
|
|
+ UfoFilterLaminoFTConvPrivate *priv = UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE (filter);
|
|
|
+ UfoResourceManager *manager = ufo_resource_manager ();
|
|
|
+ cl_command_queue command_queue;
|
|
|
+ cl_mem filter_mem;
|
|
|
+ guint f_width, f_height;
|
|
|
+ GError *tmp_error = NULL;
|
|
|
+
|
|
|
+ priv->kernel = ufo_resource_manager_get_kernel (manager, "lamino_ft_conv.cl", "lamino_c", &tmp_error);
|
|
|
+
|
|
|
+ if (tmp_error != NULL) {
|
|
|
+ g_propagate_error (error, tmp_error);
|
|
|
+ return;
|
|
|
}
|
|
|
-}
|
|
|
-
|
|
|
-
|
|
|
-static void ufo_filter_lamino_ft_conv_process(UfoFilter *flt)
|
|
|
-{
|
|
|
-
|
|
|
- g_return_if_fail(UFO_IS_FILTER(flt));
|
|
|
- UfoFilterLaminoFTConv *self = UFO_FILTER_LAMINO_FT_CONV(flt);
|
|
|
-
|
|
|
-
|
|
|
- UfoChannel *input_channel = ufo_filter_get_input_channel_by_name(flt, "image");
|
|
|
- UfoChannel *filter_channel = ufo_filter_get_input_channel_by_name(flt, "filter");
|
|
|
- UfoChannel *output_channel = ufo_filter_get_output_channel(flt);
|
|
|
- cl_command_queue command_queue = (cl_command_queue) ufo_filter_get_command_queue(flt);
|
|
|
-
|
|
|
- UfoBuffer *input = ufo_channel_get_input_buffer(input_channel);
|
|
|
- UfoBuffer *filter = ufo_channel_get_input_buffer(filter_channel);
|
|
|
-
|
|
|
- guint num_dims = 0;
|
|
|
- guint *dim_size = NULL;
|
|
|
- ufo_buffer_get_dimensions(input, &num_dims, &dim_size);
|
|
|
- ufo_channel_allocate_output_buffers(output_channel, 2, dim_size);
|
|
|
|
|
|
- guint *fdim_size= NULL;
|
|
|
- ufo_buffer_get_dimensions(input, &num_dims, &fdim_size);
|
|
|
+ ufo_buffer_get_2d_dimensions (params[0], &priv->img_width, &priv->img_height);
|
|
|
+ ufo_buffer_get_2d_dimensions (params[1], &f_width, &f_height);
|
|
|
|
|
|
- if( (dim_size[0] != fdim_size[0]) || (dim_size[1] != fdim_size[1]))
|
|
|
- g_error("Filter and image sizes are different");
|
|
|
-
|
|
|
- size_t global_work_size[2] = { (size_t) dim_size[0], (size_t) dim_size[1] };
|
|
|
- guint width = dim_size[0];
|
|
|
- cl_kernel kernel = self->priv->kernel;
|
|
|
-
|
|
|
- cl_mem filter_mem = (cl_mem) ufo_buffer_get_device_array(filter, command_queue);
|
|
|
-
|
|
|
- while (input != NULL)
|
|
|
- {
|
|
|
- UfoBuffer *output = ufo_channel_get_output_buffer(output_channel);
|
|
|
- cl_mem input_mem = (cl_mem) ufo_buffer_get_device_array(input, command_queue);
|
|
|
- cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array(output, command_queue);
|
|
|
-
|
|
|
- CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &input_mem));
|
|
|
- CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &filter_mem));
|
|
|
- CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &output_mem));
|
|
|
- CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 3, sizeof(int), &width));
|
|
|
-
|
|
|
- cl_event event;
|
|
|
- CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel,
|
|
|
- 2, NULL, global_work_size, NULL,
|
|
|
- 0, NULL, &event));
|
|
|
- clFinish(command_queue);
|
|
|
- //ufo_buffer_attach_event(output, event);
|
|
|
-
|
|
|
-
|
|
|
- ufo_channel_finalize_input_buffer(input_channel, input);
|
|
|
- ufo_channel_finalize_output_buffer(output_channel, output);
|
|
|
- input = ufo_channel_get_input_buffer(input_channel);
|
|
|
+ if ((priv->img_width != f_width ) || (priv->img_height != f_height)) {
|
|
|
+ g_set_error (error, UFO_FILTER_ERROR, UFO_FILTER_ERROR_INITIALIZATION,
|
|
|
+ "Filter and image sizes are different");
|
|
|
+ return;
|
|
|
}
|
|
|
|
|
|
- ufo_channel_finalize_input_buffer(filter_channel, filter);
|
|
|
- ufo_channel_finish(output_channel);
|
|
|
+ dim_sizes[0][0] = priv->img_width;
|
|
|
+ dim_sizes[0][1] = priv->img_height;
|
|
|
+ priv->global_work_size[0] = (size_t) priv->img_width;
|
|
|
+ priv->global_work_size[1] = (size_t) priv->img_height;
|
|
|
|
|
|
- g_free(dim_size);
|
|
|
- g_free(fdim_size);
|
|
|
+ command_queue = ufo_resource_manager_get_command_queue (manager, 0);
|
|
|
+ filter_mem = ufo_buffer_get_device_array (params[INPUT_FILTER], command_queue);
|
|
|
+ priv->filter_mem = ufo_resource_manager_memdup (manager, filter_mem);
|
|
|
}
|
|
|
|
|
|
-/*static void ufo_filter_lamino_ft_conv_set_property(GObject *object,
|
|
|
- guint property_id,
|
|
|
- const GValue *value,
|
|
|
- GParamSpec *pspec)
|
|
|
+static UfoEventList *
|
|
|
+ufo_filter_lamino_ft_conv_process_gpu(UfoFilter *flt, UfoBuffer *input[], UfoBuffer *output[], gpointer cmd_queue, GError **error)
|
|
|
{
|
|
|
-
|
|
|
- switch (property_id) {
|
|
|
- default:
|
|
|
- G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
|
|
|
- break;
|
|
|
- }
|
|
|
+ g_return_val_if_fail(UFO_IS_FILTER(flt), NULL);
|
|
|
+ UfoFilterLaminoFTConvPrivate *priv = UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE (flt);
|
|
|
+ UfoEventList *event_list = ufo_event_list_new (1);
|
|
|
+ cl_event *events = ufo_event_list_get_event_array (event_list);
|
|
|
+
|
|
|
+ cl_command_queue command_queue = (cl_command_queue) cmd_queue;
|
|
|
+ cl_kernel kernel = priv->kernel;
|
|
|
+ cl_mem input_mem = (cl_mem) ufo_buffer_get_device_array(input[INPUT_IMAGE], command_queue);
|
|
|
+ cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array(output[0], command_queue);
|
|
|
+
|
|
|
+ CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &input_mem));
|
|
|
+ CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &priv->filter_mem));
|
|
|
+ CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &output_mem));
|
|
|
+ CHECK_OPENCL_ERROR(clSetKernelArg(kernel, 3, sizeof(int), &priv->img_width));
|
|
|
+
|
|
|
+ CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue, kernel,
|
|
|
+ 2, NULL, priv->global_work_size, NULL,
|
|
|
+ 0, NULL, &events[0]));
|
|
|
+
|
|
|
+ clFinish(command_queue);
|
|
|
+ return event_list;
|
|
|
}
|
|
|
|
|
|
-static void ufo_filter_lamino_ft_conv_get_property(GObject *object,
|
|
|
- guint property_id,
|
|
|
- GValue *value,
|
|
|
- GParamSpec *pspec)
|
|
|
-{
|
|
|
- switch (property_id) {
|
|
|
- default:
|
|
|
- G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
|
|
|
- break;
|
|
|
- }
|
|
|
-}*/
|
|
|
-
|
|
|
-static void ufo_filter_lamino_ft_conv_class_init(UfoFilterLaminoFTConvClass *klass)
|
|
|
+static void
|
|
|
+ufo_filter_lamino_ft_conv_class_init(UfoFilterLaminoFTConvClass *klass)
|
|
|
{
|
|
|
GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
|
|
|
UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);
|
|
|
|
|
|
- //gobject_class->set_property = ufo_filter_lamino_ft_conv_set_property;
|
|
|
- //gobject_class->get_property = ufo_filter_lamino_ft_conv_get_property;
|
|
|
filter_class->initialize = ufo_filter_lamino_ft_conv_initialize;
|
|
|
- filter_class->process = ufo_filter_lamino_ft_conv_process;
|
|
|
-
|
|
|
- /* lamino_ft_conv_properties[PROP_EXAMPLE] =
|
|
|
- g_param_spec_double("example",
|
|
|
- "This is an example property",
|
|
|
- "You should definately replace this with some meaningful property",
|
|
|
- -1.0,
|
|
|
- 1.0,
|
|
|
- 1.0,
|
|
|
- G_PARAM_READWRITE);
|
|
|
-
|
|
|
- g_object_class_install_property(gobject_class, PROP_EXAMPLE, lamino_ft_conv_properties[PROP_EXAMPLE]); */
|
|
|
+ filter_class->process_gpu = ufo_filter_lamino_ft_conv_process_gpu;
|
|
|
|
|
|
g_type_class_add_private(gobject_class, sizeof(UfoFilterLaminoFTConvPrivate));
|
|
|
-
|
|
|
}
|
|
|
|
|
|
-static void ufo_filter_lamino_ft_conv_init(UfoFilterLaminoFTConv *self)
|
|
|
+static void
|
|
|
+ufo_filter_lamino_ft_conv_init(UfoFilterLaminoFTConv *self)
|
|
|
{
|
|
|
UfoFilterLaminoFTConvPrivate *priv = self->priv = UFO_FILTER_LAMINO_FT_CONV_GET_PRIVATE(self);
|
|
|
- // priv->example = 1.0;
|
|
|
priv->kernel = NULL;
|
|
|
|
|
|
- ufo_filter_register_input(UFO_FILTER(self), "filter", 2);
|
|
|
- ufo_filter_register_input(UFO_FILTER(self), "image", 2);
|
|
|
- ufo_filter_register_output(UFO_FILTER(self), "oimage", 2);
|
|
|
+ ufo_filter_register_inputs (UFO_FILTER (self),
|
|
|
+ 2, /* filter */
|
|
|
+ 2, /* input image */
|
|
|
+ NULL);
|
|
|
+
|
|
|
+ ufo_filter_register_outputs (UFO_FILTER (self),
|
|
|
+ 2, /* output image */
|
|
|
+ NULL);
|
|
|
}
|
|
|
|
|
|
-G_MODULE_EXPORT UfoFilter *ufo_filter_plugin_new(void)
|
|
|
+G_MODULE_EXPORT UfoFilter *
|
|
|
+ufo_filter_plugin_new(void)
|
|
|
{
|
|
|
return g_object_new(UFO_TYPE_FILTER_LAMINO_FT_CONV, NULL);
|
|
|
}
|