diff --git a/app/operations/gimpoperationsetalpha.c b/app/operations/gimpoperationsetalpha.c index b41af3caa1d20bddc7b894c0f39ba50f843c9631..7909fc17a7e98c88f96be324397e7bdb2b057d5b 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/layer-modes/gimp-layer-modes.c b/app/operations/layer-modes/gimp-layer-modes.c index e7430fa6bb376e3523aeb37a3914626bca82492a..70f1cea27371cd9a6dbe7603cbbfecacb7ace5e5 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 a5e8b1b8d91d23c2d4df52df8fad4be73cbde869..ec548c5ba9a2d1aa47e87d02f60efac73c896454 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 8b129a5edf0bee285d7422cc555c088693e77552..9174381de41f5ae6d0d9ca41c198f8579d213661 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_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_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, @@ -157,12 +165,13 @@ 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; + point_composer3_class->cl_process = gimp_operation_layer_mode_point_composer3_cl_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, @@ -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"); @@ -420,11 +430,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 +468,39 @@ 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_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, + 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 +634,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 +665,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; diff --git a/app/operations/layer-modes/gimpoperationlayermode.h b/app/operations/layer-modes/gimpoperationlayermode.h index 6a2cabc76ce5299d13419a45f08c19eb19f20783..baa9f0af1d465d6347913b0c375b4d2c35f987ee 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/layer-modes/gimpoperationnormal.c b/app/operations/layer-modes/gimpoperationnormal.c index 9bd87052fcc1da58365d1cfb594019a719534629..6cd17ecd0594c50ec8f62d1926a6a0db2d5e9d54 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 011ee0380c43405e03146674c1785f056c56a3c0..e0edb700804db85c55e65a32249a40d2270b4481 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/meson.build b/app/operations/meson.build index 14f210b9893367739cc4e32df00cd2cc58050799..399546fb6cd05d6a00ce1e4c001e9843032b86b6 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 0000000000000000000000000000000000000000..ad1289479bef420f634df001d923f042b3633208 --- /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/layer-mode-normal.cl b/app/operations/opencl/layer-mode-normal.cl new file mode 100644 index 0000000000000000000000000000000000000000..5c0533b02fdda95f790a41faf85090c63935312f --- /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 new file mode 100644 index 0000000000000000000000000000000000000000..aca1d2c491007ad311d0a6593fed14fcf9f30d9c --- /dev/null +++ b/app/operations/opencl/meson.build @@ -0,0 +1,15 @@ +cltostring = files('cltostring.py') + +opencl_sources = [ + 'layer-mode-normal.cl', + '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 diff --git a/app/operations/opencl/set-alpha.cl b/app/operations/opencl/set-alpha.cl new file mode 100644 index 0000000000000000000000000000000000000000..87b6e907e8d18edf5c1c1b1405020edfda245687 --- /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; +} diff --git a/app/operations/operations-types.h b/app/operations/operations-types.h index f0cac9610f40018a60d59c6e5bb6a20463c0cc12..9e41c88b8525e01231a32a428bdacc4c0b00d913 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,