constant地址空间用于从设备内存池分配存储的缓存对象，是只读的。constant地址空间的指针或引用可以做函数的参数，向声明为常量的变量赋值会产生编译错误，声明常量但是没有赋予初始值也会产生编译错误。 在shader中，函数之外的变量（相当于全局变量），其地址空间必须是constant。

device地址空间用于从设备内存池分配出来的缓存对象，可读也可写。一个缓存对象可以被声明成一个标量、向量或是用户自定义结构体的指针或是引用。缓存对象使用的内存实际大小，应该在CPU侧调用时就确定。 纹理对象总是在device地址空间分配内存，所以纹理类型可以省略修饰符。

threadgroup地址空间用于通用计算函数变量的内存分配，变量被一个线程组的所有的线程共享，threadgroup地址空间分配的变量不能用于图形绘制函数。

thread地址空间用于每个线程内部的内存分配，被thread修饰的变量在其他线程无法访问，在图形绘制或是通用计算函数内声明的变量是thread地址空间分配。 如下一段代码，包括device、threadgroup、thread的使用：

typedef struct
{
    half3 kRec709Luma; // position的修饰符表示这个是顶点
    
} TransParam;

kernel void
sobelKernel(texture2d<half, access::read>  sourceTexture  [[texture(LYFragmentTextureIndexTextureSource)]],
                texture2d<half, access::write> destTexture [[texture(LYFragmentTextureIndexTextureDest)]],
                uint2                          grid         [[thread_position_in_grid]],
            device TransParam *param [[buffer(0)]], // param.kRec709Luma = half3(0.2126, 0.7152, 0.0722); // 把rgba转成亮度值
            threadgroup float3 *localBuffer [[threadgroup(0)]]) // threadgroup地址空间，这里并没有使用到；
{
    // 边界保护
    if(grid.x <= destTexture.get_width() && grid.y <= destTexture.get_height())
    {
        thread half4 color  = sourceTexture.read(grid); // 初始颜色
        thread half gray   = dot(color.rgb, half3(param->kRec709Luma)); // 转换成亮度
        destTexture.write(half4(gray, gray, gray, 1.0), grid); // 写回对应纹理
    }
}

数据结构

Metal中常用的数据结构有向量、矩阵、原子数据类型、缓存、纹理、采样器、数组、用户自定义结构体。

half 是16bit是浮点数 0.5h float 是32bit的浮点数 0.5f size_t 是64bit的无符号整数 通常用于sizeof的返回值 ptrdiff_t 是64bit的有符号整数 通常用于指针的差值 half2、half3、half4、float2、float3、float4等，是向量类型，表达方式为基础类型+向量维数。矩阵类似half4x4、half3x3、float4x4、float3x3。 double、long、long long不支持。

对于向量的访问，比如说vec=float4(1.0f, 1.0f, 1.0f, 1.0f)，其访问方式可以是vec[0]、vec[1]，也可以是vec.x、vec.y，也可以是vec.r、vec.g。（.xyzw和.rgba，前者对应三维坐标，后者对应RGB颜色空间） 只取部分、乱序取均可，比如说我们常用到的color=texture.bgra。

数据对齐 char3、uchar3的size是4Bytes，而不是3Bytes； 类似的，int是4Bytes，但int3是16而不是12Bytes； 矩阵是由一组向量构成，按照向量的维度对齐；float3x3由3个float3向量构成，那么每个float3的size是16Bytes； 隐式类型转换（Implicit Type Conversions） 向量到向量或是标量的隐式转换会导致编译错误，比如int4 i; float4 f = i; // compile error，无法将一个4维的整形向量转换为4维的浮点向量。 标量到向量的隐式转换，是标量被赋值给向量的每一个分量。 float4 f = 2.0f; // f = (2.0f, 2.0f, 2.0f, 2.0f) 标量到矩阵、向量到矩阵的隐式转换，矩阵到矩阵和向量及标量的隐式转换会导致编译错误。

纹理数据结构不支持指针和引用，纹理数据结构包括精度和access描述符，access修饰符描述纹理如何被访问，有三种描述符：sample、read、write，如下：

kernel void
sobelKernel(texture2d<half, access::read>  sourceTexture  [[texture(LYFragmentTextureIndexTextureSource)]],
                texture2d<half, access::write> destTexture [[texture(LYFragmentTextureIndexTextureDest)]],
                uint2                          grid         [[thread_position_in_grid]])

Sampler是采样器，决定如何对一个纹理进行采样操作。寻址模式，过滤模式，归一化坐标，比较函数。 在Metal程序里初始化的采样器必须使用constexpr修饰符声明。 采样器指针和引用是不支持的，将会导致编译错误。

    constexpr sampler textureSampler (mag_filter::linear,
                                      min_filter::linear); // sampler是采样器

运算符