前言
Metal着色语言是用于编写图形和计算内核函数的统一编程语言,是配合Metal框架编写的应用程序使用的。
Metal框架管理着色语言的运行,以及编译;Metal着色语言使用clang和LLVM,该编译器为在GPU上执行的代码提供了良好的性能。
Metal基于C++设计
Metal像素坐标系统
在metal中,纹理或是帧缓存attachment中的像素,坐标原点为左上
Metal数据类型
标量数据类型
metal不支持这些数据类型:double, long, unsigned long, long long,unsigned long long, long double
bool:值true可以扩展到整数常数1,值false可以扩展到整数常数0。
char:有符号,8bit整数
unsigned char,uchar:无符号,8bit整数
short:有符号,16bit整数
unsigned short,ushort:无符号,16bit整数
int:有符号32bit整数
unsigned int,uint:无符号32bit整数
half:16bit浮点数
float:32bit浮点数
size_t:64位无符号整数,用来配合sizeof
ptrdiff_t:64位有符号整数,表示两个指针的差
void:空
其余的:
.f,.F:表示单精度浮点型
.h,.H:表示半单精度浮点型
u,U:表示无符号整型字面量
向量和矩阵
Metal支持的向量类型名字如下:
booln, charn, shortn, intn, ucharn, ushortn, uintn, halfn, floatn
n取值2,3,4表示2维,3维,4维向量
Metal支持的矩阵类型名字如下:
halfnxm, floatnxm
n和m表示矩阵的列和行
访问向量的分量:
1.pos = float4(1.0f, 2.0f, 3.0f, 4.0f); float x = pos[0]; // x = 1.0
float z = pos[2]; // z = 3.0
2.int4 test = int4(0, 1, 2, 3);
int a = test.x; // a=0
int b = test.y; // b=1
int c = test.z; // c=2
int d = test.w; // d=3
int e = test.r; // e=0
int f = test.g; // f=1
int g = test.b; // g=2
int h = test.a; // h=3
3.float4 c;
c.xyzw = float4(1.0f, 2.0f, 3.0f, 4.0f);
c.z = 1.0f;
c.xy = float2(3.0f, 4.0f);
c.xyz = float3(3.0f, 4.0f, 5.0f);
4.float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
float4 swiz = pos.wzyx; // swiz = (4.0f, 3.0f, 2.0f, 1.0f)float4
dup = pos.xxyy; // dup = (1.0f, 1.0f, 2.0f, 2.0f)
4.float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
pos.xw = float2(5.0f, 6.0f);// pos = (5.0, 2.0, 3.0, 6.0)
pos.wx = float2(7.0f, 8.0f);// pos = (8.0, 2.0, 3.0, 7.0)
pos.xyz = float3(3.0f, 5.0f, 9.0f);// pos = (3.0, 5.0, 9.0, 7.0)
访问矩阵的分量
float4x4 m;
// sets the 2nd column to all 2.0
m[1] = float4(2.0f);
// sets the 1st element of the 1st column to 1.0
m[0][0] = 1.0f;
// sets the 4th element of the 3rd column to 3.0
m[2][3] = 3.0f;
原子数据类型
atomic_int , atomic_uint;用于多线程
Buffers类型
Metal中实现的缓存是一个指针,它指向一个由device或者constant修饰的数据块
device float4 *device_buffer;
struct my_user_data {
float4 a; float b; int2 c;
};
constant my_user_data *user_data;
Textures类型
纹理类型是一个句柄,它指向一个1维,或者二维三维的纹理数据
access结构:enum class access { sample, read, write };
texture1d<T, access a = access::sample>
texture1d_array<T, access a = access::sample>
texture2d<T, access a = access::sample>
texture2d_array<T, access a = access::sample>
texture3d<T, access a = access::sample>
texturecube<T, access a = access::sample>
texture2d_ms<T, access a = access::read>
access取值:
sample:纹理对象可以被采样,采样器可以从纹理中读取数据
read:不使用采样器,可以读取纹理对象
write:可以向纹理对象中写入数据
T是纹理数据类型
纹理不支持指针和引用
void foo (texture2d<float> imgA [[ texture(0) ]],
texture2d<float, access::read> imgB [[ texture(1) ]],
texture2d<float, access::write> imgC [[ texture(2) ]]){
...
}
Samplers类型
在着色语言中,采样器类型决定了如何对一个纹理进行采样,Metal框架中可以创意一个对应着色语言的采样器对象MTLSamplerState,当然采样器也可以着色语言中定义描述,采样器包括这些内容:
coord:纹理坐标是否归一化
address:所有纹理坐标的寻址模式
s_address,t_address,r_address:某个纹理坐标的寻址模式
filter:设置纹理的放大和缩小模式
constexpr sampler s(coord::pixel,
address::clamp_to_zero,
filter::linear);
constexpr sampler a(coord::normalized);
constexpr sampler b(address::repeat);
constexpr sampler s(address::clamp_to_zero,
filter::linear,
compare_func::less);
kernel void my_kernel(device float4 *p [[ buffer(0) ]],texture2d<float4> img [[ texture(0) ]],
sampler smp [[ sampler(3) ]]) {
...
}
采样器不支持指针和引用
Arrays and Structs
纹理和采样器类型的数组不支持
纹理和采样器类型不能在一个结构体中声明
结构体中的成员必须属于同一个地址空间