Hi again Thierry! When I told drawables I did mean some logical surface you are drawing to. Assuming the contents would not be like AAA-game or really high in polygon count, CoreGraphics would be a great choice, so there is no absolute need to do it with Metal. But if there is any good reason to do it with Metal, Metal totally supports off-screen rendering to linear textures (backed by an MTLBuffer) so you can achieve what you need. Though again, I doubt there will be any benefit assuming you are light on content. Best regards, Eugene.

Hi pgelvinaz! Please file feedback assistant ticket and post here its number, and please attach dev-signed ipa with sysdiagnose. We will look into what is going on, it is really hard to say why exactly you run out of memory just by looking at the video, so we need a repro to investigate. Thanks, Eugene.

Hi hawkiyc! Yes, we are aware of this issue and already working on a fix. We will update FB assistant ticket once the fix is included into the macOS12 beta seed. Thanks!

Hi Alecazam! Thanks for your feedback! Could you please file a feature request with Feedback Assistant and provide us with a FB ticket number? Thanks so much!

Aaargh, can you try adding #pragma unroll (1) before your loop? If it doesn't help, please file a feedback assistant ticket with a reproducer, thanks!

Sorry, formatting is insane, I am impressed how all the CR/LF just turned into nothing. But please check 6.13.3 in MSL2.4 spec.

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); }`