#!amber

SHADER compute fill GLSL
#version 450
layout(set=0, binding=0, rgba32i) uniform iimage2D im;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
  ivec2 coord = ivec2(gl_GlobalInvocationID.xy);
  ivec4 data = ivec4(gl_GlobalInvocationID.x,
                     -3,
                     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, global int4* out) {
  int gid_x = get_global_id(0);
  int gid_y = get_global_id(1);
  int linear = 2 * gid_y + gid_x;
  int2 coord = (int2)(gid_x, gid_y);
  out[linear] = read_imagei(image, coord);
}
END

BUFFER texture DATA_TYPE vec4<int32> WIDTH 2 HEIGHT 2 FILL 4
BUFFER out_buf DATA_TYPE vec4<int32> DATA
2 2 2 2
2 2 2 2
2 2 2 2
2 2 2 2
END

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
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 -3 0  0
EXPECT out_buf IDX 16 EQ 1 -3 1  0
EXPECT out_buf IDX 32 EQ 0 -3 1 -1
EXPECT out_buf IDX 48 EQ 1 -3 2 -1