amber/images/read_image2d_unsampled_rgba32ui.amber

#!amber

SHADER compute fill GLSL
#version 450
layout(set=0, binding=0, rgba32ui) uniform uimage2D im;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
  ivec2 coord = ivec2(gl_GlobalInvocationID.xy);
  uvec4 data = uvec4(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 uint4* 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_imageui(image, coord);
}
END

BUFFER texture DATA_TYPE vec4<uint32> WIDTH 2 HEIGHT 2 FILL 4
BUFFER out_buf DATA_TYPE vec4<uint32> 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