opencl: many changes

- splitting kernels from remaining filters
- removing unnecessary opencl barriers
- stop using cl_int to report GEGL errors (just opencl ones)

1 files changed, 13 insertions, 0 deletions

diff --git a/opencl/levels.cl b/opencl/levels.cl
new file mode 100644
index 00000000..80b91d2f
--- /dev/null
+++ b/opencl/levels.cl
@@ -0,0 +1,13 @@
+__kernel void kernel_levels(__global const float4     *in,
+                            __global       float4     *out,
+                            float in_offset,
+                            float out_offset,
+                            float scale)
+{
+  int gid = get_global_id(0);
+  float4 in_v  = in[gid];
+  float4 out_v;
+  out_v.xyz = (in_v.xyz - in_offset) * scale + out_offset;
+  out_v.w   =  in_v.w;
+  out[gid]  =  out_v;
+}