1 /* This file is an image processing operation for GEGL 2 * 3 * GEGL is free software; you can redistribute it and/or 4 * modify it under the terms of the GNU Lesser General Public 5 * License as published by the Free Software Foundation; either 6 * version 3 of the License, or (at your option) any later version. 7 * 8 * GEGL is distributed in the hope that it will be useful, 9 * but WITHOUT ANY WARRANTY; without even the implied warranty of 10 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU 11 * Lesser General Public License for more details. 12 * 13 * You should have received a copy of the GNU Lesser General Public 14 * License along with GEGL; if not, see <https://www.gnu.org/licenses/>. 15 * 16 * Copyright 2006 Øyvind Kolås <pippin@gimp.org> 17 * 2012 Pavel Roschin <roshin@scriptumplus.ru> 18 */ 19 20 #include "config.h" 21 #include <glib/gi18n-lib.h> 22 23 #ifdef GEGL_PROPERTIES 24 25 property_int (radius, _("Radius"), 4) 26 description(_("Radius of square pixel region, (width and height will be radius*2+1)")) 27 value_range (0, 1000) 28 ui_range (0, 100) 29 ui_gamma (1.5) 30 31 #else 32 33 #define GEGL_OP_AREA_FILTER 34 #define GEGL_OP_NAME box_blur 35 #define GEGL_OP_C_SOURCE box-blur.c 36 37 #include "gegl-op.h" 38 #include <stdio.h> 39 40 #define SRC_OFFSET (row + u + radius * 2) * 4 41 42 static void 43 hor_blur (GeglBuffer *src, 44 const GeglRectangle *src_rect, 45 GeglBuffer *dst, 46 const GeglRectangle *dst_rect, 47 gint radius, 48 const Babl *format) 49 { 50 gint u,v; 51 gint i; 52 gint offset; 53 gint src_offset; 54 gint prev_rad = radius * 4 + 4; 55 gint next_rad = radius * 4; 56 gint row; 57 gfloat *src_buf; 58 gfloat *dst_buf; 59 gfloat rad1 = 1.0 / (gfloat)(radius * 2 + 1); 60 61 src_buf = g_new0 (gfloat, src_rect->width * src_rect->height * 4); 62 dst_buf = g_new0 (gfloat, dst_rect->width * dst_rect->height * 4); 63 64 gegl_buffer_get (src, src_rect, 1.0, format, 65 src_buf, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_CLAMP); 66 67 offset = 0; 68 for (v = 0; v < dst_rect->height; v++) 69 { 70 /* here just radius, not radius * 2 as in ver_blur beacuse 71 * we enlarged dst_buf by y earlier */ 72 row = (v + radius) * src_rect->width; 73 /* prepare - set first column of pixels */ 74 for (u = -radius; u <= radius; u++) 75 { 76 src_offset = SRC_OFFSET; 77 for (i = 0; i < 4; i++) 78 dst_buf[offset + i] += src_buf[src_offset + i] * rad1; 79 } 80 offset += 4; 81 /* iterate other pixels by moving a window - very fast */ 82 for (u = 1; u < dst_rect->width; u++) 83 { 84 src_offset = SRC_OFFSET; 85 for (i = 0; i < 4; i++) 86 { 87 dst_buf[offset] = dst_buf[offset - 4] 88 - src_buf[src_offset - prev_rad] * rad1 89 + src_buf[src_offset + next_rad] * rad1; 90 src_offset++; 91 offset++; 92 } 93 } 94 } 95 96 gegl_buffer_set (dst, dst_rect, 0, format, 97 dst_buf, GEGL_AUTO_ROWSTRIDE); 98 99 g_free (src_buf); 100 g_free (dst_buf); 101 } 102 103 static void 104 ver_blur (GeglBuffer *src, 105 const GeglRectangle *src_rect, 106 GeglBuffer *dst, 107 const GeglRectangle *dst_rect, 108 gint radius, 109 const Babl *format) 110 { 111 gint u, v; 112 gint i; 113 gint offset; 114 gint src_offset; 115 gint prev_rad = (radius * 4 + 4) * src_rect->width; 116 gint next_rad = (radius * 4) * src_rect->width; 117 gint row; 118 gfloat *src_buf; 119 gfloat *dst_buf; 120 gfloat rad1 = 1.0 / (gfloat)(radius * 2 + 1); 121 122 src_buf = g_new0 (gfloat, src_rect->width * src_rect->height * 4); 123 dst_buf = g_new0 (gfloat, dst_rect->width * dst_rect->height * 4); 124 125 gegl_buffer_get (src, src_rect, 1.0, format, 126 src_buf, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_CLAMP); 127 128 /* prepare: set first row of pixels */ 129 for (v = -radius; v <= radius; v++) 130 { 131 row = (v + radius * 2) * src_rect->width; 132 for (u = 0; u < dst_rect->width; u++) 133 { 134 src_offset = SRC_OFFSET; 135 for (i = 0; i < 4; i++) 136 dst_buf[u * 4 + i] += src_buf[src_offset + i] * rad1; 137 } 138 } 139 /* skip first row */ 140 offset = dst_rect->width * 4; 141 for (v = 1; v < dst_rect->height; v++) 142 { 143 row = (v + radius * 2) * src_rect->width; 144 for (u = 0; u < dst_rect->width; u++) 145 { 146 src_offset = SRC_OFFSET; 147 for (i = 0; i < 4; i++) 148 { 149 dst_buf[offset] = dst_buf[offset - 4 * dst_rect->width] 150 - src_buf[src_offset - prev_rad] * rad1 151 + src_buf[src_offset + next_rad] * rad1; 152 src_offset++; 153 offset++; 154 } 155 } 156 } 157 158 gegl_buffer_set (dst, dst_rect, 0, format, 159 dst_buf, GEGL_AUTO_ROWSTRIDE); 160 161 g_free (src_buf); 162 g_free (dst_buf); 163 } 164 165 #undef SRC_OFFSET 166 167 static void prepare (GeglOperation *operation) 168 { 169 GeglProperties *o; 170 GeglOperationAreaFilter *op_area; 171 const Babl *space = gegl_operation_get_source_space (operation, "input"); 172 const Babl *format = babl_format_with_space ("RaGaBaA float", space); 173 174 op_area = GEGL_OPERATION_AREA_FILTER (operation); 175 o = GEGL_PROPERTIES (operation); 176 177 op_area->left = 178 op_area->right = 179 op_area->top = 180 op_area->bottom = o->radius; 181 182 gegl_operation_set_format (operation, "input", format); 183 gegl_operation_set_format (operation, "output", format); 184 } 185 186 #include "opencl/gegl-cl.h" 187 #include "gegl-buffer-cl-iterator.h" 188 189 #include "opencl/box-blur.cl.h" 190 static GeglClRunData *cl_data = NULL; 191 static gboolean 192 cl_box_blur (cl_mem in_tex, 193 cl_mem aux_tex, 194 cl_mem out_tex, 195 size_t global_worksize, 196 const GeglRectangle *roi, 197 gint radius) 198 { 199 cl_int cl_err = 0; 200 201 if (!cl_data) 202 { 203 const char *kernel_name[] = { "kernel_blur_hor", "kernel_blur_ver","kernel_box_blur_fast", NULL}; 204 cl_data = gegl_cl_compile_and_build (box_blur_cl_source, kernel_name); 205 } 206 207 if (!cl_data) 208 return TRUE; 209 210 #if 0 211 /* XXX: this code path is commented out, because it introduces vertical stripes/gaps, 212 wheras the generic code path does not. */ 213 if( radius <=110 ) 214 { 215 size_t global_ws[2]; 216 local_ws[2]; 217 size_t step_size = 64; 218 local_ws[0]=256; 219 local_ws[1]=1; 220 global_ws[0] = (roi->width + local_ws[0] - 2 * radius - 1) / ( local_ws[0] - 2 * radius ) * local_ws[0]; 221 global_ws[1] = (roi->height + step_size - 1) / step_size; 222 cl_err = gegl_cl_set_kernel_args(cl_data->kernel[2], 223 sizeof(cl_mem), (void *)&in_tex, 224 sizeof(cl_mem), (void *)&out_tex, 225 sizeof(cl_float4)*local_ws[0], (void *)NULL, 226 sizeof(cl_int), (void *)&roi->width, 227 sizeof(cl_int), (void *)&roi->height, 228 sizeof(cl_int), (void *)&radius, 229 sizeof(cl_int), (void *)&step_size, NULL); 230 CL_CHECK; 231 cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), 232 cl_data->kernel[2], 2, 233 NULL, global_ws, local_ws, 0, NULL, NULL ); 234 CL_CHECK; 235 236 } 237 else 238 #endif 239 { 240 size_t global_ws_hor[2], global_ws_ver[2]; 241 size_t local_ws_hor[2], local_ws_ver[2]; 242 243 local_ws_hor[0] = 1; 244 local_ws_hor[1] = 256; 245 global_ws_hor[0] = roi->height + 2 * radius; 246 global_ws_hor[1] = ((roi->width + local_ws_hor[1] -1)/local_ws_hor[1]) * local_ws_hor[1]; 247 248 local_ws_ver[0] = 1; 249 local_ws_ver[1] = 256; 250 global_ws_ver[0] = roi->height; 251 global_ws_ver[1] = ((roi->width + local_ws_ver[1] -1)/local_ws_ver[1]) * local_ws_ver[1]; 252 253 254 cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0], 255 sizeof(cl_mem), (void*)&in_tex, 256 sizeof(cl_mem), (void*)&aux_tex, 257 sizeof(cl_int), (void*)&roi->width, 258 sizeof(cl_int), (void*)&radius, 259 NULL); 260 CL_CHECK; 261 cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), 262 cl_data->kernel[0], 2, 263 NULL, global_ws_hor, local_ws_hor, 264 0, NULL, NULL); 265 CL_CHECK; 266 267 268 cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1], 269 sizeof(cl_mem), (void*)&aux_tex, 270 sizeof(cl_mem), (void*)&out_tex, 271 sizeof(cl_int), (void*)&roi->width, 272 sizeof(cl_int), (void*)&radius, 273 NULL); 274 CL_CHECK; 275 276 cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), 277 cl_data->kernel[1], 2, 278 NULL, global_ws_ver, local_ws_ver, 279 0, NULL, NULL); 280 CL_CHECK; 281 } 282 283 return FALSE; 284 285 error: 286 return TRUE; 287 } 288 289 static gboolean 290 cl_process (GeglOperation *operation, 291 GeglBuffer *input, 292 GeglBuffer *output, 293 const GeglRectangle *result) 294 { 295 const Babl *in_format = gegl_operation_get_format (operation, "input"); 296 const Babl *out_format = gegl_operation_get_format (operation, "output"); 297 298 gint err = 0; 299 300 GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation); 301 GeglProperties *o = GEGL_PROPERTIES (operation); 302 303 GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, 304 result, 305 out_format, 306 GEGL_CL_BUFFER_WRITE); 307 308 gint read = gegl_buffer_cl_iterator_add_2 (i, 309 input, 310 result, 311 in_format, 312 GEGL_CL_BUFFER_READ, 313 op_area->left, 314 op_area->right, 315 op_area->top, 316 op_area->bottom, 317 GEGL_ABYSS_CLAMP); 318 319 gint aux = gegl_buffer_cl_iterator_add_aux (i, 320 result, 321 in_format, 322 0, 323 0, 324 op_area->top, 325 op_area->bottom); 326 327 while (gegl_buffer_cl_iterator_next (i, &err) && !err) 328 { 329 err = cl_box_blur (i->tex[read], 330 i->tex[aux], 331 i->tex[0], 332 i->size[0], 333 &i->roi[0], 334 o->radius); 335 336 if (err) 337 { 338 gegl_buffer_cl_iterator_stop (i); 339 break; 340 } 341 } 342 343 return !err; 344 } 345 346 static gboolean 347 process (GeglOperation *operation, 348 GeglBuffer *input, 349 GeglBuffer *output, 350 const GeglRectangle *result, 351 gint level) 352 { 353 GeglRectangle rect; 354 GeglRectangle tmprect; 355 GeglProperties *o = GEGL_PROPERTIES (operation); 356 GeglBuffer *temp; 357 GeglOperationAreaFilter *op_area; 358 const Babl *out_format = gegl_operation_get_format (operation, "output"); 359 360 op_area = GEGL_OPERATION_AREA_FILTER (operation); 361 362 if (gegl_operation_use_opencl (operation)) 363 if (cl_process (operation, input, output, result)) 364 return TRUE; 365 366 rect = *result; 367 tmprect = *result; 368 369 rect.x -= op_area->left * 2; 370 rect.y -= op_area->top * 2; 371 rect.width += (op_area->left + op_area->right) * 2; 372 rect.height += (op_area->top + op_area->bottom) * 2; 373 /* very tricky: enlarge temp buffer to avoid seams in second pass */ 374 tmprect.y -= o->radius; 375 tmprect.height += o->radius * 2; 376 377 temp = gegl_buffer_new (&tmprect, out_format); 378 379 /* doing second pass in separate gegl op may be significantly faster */ 380 hor_blur (input, &rect, temp, &tmprect, o->radius, out_format); 381 ver_blur (temp, &rect, output, result, o->radius, out_format); 382 383 g_object_unref (temp); 384 return TRUE; 385 } 386 387 388 static void 389 gegl_op_class_init (GeglOpClass *klass) 390 { 391 GeglOperationClass *operation_class; 392 GeglOperationFilterClass *filter_class; 393 394 operation_class = GEGL_OPERATION_CLASS (klass); 395 filter_class = GEGL_OPERATION_FILTER_CLASS (klass); 396 397 filter_class->process = process; 398 operation_class->prepare = prepare; 399 400 operation_class->opencl_support = TRUE; 401 402 gegl_operation_class_set_keys (operation_class, 403 "name", "gegl:box-blur", 404 "title", _("Box Blur"), 405 "categories", "blur", 406 "reference-hash", "a1373d1ad34431271aeec3c289f6ca10", 407 "description", _("Blur resulting from averaging the colors of a square neighbourhood."), 408 NULL); 409 } 410 411 #endif 412