From ec6e291a3e3cc2179fa45d475d1aaf28d4b30230 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ond=C5=99ej=20M=C3=ADchal?= Date: Fri, 15 Aug 2025 21:02:56 +0200 Subject: [PATCH 1/5] app/operations: Add central point for OpenCL implementations This brings in the same basic infrastructure for writing OpenCL implementations of operations from GEGL. GEGL operations expect CL implementations in source form which is then passed to the driver for compilation. The source is passed as a string and this infrastructure allows for writing CL code in plain text files which are automatically "stringified" and made available via header files. --- app/operations/meson.build | 4 +- app/operations/opencl/cltostring.py | 66 +++++++++++++++++++++++++++++ app/operations/opencl/meson.build | 14 ++++++ 3 files changed, 83 insertions(+), 1 deletion(-) create mode 100755 app/operations/opencl/cltostring.py create mode 100644 app/operations/opencl/meson.build diff --git a/app/operations/meson.build b/app/operations/meson.build index 14f210b9893..399546fb6cd 100644 --- a/app/operations/meson.build +++ b/app/operations/meson.build @@ -17,6 +17,8 @@ stamp_operations_enums = custom_target('stamp-operations-enums.h', build_by_default: true ) +subdir('opencl') + libappoperations_sources = [ 'gimp-operation-config.c', 'gimp-operations.c', @@ -62,7 +64,7 @@ libappoperations_sources = [ ] libappoperations = static_library('appoperations', - libappoperations_sources, + [libappoperations_sources, opencl_headers], include_directories: [ rootInclude, rootAppInclude, ], c_args: '-DG_LOG_DOMAIN="Gimp-Operations"', dependencies: [ diff --git a/app/operations/opencl/cltostring.py b/app/operations/opencl/cltostring.py new file mode 100755 index 00000000000..ad1289479be --- /dev/null +++ b/app/operations/opencl/cltostring.py @@ -0,0 +1,66 @@ +#!/usr/bin/env python +from __future__ import print_function +from __future__ import unicode_literals + +import os +import sys +import io + +# Search for lines that look like #include "blah.h" and replace them +# with the contents of blah.h. +def do_includes (source): + result = list() + for line in source.split("\n"): + if line.lstrip().startswith("#include"): + splitstr = line.split('"') + if len(splitstr) != 3: + raise RuntimeError("Invalid include: %s" % line) + include_path = splitstr[1] + if not os.path.isfile(include_path): + raise RuntimeError("Could not find include: %s" % line) + with open(include_path, "r") as inc_file: + result += do_includes(inc_file.read()).split("\n") + else: + result.append(line) + return "\n".join(result) + +# From http://stackoverflow.com/questions/14945095/how-to-escape-string-for-generated-c +def escape_string(s): + result = '' + for c in s: + if not (32 <= ord(c) < 127) or c in ('\\', '"'): + result += '\\%03o' % ord(c) + else: + result += c + return result + + +if len(sys.argv) == 2: + infile = io.open(sys.argv[1], "r", encoding="utf-8") + outfile = io.open(sys.argv[1] + '.h', "w", encoding="utf-8") + +elif len(sys.argv) == 3: + infile = io.open(sys.argv[1], "r", encoding="utf-8") + outfile = io.open(sys.argv[2], "w", encoding="utf-8") + +else: + print("Usage: %s input [output]" % sys.argv[0]) + sys.exit(1) + + +cl_source = infile.read() +cl_source = do_includes(cl_source) +infile.close() + +string_var_name = os.path.basename(sys.argv[1]).replace("-", "_").replace(":", "_") +if string_var_name.endswith(".cl"): + string_var_name = string_var_name[:-3] + +outfile.write("static const char* %s_cl_source =\n" % string_var_name) +for line in cl_source.rstrip().split("\n"): + line = line.rstrip() + line = escape_string(line) + line = '"%-78s\\n"\n' % line + outfile.write(line) +outfile.write(";\n") +outfile.close() diff --git a/app/operations/opencl/meson.build b/app/operations/opencl/meson.build new file mode 100644 index 00000000000..31967af7d2a --- /dev/null +++ b/app/operations/opencl/meson.build @@ -0,0 +1,14 @@ +cltostring = files('cltostring.py') + +opencl_sources = [ + 'set-alpha.cl', +] + +opencl_headers = files() +foreach source : opencl_sources + opencl_headers += custom_target(source + '.h', + input : source, + output: source +'.h', + command: [ python, cltostring, '@INPUT@', '@OUTPUT@' ], + ) +endforeach -- GitLab From bef8dc9d50e60bef81c82454fdfedc8b0e2d0092 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ond=C5=99ej=20M=C3=ADchal?= Date: Sat, 5 Jul 2025 19:49:25 +0200 Subject: [PATCH 2/5] app/operations: Add OpenCL implementation for set-alpha This is the MVC to try out the introduced basic OpenCL infrastructure in GIMP. The operation itself is very simple and thus was the perfect candidate for a quick implementation. As a note for myself and others writing OpenCL implementations in the future, be mindful of parameters with double type. Support for double is not present in all drivers and so using floats is something I'd recommend. For that to work, convert all double parameters to floats before setting kernel arguments! --- app/operations/gimpoperationsetalpha.c | 70 ++++++++++++++++++++++++++ app/operations/opencl/set-alpha.cl | 24 +++++++++ 2 files changed, 94 insertions(+) create mode 100644 app/operations/opencl/set-alpha.cl diff --git a/app/operations/gimpoperationsetalpha.c b/app/operations/gimpoperationsetalpha.c index b41af3caa1d..7909fc17a7e 100644 --- a/app/operations/gimpoperationsetalpha.c +++ b/app/operations/gimpoperationsetalpha.c @@ -21,8 +21,10 @@ #include "config.h" #include +#include "opencl/gegl-cl.h" #include "operations-types.h" +#include "app/operations/opencl/set-alpha.cl.h" #include "gimpoperationsetalpha.h" @@ -51,6 +53,13 @@ static gboolean gimp_operation_set_alpha_process (GeglOperation *op glong samples, const GeglRectangle *roi, gint level); +static gboolean gimp_operation_set_alpha_cl_process (GeglOperation *self, + cl_mem in_tex, + cl_mem aux_tex, + cl_mem out_tex, + size_t global_worksize, + const GeglRectangle *roi, + gint level); G_DEFINE_TYPE (GimpOperationSetAlpha, gimp_operation_set_alpha, @@ -78,6 +87,7 @@ gimp_operation_set_alpha_class_init (GimpOperationSetAlphaClass *klass) operation_class->prepare = gimp_operation_set_alpha_prepare; point_class->process = gimp_operation_set_alpha_process; + point_class->cl_process = gimp_operation_set_alpha_cl_process; g_object_class_install_property (object_class, PROP_VALUE, g_param_spec_double ("value", @@ -186,3 +196,63 @@ gimp_operation_set_alpha_process (GeglOperation *operation, return TRUE; } + +static gboolean +gimp_operation_set_alpha_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) +{ + static GeglClRunData *cl_data = NULL; + GimpOperationSetAlpha *self = GIMP_OPERATION_SET_ALPHA (operation); + /* self->value needs to be cast down to float to properly fit into cl_float */ + float value = (float)self->value; + gint cl_kernel = 0; + gint cl_err = 0; + + if (cl_data == NULL) + { + const char *kernel_name[] = { "kernel_gimp_operation_set_alpha_with_aux", + "kernel_gimp_operation_set_alpha", NULL }; + cl_data = gegl_cl_compile_and_build (set_alpha_cl_source, + kernel_name); + if (cl_data == NULL) + goto error; + } + + if (aux_tex != NULL) + { + cl_kernel = 0; + cl_err = gegl_cl_set_kernel_args (cl_data->kernel[cl_kernel], + sizeof(cl_mem), &in_tex, + sizeof(cl_mem), &aux_tex, + sizeof(cl_mem), &out_tex, + sizeof(cl_float), &value, + NULL); + CL_CHECK; + } + else + { + cl_kernel = 1; + cl_err = gegl_cl_set_kernel_args (cl_data->kernel[cl_kernel], + sizeof(cl_mem), &in_tex, + sizeof(cl_mem), &out_tex, + sizeof(cl_float), &value, + NULL); + CL_CHECK; + } + + cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (), + cl_data->kernel[cl_kernel], 1, + NULL, &global_worksize, NULL, + 0, NULL, NULL); + CL_CHECK; + + return FALSE; + +error: + return TRUE; +} diff --git a/app/operations/opencl/set-alpha.cl b/app/operations/opencl/set-alpha.cl new file mode 100644 index 00000000000..87b6e907e8d --- /dev/null +++ b/app/operations/opencl/set-alpha.cl @@ -0,0 +1,24 @@ +__kernel void kernel_gimp_operation_set_alpha_with_aux (__global const float4 *in, + __global const float *aux, + __global float4 *out, + float value) +{ + int gid = get_global_id(0); + float4 in_v = in[gid]; + + in_v.w = value * aux[gid]; + + out[gid] = in_v; +} + +__kernel void kernel_gimp_operation_set_alpha (__global const float4 *in, + __global float4 *out, + float value) +{ + int gid = get_global_id(0); + float4 in_v = in[gid]; + + in_v.w = value; + + out[gid] = in_v; +} -- GitLab From 7d66d981a47ca202e521032d2bcf52a4facba6e3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ond=C5=99ej=20M=C3=ADchal?= Date: Mon, 18 Aug 2025 15:12:46 +0300 Subject: [PATCH 3/5] app:operations/layer-modes: Clarify names of GEGL vfuncs Personally, I find names like "parent", "real_parent" rather confusing when the class tree is not exactly simple. In the case of layer modes it is GObject -> GeglOperation -> GeglOperationPointCompose3 -> LayerMode. And 3 of these have a vfunc called 'process'. I find prepending the class name to the name of a vfunc implementation to be much more helpful in understanding which function is which. --- .../layer-modes/gimpoperationlayermode.c | 178 +++++++++--------- 1 file changed, 89 insertions(+), 89 deletions(-) diff --git a/app/operations/layer-modes/gimpoperationlayermode.c b/app/operations/layer-modes/gimpoperationlayermode.c index 8b129a5edf0..7c50405e667 100644 --- a/app/operations/layer-modes/gimpoperationlayermode.c +++ b/app/operations/layer-modes/gimpoperationlayermode.c @@ -68,61 +68,61 @@ typedef void (* CompositeFunc) (const gfloat *in, gint samples); -static void gimp_operation_layer_mode_finalize (GObject *object); -static void gimp_operation_layer_mode_set_property (GObject *object, - guint property_id, - const GValue *value, - GParamSpec *pspec); -static void gimp_operation_layer_mode_get_property (GObject *object, - guint property_id, - GValue *value, - GParamSpec *pspec); - -static void gimp_operation_layer_mode_prepare (GeglOperation *operation); -static GeglRectangle gimp_operation_layer_mode_get_bounding_box (GeglOperation *operation); -static gboolean gimp_operation_layer_mode_parent_process (GeglOperation *operation, - GeglOperationContext *context, - const gchar *output_prop, - const GeglRectangle *result, - gint level); - -static gboolean gimp_operation_layer_mode_process (GeglOperation *operation, - void *in, - void *layer, - void *mask, - void *out, - glong samples, - const GeglRectangle *roi, - gint level); - -static gboolean gimp_operation_layer_mode_real_parent_process (GeglOperation *operation, - GeglOperationContext *context, - const gchar *output_prop, - const GeglRectangle *result, - gint level); -static gboolean gimp_operation_layer_mode_real_process (GeglOperation *operation, - void *in, - void *layer, - void *mask, - void *out, - glong samples, - const GeglRectangle *roi, - gint level); - -static gboolean process_last_node (GeglOperation *operation, - void *in, - void *layer, - void *mask, - void *out, - glong samples, - const GeglRectangle *roi, - gint level); - -static void gimp_operation_layer_mode_cache_fishes (GimpOperationLayerMode *op, - const Babl *preferred_format, - const Babl **out_format, - const Babl **composite_to_blend_fish, - const Babl **blend_to_composite_fish); +static void gimp_operation_layer_mode_finalize (GObject *object); +static void gimp_operation_layer_mode_set_property (GObject *object, + guint property_id, + const GValue *value, + GParamSpec *pspec); +static void gimp_operation_layer_mode_get_property (GObject *object, + guint property_id, + GValue *value, + GParamSpec *pspec); + +static void gimp_operation_layer_mode_prepare (GeglOperation *operation); +static GeglRectangle gimp_operation_layer_mode_get_bounding_box (GeglOperation *operation); +static gboolean gimp_operation_layer_mode_operation_process (GeglOperation *operation, + GeglOperationContext *context, + const gchar *output_prop, + const GeglRectangle *result, + gint level); + +static gboolean gimp_operation_layer_mode_point_composer3_process (GeglOperation *operation, + void *in, + void *layer, + void *mask, + void *out, + glong samples, + const GeglRectangle *roi, + gint level); + +static gboolean gimp_operation_layer_mode_parent_process (GeglOperation *operation, + GeglOperationContext *context, + const gchar *output_prop, + const GeglRectangle *result, + gint level); +static gboolean gimp_operation_layer_mode_process (GeglOperation *operation, + void *in, + void *layer, + void *mask, + void *out, + glong samples, + const GeglRectangle *roi, + gint level); + +static gboolean process_last_node (GeglOperation *operation, + void *in, + void *layer, + void *mask, + void *out, + glong samples, + const GeglRectangle *roi, + gint level); + +static void gimp_operation_layer_mode_cache_fishes (GimpOperationLayerMode *op, + const Babl *preferred_format, + const Babl **out_format, + const Babl **composite_to_blend_fish, + const Babl **blend_to_composite_fish); G_DEFINE_TYPE (GimpOperationLayerMode, gimp_operation_layer_mode, @@ -157,12 +157,12 @@ gimp_operation_layer_mode_class_init (GimpOperationLayerModeClass *klass) operation_class->prepare = gimp_operation_layer_mode_prepare; operation_class->get_bounding_box = gimp_operation_layer_mode_get_bounding_box; - operation_class->process = gimp_operation_layer_mode_parent_process; + operation_class->process = gimp_operation_layer_mode_operation_process; - point_composer3_class->process = gimp_operation_layer_mode_process; + point_composer3_class->process = gimp_operation_layer_mode_point_composer3_process; - klass->parent_process = gimp_operation_layer_mode_real_parent_process; - klass->process = gimp_operation_layer_mode_real_process; + klass->parent_process = gimp_operation_layer_mode_parent_process; + klass->process = gimp_operation_layer_mode_process; klass->get_affected_region = NULL; g_object_class_install_property (object_class, PROP_LAYER_MODE, @@ -420,11 +420,11 @@ gimp_operation_layer_mode_get_bounding_box (GeglOperation *op) } static gboolean -gimp_operation_layer_mode_parent_process (GeglOperation *operation, - GeglOperationContext *context, - const gchar *output_prop, - const GeglRectangle *result, - gint level) +gimp_operation_layer_mode_operation_process (GeglOperation *operation, + GeglOperationContext *context, + const gchar *output_prop, + const GeglRectangle *result, + gint level) { GimpOperationLayerMode *point = GIMP_OPERATION_LAYER_MODE (operation); @@ -458,25 +458,25 @@ gimp_operation_layer_mode_parent_process (GeglOperation *operation, } static gboolean -gimp_operation_layer_mode_process (GeglOperation *operation, - void *in, - void *layer, - void *mask, - void *out, - glong samples, - const GeglRectangle *roi, - gint level) +gimp_operation_layer_mode_point_composer3_process (GeglOperation *operation, + void *in, + void *layer, + void *mask, + void *out, + glong samples, + const GeglRectangle *roi, + gint level) { return ((GimpOperationLayerMode *) operation)->function ( operation, in, layer, mask, out, samples, roi, level); } static gboolean -gimp_operation_layer_mode_real_parent_process (GeglOperation *operation, - GeglOperationContext *context, - const gchar *output_prop, - const GeglRectangle *result, - gint level) +gimp_operation_layer_mode_parent_process (GeglOperation *operation, + GeglOperationContext *context, + const gchar *output_prop, + const GeglRectangle *result, + gint level) { GimpOperationLayerMode *point = GIMP_OPERATION_LAYER_MODE (operation); GObject *input; @@ -610,14 +610,14 @@ gimp_operation_layer_mode_real_parent_process (GeglOperation *operation, } static gboolean -gimp_operation_layer_mode_real_process (GeglOperation *operation, - void *in_p, - void *layer_p, - void *mask_p, - void *out_p, - glong samples, - const GeglRectangle *roi, - gint level) +gimp_operation_layer_mode_process (GeglOperation *operation, + void *in_p, + void *layer_p, + void *mask_p, + void *out_p, + glong samples, + const GeglRectangle *roi, + gint level) { GimpOperationLayerMode *layer_mode = (gpointer) operation; gfloat *in = in_p; @@ -641,10 +641,10 @@ gimp_operation_layer_mode_real_process (GeglOperation *operation, */ while (samples > GIMP_COMPOSITE_BLEND_MAX_SAMPLES) { - gimp_operation_layer_mode_real_process (operation, - in, layer, mask, out, - GIMP_COMPOSITE_BLEND_MAX_SAMPLES, - roi, level); + gimp_operation_layer_mode_process (operation, + in, layer, mask, out, + GIMP_COMPOSITE_BLEND_MAX_SAMPLES, + roi, level); in += 4 * GIMP_COMPOSITE_BLEND_MAX_SAMPLES; layer += 4 * GIMP_COMPOSITE_BLEND_MAX_SAMPLES; -- GitLab From 41c0ce688bb739ede3fdaa0d18d3b9557c0ac83f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ond=C5=99ej=20M=C3=ADchal?= Date: Mon, 18 Aug 2025 15:14:15 +0300 Subject: [PATCH 4/5] app/operations/layer-modes: Enable OpenCL for layer modes With this layer mode implementations can register a 'cl_process' vfunc that will automatically be called when GIMP has OpenCL support enabled. --- app/operations/layer-modes/gimp-layer-modes.c | 10 ++ app/operations/layer-modes/gimp-layer-modes.h | 1 + .../layer-modes/gimpoperationlayermode.c | 134 +++++++++++------- .../layer-modes/gimpoperationlayermode.h | 9 ++ app/operations/operations-types.h | 11 ++ 5 files changed, 110 insertions(+), 55 deletions(-) diff --git a/app/operations/layer-modes/gimp-layer-modes.c b/app/operations/layer-modes/gimp-layer-modes.c index e7430fa6bb3..70f1cea2737 100644 --- a/app/operations/layer-modes/gimp-layer-modes.c +++ b/app/operations/layer-modes/gimp-layer-modes.c @@ -1360,6 +1360,16 @@ gimp_layer_mode_get_function (GimpLayerMode mode) return GIMP_OPERATION_LAYER_MODE_GET_CLASS (operation)->process; } +GimpLayerModeCLFunc +gimp_layer_mode_get_cl_function (GimpLayerMode mode) +{ + GeglOperation *operation; + + operation = gimp_layer_mode_get_operation (mode); + + return GIMP_OPERATION_LAYER_MODE_GET_CLASS (operation)->cl_process; +} + GimpLayerModeBlendFunc gimp_layer_mode_get_blend_function (GimpLayerMode mode) { diff --git a/app/operations/layer-modes/gimp-layer-modes.h b/app/operations/layer-modes/gimp-layer-modes.h index a5e8b1b8d91..ec548c5ba9a 100644 --- a/app/operations/layer-modes/gimp-layer-modes.h +++ b/app/operations/layer-modes/gimp-layer-modes.h @@ -44,6 +44,7 @@ const gchar * gimp_layer_mode_get_operation_name (GimpLayer GeglOperation * gimp_layer_mode_get_operation (GimpLayerMode mode); GimpLayerModeFunc gimp_layer_mode_get_function (GimpLayerMode mode); +GimpLayerModeCLFunc gimp_layer_mode_get_cl_function (GimpLayerMode mode); GimpLayerModeBlendFunc gimp_layer_mode_get_blend_function (GimpLayerMode mode); GimpLayerModeContext gimp_layer_mode_get_context (GimpLayerMode mode); diff --git a/app/operations/layer-modes/gimpoperationlayermode.c b/app/operations/layer-modes/gimpoperationlayermode.c index 7c50405e667..9174381de41 100644 --- a/app/operations/layer-modes/gimpoperationlayermode.c +++ b/app/operations/layer-modes/gimpoperationlayermode.c @@ -68,61 +68,69 @@ typedef void (* CompositeFunc) (const gfloat *in, gint samples); -static void gimp_operation_layer_mode_finalize (GObject *object); -static void gimp_operation_layer_mode_set_property (GObject *object, - guint property_id, - const GValue *value, - GParamSpec *pspec); -static void gimp_operation_layer_mode_get_property (GObject *object, - guint property_id, - GValue *value, - GParamSpec *pspec); - -static void gimp_operation_layer_mode_prepare (GeglOperation *operation); -static GeglRectangle gimp_operation_layer_mode_get_bounding_box (GeglOperation *operation); -static gboolean gimp_operation_layer_mode_operation_process (GeglOperation *operation, - GeglOperationContext *context, - const gchar *output_prop, - const GeglRectangle *result, - gint level); - -static gboolean gimp_operation_layer_mode_point_composer3_process (GeglOperation *operation, - void *in, - void *layer, - void *mask, - void *out, - glong samples, - const GeglRectangle *roi, - gint level); - -static gboolean gimp_operation_layer_mode_parent_process (GeglOperation *operation, - GeglOperationContext *context, - const gchar *output_prop, - const GeglRectangle *result, - gint level); -static gboolean gimp_operation_layer_mode_process (GeglOperation *operation, - void *in, - void *layer, - void *mask, - void *out, - glong samples, - const GeglRectangle *roi, - gint level); - -static gboolean process_last_node (GeglOperation *operation, - void *in, - void *layer, - void *mask, - void *out, - glong samples, - const GeglRectangle *roi, - gint level); - -static void gimp_operation_layer_mode_cache_fishes (GimpOperationLayerMode *op, - const Babl *preferred_format, - const Babl **out_format, - const Babl **composite_to_blend_fish, - const Babl **blend_to_composite_fish); +static void gimp_operation_layer_mode_finalize (GObject *object); +static void gimp_operation_layer_mode_set_property (GObject *object, + guint property_id, + const GValue *value, + GParamSpec *pspec); +static void gimp_operation_layer_mode_get_property (GObject *object, + guint property_id, + GValue *value, + GParamSpec *pspec); + +static void gimp_operation_layer_mode_prepare (GeglOperation *operation); +static GeglRectangle gimp_operation_layer_mode_get_bounding_box (GeglOperation *operation); +static gboolean gimp_operation_layer_mode_operation_process (GeglOperation *operation, + GeglOperationContext *context, + const gchar *output_prop, + const GeglRectangle *result, + gint level); + +static gboolean gimp_operation_layer_mode_point_composer3_process (GeglOperation *operation, + void *in, + void *layer, + void *mask, + void *out, + glong samples, + const GeglRectangle *roi, + gint level); +static gboolean gimp_operation_layer_mode_point_composer3_cl_process (GeglOperation *self, + cl_mem in_tex, + cl_mem aux_tex, + cl_mem mask_tex, + cl_mem out_tex, + size_t global_worksize, + const GeglRectangle *roi, + gint level); + +static gboolean gimp_operation_layer_mode_parent_process (GeglOperation *operation, + GeglOperationContext *context, + const gchar *output_prop, + const GeglRectangle *result, + gint level); +static gboolean gimp_operation_layer_mode_process (GeglOperation *operation, + void *in, + void *layer, + void *mask, + void *out, + glong samples, + const GeglRectangle *roi, + gint level); + +static gboolean process_last_node (GeglOperation *operation, + void *in, + void *layer, + void *mask, + void *out, + glong samples, + const GeglRectangle *roi, + gint level); + +static void gimp_operation_layer_mode_cache_fishes (GimpOperationLayerMode *op, + const Babl *preferred_format, + const Babl **out_format, + const Babl **composite_to_blend_fish, + const Babl **blend_to_composite_fish); G_DEFINE_TYPE (GimpOperationLayerMode, gimp_operation_layer_mode, @@ -160,6 +168,7 @@ gimp_operation_layer_mode_class_init (GimpOperationLayerModeClass *klass) operation_class->process = gimp_operation_layer_mode_operation_process; point_composer3_class->process = gimp_operation_layer_mode_point_composer3_process; + point_composer3_class->cl_process = gimp_operation_layer_mode_point_composer3_cl_process; klass->parent_process = gimp_operation_layer_mode_parent_process; klass->process = gimp_operation_layer_mode_process; @@ -320,6 +329,7 @@ gimp_operation_layer_mode_prepare (GeglOperation *operation) } self->function = gimp_layer_mode_get_function (self->layer_mode); + self->cl_function = gimp_layer_mode_get_cl_function (self->layer_mode); self->blend_function = gimp_layer_mode_get_blend_function (self->layer_mode); input_extent = gegl_operation_source_get_bounding_box (operation, "input"); @@ -457,6 +467,20 @@ gimp_operation_layer_mode_operation_process (GeglOperation *operation, operation, context, output_prop, result, level); } +static gboolean +gimp_operation_layer_mode_point_composer3_cl_process (GeglOperation *operation, + cl_mem in_tex, + cl_mem aux_tex, + cl_mem mask_tex, + cl_mem out_tex, + size_t global_worksize, + const GeglRectangle *roi, + gint level) +{ + return ((GimpOperationLayerMode *) operation)->cl_function ( + operation, in_tex, aux_tex, mask_tex, out_tex, global_worksize, roi, level); +} + static gboolean gimp_operation_layer_mode_point_composer3_process (GeglOperation *operation, void *in, diff --git a/app/operations/layer-modes/gimpoperationlayermode.h b/app/operations/layer-modes/gimpoperationlayermode.h index 6a2cabc76ce..baa9f0af1d4 100644 --- a/app/operations/layer-modes/gimpoperationlayermode.h +++ b/app/operations/layer-modes/gimpoperationlayermode.h @@ -50,6 +50,7 @@ struct _GimpOperationLayerMode GimpLayerCompositeMode prop_composite_mode; GimpLayerModeFunc function; + GimpLayerModeCLFunc cl_function; GimpLayerModeBlendFunc blend_function; gboolean is_last_node; gboolean has_mask; @@ -73,6 +74,14 @@ struct _GimpOperationLayerModeClass glong samples, const GeglRectangle *roi, gint level); + gboolean (* cl_process) (GeglOperation *self, + cl_mem in_tex, + cl_mem aux_tex, + cl_mem mask_tex, + cl_mem out_tex, + size_t global_worksize, + const GeglRectangle *roi, + gint level); /* Returns the composite region (any combination of the layer and the * backdrop) that the layer mode affects. Most modes only affect the diff --git a/app/operations/operations-types.h b/app/operations/operations-types.h index f0cac9610f4..9e41c88b852 100644 --- a/app/operations/operations-types.h +++ b/app/operations/operations-types.h @@ -22,6 +22,7 @@ #include #include "gegl/gimp-gegl-types.h" +#include "opencl/gegl-cl.h" #include "operations-enums.h" @@ -64,6 +65,16 @@ typedef gboolean (* GimpLayerModeFunc) (GeglOperation *operation, const GeglRectangle *roi, gint level); +typedef gboolean (* GimpLayerModeCLFunc) (GeglOperation *operation, + cl_mem in_tex, + cl_mem aux_tex, + cl_mem mask_tex, + cl_mem out_tex, + size_t global_worksize, + const GeglRectangle *roi, + gint level); + + typedef void (* GimpLayerModeBlendFunc) (GeglOperation *operation, const gfloat *in, const gfloat *layer, -- GitLab From a2766fe77e364eaf076f82d3313516766ecdc7d5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ond=C5=99ej=20M=C3=ADchal?= Date: Mon, 18 Aug 2025 15:15:09 +0300 Subject: [PATCH 5/5] app/operations: Add OpenCL implementation for normal layer mode Initial test to stress test support for OpenCL in GEGL's GeglOperationPointComposer3. Supports all composite modes and masks. --- .../layer-modes/gimpoperationnormal.c | 89 +++++++- .../layer-modes/gimpoperationnormal.h | 9 + app/operations/opencl/layer-mode-normal.cl | 201 ++++++++++++++++++ app/operations/opencl/meson.build | 1 + 4 files changed, 299 insertions(+), 1 deletion(-) create mode 100644 app/operations/opencl/layer-mode-normal.cl diff --git a/app/operations/layer-modes/gimpoperationnormal.c b/app/operations/layer-modes/gimpoperationnormal.c index 9bd87052fcc..6cd17ecd059 100644 --- a/app/operations/layer-modes/gimpoperationnormal.c +++ b/app/operations/layer-modes/gimpoperationnormal.c @@ -24,8 +24,10 @@ #include #include "libgimpbase/gimpbase.h" +#include "opencl/gegl-cl.h" #include "../operations-types.h" +#include "app/operations/opencl/layer-mode-normal.cl.h" #include "gimpoperationnormal.h" @@ -64,7 +66,8 @@ gimp_operation_normal_class_init (GimpOperationNormalClass *klass) "reference-composition", reference_xml, NULL); - layer_mode_class->process = gimp_operation_normal_process; + layer_mode_class->process = gimp_operation_normal_process; + layer_mode_class->cl_process = gimp_operation_normal_cl_process; #if COMPILE_SSE2_INTRINISICS if (gimp_cpu_accel_get_support() & GIMP_CPU_ACCEL_X86_SSE2) @@ -82,6 +85,90 @@ gimp_operation_normal_init (GimpOperationNormal *self) { } +gboolean +gimp_operation_normal_cl_process (GeglOperation *operation, + cl_mem in_tex, + cl_mem layer_tex, + cl_mem mask_tex, + cl_mem out_tex, + size_t global_worksize, + const GeglRectangle *roi, + gint level) +{ + static GeglClRunData *cl_data = NULL; + GimpOperationLayerMode *layer_mode = GIMP_OPERATION_LAYER_MODE (operation); + /* layer_mode->opacity needs to be cast down to float to properly fit into cl_float */ + float opacity = (float)layer_mode->opacity; + cl_int cl_err = 0; + gint kernel = 0; + + if (!cl_data) + { + const char *kernel_names[] = { "kernel_gimp_operation_normal_union_with_mask", + "kernel_gimp_operation_normal_union", + "kernel_gimp_operation_normal_clip_to_backdrop_with_mask", + "kernel_gimp_operation_normal_clip_to_backdrop", + "kernel_gimp_operation_normal_clip_to_layer_with_mask", + "kernel_gimp_operation_normal_clip_to_layer", + "kernel_gimp_operation_normal_intersection_with_mask", + "kernel_gimp_operation_normal_intersection", NULL }; + cl_data = gegl_cl_compile_and_build (layer_mode_normal_cl_source, kernel_names); + if (!cl_data) + return TRUE; + } + + switch (layer_mode->composite_mode) + { + case GIMP_LAYER_COMPOSITE_UNION: + case GIMP_LAYER_COMPOSITE_AUTO: + kernel = 0; + break; + case GIMP_LAYER_COMPOSITE_CLIP_TO_BACKDROP: + kernel = 2; + break; + case GIMP_LAYER_COMPOSITE_CLIP_TO_LAYER: + kernel = 4; + break; + case GIMP_LAYER_COMPOSITE_INTERSECTION: + kernel = 6; + break; + } + + if (mask_tex != NULL) + { + cl_err = gegl_cl_set_kernel_args (cl_data->kernel[kernel], + sizeof(cl_mem), &in_tex, + sizeof(cl_mem), &layer_tex, + sizeof(cl_mem), &mask_tex, + sizeof(cl_mem), &out_tex, + sizeof(cl_float), &opacity, + NULL); + CL_CHECK; + } + else + { + kernel += 1; + cl_err = gegl_cl_set_kernel_args (cl_data->kernel[kernel], + sizeof(cl_mem), &in_tex, + sizeof(cl_mem), &layer_tex, + sizeof(cl_mem), &out_tex, + sizeof(cl_float), &opacity, + NULL); + 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; +} + gboolean gimp_operation_normal_process (GeglOperation *op, void *in_p, diff --git a/app/operations/layer-modes/gimpoperationnormal.h b/app/operations/layer-modes/gimpoperationnormal.h index 011ee0380c4..e0edb700804 100644 --- a/app/operations/layer-modes/gimpoperationnormal.h +++ b/app/operations/layer-modes/gimpoperationnormal.h @@ -84,3 +84,12 @@ gboolean gimp_operation_normal_process_sse4 (GeglOperation *op, gint level); #endif /* COMPILE_SSE4_1_INTRINISICS */ + +gboolean gimp_operation_normal_cl_process (GeglOperation *operation, + cl_mem in_tex, + cl_mem layer_tex, + cl_mem mask_tex, + cl_mem out_tex, + size_t global_worksize, + const GeglRectangle *roi, + gint level); diff --git a/app/operations/opencl/layer-mode-normal.cl b/app/operations/opencl/layer-mode-normal.cl new file mode 100644 index 00000000000..5c0533b02fd --- /dev/null +++ b/app/operations/opencl/layer-mode-normal.cl @@ -0,0 +1,201 @@ +__kernel void kernel_gimp_operation_normal_union_with_mask (__global const float4 *in, + __global const float4 *layer, + __global const float *mask, + __global float4 *out, + float opacity) +{ + int gid = get_global_id (0); + + float4 in_v = in[gid]; + float4 layer_v = layer[gid]; + float4 out_v = in_v; + + float layer_alpha = layer_v.w * opacity * mask[gid]; + out_v.w = layer_alpha + in_v.w - layer_alpha * in_v.w; + + float opaque = (out_v.w == 0.f); + float translucent = 1.f - opaque; + + float layer_weight = (opaque * 0.f) + (translucent * (layer_alpha / out_v.w)); + float in_weight = (opaque * 0.f) + (translucent * (1.f - layer_weight)); + + out_v.x = (opaque * in_v.x) + (translucent * (layer_v.x * layer_weight + in_v.x * in_weight)); + out_v.y = (opaque * in_v.y) + (translucent * (layer_v.y * layer_weight + in_v.y * in_weight)); + out_v.z = (opaque * in_v.z) + (translucent * (layer_v.z * layer_weight + in_v.z * in_weight)); + + out[gid] = out_v; +} + +__kernel void kernel_gimp_operation_normal_union (__global const float4 *in, + __global const float4 *layer, + __global float4 *out, + float opacity) +{ + int gid = get_global_id (0); + + float4 in_v = in[gid]; + float4 layer_v = layer[gid]; + float4 out_v = in_v; + + float layer_alpha = layer_v.w * opacity; + out_v.w = layer_alpha + in_v.w - layer_alpha * in_v.w; + + float opaque = (out_v.w == 0.f); + float translucent = 1.f - opaque; + + float layer_weight = (opaque * 0.f) + (translucent * (layer_alpha / out_v.w)); + float in_weight = (opaque * 0.f) + (translucent * (1.f - layer_weight)); + + out_v.x = (opaque * in_v.x) + (translucent * (layer_v.x * layer_weight + in_v.x * in_weight)); + out_v.y = (opaque * in_v.y) + (translucent * (layer_v.y * layer_weight + in_v.y * in_weight)); + out_v.z = (opaque * in_v.z) + (translucent * (layer_v.z * layer_weight + in_v.z * in_weight)); + + out[gid] = out_v; +} + +__kernel void kernel_gimp_operation_normal_clip_to_backdrop_with_mask (__global const float4 *in, + __global const float4 *layer, + __global const float *mask, + __global float4 *out, + float opacity) +{ + int gid = get_global_id (0); + + float4 in_v = in[gid]; + float4 layer_v = layer[gid]; + float4 out_v = in_v; + + float layer_alpha = layer_v.w * opacity * mask[gid]; + + float opaque = (out_v.w == 0.f); + float translucent = 1.f - opaque; + + out_v.x = (opaque * in_v.x) + (translucent * (in_v.x + (layer_v.x - in_v.x) * layer_alpha)); + out_v.y = (opaque * in_v.y) + (translucent * (in_v.y + (layer_v.y - in_v.y) * layer_alpha)); + out_v.z = (opaque * in_v.z) + (translucent * (in_v.z + (layer_v.z - in_v.z) * layer_alpha)); + + out[gid] = out_v; +} + +__kernel void kernel_gimp_operation_normal_clip_to_backdrop (__global const float4 *in, + __global const float4 *layer, + __global float4 *out, + float opacity) +{ + int gid = get_global_id (0); + + float4 in_v = in[gid]; + float4 layer_v = layer[gid]; + float4 out_v = in_v; + + float layer_alpha = layer_v.w * opacity; + + float opaque = (out_v.w == 0.f); + float translucent = 1.f - opaque; + + out_v.x = (opaque * in_v.x) + (translucent * (in_v.x + (layer_v.x - in_v.x) * layer_alpha)); + out_v.y = (opaque * in_v.y) + (translucent * (in_v.y + (layer_v.y - in_v.y) * layer_alpha)); + out_v.z = (opaque * in_v.z) + (translucent * (in_v.z + (layer_v.z - in_v.z) * layer_alpha)); + + out[gid] = out_v; +} + +__kernel void kernel_gimp_operation_normal_clip_to_layer_with_mask (__global const float4 *in, + __global const float4 *layer, + __global const float *mask, + __global float4 *out, + float opacity) +{ + int gid = get_global_id (0); + + float4 in_v = in[gid]; + float4 layer_v = layer[gid]; + float4 out_v = in_v; + + float layer_alpha = layer_v.w * opacity * mask[gid]; + out_v.w = layer_alpha; + + if (out_v.w != 0.f) + { + out_v.x = layer_v.x; + out_v.y = layer_v.y; + out_v.z = layer_v.z; + } + + out[gid] = out_v; +} + + +__kernel void kernel_gimp_operation_normal_clip_to_layer (__global const float4 *in, + __global const float4 *layer, + __global float4 *out, + float opacity) +{ + int gid = get_global_id (0); + + float4 in_v = in[gid]; + float4 layer_v = layer[gid]; + float4 out_v = in_v; + + float layer_alpha = layer_v.w * opacity; + out_v.w = layer_alpha; + + if (out_v.w != 0.f) + { + out_v.x = layer_v.x; + out_v.y = layer_v.y; + out_v.z = layer_v.z; + } + + out[gid] = out_v; +} + +__kernel void kernel_gimp_operation_normal_intersection_with_mask (__global const float4 *in, + __global const float4 *layer, + __global const float *mask, + __global float4 *out, + float opacity) +{ + int gid = get_global_id (0); + + float4 in_v = in[gid]; + float4 layer_v = layer[gid]; + float4 out_v = in_v; + + float layer_alpha = layer_v.w * opacity * mask[gid]; + out_v.w = in_v.w * layer_alpha; + + if (out_v.w != 0.f) + { + out_v.x = layer_v.x; + out_v.y = layer_v.y; + out_v.z = layer_v.z; + } + + out[gid] = out_v; + +} + +__kernel void kernel_gimp_operation_normal_intersection (__global const float4 *in, + __global const float4 *layer, + __global float4 *out, + float opacity) +{ + int gid = get_global_id (0); + + float4 in_v = in[gid]; + float4 layer_v = layer[gid]; + float4 out_v = in[gid]; + + float layer_alpha = layer_v.w * opacity; + out_v.w = in_v.w * layer_alpha; + + if (out_v.w != 0.f) + { + out_v.x = layer_v.x; + out_v.y = layer_v.y; + out_v.z = layer_v.z; + } + + out[gid] = out_v; +} diff --git a/app/operations/opencl/meson.build b/app/operations/opencl/meson.build index 31967af7d2a..aca1d2c4910 100644 --- a/app/operations/opencl/meson.build +++ b/app/operations/opencl/meson.build @@ -1,6 +1,7 @@ cltostring = files('cltostring.py') opencl_sources = [ + 'layer-mode-normal.cl', 'set-alpha.cl', ] -- GitLab