Metal Shading Language Specification

The following 
slice
 member functions are used to get a region of a slice for a given data 
member in the imageblock struct. This is used to write data associated with a specific data 
member described in the imageblock struct for all threads in the threadgroup to a specified 
region in a texture.  
data_member
 is a data member declared in the struct type specified in 
imageblock
. 
origin
 is the (x, y) coordinate in the slice. 
size
 is the actual size of the slice to be copied.  
const imageblock_slice 
slice(const threadgroup_imageblock E& data_member) const;  
const imageblock_slice 
slice(const threadgroup_imageblock E& data_member, ushort2 origin, ushort2 
size ) const;  
The 
slice(…)
 member function that does not take 
origin
 and 
size
 as arguments specifies 
the region to copy with an origin of (0,0) and the width and height of the imageblock.  
5.15.3  
Writing an Imageblock Slice to a Region in a Texture  
The following 
write(…)
 member function in these texture types are used to write pixels 
associated with a slice in the imageblock to a texture starting at location given by 
coord
. 
For 1D texture:  
void write(imageblock_slice slice, uint 
coord, uint lod = 0);  
void write(imageblock_slice slice, ushort 
coord, ushort lod = 0);  
void write(imageblock_slice slice, uint 
coord, uint lod = 0);  
void write(imageblock_slice slice, ushort 
coord, ushort lod = 0); 
For 1D texture array:  
void write(imageblock_slice slice, uint 
coord, uint array, uint lod = 0);  
void write(imageblock_slice slice, ushort 
coord, ushort array, ushort lod = 0);  
void write(imageblock_slice slice, uint 
coord, uint array, uint lod = 0);  
void write(imageblock_slice slice, ushort 
coord, ushort array, ushort lod = 0);  
For 2D texture:  
void write(imageblock_slice slice, uint2 
coord, uint lod = 0);  
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page  
 of  
154
174

---

void write(imageblock_slice slice, ushort2 
coord, ushort lod = 0);  
void write(imageblock_slice slice, uint2 
coord, uint lod = 0);  
void write(imageblock_slice slice, ushort2 
coord, ushort lod = 0); 
For 2D MSAA texture:  
void write(imageblock_slice slice, uint2 
coord, uint lod = 0);  
void write(imageblock_slice slice, ushort2 
coord, ushort lod = 0);  
void write(imageblock_slice slice, uint2 
coord, uint lod = 0);  
void write(imageblock_slice slice, ushort2 
coord, ushort lod = 0); 
For 2D texture array:  
void write(imageblock_slice slice, uint2 
coord, uint array, uint lod = 0);  
void write(imageblock_slice slice, ushort2 
coord, ushort array, ushort lod = 0);  
void write(imageblock_slice slice, uint2 
coord, uint array, uint lod = 0);  
void write(imageblock_slice slice, ushort2 
coord, ushort array, ushort lod = 0); 
For cube texture:  
void write(imageblock_slice slice, uint2 
coord, uint face, uint lod = 0);  
void write(imageblock_slice slice, ushort2 
coord, ushort face, ushort lod = 0);  
void write(imageblock_slice slice, uint2 
coord, uint face, uint lod = 0);  
void write(imageblock_slice slice, ushort2 
coord, ushort face, ushort lod = 0); 
For cube array texture:  
void write(imageblock_slice slice, uint2 
coord, uint face, uint array, uint lod = 0);  
void write(imageblock_slice slice, ushort2 
coord, ushort face, ushort array, ushort lod = 0);  
void write(imageblock_slice slice, uint2 
coord, uint face, uint array, uint lod = 0);  
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page  
 of  
155
174

---

void write(imageblock_slice slice, ushort2 
coord, ushort face, ushort array, ushort lod = 0); 
For 3D texture:  
void write(imageblock_slice slice, uint3 
coord, uint lod = 0);  
void write(imageblock_slice slice, ushort3 
coord, ushort lod = 0);  
void write(imageblock_slice slice, uint3 
coord, uint lod = 0);  
void write(imageblock_slice slice, ushort3 
coord, ushort lod = 0); 
Example  
struct Foo { 
half4 a; 
int b;  
float c; 
};  
kernel void 
my_kernel(texture2d src [[ texture(0) ]],  
     texture2d dst [[ texture(1) ]], 
imageblock img_blk, 
ushort2 lid [[ thread_position_in_threadgroup ]], 
ushort2 gid [[ thread_position_in_grid ]])  
{ 
// read 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 have finished writing to the 
imageblock 
 
//  in this case each thread writes to its location in the imageblock 
so a barrier is not necessary 
 
2017-9-12   |  Copyright © 2017 Apple Inc. All Rights Reserved.  
Page  
 of  
156
174

---