#!amber

SHADER compute write_imagef OPENCL-C
kernel void foo(write_only image2d_t image, global uint4* data) {
  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);
  write_imageui(image, coord, data[linear]);
}
END

BUFFER texture DATA_TYPE vec2<uint32> WIDTH 2 HEIGHT 2 FILL 0
BUFFER data DATA_TYPE vec4<uint32> DATA
1 2 3 4
2 3 4 1
3 4 1 2
4 1 2 3
END

PIPELINE compute write_pipe
  ATTACH write_imagef ENTRY_POINT foo
  BIND BUFFER data KERNEL ARG_NAME data
  BIND BUFFER texture KERNEL ARG_NAME image
END

RUN write_pipe 2 2 1

EXPECT texture IDX 0  EQ 1 2
EXPECT texture IDX 8  EQ 2 3
EXPECT texture IDX 16 EQ 3 4
EXPECT texture IDX 24 EQ 4 1