CompilerMSL map many GLSL functions to MSL functions.

Add bool members is_read and is_written to SPIRType::Image.
Output correct texture read/write access by marking whether textures
are read from and written to by the shader.
Override bitcast_glsl_op() to use Metal as_type<type> functions.
Add implementations of SPIR-V functions inverse(), degrees() & radians().
Map inverseSqrt() to rsqrt().
Map roundEven() to rint().
GLSL functions imageSize() and textureSize() map to equivalent
expression using MSL get_width() & get_height() functions.
Map several SPIR-V integer bitfield functions to MSL equivalents.
Map SPIR-V atomic functions to MSL equivalents.
Map texture packing and unpacking functions to MSL equivalents.
Refactor existing, and add new, image query functions.
Reorganize header lines into includes and pragmas.
Simplify type_to_glsl() logic.
Add MSL test case vert/functions.vert for added function implementations.
Add MSL test case comp/atomic.comp for added function implementations.
test_shaders.py use macOS compilation for MSL shader compilation validations.

reference/shaders-msl/comp/atomic.comp

@@ -0,0 +1,36 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    uint u32;
+    int i32;
+};
+
+kernel void main0(device SSBO& ssbo [[buffer(0)]])
+{
+    uint _16 = atomic_fetch_add_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
+    uint _18 = atomic_fetch_or_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
+    uint _20 = atomic_fetch_xor_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
+    uint _22 = atomic_fetch_and_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
+    uint _24 = atomic_fetch_min_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
+    uint _26 = atomic_fetch_max_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
+    uint _28 = atomic_exchange_explicit((volatile device atomic_uint*)&(ssbo.u32), 1u, memory_order_relaxed);
+    uint _30 = 10u;
+    uint _32 = atomic_compare_exchange_weak_explicit((volatile device atomic_uint*)&(ssbo.u32), &(_30), 2u, memory_order_relaxed, memory_order_relaxed);
+    int _36 = atomic_fetch_add_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
+    int _38 = atomic_fetch_or_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
+    int _40 = atomic_fetch_xor_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
+    int _42 = atomic_fetch_and_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
+    int _44 = atomic_fetch_min_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
+    int _46 = atomic_fetch_max_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
+    int _48 = atomic_exchange_explicit((volatile device atomic_int*)&(ssbo.i32), 1, memory_order_relaxed);
+    int _50 = 10;
+    int _52 = atomic_compare_exchange_weak_explicit((volatile device atomic_int*)&(ssbo.i32), &(_50), 2, memory_order_relaxed, memory_order_relaxed);
+}
+

reference/shaders-msl/comp/global-invocation-id-writable-ssbo-in-function.comp

@@ -11,7 +11,7 @@ struct myBlock
     float b[1];
 };
 
-// Support GLSL mod(), which is slightly different than Metal fmod()
+// Implementation of the GLSL mod() function, which is slightly different than Metal fmod()
 template<typename Tx, typename Ty>
 Tx mod(Tx x, Ty y)
 {

reference/shaders-msl/comp/global-invocation-id.comp

@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -9,7 +11,7 @@ struct myBlock
     float b[1];
 };
 
-// Support GLSL mod(), which is slightly different than Metal fmod()
+// Implementation of the GLSL mod() function, which is slightly different than Metal fmod()
 template<typename Tx, typename Ty>
 Tx mod(Tx x, Ty y)
 {

reference/shaders-msl/comp/local-invocation-id.comp

@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -9,7 +11,7 @@ struct myBlock
     float b[1];
 };
 
-// Support GLSL mod(), which is slightly different than Metal fmod()
+// Implementation of the GLSL mod() function, which is slightly different than Metal fmod()
 template<typename Tx, typename Ty>
 Tx mod(Tx x, Ty y)
 {

reference/shaders-msl/comp/local-invocation-index.comp

@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -9,7 +11,7 @@ struct myBlock
     float b[1];
 };
 
-// Support GLSL mod(), which is slightly different than Metal fmod()
+// Implementation of the GLSL mod() function, which is slightly different than Metal fmod()
 template<typename Tx, typename Ty>
 Tx mod(Tx x, Ty y)
 {

reference/shaders-msl/comp/writable-ssbo.comp

@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -9,7 +11,7 @@ struct myBlock
     float b;
 };
 
-// Support GLSL mod(), which is slightly different than Metal fmod()
+// Implementation of the GLSL mod() function, which is slightly different than Metal fmod()
 template<typename Tx, typename Ty>
 Tx mod(Tx x, Ty y)
 {

reference/shaders-msl/desktop-only/frag/image-ms.desktop.frag

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+fragment void main0(texture2d_ms<float> uImageMS [[texture(0)]], texture2d_array<float, access::read_write> uImageArray [[texture(1)]], texture2d<float, access::write> uImage [[texture(2)]])
+{
+    float4 a = uImageMS.read(uint2(int2(1, 2)), 2);
+    float4 b = uImageArray.read(uint2(int3(1, 2, 4).xy), uint(int3(1, 2, 4).z));
+    uImage.write(a, uint2(int2(2, 3)));
+    uImageArray.write(b, uint2(int3(2, 3, 7).xy), uint(int3(2, 3, 7).z));
+}
+

reference/shaders-msl/desktop-only/frag/query-levels.desktop.frag

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 FragColor [[color(0)]];
+};
+
+fragment main0_out main0(texture2d<float> uSampler [[texture(0)]], sampler uSamplerSmplr [[sampler(0)]])
+{
+    main0_out out = {};
+    out.FragColor = float4(float(int(uSampler.get_num_mip_levels())));
+    return out;
+}
+

reference/shaders-msl/desktop-only/frag/sampler-ms-query.desktop.frag

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 FragColor [[color(0)]];
+};
+
+fragment main0_out main0(texture2d_ms<float> uSampler [[texture(0)]], sampler uSamplerSmplr [[sampler(0)]], texture2d_ms<float> uSamplerArray [[texture(1)]], sampler uSamplerArraySmplr [[sampler(1)]], texture2d_ms<float> uImage [[texture(2)]], texture2d_ms<float> uImageArray [[texture(3)]])
+{
+    main0_out out = {};
+    out.FragColor = float4(float(((int(uSampler.get_num_samples()) + int(uSamplerArray.get_num_samples())) + int(uImage.get_num_samples())) + int(uImageArray.get_num_samples())));
+    return out;
+}
+

reference/shaders-msl/frag/bitcasting.frag

@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_in
+{
+    float4 VertGeom [[user(locn0)]];
+};
+
+struct main0_out
+{
+    float4 FragColor0 [[color(0)]];
+    float4 FragColor1 [[color(1)]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]], texture2d<float> TextureBase [[texture(0)]], sampler TextureBaseSmplr [[sampler(0)]], texture2d<float> TextureDetail [[texture(1)]], sampler TextureDetailSmplr [[sampler(1)]])
+{
+    main0_out out = {};
+    float4 texSample0 = TextureBase.sample(TextureBaseSmplr, in.VertGeom.xy);
+    float4 texSample1 = TextureDetail.sample(TextureDetailSmplr, in.VertGeom.xy, int2(3, 2));
+    int4 iResult0 = as_type<int4>(texSample0);
+    int4 iResult1 = as_type<int4>(texSample1);
+    out.FragColor0 = as_type<float4>(iResult0) * as_type<float4>(iResult1);
+    uint4 uResult0 = as_type<uint4>(texSample0);
+    uint4 uResult1 = as_type<uint4>(texSample1);
+    out.FragColor1 = as_type<float4>(uResult0) * as_type<float4>(uResult1);
+    return out;
+}
+

reference/shaders-msl/vert/functions.vert

@@ -0,0 +1,120 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct UBO
+{
+    float4x4 uMVP;
+    float3 rotDeg;
+    float3 rotRad;
+    int2 bits;
+};
+
+struct main0_in
+{
+    float3 aNormal [[attribute(1)]];
+    float4 aVertex [[attribute(0)]];
+};
+
+struct main0_out
+{
+    float3 vRotRad [[user(locn0)]];
+    float3 vRotDeg [[user(locn1)]];
+    float3 vNormal [[user(locn2)]];
+    int2 vMSB [[user(locn3)]];
+    int2 vLSB [[user(locn4)]];
+    float4 gl_Position [[position]];
+    float gl_PointSize [[point_size]];
+};
+
+// Implementation of the GLSL radians() function
+template<typename T>
+T radians(T d)
+{
+    return d * 0.01745329251;
+}
+
+// Implementation of the GLSL degrees() function
+template<typename T>
+T degrees(T r)
+{
+    return r * 57.2957795131;
+}
+
+// Implementation of the GLSL findLSB() function
+template<typename T>
+T findLSB(T x)
+{
+    return select(ctz(x), -1, x == 0);
+}
+
+// Implementation of the signed GLSL findMSB() function
+template<typename T>
+T findSMSB(T x)
+{
+    T v = select(x, -1 - x, x < 0);
+    return select(clz(0) - (clz(v) + 1), -1, v == 0);
+}
+
+// Returns the determinant of a 2x2 matrix.
+inline float spvDet2x2(float a1, float a2, float b1, float b2)
+{
+    return a1 * b2 - b1 * a2;
+}
+
+// Returns the determinant of a 3x3 matrix.
+inline float spvDet3x3(float a1, float a2, float a3, float b1, float b2, float b3, float c1, float c2, float c3)
+{
+    return a1 * spvDet2x2(b2, b3, c2, c3) - b1 * spvDet2x2(a2, a3, c2, c3) + c1 * spvDet2x2(a2, a3, b2, b3);
+}
+
+// Returns the inverse of a matrix, by using the algorithm of calculating the classical
+// adjoint and dividing by the determinant. The contents of the matrix are changed.
+float4x4 spvInverse4x4(float4x4 m)
+{
+    float4x4 adj;	// The adjoint matrix (inverse after dividing by determinant)
+    
+    // Create the transpose of the cofactors, as the classical adjoint of the matrix.
+    adj[0][0] =  spvDet3x3(m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]);
+    adj[0][1] = -spvDet3x3(m[0][1], m[0][2], m[0][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]);
+    adj[0][2] =  spvDet3x3(m[0][1], m[0][2], m[0][3], m[1][1], m[1][2], m[1][3], m[3][1], m[3][2], m[3][3]);
+    adj[0][3] = -spvDet3x3(m[0][1], m[0][2], m[0][3], m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3]);
+    
+    adj[1][0] = -spvDet3x3(m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]);
+    adj[1][1] =  spvDet3x3(m[0][0], m[0][2], m[0][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]);
+    adj[1][2] = -spvDet3x3(m[0][0], m[0][2], m[0][3], m[1][0], m[1][2], m[1][3], m[3][0], m[3][2], m[3][3]);
+    adj[1][3] =  spvDet3x3(m[0][0], m[0][2], m[0][3], m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3]);
+    
+    adj[2][0] =  spvDet3x3(m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]);
+    adj[2][1] = -spvDet3x3(m[0][0], m[0][1], m[0][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]);
+    adj[2][2] =  spvDet3x3(m[0][0], m[0][1], m[0][3], m[1][0], m[1][1], m[1][3], m[3][0], m[3][1], m[3][3]);
+    adj[2][3] = -spvDet3x3(m[0][0], m[0][1], m[0][3], m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3]);
+    
+    adj[3][0] = -spvDet3x3(m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]);
+    adj[3][1] =  spvDet3x3(m[0][0], m[0][1], m[0][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]);
+    adj[3][2] = -spvDet3x3(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], m[3][0], m[3][1], m[3][2]);
+    adj[3][3] =  spvDet3x3(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2]);
+    
+    // Calculate the determinant as a combination of the cofactors of the first row.
+    float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]) + (adj[0][2] * m[2][0]) + (adj[0][3] * m[3][0]);
+    
+    // Divide the classical adjoint matrix by the determinant.
+    // If determinant is zero, matrix is not invertable, so leave it unchanged.
+    return (det != 0.0f) ? (adj * (1.0f / det)) : m;
+}
+
+vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _18 [[buffer(0)]])
+{
+    main0_out out = {};
+    out.gl_Position = spvInverse4x4(_18.uMVP) * in.aVertex;
+    out.vNormal = in.aNormal;
+    out.vRotDeg = degrees(_18.rotRad);
+    out.vRotRad = radians(_18.rotDeg);
+    out.vLSB = findLSB(_18.bits);
+    out.vMSB = findSMSB(_18.bits);
+    return out;
+}
+

reference/shaders-msl/vert/texture_buffer.vert

@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 gl_Position [[position]];
+    float gl_PointSize [[point_size]];
+};
+
+vertex main0_out main0(texture2d<float> uSamp [[texture(0)]], texture2d<float> uSampo [[texture(1)]])
+{
+    main0_out out = {};
+    out.gl_Position = uSamp.read(uint2(10, 0)) + uSampo.read(uint2(100, 0));
+    return out;
+}
+

shaders-msl/comp/atomic.comp

@@ -0,0 +1,33 @@
+#version 310 es
+#extension GL_OES_shader_image_atomic : require
+layout(local_size_x = 1) in;
+
+layout(r32ui, binding = 0) uniform highp uimage2D uImage;
+layout(r32i, binding = 1) uniform highp iimage2D iImage;
+layout(binding = 2, std430) buffer SSBO
+{
+    uint u32;
+    int  i32;
+} ssbo;
+
+void main()
+{
+    atomicAdd(ssbo.u32, 1u);
+    atomicOr(ssbo.u32, 1u);
+    atomicXor(ssbo.u32, 1u);
+    atomicAnd(ssbo.u32, 1u);
+    atomicMin(ssbo.u32, 1u);
+    atomicMax(ssbo.u32, 1u);
+    atomicExchange(ssbo.u32, 1u);
+    atomicCompSwap(ssbo.u32, 10u, 2u);
+
+    atomicAdd(ssbo.i32, 1);
+    atomicOr(ssbo.i32, 1);
+    atomicXor(ssbo.i32, 1);
+    atomicAnd(ssbo.i32, 1);
+    atomicMin(ssbo.i32, 1);
+    atomicMax(ssbo.i32, 1);
+    atomicExchange(ssbo.i32, 1);
+    atomicCompSwap(ssbo.i32, 10, 2);
+}
+

shaders-msl/desktop-only/frag/image-ms.desktop.frag

@@ -0,0 +1,13 @@
+#version 450
+
+layout(rgba8, binding = 0) uniform image2D uImage;
+layout(rgba8, binding = 1) uniform image2DArray uImageArray;
+layout(rgba8, binding = 2) uniform image2DMS uImageMS;
+
+void main()
+{
+	vec4 a = imageLoad(uImageMS, ivec2(1, 2), 2);
+	vec4 b = imageLoad(uImageArray, ivec3(1, 2, 4));
+	imageStore(uImage, ivec2(2, 3), a);
+	imageStore(uImageArray, ivec3(2, 3, 7), b);
+}

shaders-msl/desktop-only/frag/query-levels.desktop.frag

@@ -0,0 +1,11 @@
+#version 450
+
+layout(binding = 0) uniform sampler2D uSampler;
+
+layout(location = 0) out vec4 FragColor;
+
+void main()
+{
+    FragColor = vec4(float(textureQueryLevels(uSampler)));
+}
+

shaders-msl/desktop-only/frag/sampler-ms-query.desktop.frag

@@ -0,0 +1,14 @@
+#version 450
+
+layout(binding = 0) uniform sampler2DMS uSampler;
+layout(binding = 1) uniform sampler2DMSArray uSamplerArray;
+layout(binding = 2, rgba8) uniform readonly writeonly image2DMS uImage;
+layout(binding = 3, rgba8) uniform readonly writeonly image2DMSArray uImageArray;
+
+layout(location = 0) out vec4 FragColor;
+
+void main()
+{
+    FragColor = vec4(float(((textureSamples(uSampler) + textureSamples(uSamplerArray)) + imageSamples(uImage)) + imageSamples(uImageArray)));
+}
+

shaders-msl/frag/bitcasting.frag

@@ -0,0 +1,24 @@
+#version 310 es
+precision mediump float;
+
+layout(binding = 0) uniform sampler2D TextureBase;
+layout(binding = 1) uniform sampler2D TextureDetail;
+
+layout(location = 0) in vec4 VertGeom;
+
+layout(location = 0) out vec4 FragColor0;
+layout(location = 1) out vec4 FragColor1;
+
+void main()
+{
+	vec4 texSample0 = texture(TextureBase, VertGeom.xy);
+	vec4 texSample1 = textureOffset(TextureDetail, VertGeom.xy, ivec2(3, 2));
+	
+    ivec4 iResult0 = floatBitsToInt(texSample0);
+    ivec4 iResult1 = floatBitsToInt(texSample1);
+    FragColor0 = (intBitsToFloat(iResult0) * intBitsToFloat(iResult1));    
+    
+    uvec4 uResult0 = floatBitsToUint(texSample0);
+    uvec4 uResult1 = floatBitsToUint(texSample1);
+    FragColor1 = (uintBitsToFloat(uResult0) * uintBitsToFloat(uResult1));    
+}

shaders-msl/vert/functions.vert

@@ -0,0 +1,28 @@
+#version 310 es
+
+layout(std140) uniform UBO
+{
+    uniform mat4 uMVP;
+    uniform vec3 rotDeg;
+    uniform vec3 rotRad;
+    uniform ivec2 bits;
+};
+
+layout(location = 0) in vec4 aVertex;
+layout(location = 1) in vec3 aNormal;
+
+out vec3 vNormal;
+out vec3 vRotDeg;
+out vec3 vRotRad;
+out ivec2 vLSB;
+out ivec2 vMSB;
+
+void main()
+{
+    gl_Position = inverse(uMVP) * aVertex;
+    vNormal = aNormal;
+    vRotDeg = degrees(rotRad);
+    vRotRad = radians(rotDeg);
+    vLSB = findLSB(bits);
+    vMSB = findMSB(bits);
+}

shaders-msl/vert/texture_buffer.vert

@@ -0,0 +1,10 @@
+#version 310 es
+#extension GL_OES_texture_buffer : require
+
+layout(binding = 4) uniform highp samplerBuffer uSamp;
+layout(rgba32f, binding = 5) uniform readonly highp imageBuffer uSampo;
+
+void main()
+{
+   gl_Position = texelFetch(uSamp, 10) + imageLoad(uSampo, 100);
+}

spirv_common.hpp

@@ -276,6 +276,8 @@ struct SPIRType : IVariant
 		bool depth;
 		bool arrayed;
 		bool ms;
+		bool is_read = false;
+		bool is_written = false;
 		uint32_t sampled;
 		spv::ImageFormat format;
 	} image;

spirv_cross.cpp

@@ -127,6 +127,7 @@ bool Compiler::block_is_pure(const SPIRBlock &block)
 		case OpAtomicStore:
 		case OpAtomicExchange:
 		case OpAtomicCompareExchange:
+		case OpAtomicCompareExchangeWeak:
 		case OpAtomicIIncrement:
 		case OpAtomicIDecrement:
 		case OpAtomicIAdd:
@@ -538,6 +539,7 @@ bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t
 	case OpAtomicLoad:
 	case OpAtomicExchange:
 	case OpAtomicCompareExchange:
+	case OpAtomicCompareExchangeWeak:
 	case OpAtomicIIncrement:
 	case OpAtomicIDecrement:
 	case OpAtomicIAdd:

spirv_glsl.cpp

@@ -2769,6 +2769,7 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
 		break;
 
 	case OpImageFetch:
+	case OpImageRead: // Reads == fetches in Metal (other langs will not get here)
 		opt = &ops[4];
 		length -= 4;
 		fetch = true;
@@ -2787,7 +2788,13 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
 		break;
 	}
 
-	auto &imgtype = expression_type(img);
+	// Bypass pointers because we need the real image struct
+	auto &type = expression_type(img);
+	auto &imgtype = get<SPIRType>(type.self);
+
+	// Mark that this shader reads from this image
+	((SPIRType &)imgtype).image.is_read = true;
+
 	uint32_t coord_components = 0;
 	switch (imgtype.image.dim)
 	{

spirv_msl.hpp

@@ -92,6 +92,22 @@ public:
 		options = opts;
 	}
 
+	// An enum of SPIR-V functions that are implemented in additional
+	// source code that is added to the shader if necessary.
+	enum SPVFuncImpl
+	{
+		SPVFuncImplNone,
+		SPVFuncImplMod,
+		SPVFuncImplRadians,
+		SPVFuncImplDegrees,
+		SPVFuncImplFindILsb,
+		SPVFuncImplFindSMsb,
+		SPVFuncImplFindUMsb,
+		SPVFuncImplInverse2x2,
+		SPVFuncImplInverse3x3,
+		SPVFuncImplInverse4x4,
+	};
+
 	// Constructs an instance to compile the SPIR-V code into Metal Shading Language,
 	// using the configuration parameters, if provided:
 	//  - p_vtx_attrs is an optional list of vertex attribute bindings used to match
@@ -149,8 +165,7 @@ protected:
 	                             uint32_t grad_y, uint32_t lod, uint32_t coffset, uint32_t offset, uint32_t bias,
 	                             uint32_t comp, uint32_t sample, bool *p_forward) override;
 	std::string unpack_expression_type(std::string expr_str, const SPIRType &type) override;
-
-	std::string get_argument_address_space(const SPIRVariable &argument);
+	std::string bitcast_glsl_op(const SPIRType &result_type, const SPIRType &argument_type) override;
 
 	void preprocess_op_codes();
 	void emit_custom_functions();
@@ -196,14 +211,22 @@ protected:
 	void align_struct(SPIRType &ib_type);
 	bool is_member_packable(SPIRType &ib_type, uint32_t index);
 	MSLStructMemberKey get_struct_member_key(uint32_t type_id, uint32_t index);
+	SPVFuncImpl get_spv_func_impl(spv::Op opcode, const uint32_t *args);
+	std::string get_argument_address_space(const SPIRVariable &argument);
+	void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, uint32_t mem_order_1,
+	                         uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0,
+	                         bool op1_is_pointer = false, uint32_t op2 = 0);
+	const char *get_memory_order(uint32_t spv_mem_sem);
+	void add_pragma_line(const std::string &line);
 
 	Options options;
 	std::unordered_map<std::string, std::string> func_name_overrides;
 	std::unordered_map<std::string, std::string> var_name_overrides;
-	std::set<uint32_t> custom_function_ops;
+	std::set<SPVFuncImpl> spv_function_implementations;
 	std::unordered_map<uint32_t, MSLVertexAttr *> vtx_attrs_by_location;
 	std::map<uint32_t, uint32_t> non_stage_in_input_var_ids;
 	std::unordered_map<MSLStructMemberKey, uint32_t> struct_member_padding;
+	std::vector<std::string> pragma_lines;
 	std::vector<MSLResourceBinding *> resource_bindings;
 	MSLResourceBinding next_metal_resource_index;
 	uint32_t stage_in_var_id = 0;
@@ -229,6 +252,7 @@ protected:
 
 		CompilerMSL &compiler;
 		bool suppress_missing_prototypes = false;
+		bool uses_atomics = false;
 	};
 
 	// Sorts the members of a SPIRType and associated Meta info based on a settable sorting

test_shaders.py

@@ -75,7 +75,8 @@ def print_msl_compiler_version():
 def validate_shader_msl(shader):
     msl_path = reference_path(shader[0], shader[1])
     try:
-        subprocess.check_call(['xcrun', '--sdk', 'iphoneos', 'metal', '-x', 'metal', '-std=ios-metal1.2', '-Werror', msl_path])
+        subprocess.check_call(['xcrun', '--sdk', 'macosx', 'metal', '-x', 'metal', '-std=osx-metal1.2', '-Werror', msl_path])
+#        subprocess.check_call(['xcrun', '--sdk', 'iphoneos', 'metal', '-x', 'metal', '-std=ios-metal1.2', '-Werror', msl_path])
         print('Compiled Metal shader: ' + msl_path)   # display after so xcrun FNF is silent
     except OSError as oe:
         if (oe.errno != os.errno.ENOENT):   # Ignore xcrun not found error