[gegl] operations/common/exposure: Fix unpacking of floats into float3

[Øyvind Kolås <ok src gnome org>]

commit 813162c8768b2aca1490cabdcc2f59005a4d94ed
Author: Debarshi Ray <debarshir gnome org>
Date:   Sat Nov 11 18:30:04 2017 +0100

    operations/common/exposure: Fix unpacking of floats into float3
    
    Fallout from 55cf906717510a56a65daff01919fc11a255ad45
    
    https://bugzilla.gnome.org/show_bug.cgi?id=789807

 operations/common/exposure.c |   15 +++++++++------
 1 files changed, 9 insertions(+), 6 deletions(-)
---
diff --git a/operations/common/exposure.c b/operations/common/exposure.c
index 114c937..38d2bb0 100644
--- a/operations/common/exposure.c
+++ b/operations/common/exposure.c
@@ -66,16 +66,19 @@ static GeglClRunData *cl_data_y = NULL;
 static GeglClRunData *cl_data_ya = NULL;
 
 static const char* kernel_source_rgb =
-"__kernel void kernel_exposure_rgb(__global const float3 *in,          \n"
-"                                  __global       float3 *out,         \n"
-"                                  float                  black_level, \n"
-"                                  float                  gain)        \n"
+"__kernel void kernel_exposure_rgb(__global const float *in,           \n"
+"                                  __global       float *out,          \n"
+"                                  float                 black_level,  \n"
+"                                  float                 gain)         \n"
 "{                                                                     \n"
 "  int gid = get_global_id(0);                                         \n"
-"  float3 in_v  = in[gid];                                             \n"
+"  int offset  = 3 * gid;                                              \n"
+"  float3 in_v = (float3) (in[offset], in[offset + 1], in[offset+2]);  \n"
 "  float3 out_v;                                                       \n"
 "  out_v.xyz =  ((in_v.xyz - black_level) * gain);                     \n"
-"  out[gid]  =  out_v;                                                 \n"
+"  out[offset]     = out_v.x;                                          \n"
+"  out[offset + 1] = out_v.y;                                          \n"
+"  out[offset + 2] = out_v.z;                                          \n"
 "}                                                                     \n";
 
 static const char* kernel_source_rgba =