static gboolean 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_in[2]; size_t global_ws_aux[2]; if (!cl_data) { const char *kernel_name[] = {"pre_edgelaplace", "knl_edgelaplace", NULL}; cl_data = gegl_cl_compile_and_build (edge_laplace_cl_source, kernel_name); } if (!cl_data) return TRUE; global_ws_in[0] = roi->width + LAPLACE_RADIUS; global_ws_in[1] = roi->height + LAPLACE_RADIUS; global_ws_aux[0] = roi->width; global_ws_aux[1] = roi->height; cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof (cl_mem), &in_tex, sizeof (cl_mem), &aux_tex, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, global_ws_in, NULL, 0, NULL, NULL); CL_CHECK; cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1], sizeof (cl_mem), &aux_tex, sizeof (cl_mem), &out_tex, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[1], 2, NULL, global_ws_aux, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
static gboolean cl_gaussian_blur (cl_mem in_tex, cl_mem out_tex, const GeglRectangle *roi, cl_mem cl_cmatrix, gint clen, GeglOrientation orientation) { cl_int cl_err = 0; size_t global_ws[2]; gint kernel_num; if (!cl_data) { const char *kernel_name[] = {"fir_ver_blur", "fir_hor_blur", NULL}; cl_data = gegl_cl_compile_and_build (gblur_1d_cl_source, kernel_name); } if (!cl_data) return TRUE; if (orientation == GEGL_ORIENTATION_VERTICAL) kernel_num = 0; else kernel_num = 1; global_ws[0] = roi->width; global_ws[1] = roi->height; cl_err = gegl_cl_set_kernel_args (cl_data->kernel[kernel_num], sizeof(cl_mem), (void*)&in_tex, sizeof(cl_mem), (void*)&out_tex, sizeof(cl_mem), (void*)&cl_cmatrix, sizeof(cl_int), (void*)&clen, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[kernel_num], 2, NULL, global_ws, NULL, 0, NULL, NULL); CL_CHECK; cl_err = gegl_clFinish (gegl_cl_get_command_queue ()); CL_CHECK; return FALSE; error: return TRUE; }
static gboolean cl_process (GeglOperation *operation, cl_mem in_buf, cl_mem out_buf, const size_t n_pixels, const GeglRectangle *roi, gint level) { GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation); GeglClRunData *cl_data = operation_class->cl_data; GeglProperties *o = GEGL_PROPERTIES (operation); const size_t gbl_size[2] = {roi->width, roi->height}; const size_t gbl_off[2] = {roi->x, roi->y}; cl_int cl_err = 0; cl_mem filter_pat = NULL; if (!cl_data) goto error; filter_pat = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, pattern_width[o->pattern] * pattern_height[o->pattern] * sizeof(cl_int), (void*)pattern[o->pattern], &cl_err); CL_CHECK; cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), &in_buf, sizeof(cl_mem), &out_buf, sizeof(cl_mem), &filter_pat, sizeof(cl_int), &pattern_width[o->pattern], sizeof(cl_int), &pattern_height[o->pattern], sizeof(cl_int), &o->additive, sizeof(cl_int), &o->rotated, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[0], 2, gbl_off, gbl_size, NULL, 0, NULL, NULL); CL_CHECK; cl_err = gegl_clFinish (gegl_cl_get_command_queue ()); CL_CHECK; cl_err = gegl_clReleaseMemObject (filter_pat); CL_CHECK; return FALSE; error: if (filter_pat) gegl_clReleaseMemObject (filter_pat); return TRUE; }
static gboolean cl_process (GeglOperation *operation, cl_mem in_tex, cl_mem aux_tex, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint level) { GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation); cl_int cl_err = 0; /* The kernel will have been compiled by our parent class */ if (!operation_class->cl_data) return TRUE; cl_err = gegl_cl_set_kernel_args (operation_class->cl_data->kernel[0], sizeof(cl_mem), &in_tex, sizeof(cl_mem), &aux_tex, sizeof(cl_mem), &out_tex, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), operation_class->cl_data->kernel[0], 1, NULL, &global_worksize, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
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; }
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; }
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); cl_float3 freq; cl_float3 phaseshift; cl_int3 keep; cl_int cl_err = 0; if (!cl_data) { const char *kernel_name[] = {"cl_alien_map", NULL}; cl_data = gegl_cl_compile_and_build (alien_map_cl_source, kernel_name); } if (!cl_data) return TRUE; freq.s[0] = o->cpn_1_frequency * G_PI; freq.s[1] = o->cpn_2_frequency * G_PI; freq.s[2] = o->cpn_3_frequency * G_PI; phaseshift.s[0] = G_PI * o->cpn_1_phaseshift / 180.0; phaseshift.s[1] = G_PI * o->cpn_2_phaseshift / 180.0; phaseshift.s[2] = G_PI * o->cpn_3_phaseshift / 180.0; keep.s[0] = (cl_int)o->cpn_1_keep; keep.s[1] = (cl_int)o->cpn_2_keep; keep.s[2] = (cl_int)o->cpn_3_keep; cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), &in, sizeof(cl_mem), &out, sizeof(cl_float3), &freq, sizeof(cl_float3), &phaseshift, sizeof(cl_int3), &keep, NULL); 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 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; }
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); CeParamsType *params = (CeParamsType*) o->user_data; cl_float3 color_diff; cl_float3 min; cl_float3 max; cl_int cl_err = 0; gint i; if (!cl_data) { const char *kernel_name[] = {"cl_color_exchange", NULL}; cl_data = gegl_cl_compile_and_build (color_exchange_cl_source, kernel_name); } if (!cl_data) return TRUE; for (i = 0; i < 3; i++) { color_diff.s[i] = params->color_diff[i]; min.s[i] = params->min[i]; max.s[i] = params->max[i]; } cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), &in, sizeof(cl_mem), &out, sizeof(cl_float3), &color_diff, sizeof(cl_float3), &min, sizeof(cl_float3), &max, NULL); 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 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_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 checkerboard_cl_process (GeglOperation *operation, cl_mem out_tex, size_t global_worksize, const GeglRectangle *roi, gint level) { GeglProperties *o = GEGL_PROPERTIES (operation); const Babl *out_format = gegl_operation_get_format (operation, "output"); const size_t gbl_size[2] = {roi->width, roi->height}; const size_t gbl_offs[2] = {roi->x, roi->y}; cl_int cl_err = 0; float color1[4]; float color2[4]; if (!cl_data) { const char *kernel_name[] = {"kernel_checkerboard", NULL}; cl_data = gegl_cl_compile_and_build (checkerboard_cl_source, kernel_name); if (!cl_data) return TRUE; } gegl_color_get_pixel (o->color1, out_format, color1); gegl_color_get_pixel (o->color2, out_format, color2); cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), &out_tex, sizeof(color1), &color1, sizeof(color2), &color2, sizeof(cl_int), &o->x, sizeof(cl_int), &o->y, sizeof(cl_int), &o->x_offset, sizeof(cl_int), &o->y_offset, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[0], 2, gbl_offs, gbl_size, 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 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; }
/* 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 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 *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; }
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); cl_float threshold = o->threshold; if (!cl_data) { const char *kernel_name[] = {"cl_red_eye_removal", NULL}; cl_data = gegl_cl_compile_and_build(red_eye_removal_cl_source, kernel_name); } if (!cl_data) return TRUE; { cl_int cl_err = 0; gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), &in, sizeof(cl_mem), &out, sizeof(cl_float), &threshold, NULL); 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 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_pixelize (cl_mem in_tex, cl_mem aux_tex, cl_mem out_tex, const GeglRectangle *src_rect, const GeglRectangle *roi, gint xsize, gint ysize, gfloat xratio, gfloat yratio, gfloat bg_color[4], gint norm, GeglRectangle *image_extent) { cl_int cl_err = 0; const size_t gbl_size[2]= {roi->width, roi->height}; gint cx0 = block_index (roi->x, xsize); gint cy0 = block_index (roi->y, ysize); gint block_count_x = block_index (roi->x + roi->width + xsize - 1, xsize) - cx0; gint block_count_y = block_index (roi->y + roi->height + ysize - 1, ysize) - cy0; cl_int4 bbox = {{ image_extent->x, image_extent->y, image_extent->x + image_extent->width, image_extent->y + image_extent->height }}; cl_int line_width = roi->width + 2 * xsize; size_t gbl_size_tmp[2] = {block_count_x, block_count_y}; if (!cl_data) { const char *kernel_name[] = {"calc_block_color", "kernel_pixelize", NULL}; cl_data = gegl_cl_compile_and_build (pixelize_cl_source, kernel_name); } if (!cl_data) return 1; cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), (void*)&in_tex, sizeof(cl_mem), (void*)&aux_tex, sizeof(cl_int), (void*)&xsize, sizeof(cl_int), (void*)&ysize, sizeof(cl_int), (void*)&roi->x, sizeof(cl_int), (void*)&roi->y, sizeof(cl_int4), &bbox, sizeof(cl_int), (void*)&line_width, sizeof(cl_int), (void*)&block_count_x, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, gbl_size_tmp, NULL, 0, NULL, NULL); CL_CHECK; cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1], sizeof(cl_mem), (void*)&aux_tex, sizeof(cl_mem), (void*)&out_tex, sizeof(cl_int), (void*)&xsize, sizeof(cl_int), (void*)&ysize, sizeof(cl_float), (void*)&xratio, sizeof(cl_float), (void*)&yratio, sizeof(cl_int), (void*)&roi->x, sizeof(cl_int), (void*)&roi->y, sizeof(cl_float4),(void*)bg_color, sizeof(cl_int), (void*)&norm, sizeof(cl_int), (void*)&block_count_x, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[1], 2, NULL, gbl_size, NULL, 0, NULL, NULL); CL_CHECK; return FALSE; error: return TRUE; }
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; }
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 inline gboolean _gegl_buffer_cl_cache_flush2 (GeglTileHandlerCache *cache, const GeglRectangle *roi) { size_t size; GList *elem; GeglRectangle tmp; cl_int cl_err = 0; gpointer data; gboolean need_cl = FALSE; for (elem=cache_entries; elem; elem=elem->next) { CacheEntry *entry = elem->data; if (entry->valid && entry->tile_storage->cache == cache && (!roi || gegl_rectangle_intersect (&tmp, roi, &entry->roi))) { entry->valid = FALSE; entry->used ++; gegl_cl_color_babl (entry->buffer->soft_format, &size); data = g_malloc(entry->roi.width * entry->roi.height * size); cl_err = gegl_clEnqueueReadBuffer(gegl_cl_get_command_queue(), entry->tex, CL_TRUE, 0, entry->roi.width * entry->roi.height * size, data, 0, NULL, NULL); /* tile-ize */ gegl_buffer_set (entry->buffer, &entry->roi, 0, entry->buffer->soft_format, data, GEGL_AUTO_ROWSTRIDE); entry->used --; need_cl = TRUE; g_free(data); CL_CHECK; } } if (need_cl) { cl_err = gegl_clFinish (gegl_cl_get_command_queue ()); CL_CHECK; g_mutex_lock (&cache_mutex); while (cache_entry_find_invalid (&data)) { CacheEntry *entry = data; #if 1 GEGL_NOTE (GEGL_DEBUG_OPENCL, "Removing from cl-cache: %p %s {%d %d %d %d}", entry->buffer, babl_get_name(entry->buffer->soft_format), entry->roi.x, entry->roi.y, entry->roi.width, entry->roi.height); #endif gegl_clReleaseMemObject(entry->tex); memset (entry, 0x0, sizeof (CacheEntry)); g_slice_free (CacheEntry, data); cache_entries = g_list_remove (cache_entries, data); } g_mutex_unlock (&cache_mutex); } return TRUE; error: g_mutex_lock (&cache_mutex); while (cache_entry_find_invalid (&data)) { g_slice_free (CacheEntry, data); cache_entries = g_list_remove (cache_entries, data); } g_mutex_unlock (&cache_mutex); /* XXX : result is corrupted */ return FALSE; }
gboolean gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err) { GeglBufferClIterators *i = (gpointer)iterator; gboolean result = FALSE; gint no, j; cl_int cl_err = 0; if (i->is_finished) g_error ("%s called on finished buffer iterator", G_STRFUNC); if (i->iteration_no == 0) { for (no=0; no<i->iterators;no++) { if (i->buffer[no]) { gint j; gboolean found = FALSE; for (j=0; j<no; j++) if (i->buffer[no]==i->buffer[j]) { found = TRUE; break; } if (!found) gegl_buffer_lock (i->buffer[no]); if (i->flags[no] == GEGL_CL_BUFFER_WRITE || (i->flags[no] == GEGL_CL_BUFFER_READ && (i->area[no][0] > 0 || i->area[no][1] > 0 || i->area[no][2] > 0 || i->area[no][3] > 0))) { gegl_buffer_cl_cache_flush (i->buffer[no], &i->rect[no]); } } } } else { /* complete pending write work */ for (no=0; no<i->iterators;no++) { if (i->flags[no] == GEGL_CL_BUFFER_WRITE) { /* Wait Processing */ cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue()); if (cl_err != CL_SUCCESS) CL_ERROR; /* color conversion in the GPU (output) */ if (i->conv[no] == GEGL_CL_COLOR_CONVERT) for (j=0; j < i->n; j++) { cl_err = gegl_cl_color_conv (i->tex_op[no][j], i->tex_buf[no][j], i->size[no][j], i->format[no], i->buffer[no]->soft_format); if (cl_err == FALSE) CL_ERROR; } /* Wait Processing */ cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue()); if (cl_err != CL_SUCCESS) CL_ERROR; /* GPU -> CPU */ for (j=0; j < i->n; j++) { gpointer data; /* tile-ize */ if (i->conv[no] == GEGL_CL_COLOR_NOT_SUPPORTED) { data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE, CL_MAP_READ, 0, i->size[no][j] * i->op_cl_format_size [no], 0, NULL, NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* color conversion using BABL */ gegl_buffer_set (i->buffer[no], &i->roi[no][j], 0, i->format[no], data, GEGL_AUTO_ROWSTRIDE); cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_op[no][j], data, 0, NULL, NULL); if (cl_err != CL_SUCCESS) CL_ERROR; } else #ifdef OPENCL_USE_CACHE { gegl_buffer_cl_cache_new (i->buffer[no], &i->roi[no][j], i->tex_buf[no][j]); /* don't release this texture */ i->tex_buf[no][j] = NULL; } #else { data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE, CL_MAP_READ, 0, i->size[no][j] * i->buf_cl_format_size [no], 0, NULL, NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* color conversion using BABL */ gegl_buffer_set (i->buffer[no], &i->roi[no][j], i->format[no], data, GEGL_AUTO_ROWSTRIDE); cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data, 0, NULL, NULL); if (cl_err != CL_SUCCESS) CL_ERROR; } #endif } } } /* Run! */ cl_err = gegl_clFinish(gegl_cl_get_command_queue()); if (cl_err != CL_SUCCESS) CL_ERROR; for (no=0; no < i->iterators; no++) for (j=0; j < i->n; j++) { if (i->tex_buf_from_cache [no][j]) { gboolean ok = gegl_buffer_cl_cache_release (i->tex_buf[no][j]); g_assert (ok); } if (i->tex_buf[no][j] && !i->tex_buf_from_cache [no][j]) gegl_clReleaseMemObject (i->tex_buf[no][j]); if (i->tex_op [no][j]) gegl_clReleaseMemObject (i->tex_op [no][j]); i->tex [no][j] = NULL; i->tex_buf[no][j] = NULL; i->tex_op [no][j] = NULL; } } g_assert (i->iterators > 0); result = (i->roi_no >= i->rois)? FALSE : TRUE; i->n = MIN(GEGL_CL_NTEX, i->rois - i->roi_no); /* then we iterate all */ for (no=0; no<i->iterators;no++) { for (j = 0; j < i->n; j++) { GeglRectangle r = {i->rect[no].x + i->roi_all[i->roi_no+j].x - i->area[no][0], i->rect[no].y + i->roi_all[i->roi_no+j].y - i->area[no][2], i->roi_all[i->roi_no+j].width + i->area[no][0] + i->area[no][1], i->roi_all[i->roi_no+j].height + i->area[no][2] + i->area[no][3]}; i->roi [no][j] = r; i->size[no][j] = r.width * r.height; } if (i->flags[no] == GEGL_CL_BUFFER_READ) { for (j=0; j < i->n; j++) { gpointer data; /* un-tile */ switch (i->conv[no]) { case GEGL_CL_COLOR_NOT_SUPPORTED: { gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]); g_assert (i->tex_op[no][j] == NULL); i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY, i->size[no][j] * i->op_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* pre-pinned memory */ data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE, CL_MAP_WRITE, 0, i->size[no][j] * i->op_cl_format_size [no], 0, NULL, NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* color conversion using BABL */ gegl_buffer_get (i->buffer[no], &i->roi[no][j], 1.0, i->format[no], data, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE); cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_op[no][j], data, 0, NULL, NULL); if (cl_err != CL_SUCCESS) CL_ERROR; i->tex[no][j] = i->tex_op[no][j]; break; } case GEGL_CL_COLOR_EQUAL: { i->tex_buf[no][j] = gegl_buffer_cl_cache_get (i->buffer[no], &i->roi[no][j]); if (i->tex_buf[no][j]) i->tex_buf_from_cache [no][j] = TRUE; /* don't free texture from cache */ else { gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]); g_assert (i->tex_buf[no][j] == NULL); i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY, i->size[no][j] * i->buf_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* pre-pinned memory */ data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE, CL_MAP_WRITE, 0, i->size[no][j] * i->buf_cl_format_size [no], 0, NULL, NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* color conversion will be performed in the GPU later */ gegl_buffer_cl_worker_transf (i->buffer[no], data, i->buf_cl_format_size [no], i->roi[no][j], FALSE); cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data, 0, NULL, NULL); if (cl_err != CL_SUCCESS) CL_ERROR; } i->tex[no][j] = i->tex_buf[no][j]; break; } case GEGL_CL_COLOR_CONVERT: { i->tex_buf[no][j] = gegl_buffer_cl_cache_get (i->buffer[no], &i->roi[no][j]); if (i->tex_buf[no][j]) i->tex_buf_from_cache [no][j] = TRUE; /* don't free texture from cache */ else { gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]); g_assert (i->tex_buf[no][j] == NULL); i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY, i->size[no][j] * i->buf_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* pre-pinned memory */ data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE, CL_MAP_WRITE, 0, i->size[no][j] * i->buf_cl_format_size [no], 0, NULL, NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* color conversion will be performed in the GPU later */ /* get buffer data using multiple worker threads to increase bandwidth */ gegl_buffer_cl_worker_transf (i->buffer[no], data, i->buf_cl_format_size [no], i->roi[no][j], FALSE); cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data, 0, NULL, NULL); if (cl_err != CL_SUCCESS) CL_ERROR; } g_assert (i->tex_op[no][j] == NULL); i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_READ_WRITE, i->size[no][j] * i->op_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; /* color conversion in the GPU (input) */ g_assert (i->tex_buf[no][j] && i->tex_op[no][j]); cl_err = gegl_cl_color_conv (i->tex_buf[no][j], i->tex_op[no][j], i->size[no][j], i->buffer[no]->soft_format, i->format[no]); if (cl_err == FALSE) CL_ERROR; i->tex[no][j] = i->tex_op[no][j]; break; } } } /* Wait Processing */ cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue()); if (cl_err != CL_SUCCESS) CL_ERROR; } else if (i->flags[no] == GEGL_CL_BUFFER_WRITE) { for (j=0; j < i->n; j++) { switch (i->conv[no]) { case GEGL_CL_COLOR_NOT_SUPPORTED: { g_assert (i->tex_op[no][j] == NULL); i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_ALLOC_HOST_PTR | CL_MEM_WRITE_ONLY, i->size[no][j] * i->op_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; i->tex[no][j] = i->tex_op[no][j]; break; } case GEGL_CL_COLOR_EQUAL: { g_assert (i->tex_buf[no][j] == NULL); i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE, /* cache */ i->size[no][j] * i->buf_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; i->tex[no][j] = i->tex_buf[no][j]; break; } case GEGL_CL_COLOR_CONVERT: { g_assert (i->tex_buf[no][j] == NULL); i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE, /* cache */ i->size[no][j] * i->buf_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; g_assert (i->tex_op[no][j] == NULL); i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_READ_WRITE, i->size[no][j] * i->op_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; i->tex[no][j] = i->tex_op[no][j]; break; } } } } else if (i->flags[no] == GEGL_CL_BUFFER_AUX) { for (j=0; j < i->n; j++) { g_assert (i->tex_op[no][j] == NULL); i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_READ_WRITE, i->size[no][j] * i->op_cl_format_size [no], NULL, &cl_err); if (cl_err != CL_SUCCESS) CL_ERROR; i->tex[no][j] = i->tex_op[no][j]; } } } i->roi_no += i->n; i->iteration_no++; if (result == FALSE) { for (no=0; no<i->iterators;no++) { if (i->buffer[no]) { gint j; gboolean found = FALSE; for (j=0; j<no; j++) if (i->buffer[no]==i->buffer[j]) { found = TRUE; break; } if (!found) gegl_buffer_unlock (i->buffer[no]); g_object_unref (i->buffer[no]); } } i->is_finished = TRUE; g_free (i->roi_all); g_slice_free (GeglBufferClIterators, i); } *err = FALSE; return result; error: for (no=0; no<i->iterators;no++) for (j=0; j < i->n; j++) { if (i->tex_buf[no][j]) gegl_clReleaseMemObject (i->tex_buf[no][j]); if (i->tex_op [no][j]) gegl_clReleaseMemObject (i->tex_op [no][j]); i->tex [no][j] = NULL; i->tex_buf[no][j] = NULL; i->tex_op [no][j] = NULL; } *err = TRUE; return FALSE; }
static gboolean cl_bilateral (cl_mem in_tex, cl_mem out_tex, const GeglRectangle *roi, const GeglRectangle *src_rect, gint s_sigma, gfloat r_sigma) { cl_int cl_err = 0; gint c; const gint width = src_rect->width; const gint height = src_rect->height; const gint sw = (width -1) / s_sigma + 1; const gint sh = (height-1) / s_sigma + 1; const gint depth = (int)(1.0f / r_sigma) + 1; size_t global_ws[2]; size_t local_ws[2]; cl_mem grid = NULL; cl_mem blur[4] = {NULL, NULL, NULL, NULL}; if (!cl_data) { const char *kernel_name[] = {"bilateral_downsample", "bilateral_blur", "bilateral_interpolate", NULL}; cl_data = gegl_cl_compile_and_build (bilateral_filter_fast_cl_source, kernel_name); } if (!cl_data) return 1; grid = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_READ_WRITE, sw * sh * depth * sizeof(cl_float8), NULL, &cl_err); CL_CHECK; for(c = 0; c < 4; c++) { blur[c] = gegl_clCreateBuffer (gegl_cl_get_context (), CL_MEM_READ_WRITE, sw * sh * depth * sizeof(cl_float2), NULL, &cl_err); CL_CHECK; } local_ws[0] = 8; local_ws[1] = 8; global_ws[0] = ((sw + local_ws[0] - 1)/local_ws[0])*local_ws[0]; global_ws[1] = ((sh + local_ws[1] - 1)/local_ws[1])*local_ws[1]; gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), &in_tex, sizeof(cl_mem), &grid, sizeof(cl_int), &width, sizeof(cl_int), &height, sizeof(cl_int), &sw, sizeof(cl_int), &sh, sizeof(cl_int), &depth, sizeof(cl_int), &s_sigma, sizeof(cl_float), &r_sigma, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[0], 2, NULL, global_ws, local_ws, 0, NULL, NULL); CL_CHECK; local_ws[0] = 16; local_ws[1] = 16; global_ws[0] = ((sw + local_ws[0] - 1)/local_ws[0])*local_ws[0]; global_ws[1] = ((sh + local_ws[1] - 1)/local_ws[1])*local_ws[1]; gegl_cl_set_kernel_args (cl_data->kernel[1], sizeof(cl_mem), &grid, sizeof(cl_mem), &blur[0], sizeof(cl_mem), &blur[1], sizeof(cl_mem), &blur[2], sizeof(cl_mem), &blur[3], sizeof(cl_int), &sw, sizeof(cl_int), &sh, sizeof(cl_int), &depth, NULL); cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), cl_data->kernel[1], 2, NULL, global_ws, local_ws, 0, NULL, NULL); CL_CHECK; global_ws[0] = width; global_ws[1] = height; gegl_cl_set_kernel_args (cl_data->kernel[2], sizeof(cl_mem), &in_tex, sizeof(cl_mem), &blur[0], sizeof(cl_mem), &blur[1], sizeof(cl_mem), &blur[2], sizeof(cl_mem), &blur[3], sizeof(cl_mem), &out_tex, sizeof(cl_int), &width, sizeof(cl_int), &sw, sizeof(cl_int), &sh, sizeof(cl_int), &depth, sizeof(cl_int), &s_sigma, sizeof(cl_float), &r_sigma, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (), cl_data->kernel[2], 2, NULL, global_ws, NULL, 0, NULL, NULL); CL_CHECK; cl_err = gegl_clFinish (gegl_cl_get_command_queue ()); CL_CHECK; cl_err = gegl_clReleaseMemObject (grid); CL_CHECK_ONLY (cl_err); for(c = 0; c < 4; c++) { cl_err = gegl_clReleaseMemObject (blur[c]); CL_CHECK_ONLY (cl_err); } return FALSE; error: if (grid) gegl_clReleaseMemObject (grid); for (c = 0; c < 4; c++) { if (blur[c]) gegl_clReleaseMemObject (blur[c]); } return TRUE; }
static gboolean cl_process (GeglOperation *operation, cl_mem out_tex, const GeglRectangle *roi) { GeglProperties *o = GEGL_PROPERTIES (operation); const size_t gbl_size[] = { roi->width, roi->height }; cl_int cl_err = 0; cl_int offset_x; cl_int offset_y; cl_int width; cl_int height; cl_float3 sedges; cl_float3 contours; cl_float3 frequency; cl_float brightness; cl_float polarization; cl_float scattering; cl_int iterations; cl_float weird_factor; if (!cl_data) { const char *kernel_name[] = { "cl_diffraction_patterns", NULL }; cl_data = gegl_cl_compile_and_build (diffraction_patterns_cl_source, kernel_name); if (!cl_data) return TRUE; } offset_x = roi->x; offset_y = roi->y; width = o->width; height = o->height; sedges.s[0] = o->red_sedges; sedges.s[1] = o->green_sedges; sedges.s[2] = o->blue_sedges; contours.s[0] = o->red_contours; contours.s[1] = o->green_contours; contours.s[2] = o->blue_contours; frequency.s[0] = o->red_frequency; frequency.s[1] = o->green_frequency; frequency.s[2] = o->blue_frequency; brightness = o->brightness; polarization = o->polarization; scattering = o->scattering; iterations = ITERATIONS; weird_factor = WEIRD_FACTOR; cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), &out_tex, sizeof(cl_int), &offset_x, sizeof(cl_int), &offset_y, sizeof(cl_int), &width, sizeof(cl_int), &height, sizeof(cl_float3), &sedges, sizeof(cl_float3), &contours, sizeof(cl_float3), &frequency, sizeof(cl_float), &brightness, sizeof(cl_float), &polarization, sizeof(cl_float), &scattering, sizeof(cl_int), &iterations, sizeof(cl_float), &weird_factor, NULL); 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; cl_err = gegl_clFinish (gegl_cl_get_command_queue ()); CL_CHECK; return FALSE; error: return TRUE; }
static gboolean cl_box_blur (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], global_ws[2]; size_t local_ws_hor[2], local_ws_ver[2], local_ws[2]; size_t step_size ; if (!cl_data) { const char *kernel_name[] = { "kernel_blur_hor", "kernel_blur_ver","kernel_box_blur_fast", NULL}; cl_data = gegl_cl_compile_and_build (box_blur_cl_source, kernel_name); } if (!cl_data) return TRUE; step_size = 64; local_ws[0]=256; local_ws[1]=1; if( radius <=110 ) { global_ws[0] = (roi->width + local_ws[0] - 2 * radius - 1) / ( local_ws[0] - 2 * radius ) * local_ws[0]; global_ws[1] = (roi->height + step_size - 1) / step_size; cl_err = gegl_cl_set_kernel_args(cl_data->kernel[2], sizeof(cl_mem), (void *)&in_tex, sizeof(cl_mem), (void *)&out_tex, sizeof(cl_float4)*local_ws[0], (void *)NULL, sizeof(cl_int), (void *)&roi->width, sizeof(cl_int), (void *)&roi->height, sizeof(cl_int), (void *)&radius, sizeof(cl_int), (void *)&step_size, NULL); CL_CHECK; cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[2], 2, NULL, global_ws, local_ws, 0, NULL, NULL ); CL_CHECK; } else { 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_cl_set_kernel_args (cl_data->kernel[0], sizeof(cl_mem), (void*)&in_tex, sizeof(cl_mem), (void*)&aux_tex, sizeof(cl_int), (void*)&roi->width, sizeof(cl_int), (void*)&radius, NULL); 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_cl_set_kernel_args (cl_data->kernel[1], sizeof(cl_mem), (void*)&aux_tex, sizeof(cl_mem), (void*)&out_tex, sizeof(cl_int), (void*)&roi->width, sizeof(cl_int), (void*)&radius, NULL); 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; }