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; }
static gboolean process (GeglOperation *operation, GeglBuffer *input, const GeglRectangle *result, gint level) { GeglChantO *o = GEGL_CHANT_PROPERTIES (operation); if (o->buffer) { GeglBuffer *output = GEGL_BUFFER (o->buffer); const Babl *in_format = gegl_buffer_get_format (input); const Babl *out_format = gegl_buffer_get_format (output); if (gegl_operation_use_opencl (operation) && gegl_cl_color_supported (in_format, out_format) == GEGL_CL_COLOR_CONVERT) { size_t size; gboolean err; cl_int cl_err = 0; GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE); gint read = gegl_buffer_cl_iterator_add (i, input, result, out_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE); gegl_cl_color_babl (out_format, &size); GEGL_NOTE (GEGL_DEBUG_OPENCL, "write-buffer: " "%p %p %s %s {%d %d %d %d}", input, output, babl_get_name (in_format), babl_get_name (out_format), result->x, result->y, result->width, result->height); while (gegl_buffer_cl_iterator_next (i, &err)) { if (err) break; cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue (), i->tex[read], i->tex[0], 0, 0, i->size[0] * size, 0, NULL, NULL); if (cl_err != CL_SUCCESS) { GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring (cl_err)); break; } } if (cl_err || err) gegl_buffer_copy (input, result, output, result); } else gegl_buffer_copy (input, result, output, result); gegl_buffer_flush (output); } return TRUE; }
gint gegl_buffer_cl_iterator_add_2 (GeglBufferClIterator *iterator, GeglBuffer *buffer, const GeglRectangle *result, const Babl *format, guint flags, gint left, gint right, gint top, gint bottom, GeglAbyssPolicy abyss_policy) { GeglBufferClIterators *i = (gpointer)iterator; gint self = 0; if (i->iterators+1 > GEGL_CL_BUFFER_MAX_ITERATORS) { g_error ("too many iterators (%i)", i->iterators+1); } if (i->iterators == 0) /* for sanity, we zero at init */ { memset (i, 0, sizeof (GeglBufferClIterators)); } self = i->iterators++; if (!result) result = self==0?&(buffer->extent):&(i->rect[0]); i->rect[self]=*result; i->flags[self]=flags; if (flags == GEGL_CL_BUFFER_WRITE || flags == GEGL_CL_BUFFER_READ) { g_assert (buffer); i->buffer[self]= g_object_ref (buffer); if (format) i->format[self]=format; else i->format[self]=buffer->soft_format; if (flags == GEGL_CL_BUFFER_WRITE) i->conv[self] = gegl_cl_color_supported (format, buffer->soft_format); else i->conv[self] = gegl_cl_color_supported (buffer->soft_format, format); gegl_cl_color_babl (buffer->soft_format, &i->buf_cl_format_size[self]); gegl_cl_color_babl (format, &i->op_cl_format_size [self]); } else /* GEGL_CL_BUFFER_AUX */ { g_assert (buffer == NULL); i->buffer[self] = NULL; i->format[self] = NULL; i->conv[self] = -1; i->buf_cl_format_size[self] = SIZE_MAX; gegl_cl_color_babl (format, &i->op_cl_format_size [self]); } i->area[self][0] = left; i->area[self][1] = right; i->area[self][2] = top; i->area[self][3] = bottom; if (flags == GEGL_CL_BUFFER_WRITE && (left > 0 || right > 0 || top > 0 || bottom > 0)) g_assert(FALSE); if (self!=0) { /* we make all subsequently added iterators share the width and height of the first one */ i->rect[self].width = i->rect[0].width; i->rect[self].height = i->rect[0].height; } else { gint x, y, j; i->rois = 0; for (y=result->y; y < result->y + result->height; y += gegl_cl_get_iter_height ()) for (x=result->x; x < result->x + result->width; x += gegl_cl_get_iter_width ()) i->rois++; i->roi_no = 0; i->roi_all = g_new0 (GeglRectangle, i->rois); j = 0; for (y=0; y < result->height; y += gegl_cl_get_iter_height ()) for (x=0; x < result->width; x += gegl_cl_get_iter_width ()) { GeglRectangle r = {x, y, MIN(gegl_cl_get_iter_width (), result->width - x), MIN(gegl_cl_get_iter_height (), result->height - y)}; i->roi_all[j] = r; j++; } } return self; }
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; }