static gboolean cl_stretch_contrast (cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, cl_float4 min, cl_float4 diff) { cl_int cl_err = 0; cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 1, sizeof(cl_mem), (void*)&out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 2, sizeof(cl_float4), (void*)&min); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 3, sizeof(cl_float4), (void*)&diff); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[2], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
/* OpenCL processing function */ static cl_int cl_process (GeglOperation *op, cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint level) { /* Retrieve a pointer to GeglChantO structure which contains all the * chanted properties */ GeglChantO *o = GEGL_CHANT_PROPERTIES (op); gfloat in_range; gfloat out_range; gfloat in_offset; gfloat out_offset; gfloat scale; cl_int cl_err = 0; in_offset = o->in_low * 1.0; out_offset = o->out_low * 1.0; in_range = o->in_high-o->in_low; out_range = o->out_high-o->out_low; if (in_range == 0.0) in_range = 0.00000001; scale = out_range/in_range; if (!cl_data) { const char *kernel_name[] = {"kernel_levels", NULL}; cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name); } if (!cl_data) return 1; cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&in_offset); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&out_offset); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&scale); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); if (cl_err != CL_SUCCESS) return cl_err; return cl_err; }
/* OpenCL processing function */ static gboolean cl_process (GeglOperation *op, cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, int level) { /* Retrieve a pointer to GeglProperties structure which contains all the * chanted properties */ GeglProperties *o = GEGL_PROPERTIES (op); const gfloat *coeffs = o->user_data; cl_int cl_err = 0; if (! coeffs) { coeffs = o->user_data = preprocess (o); } if (!cl_data) { const char *kernel_name[] = {"gegl_color_temperature", NULL}; cl_data = gegl_cl_compile_and_build (color_temperature_cl_source, kernel_name); } if (!cl_data) return 1; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&coeffs[0]); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&coeffs[1]); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&coeffs[2]); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
static gboolean cl_process (GeglOperation *operation, cl_mem in, cl_mem out, size_t global_worksize, const GeglRectangle *roi, gint level) { GeglProperties *o = GEGL_PROPERTIES (operation); gfloat color[4]; gegl_color_get_pixel (o->color, babl_format ("R'G'B'A float"), color); if (!cl_data) { const char *kernel_name[] = {"cl_color_to_alpha",NULL}; cl_data = gegl_cl_compile_and_build (color_to_alpha_cl_source, kernel_name); } if (!cl_data) return TRUE; else { cl_int cl_err = 0; cl_float4 f_color; f_color.s[0] = color[0]; f_color.s[1] = color[1]; f_color.s[2] = color[2]; f_color.s[3] = color[3]; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float4),(void*)&f_color); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); CL_CHECK; } return FALSE; error: return TRUE; }
/* OpenCL processing function */ static cl_int cl_process (GeglOperation *op, cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint level) { /* Retrieve a pointer to GeglProperties structure which contains all the * chanted properties */ GeglProperties *o = GEGL_PROPERTIES (op); gfloat black_level = (gfloat) o->black_level; gfloat diff; gfloat exposure_negated = (gfloat) -o->exposure; gfloat gain; gfloat white; cl_int cl_err = 0; if (!cl_data) { const char *kernel_name[] = {"kernel_exposure", NULL}; cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name); } if (!cl_data) return 1; white = exp2f (exposure_negated); diff = MAX (white - black_level, 0.01); gain = 1.0f / diff; cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&black_level); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&gain); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); if (cl_err != CL_SUCCESS) return cl_err; return cl_err; }
static cl_int cl_edge_laplace (cl_mem in_tex, cl_mem aux_tex, cl_mem out_tex, const GeglRectangle *src_rect, const GeglRectangle *roi, gint radius) { cl_int cl_err = 0; size_t global_ws[2]; if (!cl_data) { const char *kernel_name[] = {"pre_edgelaplace", "knl_edgelaplace", NULL}; cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name); } if (!cl_data) return 1; global_ws[0] = roi->width; global_ws[1] = roi->height; cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&aux_tex); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, global_ws, NULL, 0, NULL, NULL); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue()); if (CL_SUCCESS != cl_err) return cl_err; cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&aux_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&out_tex); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[1], 2, NULL, global_ws, NULL, 0, NULL, NULL); if (cl_err != CL_SUCCESS) return cl_err; return cl_err; }
/* OpenCL processing function */ static cl_int cl_process (GeglOperation *op, cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, int level) { /* Retrieve a pointer to GeglChantO structure which contains all the * chanted properties */ GeglChantO *o = GEGL_CHANT_PROPERTIES (op); const gfloat *coeffs = o->chant_data; cl_int cl_err = 0; if (! coeffs) { coeffs = o->chant_data = preprocess (o); } if (!cl_data) { const char *kernel_name[] = {"kernel_temp", NULL}; cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name); } if (!cl_data) return 1; cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&coeffs[0]); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&coeffs[1]); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&coeffs[2]); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); if (cl_err != CL_SUCCESS) return cl_err; return cl_err; }
static gboolean cl_process (GeglOperation *op, cl_mem in_tex, cl_mem aux_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint level) { cl_int cl_err = 0; int kernel; gfloat value; if (!cl_data) { const char *kernel_name[] = {"gegl_opacity_RaGaBaA_float", "gegl_opacity_RGBA_float", NULL}; cl_data = gegl_cl_compile_and_build (opacity_cl_source, kernel_name); } if (!cl_data) return TRUE; value = GEGL_CHANT_PROPERTIES (op)->value; kernel = (GEGL_CHANT_PROPERTIES (op)->chant_data == NULL)? 0 : 1; cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 1, sizeof(cl_mem), (aux_tex)? (void*)&aux_tex : NULL); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 2, sizeof(cl_mem), (void*)&out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 3, sizeof(cl_float), (void*)&value); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[kernel], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
/* OpenCL processing function */ static cl_int cl_process (GeglOperation *op, cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint level) { /* Retrieve a pointer to GeglChantO structure which contains all the * chanted properties */ GeglChantO *o = GEGL_CHANT_PROPERTIES (op); gfloat gain = powf(2.0, o->exposure); gfloat offset = o->offset; gfloat gamma = 1.0 / o->gamma; cl_int cl_err = 0; if (!cl_data) { const char *kernel_name[] = {"kernel_exposure", NULL}; cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name); } if (!cl_data) return 1; cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&gain); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&offset); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&gamma); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); if (cl_err != CL_SUCCESS) return cl_err; return cl_err; }
static gboolean cl_bilateral_filter (cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gfloat radius, gfloat preserve) { cl_int cl_err = 0; size_t global_ws[2]; if (!cl_data) { const char *kernel_name[] = {"bilateral_filter", NULL}; cl_data = gegl_cl_compile_and_build (bilateral_filter_cl_source, kernel_name); } if (!cl_data) return TRUE; global_ws[0] = roi->width; global_ws[1] = roi->height; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&radius); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&preserve); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, global_ws, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
static cl_int cl_mono_mixer(cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gfloat red, gfloat green, gfloat blue) { cl_int cl_err = 0; if (!cl_data) { const char *kernel_name[] = {"Mono_mixer_cl", NULL}; cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name); } if (!cl_data) return 0; { cl_float4 color; color.s[0] = red; color.s[1] = green; color.s[2] = blue; color.s[3] = 1.0f; cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_float4), (void*)&color); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&out_tex); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); } return cl_err; }
static gboolean cl_edge_sobel (cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gboolean horizontal, gboolean vertical, gboolean keep_signal, gboolean has_alpha) { const size_t gbl_size[2] = {roi->width, roi->height}; cl_int n_horizontal = horizontal; cl_int n_vertical = vertical; cl_int n_keep_signal = keep_signal; cl_int n_has_alpha = has_alpha; cl_int cl_err = 0; if (!cl_data) { const char *kernel_name[] = {"kernel_edgesobel", NULL}; cl_data = gegl_cl_compile_and_build (edge_sobel_cl_source, kernel_name); } if (!cl_data) return TRUE; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&n_horizontal); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&n_vertical); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&n_keep_signal); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&n_has_alpha); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[0], 2, NULL, gbl_size, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
static gboolean cl_snn_mean (cl_mem in_tex, cl_mem out_tex, const GeglRectangle *src_rect, const GeglRectangle *roi, gint radius, gint pairs) { cl_int cl_err = 0; size_t global_ws[2]; if (!cl_data) { const char *kernel_name[] = {"snn_mean", NULL}; cl_data = gegl_cl_compile_and_build (snn_mean_cl_source, kernel_name); } if (!cl_data) return TRUE; global_ws[0] = roi->width; global_ws[1] = roi->height; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int), (void*)&src_rect->width); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&src_rect->height); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_mem), (void*)&out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&radius); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&pairs); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, global_ws, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
static gboolean cl_process (GeglOperation *self, cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint level) { GeglChantO *o = GEGL_CHANT_PROPERTIES (self); gint num_sampling_points; gdouble *xs, *ys; gfloat *ysf = NULL; cl_mem cl_curve = NULL; cl_ulong cl_max_constant_size; cl_int cl_err = 0; num_sampling_points = o->sampling_points; if (!cl_data) { const char *kernel_name[] = {"cl_contrast_curve",NULL}; cl_data = gegl_cl_compile_and_build (contrast_curve_cl_source, kernel_name); } if (!cl_data) return TRUE; if (num_sampling_points > 0) { xs = g_new (gdouble, num_sampling_points); ys = g_new (gdouble, num_sampling_points); gegl_curve_calc_values (o->curve, 0.0, 1.0, num_sampling_points, xs, ys); g_free (xs); /*We need to downscale the array to pass it to the GPU*/ ysf = g_new (gfloat, num_sampling_points); copy_double_array_to_float_array (ys, ysf, num_sampling_points); g_free (ys); cl_err = gegl_clGetDeviceInfo (gegl_cl_get_device (), CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof (cl_ulong), &cl_max_constant_size, NULL); CL_CHECK; GEGL_NOTE (GEGL_DEBUG_OPENCL, "Max Constant Mem Size: %lu bytes", (unsigned long) cl_max_constant_size); if (sizeof (cl_float) * num_sampling_points < cl_max_constant_size) { cl_curve = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY, num_sampling_points * sizeof (cl_float), ysf, &cl_err); CL_CHECK; cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 0, sizeof (cl_mem), (void*) &in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 1, sizeof (cl_mem), (void*) &out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 2, sizeof (cl_mem), (void*) &cl_curve); CL_CHECK; cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 3, sizeof (gint), (void*) &num_sampling_points); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); CL_CHECK; cl_err = gegl_clFinish (gegl_cl_get_command_queue ()); CL_CHECK; cl_err = gegl_clReleaseMemObject (cl_curve); CL_CHECK_ONLY (cl_err); } else { /*If the curve size doesn't fit constant memory is better to use CPU*/ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Not enough constant memory for the curve"); g_free (ysf); return TRUE; } g_free (ysf); return FALSE; error: if (ysf) g_free (ysf); if (cl_curve) gegl_clReleaseMemObject (cl_curve); return TRUE; } else /*If the curve doesn't have a lookup table is better to use CPU*/ { GEGL_NOTE (GEGL_DEBUG_OPENCL, "Curve not suitable to be computed in the GPU"); return TRUE; } }
static void motion_blur_cl (GeglBuffer *src, const GeglRectangle *src_rect, GeglBuffer *dst, const GeglRectangle *dst_rect, const int num_steps, const float offset_x, const float offset_y) { const Babl * in_format = babl_format("RaGaBaA float"); const Babl *out_format = babl_format("RaGaBaA float"); /* AreaFilter general processing flow. Loading data and making the necessary color space conversion. */ #include "gegl-cl-operation-area-filter-fw1.h" /////////////////////////////////////////////////////////////////////////// /* Algorithm specific processing flow. Build kernels, setting parameters, and running them. */ if (!cl_data) { const char *kernel_name[] = { "motion_blur_CL", NULL }; cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name); } if (!cl_data) CL_ERROR; cl_int cl_src_width = src_rect->width; cl_int cl_src_height = src_rect->height; cl_int cl_src_x = src_rect->x; cl_int cl_src_y = src_rect->y; cl_int cl_dst_x = dst_rect->x; cl_int cl_dst_y = dst_rect->y; cl_int cl_num_steps = num_steps; cl_float cl_offset_x = offset_x; cl_float cl_offset_y = offset_y; CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&src_mem)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 1, sizeof(cl_int), (void*)&cl_src_width)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 2, sizeof(cl_int), (void*)&cl_src_height)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 3, sizeof(cl_int), (void*)&cl_src_x)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 4, sizeof(cl_int), (void*)&cl_src_y)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 5, sizeof(cl_mem), (void*)&dst_mem)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 6, sizeof(cl_int), (void*)&cl_dst_x)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 7, sizeof(cl_int), (void*)&cl_dst_y)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 8, sizeof(cl_int), (void*)&cl_num_steps)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 9, sizeof(cl_float), (void*)&cl_offset_x)); CL_SAFE_CALL(errcode = gegl_clSetKernelArg( cl_data->kernel[0], 10, sizeof(cl_float), (void*)&cl_offset_y)); CL_SAFE_CALL(errcode = gegl_clEnqueueNDRangeKernel( gegl_cl_get_command_queue(), cl_data->kernel[0], 2, NULL, gbl_size, NULL, 0, NULL, NULL)); errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue()); if (CL_SUCCESS != errcode) CL_ERROR; /////////////////////////////////////////////////////////////////////////// /* AreaFilter general processing flow. Making the necessary color space conversion and Saving data. */ #include "gegl-cl-operation-area-filter-fw2.h" }
static cl_int cl_motion_blur (cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, const GeglRectangle *src_rect, gint num_steps, gfloat offset_x, gfloat offset_y) { cl_int cl_err = 0; size_t global_ws[2]; if (!cl_data) { const char *kernel_name[] = {"motion_blur_CL", NULL}; cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name); } if (!cl_data) return 1; global_ws[0] = roi->width; global_ws[1] = roi->height; cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int), (void*)&src_rect->width); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&src_rect->height); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&src_rect->x); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&src_rect->y); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_mem), (void*)&out_tex); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 6, sizeof(cl_int), (void*)&roi->x); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 7, sizeof(cl_int), (void*)&roi->y); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 8, sizeof(cl_int), (void*)&num_steps); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_float), (void*)&offset_x); cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 10, sizeof(cl_float), (void*)&offset_y); if (cl_err != CL_SUCCESS) return cl_err; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, global_ws, NULL, 0, NULL, NULL); if (cl_err != CL_SUCCESS) return cl_err; return cl_err; }
static gboolean cl_process (GeglOperation *operation, cl_mem in_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint level) { GeglChantO *o = GEGL_CHANT_PROPERTIES (operation); gfloat scale; gfloat radius0, radius1; gint roi_x, roi_y,x; gint midx, midy; GeglRectangle *bounds = gegl_operation_source_get_bounding_box (operation, "input"); gfloat length = hypot (bounds->width, bounds->height)/2; gfloat rdiff; gfloat cost, sint; gfloat color[4]; scale = bounds->width / (1.0 * bounds->height); scale = scale * (o->proportion) + 1.0 * (1.0-o->proportion); scale *= aspect_to_scale (o->squeeze); length = (bounds->width/2.0); if (scale > 1.0) length /= scale; gegl_color_get_pixel (o->color, babl_format ("RGBA float"), color); for (x=0; x<3; x++) /* premultiply */ color[x] *= color[3]; radius0 = o->radius * (1.0-o->softness); radius1 = o->radius; rdiff = radius1-radius0; if (fabs (rdiff) < 0.0001) rdiff = 0.0001; midx = bounds->x + bounds->width * o->x; midy = bounds->y + bounds->height * o->y; roi_x = roi->x; roi_y = roi->y; /* constant for all pixels */ cost = cos(-o->rotation * (G_PI*2/360.0)); sint = sin(-o->rotation * (G_PI*2/360.0)); if (!cl_data) { const char *kernel_name[] = {"vignette_cl",NULL}; cl_data = gegl_cl_compile_and_build (vignette_cl_source, kernel_name); } if (!cl_data) return TRUE; { const size_t gbl_size[2] = {roi->width, roi->height}; gint shape = (gint) o->shape; gfloat gamma = o->gamma; cl_int cl_err = 0; cl_float4 f_color; f_color.s[0] = color[0]; f_color.s[1] = color[1]; f_color.s[2] = color[2]; f_color.s[3] = color[3]; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float4),(void*)&f_color); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&scale); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&cost); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_float), (void*)&sint); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 6, sizeof(cl_int), (void*)&roi_x); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 7, sizeof(cl_int), (void*)&roi_y); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 8, sizeof(cl_int), (void*)&midx); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_int), (void*)&midy); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 10, sizeof(cl_int), (void*)&shape); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 11, sizeof(cl_float), (void*)&gamma); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 12, sizeof(cl_float), (void*)&length); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 13, sizeof(cl_float), (void*)&radius0); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 14, sizeof(cl_float), (void*)&rdiff); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, gbl_size, NULL, 0, NULL, NULL); CL_CHECK; } return FALSE; error: return TRUE; }
gboolean gegl_operation_cl_set_kernel_args (GeglOperation *operation, cl_kernel kernel, gint *p, cl_int *err) { GParamSpec **self; GParamSpec **parent; guint n_self; guint n_parent; gint prop_no; self = g_object_class_list_properties ( G_OBJECT_CLASS (g_type_class_ref (G_OBJECT_CLASS_TYPE (GEGL_OPERATION_GET_CLASS(operation)))), &n_self); parent = g_object_class_list_properties ( G_OBJECT_CLASS (g_type_class_ref (GEGL_TYPE_OPERATION)), &n_parent); for (prop_no=0;prop_no<n_self;prop_no++) { gint parent_no; gboolean found=FALSE; for (parent_no=0;parent_no<n_parent;parent_no++) if (self[prop_no]==parent[parent_no]) found=TRUE; /* only print properties if we are an addition compared to * GeglOperation */ /* Removing pads */ if (!strcmp(g_param_spec_get_name (self[prop_no]), "input") || !strcmp(g_param_spec_get_name (self[prop_no]), "output") || !strcmp(g_param_spec_get_name (self[prop_no]), "aux")) continue; if (!found) { if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_DOUBLE)) { gdouble value; cl_float v; g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL); v = value; *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_float), (void*)&v); } else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_FLOAT)) { gfloat value; cl_float v; g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL); v = value; *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_float), (void*)&v); } else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_INT)) { gint value; cl_int v; g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL); v = value; *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_int), (void*)&v); } else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_BOOLEAN)) { gboolean value; cl_bool v; g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL); v = value; *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_bool), (void*)&v); } else { g_error ("Unsupported OpenCL kernel argument"); return FALSE; } } } if (self) g_free (self); if (parent) g_free (parent); return TRUE; }
static gboolean cl_buffer_get_min_max (cl_mem in_tex, size_t global_worksize, const GeglRectangle *roi, gfloat min[4], gfloat max[4]) { cl_int cl_err = 0; size_t local_ws, max_local_ws; size_t work_groups; size_t global_ws; cl_mem cl_aux_min = NULL; cl_mem cl_aux_max = NULL; cl_mem cl_min_max = NULL; cl_int n_pixels = (cl_int)global_worksize; cl_float4 min_max_buf[2]; if (global_worksize < 1) { min[0] = min[1] = min[2] = min[3] = G_MAXFLOAT; max[0] = max[1] = max[2] = max[3] = -G_MAXFLOAT; return FALSE; } cl_err = gegl_clGetDeviceInfo (gegl_cl_get_device (), CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof (size_t), &max_local_ws, NULL); CL_CHECK; max_local_ws = MIN (max_local_ws, MIN (cl_data->work_group_size[0], cl_data->work_group_size[1])); /* Needs to be a power of two */ local_ws = 256; while (local_ws > max_local_ws) local_ws /= 2; work_groups = MIN ((global_worksize + local_ws - 1) / local_ws, local_ws); global_ws = work_groups * local_ws; cl_aux_min = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_READ_WRITE, local_ws * sizeof(cl_float4), NULL, &cl_err); CL_CHECK; cl_aux_max = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_READ_WRITE, local_ws * sizeof(cl_float4), NULL, &cl_err); CL_CHECK; cl_min_max = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_WRITE_ONLY, 2 * sizeof(cl_float4), NULL, &cl_err); CL_CHECK; /* The full initialization is done in the two_stages_local_min_max_reduce kernel */ #if 0 cl_err = gegl_clSetKernelArg(cl_data->kernel[3], 0, sizeof(cl_mem), (void*)&cl_aux_min); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[3], 1, sizeof(cl_mem), (void*)&cl_aux_max); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[3], 1, NULL, &local_ws, &local_ws, 0, NULL, NULL); CL_CHECK; #endif cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&cl_aux_min); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&cl_aux_max); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float4) * local_ws, NULL); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float4) * local_ws, NULL); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&n_pixels); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 1, NULL, &global_ws, &local_ws, 0, NULL, NULL); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&cl_aux_min); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&cl_aux_max); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_mem), (void*)&cl_min_max); CL_CHECK; /* Only one work group */ cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[1], 1, NULL, &local_ws, &local_ws, 0, NULL, NULL); CL_CHECK; /* Read the memory buffer, probably better to keep it in GPU memory */ cl_err = gegl_clEnqueueReadBuffer (gegl_cl_get_command_queue (), cl_min_max, CL_TRUE, 0, 2 * sizeof (cl_float4), &min_max_buf, 0, NULL, NULL); CL_CHECK; min[0] = min_max_buf[0].x; min[1] = min_max_buf[0].y; min[2] = min_max_buf[0].z; min[3] = min_max_buf[0].w; max[0] = min_max_buf[1].x; max[1] = min_max_buf[1].y; max[2] = min_max_buf[1].z; max[3] = min_max_buf[1].w; cl_err = gegl_clReleaseMemObject (cl_aux_min); CL_CHECK_ONLY (cl_err); cl_err = gegl_clReleaseMemObject (cl_aux_max); CL_CHECK_ONLY (cl_err); cl_err = gegl_clReleaseMemObject (cl_min_max); CL_CHECK_ONLY (cl_err); return FALSE; error: if (cl_aux_min) gegl_clReleaseMemObject (cl_aux_min); if (cl_aux_max) gegl_clReleaseMemObject (cl_aux_max); if (cl_min_max) gegl_clReleaseMemObject (cl_min_max); 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; }
static gboolean cl_box_max (cl_mem in_tex, cl_mem aux_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint radius) { cl_int cl_err = 0; size_t global_ws_hor[2], global_ws_ver[2]; size_t local_ws_hor[2], local_ws_ver[2]; if (!cl_data) { const char *kernel_name[] = {"kernel_max_hor", "kernel_max_ver", NULL}; cl_data = gegl_cl_compile_and_build (box_max_cl_source, kernel_name); } if (!cl_data) return TRUE; local_ws_hor[0] = 1; local_ws_hor[1] = 256; global_ws_hor[0] = roi->height + 2 * radius; global_ws_hor[1] = ((roi->width + local_ws_hor[1] -1)/local_ws_hor[1]) * local_ws_hor[1]; local_ws_ver[0] = 1; local_ws_ver[1] = 256; global_ws_ver[0] = roi->height; global_ws_ver[1] = ((roi->width + local_ws_ver[1] -1)/local_ws_ver[1]) * local_ws_ver[1]; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&aux_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&roi->width); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&radius); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, global_ws_hor, local_ws_hor, 0, NULL, NULL); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&aux_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&out_tex); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_int), (void*)&roi->width); CL_CHECK; cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 3, sizeof(cl_int), (void*)&radius); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[1], 2, NULL, global_ws_ver, local_ws_ver, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }