#!amber

SHADER compute fill GLSL
#version 450
layout(set=0, binding=0, rgba32f) uniform image2D im;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
  ivec2 coord = ivec2(gl_GlobalInvocationID.xy);
  vec4 data = vec4(gl_GlobalInvocationID.x,
                   3.0,
                   gl_GlobalInvocationID.x + gl_GlobalInvocationID.y,
                   gl_GlobalInvocationID.y);
  imageStore(im, coord, data);
}
END

SHADER compute read_imagef OPENCL-C
kernel void foo(read_only image2d_t image, sampler_t sampler, global float4* out) {
  int gid_x = get_global_id(0);
  int gid_y = get_global_id(1);
  int linear = 2 * gid_y + gid_x;
  float2 coord = (float2)(gid_x, gid_y);
  out[linear] = read_imagef(image, sampler, coord);
}
END

BUFFER texture DATA_TYPE vec4<float> WIDTH 2 HEIGHT 2 FILL 0.0
BUFFER out_buf DATA_TYPE vec4<float> DATA
2.0 2.0 2.0 2.0
2.0 2.0 2.0 2.0
2.0 2.0 2.0 2.0
2.0 2.0 2.0 2.0
END
SAMPLER sampler  ADDRESS_MODE_U clamp_to_edge ADDRESS_MODE_V clamp_to_edge

PIPELINE compute read_pipe
  ATTACH read_imagef ENTRY_POINT foo
  BIND BUFFER out_buf KERNEL ARG_NAME out
  BIND BUFFER texture KERNEL ARG_NAME image
  BIND SAMPLER sampler KERNEL ARG_NAME sampler
END

PIPELINE compute fill_pipe
  ATTACH fill
  BIND BUFFER texture AS storage_image DESCRIPTOR_SET 0 BINDING 0
END

RUN fill_pipe 2 2 1
RUN read_pipe 2 2 1

EXPECT out_buf IDX 0  EQ 0.0 3.0 0.0 0.0
EXPECT out_buf IDX 16 EQ 1.0 3.0 1.0 0.0
EXPECT out_buf IDX 32 EQ 0.0 3.0 1.0 1.0
EXPECT out_buf IDX 48 EQ 1.0 3.0 2.0 1.0