static gboolean cl_process (GeglOperation *operation, GeglBuffer *input, GeglBuffer *output, const GeglRectangle *result) { const Babl *in_format = gegl_operation_get_format (operation, "input"); const Babl *out_format = gegl_operation_get_format (operation, "output"); gint err; gint j; cl_int cl_err; GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation); GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE); gint read = gegl_buffer_cl_iterator_add_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ, op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE); gint aux = gegl_buffer_cl_iterator_add_2 (i, NULL, result, in_format, GEGL_CL_BUFFER_AUX, op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE); while (gegl_buffer_cl_iterator_next (i, &err)) { if (err) return FALSE; for (j=0; j < i->n; j++) { cl_err = cl_edge_laplace(i->tex[read][j], i->tex[aux][j], i->tex[0][j], &i->roi[read][j], &i->roi[0][j], LAPLACE_RADIUS); if (cl_err != CL_SUCCESS) { g_warning("[OpenCL] Error in gegl:edge-laplace: %s", gegl_cl_errstring(cl_err)); return FALSE; } } } return TRUE; }
static gboolean cl_process (GeglOperation *operation, GeglBuffer *input, GeglBuffer *output, const GeglRectangle *result) { const Babl *in_format = gegl_operation_get_format (operation, "input"); const Babl *out_format = gegl_operation_get_format (operation, "output"); gint err; gint j; cl_int cl_err; GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation); GeglChantO *o = GEGL_CHANT_PROPERTIES (operation); GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE); gint read = gegl_buffer_cl_iterator_add_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ, op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE); while (gegl_buffer_cl_iterator_next (i, &err)) { if (err) return FALSE; for (j=0; j < i->n; j++) { cl_err = cl_bilateral_filter(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], ceil(o->blur_radius), o->edge_preservation); if (cl_err != CL_SUCCESS) { g_warning("[OpenCL] Error in gegl:bilateral-filter: %s", gegl_cl_errstring(cl_err)); return FALSE; } } } return TRUE; }
static gboolean cl_process (GeglOperation *operation, GeglBuffer *input, GeglBuffer *output, const GeglRectangle *result) { const Babl *in_format = gegl_operation_get_format (operation, "input"); const Babl *out_format = gegl_operation_get_format (operation, "output"); gint err; gint j; cl_int cl_err; GeglChantO *o = GEGL_CHANT_PROPERTIES (operation); GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE); gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE); while (gegl_buffer_cl_iterator_next (i, &err)) { if (err) return FALSE; for (j=0; j < i->n; j++) { cl_err = cl_mono_mixer(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], o->red ,o->green , o->blue); if (cl_err != CL_SUCCESS) { g_warning("[OpenCL] Error in gegl:mono-mixer: %s", gegl_cl_errstring(cl_err)); return FALSE; } } } return TRUE; }
/* XXX: same program_source with different kernel_name[], context or device * will retrieve the same key */ gegl_cl_run_data * gegl_cl_compile_and_build (const char *program_source, const char *kernel_name[]) { gint errcode; gegl_cl_run_data *cl_data = NULL; if ((cl_data = (gegl_cl_run_data *)g_hash_table_lookup(cl_program_hash, program_source)) == NULL) { size_t length = strlen(program_source); gint i; guint kernel_n = 0; while (kernel_name[++kernel_n] != NULL); cl_data = (gegl_cl_run_data *) g_malloc(sizeof(gegl_cl_run_data)+sizeof(cl_kernel)*kernel_n); CL_SAFE_CALL( cl_data->program = gegl_clCreateProgramWithSource(gegl_cl_get_context(), 1, &program_source, &length, &errcode) ); errcode = gegl_clBuildProgram(cl_data->program, 0, NULL, NULL, NULL, NULL); if (errcode != CL_SUCCESS) { char buffer[2000]; CL_SAFE_CALL( errcode = gegl_clGetProgramBuildInfo(cl_data->program, gegl_cl_get_device_id(), CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, NULL) ); g_warning("OpenCL Build Error:%s\n%s", gegl_cl_errstring(errcode), buffer); return NULL; } else { g_printf("[OpenCL] Compiling successful\n"); } for (i=0; i<kernel_n; i++) CL_SAFE_CALL( cl_data->kernel[i] = gegl_clCreateKernel(cl_data->program, kernel_name[i], &errcode) ); g_hash_table_insert(cl_program_hash, g_strdup (program_source), (void*)cl_data); } return cl_data; }
static gboolean cl_process (GeglOperation *operation, GeglBuffer *input, GeglBuffer *output, const GeglRectangle *result, const GeglRectangle *src_rect) { const Babl *in_format = gegl_operation_get_format (operation, "input"); const Babl *out_format = gegl_operation_get_format (operation, "output"); gint err; gint j; cl_int cl_err; GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation); GeglChantO *o = GEGL_CHANT_PROPERTIES (operation); gdouble theta = o->angle * G_PI / 180.0; gfloat offset_x = (gfloat)(o->length * cos(theta)); gfloat offset_y = (gfloat)(o->length * sin(theta)); gint num_steps = (gint)ceil(o->length) + 1; GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE); gint read = gegl_buffer_cl_iterator_add_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ, op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE); while (gegl_buffer_cl_iterator_next (i, &err)) { if (err) return FALSE; for (j=0; j < i->n; j++) { cl_err = cl_motion_blur(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], &i->roi[read][j], num_steps, offset_x, offset_y); if (cl_err != CL_SUCCESS) { g_warning("[OpenCL] Error in gegl:motion-blur: %s", gegl_cl_errstring(cl_err)); return FALSE; } } } return TRUE; }
static gboolean process (GeglOperation *operation, GeglBuffer *input, const GeglRectangle *result, gint level) { GeglChantO *o = GEGL_CHANT_PROPERTIES (operation); if (o->buffer) { GeglBuffer *output = GEGL_BUFFER (o->buffer); const Babl *in_format = gegl_buffer_get_format (input); const Babl *out_format = gegl_buffer_get_format (output); if (gegl_operation_use_opencl (operation) && gegl_cl_color_supported (in_format, out_format) == GEGL_CL_COLOR_CONVERT) { size_t size; gboolean err; cl_int cl_err = 0; GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE); gint read = gegl_buffer_cl_iterator_add (i, input, result, out_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE); gegl_cl_color_babl (out_format, &size); GEGL_NOTE (GEGL_DEBUG_OPENCL, "write-buffer: " "%p %p %s %s {%d %d %d %d}", input, output, babl_get_name (in_format), babl_get_name (out_format), result->x, result->y, result->width, result->height); while (gegl_buffer_cl_iterator_next (i, &err)) { if (err) break; cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue (), i->tex[read], i->tex[0], 0, 0, i->size[0] * size, 0, NULL, NULL); if (cl_err != CL_SUCCESS) { GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring (cl_err)); break; } } if (cl_err || err) gegl_buffer_copy (input, result, output, result); } else gegl_buffer_copy (input, result, output, result); gegl_buffer_flush (output); } return TRUE; }
static gboolean gegl_operation_point_filter_cl_process (GeglOperation *operation, GeglBuffer *input, GeglBuffer *output, const GeglRectangle *result, gint level) { const Babl *in_format = gegl_operation_get_format (operation, "input"); const Babl *out_format = gegl_operation_get_format (operation, "output"); GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation); GeglOperationPointFilterClass *point_filter_class = GEGL_OPERATION_POINT_FILTER_GET_CLASS (operation); GeglBufferClIterator *iter = NULL; cl_int cl_err = 0; gboolean err; /* non-texturizable format! */ if (!gegl_cl_color_babl (in_format, NULL) || !gegl_cl_color_babl (out_format, NULL)) { GEGL_NOTE (GEGL_DEBUG_OPENCL, "Non-texturizable format!"); return FALSE; } GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_FILTER: %s", operation_class->name); /* Process */ iter = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE); gegl_buffer_cl_iterator_add (iter, input, result, in_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE); while (gegl_buffer_cl_iterator_next (iter, &err)) { if (err) return FALSE; if (point_filter_class->cl_process) { err = point_filter_class->cl_process (operation, iter->tex[1], iter->tex[0], iter->size[0], &iter->roi[0], level); if (err) { GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", operation_class->name); gegl_buffer_cl_iterator_stop (iter); return FALSE; } } else if (operation_class->cl_data) { gint p = 0; GeglClRunData *cl_data = operation_class->cl_data; cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&iter->tex[1]); CL_CHECK; cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&iter->tex[0]); CL_CHECK; gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &iter->size[0], NULL, 0, NULL, NULL); CL_CHECK; } else { g_warning ("OpenCL support enabled, but no way to execute"); gegl_buffer_cl_iterator_stop (iter); return FALSE; } } return TRUE; error: GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring (cl_err)); if (iter) gegl_buffer_cl_iterator_stop (iter); return FALSE; }