Hi jcookie!
Is this sample from Metal Shading Language v 2.4 is what you are looking for?
6.13.3 Writing an Imageblock Slice to a Region in a Texture
Example:
`struct Foo {
half4 a;
int b;
float c;
};
kernel void
my_kernel(texture2d src [[ texture(0) ]],
texture2d<half, access::write> dst [[ texture(1) ]], imageblock img_blk,
ushort2 lid [[ thread_position_in_threadgroup ]], ushort2 gid [[ thread_position_in_grid ]])
{
// Read the pixel from the input image using the thread ID.
half4 clr = src.read(gid);
// Get the image slice.
threadgroup_imageblock Foo* f = img_blk.data(lid);
// Write the pixel in the imageblock using the thread ID in threadgroup.
f->a = clr;
// A barrier to make sure all threads finish writing to the imageblock. // In this case, each thread writes to its location in the imageblock // so a barrier is not necessary.
threadgroup_barrier(mem_flags::mem_threadgroup_imageblock);
// Process the pixels in imageblock, and update the elements in slice.
process_pixels_in_imageblock(img_blk, gid, lid);
// A barrier to make sure all threads finish writing to the elements in the // imageblock.
threadgroup_barrier(mem_flags::mem_threadgroup_imageblock);
// Write a specific element in an imageblock to the output image.
// Only one thread in the threadgroup performs the imageblock write.
if (lid.x == 0 && lid.y == 0)
dst.write(img_blk.slice(f->a), gid);
}`
Topic:
Graphics & Games
SubTopic:
General
Tags: