[gegl] opencl: new opacity kernels

[Victor Matheus de Araujo Oliveira <vmaolive src gnome org>]

commit 4261903818e384f5605c4c0c7b84e073864e837d
Author: Victor Oliveira <victormatheus gmail com>
Date:   Sun Jan 13 21:01:14 2013 -0200

    opencl: new opacity kernels

 opencl/opacity.cl           |   25 +++++++++++++++++++++++++
 opencl/opacity.cl.h         |   27 +++++++++++++++++++++++++++
 operations/common/opacity.c |    6 +++---
 3 files changed, 55 insertions(+), 3 deletions(-)
---
diff --git a/opencl/opacity.cl b/opencl/opacity.cl
new file mode 100644
index 0000000..ef17d7d
--- /dev/null
+++ b/opencl/opacity.cl
@@ -0,0 +1,25 @@
+__kernel void gegl_opacity_RaGaBaA_float (__global const float4     *in,
+                                          __global const float      *aux,
+                                          __global       float4     *out,
+                                          float value)
+{
+  int gid = get_global_id(0);
+  float4 in_v  = in [gid];
+  float  aux_v = (aux)? aux[gid] : 1.0f;
+  float4 out_v;
+  out_v = in_v * aux_v * value;
+  out[gid]  =  out_v;
+}
+__kernel void gegl_opacity_RGBA_float (__global const float4     *in,
+                                       __global const float      *aux,
+                                       __global       float4     *out,
+                                       float value)
+{
+  int gid = get_global_id(0);
+  float4 in_v  = in [gid];
+  float  aux_v = (aux)? aux[gid] : 1.0f;
+  float4 out_v;
+  out_v = (in_v.x, in_v.y, in_v.z, in_v.w * aux_v * value);
+  out[gid]  =  out_v;
+}
+
diff --git a/opencl/opacity.cl.h b/opencl/opacity.cl.h
new file mode 100644
index 0000000..8586e98
--- /dev/null
+++ b/opencl/opacity.cl.h
@@ -0,0 +1,27 @@
+static const char* opacity_cl_source =
+"__kernel void gegl_opacity_RaGaBaA_float (__global const float4     *in,      \n"
+"                                          __global const float      *aux,     \n"
+"                                          __global       float4     *out,     \n"
+"                                          float value)                        \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  float4 in_v  = in [gid];                                                    \n"
+"  float  aux_v = (aux)? aux[gid] : 1.0f;                                      \n"
+"  float4 out_v;                                                               \n"
+"  out_v = in_v * aux_v * value;                                               \n"
+"  out[gid]  =  out_v;                                                         \n"
+"}                                                                             \n"
+"__kernel void gegl_opacity_RGBA_float (__global const float4     *in,         \n"
+"                                       __global const float      *aux,        \n"
+"                                       __global       float4     *out,        \n"
+"                                       float value)                           \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  float4 in_v  = in [gid];                                                    \n"
+"  float  aux_v = (aux)? aux[gid] : 1.0f;                                      \n"
+"  float4 out_v;                                                               \n"
+"  out_v = (in_v.x, in_v.y, in_v.z, in_v.w * aux_v * value);                   \n"
+"  out[gid]  =  out_v;                                                         \n"
+"}                                                                             \n"
+"                                                                              \n"
+;
diff --git a/operations/common/opacity.c b/operations/common/opacity.c
index 8000c7b..dec9d80 100644
--- a/operations/common/opacity.c
+++ b/operations/common/opacity.c
@@ -219,11 +219,11 @@ cl_process (GeglOperation       *op,
 
   kernel = (GEGL_CHANT_PROPERTIES (op)->chant_data != NULL);
 
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 0, sizeof(cl_mem), (void*)&in_tex);
   CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&aux_tex);
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 1, sizeof(cl_mem), (void*)&aux_tex);
   CL_CHECK;
-  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&out_tex);
+  cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 2, sizeof(cl_mem), (void*)&out_tex);
   CL_CHECK;
 
   cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),