Updated spirv-cross.

3rdparty/spirv-cross/CMakeLists.txt

@@ -267,7 +267,7 @@ endif()
 
 if (SPIRV_CROSS_SHARED)
 	set(spirv-cross-abi-major 0)
-	set(spirv-cross-abi-minor 5)
+	set(spirv-cross-abi-minor 6)
 	set(spirv-cross-abi-patch 0)
 	set(SPIRV_CROSS_VERSION ${spirv-cross-abi-major}.${spirv-cross-abi-minor}.${spirv-cross-abi-patch})
 	set(SPIRV_CROSS_INSTALL_LIB_DIR ${CMAKE_INSTALL_PREFIX}/lib)

3rdparty/spirv-cross/main.cpp

@@ -509,6 +509,7 @@ struct CLIArguments
 	bool msl_pad_fragment_output = false;
 	bool msl_domain_lower_left = false;
 	bool msl_argument_buffers = false;
+	bool msl_texture_buffer_native = false;
 	bool glsl_emit_push_constant_as_ubo = false;
 	SmallVector<uint32_t> msl_discrete_descriptor_sets;
 	SmallVector<PLSArg> pls_in;
@@ -570,6 +571,7 @@ static void print_help()
 	                "\t[--msl-pad-fragment-output]\n"
 	                "\t[--msl-domain-lower-left]\n"
 	                "\t[--msl-argument-buffers]\n"
+	                "\t[--msl-texture-buffer-native]\n"
 	                "\t[--msl-discrete-descriptor-set <index>]\n"
 	                "\t[--hlsl]\n"
 	                "\t[--reflect]\n"
@@ -727,6 +729,7 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 		msl_opts.pad_fragment_output_components = args.msl_pad_fragment_output;
 		msl_opts.tess_domain_origin_lower_left = args.msl_domain_lower_left;
 		msl_opts.argument_buffers = args.msl_argument_buffers;
+		msl_opts.texture_buffer_native = args.msl_texture_buffer_native;
 		msl_comp->set_msl_options(msl_opts);
 		for (auto &v : args.msl_discrete_descriptor_sets)
 			msl_comp->add_discrete_descriptor_set(v);
@@ -1038,6 +1041,7 @@ static int main_inner(int argc, char *argv[])
 	cbs.add("--msl-argument-buffers", [&args](CLIParser &) { args.msl_argument_buffers = true; });
 	cbs.add("--msl-discrete-descriptor-set",
 	        [&args](CLIParser &parser) { args.msl_discrete_descriptor_sets.push_back(parser.next_uint()); });
+	cbs.add("--msl-texture-buffer-native", [&args](CLIParser &) { args.msl_texture_buffer_native = true; });
 	cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); });
 	cbs.add("--rename-entry-point", [&args](CLIParser &parser) {
 		auto old_name = parser.next_string();

3rdparty/spirv-cross/reference/opt/shaders-msl/asm/vert/copy-memory-interface.asm.vert

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 o1 [[user(locn1)]];
+    float4 gl_Position [[position]];
+};
+
+struct main0_in
+{
+    float4 v0 [[attribute(0)]];
+    float4 v1 [[attribute(1)]];
+};
+
+vertex main0_out main0(main0_in in [[stage_in]])
+{
+    main0_out out = {};
+    out.gl_Position = in.v0;
+    out.o1 = in.v1;
+    return out;
+}
+

3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read.frag → 3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read-write.frag

@@ -16,10 +16,11 @@ uint2 spvTexelBufferCoord(uint tc)
     return uint2(tc % 4096, tc / 4096);
 }
 
-fragment main0_out main0(texture2d<float> buf [[texture(0)]])
+fragment main0_out main0(texture2d<float> buf [[texture(0)]], texture2d<float, access::write> bufOut [[texture(1)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
     out.FragColor = buf.read(spvTexelBufferCoord(0));
+    bufOut.write(out.FragColor, spvTexelBufferCoord(int(gl_FragCoord.x)));
     return out;
 }
 

3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag

@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 FragColor [[color(0)]];
+};
+
+fragment main0_out main0(texture_buffer<float> buf [[texture(0)]], texture_buffer<float, access::write> bufOut [[texture(1)]], float4 gl_FragCoord [[position]])
+{
+    main0_out out = {};
+    out.FragColor = buf.read(uint(0));
+    bufOut.write(out.FragColor, uint(int(gl_FragCoord.x)));
+    return out;
+}
+

3rdparty/spirv-cross/reference/opt/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert

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

3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp.vk

@@ -0,0 +1,25 @@
+#version 450
+#extension GL_EXT_buffer_reference : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(buffer_reference) buffer Block;
+layout(buffer_reference, std430) buffer Block
+{
+    float v;
+};
+
+layout(set = 0, binding = 0, std140) uniform UBO
+{
+    Block blocks[4];
+} ubo;
+
+void main()
+{
+    Block blocks[4];
+    blocks[0] = ubo.blocks[0];
+    blocks[1] = ubo.blocks[1];
+    blocks[2] = ubo.blocks[2];
+    blocks[3] = ubo.blocks[3];
+    blocks[gl_WorkGroupID.x].v = 20.0;
+}
+

3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk

@@ -0,0 +1,45 @@
+#version 450
+#extension GL_ARB_gpu_shader_int64 : require
+#extension GL_EXT_buffer_reference : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(buffer_reference) buffer Node;
+layout(buffer_reference, std430) buffer Node
+{
+    layout(offset = 0) int value;
+    layout(offset = 16) Node next;
+    layout(offset = 32) Node prev;
+};
+
+layout(set = 0, binding = 0, std430) restrict buffer LinkedList
+{
+    Node head1;
+    Node head2;
+} _50;
+
+void main()
+{
+    Node _45;
+    if (gl_WorkGroupID.x < 4u)
+    {
+        _45 = _50.head1;
+    }
+    else
+    {
+        _45 = _50.head2;
+    }
+    restrict Node n = _45;
+    Node param = n.next;
+    Node param_1 = _50.head1;
+    Node param_2 = _50.head2;
+    param.value = param_1.value + param_2.value;
+    Node param_4 = _50.head1;
+    Node param_3 = param_4;
+    n = param_3;
+    int v = _50.head2.value;
+    n.value = 20;
+    n.value = v * 10;
+    uint64_t uptr = uint64_t(_50.head2.next);
+    Node unode = Node(uptr);
+}
+

3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp.vk

@@ -0,0 +1,147 @@
+#version 310 es
+#extension GL_EXT_scalar_block_layout : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+struct S0
+{
+    vec2 a[1];
+    float b;
+};
+
+struct S1
+{
+    vec3 a;
+    float b;
+};
+
+struct S2
+{
+    vec3 a[1];
+    float b;
+};
+
+struct S3
+{
+    vec2 a;
+    float b;
+};
+
+struct S4
+{
+    vec2 c;
+};
+
+struct Content
+{
+    S0 m0s[1];
+    S1 m1s[1];
+    S2 m2s[1];
+    S0 m0;
+    S1 m1;
+    S2 m2;
+    S3 m3;
+    float m4;
+    S4 m3s[8];
+};
+
+struct S0_1
+{
+    vec2 a[1];
+    float b;
+};
+
+struct S1_1
+{
+    vec3 a;
+    float b;
+};
+
+struct S2_1
+{
+    vec3 a[1];
+    float b;
+};
+
+struct S3_1
+{
+    vec2 a;
+    float b;
+};
+
+struct S4_1
+{
+    vec2 c;
+};
+
+struct Content_1
+{
+    S0_1 m0s[1];
+    S1_1 m1s[1];
+    S2_1 m2s[1];
+    S0_1 m0;
+    S1_1 m1;
+    S2_1 m2;
+    S3_1 m3;
+    float m4;
+    S4_1 m3s[8];
+};
+
+layout(set = 0, binding = 1, scalar) restrict buffer SSBO1
+{
+    Content content;
+    Content content1[2];
+    Content content2;
+    mat2 m0;
+    mat2 m1;
+    mat2x3 m2[4];
+    mat3x2 m3;
+    layout(row_major) mat2 m4;
+    layout(row_major) mat2 m5[9];
+    layout(row_major) mat2x3 m6[4][2];
+    layout(row_major) mat3x2 m7;
+    float array[];
+} ssbo_430;
+
+layout(set = 0, binding = 0, std140) restrict buffer SSBO0
+{
+    Content_1 content;
+    Content_1 content1[2];
+    Content_1 content2;
+    mat2 m0;
+    mat2 m1;
+    mat2x3 m2[4];
+    mat3x2 m3;
+    layout(row_major) mat2 m4;
+    layout(row_major) mat2 m5[9];
+    layout(row_major) mat2x3 m6[4][2];
+    layout(row_major) mat3x2 m7;
+    float array[];
+} ssbo_140;
+
+void main()
+{
+    ssbo_430.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
+    ssbo_430.content.m0s[0].b = ssbo_140.content.m0s[0].b;
+    ssbo_430.content.m1s[0].a = ssbo_140.content.m1s[0].a;
+    ssbo_430.content.m1s[0].b = ssbo_140.content.m1s[0].b;
+    ssbo_430.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
+    ssbo_430.content.m2s[0].b = ssbo_140.content.m2s[0].b;
+    ssbo_430.content.m0.a[0] = ssbo_140.content.m0.a[0];
+    ssbo_430.content.m0.b = ssbo_140.content.m0.b;
+    ssbo_430.content.m1.a = ssbo_140.content.m1.a;
+    ssbo_430.content.m1.b = ssbo_140.content.m1.b;
+    ssbo_430.content.m2.a[0] = ssbo_140.content.m2.a[0];
+    ssbo_430.content.m2.b = ssbo_140.content.m2.b;
+    ssbo_430.content.m3.a = ssbo_140.content.m3.a;
+    ssbo_430.content.m3.b = ssbo_140.content.m3.b;
+    ssbo_430.content.m4 = ssbo_140.content.m4;
+    ssbo_430.content.m3s[0].c = ssbo_140.content.m3s[0].c;
+    ssbo_430.content.m3s[1].c = ssbo_140.content.m3s[1].c;
+    ssbo_430.content.m3s[2].c = ssbo_140.content.m3s[2].c;
+    ssbo_430.content.m3s[3].c = ssbo_140.content.m3s[3].c;
+    ssbo_430.content.m3s[4].c = ssbo_140.content.m3s[4].c;
+    ssbo_430.content.m3s[5].c = ssbo_140.content.m3s[5].c;
+    ssbo_430.content.m3s[6].c = ssbo_140.content.m3s[6].c;
+    ssbo_430.content.m3s[7].c = ssbo_140.content.m3s[7].c;
+}
+

3rdparty/spirv-cross/reference/shaders-hlsl-no-opt/asm/comp/atomic-result-temporary.asm.comp

@@ -0,0 +1,24 @@
+RWByteAddressBuffer _5 : register(u0);
+
+static uint3 gl_GlobalInvocationID;
+struct SPIRV_Cross_Input
+{
+    uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
+};
+
+void comp_main()
+{
+    uint _24;
+    _5.InterlockedAdd(0, 1u, _24);
+    if (_24 < 1024u)
+    {
+        _5.Store(_24 * 4 + 4, gl_GlobalInvocationID.x);
+    }
+}
+
+[numthreads(1, 1, 1)]
+void main(SPIRV_Cross_Input stage_input)
+{
+    gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
+    comp_main();
+}

3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/comp/atomic-result-temporary.asm.comp

@@ -0,0 +1,23 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    uint count;
+    uint data[1];
+};
+
+kernel void main0(device SSBO& _5 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    uint _24 = atomic_fetch_add_explicit((volatile device atomic_uint*)&_5.count, 1u, memory_order_relaxed);
+    if (_24 < 1024u)
+    {
+        _5.data[_24] = gl_GlobalInvocationID.x;
+    }
+}
+

3rdparty/spirv-cross/reference/shaders-msl/asm/vert/copy-memory-interface.asm.vert

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 o1 [[user(locn1)]];
+    float4 gl_Position [[position]];
+};
+
+struct main0_in
+{
+    float4 v0 [[attribute(0)]];
+    float4 v1 [[attribute(1)]];
+};
+
+vertex main0_out main0(main0_in in [[stage_in]])
+{
+    main0_out out = {};
+    out.gl_Position = in.v0;
+    out.o1 = in.v1;
+    return out;
+}
+

3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read.frag → 3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read-write.frag

@@ -16,10 +16,11 @@ uint2 spvTexelBufferCoord(uint tc)
     return uint2(tc % 4096, tc / 4096);
 }
 
-fragment main0_out main0(texture2d<float> buf [[texture(0)]])
+fragment main0_out main0(texture2d<float> buf [[texture(0)]], texture2d<float, access::write> bufOut [[texture(1)]], float4 gl_FragCoord [[position]])
 {
     main0_out out = {};
     out.FragColor = buf.read(spvTexelBufferCoord(0));
+    bufOut.write(out.FragColor, spvTexelBufferCoord(int(gl_FragCoord.x)));
     return out;
 }
 

3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag

@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct main0_out
+{
+    float4 FragColor [[color(0)]];
+};
+
+fragment main0_out main0(texture_buffer<float> buf [[texture(0)]], texture_buffer<float, access::write> bufOut [[texture(1)]], float4 gl_FragCoord [[position]])
+{
+    main0_out out = {};
+    out.FragColor = buf.read(uint(0));
+    bufOut.write(out.FragColor, uint(int(gl_FragCoord.x)));
+    return out;
+}
+

3rdparty/spirv-cross/reference/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert

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

3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/atomic-result-temporary.asm.comp

@@ -0,0 +1,18 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0, std430) buffer SSBO
+{
+    uint count;
+    uint data[];
+} _5;
+
+void main()
+{
+    uint _24 = atomicAdd(_5.count, 1u);
+    if (_24 < 1024u)
+    {
+        _5.data[_24] = gl_GlobalInvocationID.x;
+    }
+}
+

3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer-2.asm.nocompat.vk.comp.vk

@@ -0,0 +1,21 @@
+#version 450
+#extension GL_ARB_gpu_shader_int64 : require
+#extension GL_EXT_buffer_reference : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(buffer_reference) buffer uintPointer
+{
+    uint value;
+};
+
+layout(push_constant, std430) uniform _4_12
+{
+    uint64_t _m0;
+} _12;
+
+void main()
+{
+    uintPointer _3 = uintPointer(_12._m0);
+    _3.value = 20u;
+}
+

3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer.asm.nocompat.vk.comp.vk

@@ -0,0 +1,21 @@
+#version 450
+#extension GL_ARB_gpu_shader_int64 : require
+#extension GL_EXT_buffer_reference : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(buffer_reference) buffer uint0_Pointer
+{
+    uint value[];
+};
+
+layout(push_constant, std430) uniform _6_14
+{
+    uint64_t _m0;
+} _14;
+
+void main()
+{
+    uint0_Pointer _5 = uint0_Pointer(_14._m0);
+    _5.value[10] = 20u;
+}
+

3rdparty/spirv-cross/reference/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp.vk

@@ -0,0 +1,25 @@
+#version 450
+#extension GL_EXT_buffer_reference : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(buffer_reference) buffer Block;
+layout(buffer_reference, std430) buffer Block
+{
+    float v;
+};
+
+layout(set = 0, binding = 0, std140) uniform UBO
+{
+    Block blocks[4];
+} ubo;
+
+void main()
+{
+    Block blocks[4];
+    blocks[0] = ubo.blocks[0];
+    blocks[1] = ubo.blocks[1];
+    blocks[2] = ubo.blocks[2];
+    blocks[3] = ubo.blocks[3];
+    blocks[gl_WorkGroupID.x].v = 20.0;
+}
+

3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk

@@ -0,0 +1,56 @@
+#version 450
+#extension GL_ARB_gpu_shader_int64 : require
+#extension GL_EXT_buffer_reference : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(buffer_reference) buffer Node;
+layout(buffer_reference, std430) buffer Node
+{
+    layout(offset = 0) int value;
+    layout(offset = 16) Node next;
+    layout(offset = 32) Node prev;
+};
+
+layout(set = 0, binding = 0, std430) restrict buffer LinkedList
+{
+    Node head1;
+    Node head2;
+} _50;
+
+void copy_node(restrict Node dst, restrict Node a, restrict Node b)
+{
+    dst.value = a.value + b.value;
+}
+
+void overwrite_node(out Node dst, Node src)
+{
+    dst = src;
+}
+
+void main()
+{
+    Node _45;
+    if (gl_WorkGroupID.x < 4u)
+    {
+        _45 = _50.head1;
+    }
+    else
+    {
+        _45 = _50.head2;
+    }
+    restrict Node n = _45;
+    Node param = n.next;
+    Node param_1 = _50.head1;
+    Node param_2 = _50.head2;
+    copy_node(param, param_1, param_2);
+    Node param_4 = _50.head1;
+    Node param_3;
+    overwrite_node(param_3, param_4);
+    n = param_3;
+    int v = _50.head2.value;
+    n.value = 20;
+    n.value = v * 10;
+    uint64_t uptr = uint64_t(_50.head2.next);
+    Node unode = Node(uptr);
+}
+

3rdparty/spirv-cross/reference/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp.vk

@@ -0,0 +1,147 @@
+#version 310 es
+#extension GL_EXT_scalar_block_layout : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+struct S0
+{
+    vec2 a[1];
+    float b;
+};
+
+struct S1
+{
+    vec3 a;
+    float b;
+};
+
+struct S2
+{
+    vec3 a[1];
+    float b;
+};
+
+struct S3
+{
+    vec2 a;
+    float b;
+};
+
+struct S4
+{
+    vec2 c;
+};
+
+struct Content
+{
+    S0 m0s[1];
+    S1 m1s[1];
+    S2 m2s[1];
+    S0 m0;
+    S1 m1;
+    S2 m2;
+    S3 m3;
+    float m4;
+    S4 m3s[8];
+};
+
+struct S0_1
+{
+    vec2 a[1];
+    float b;
+};
+
+struct S1_1
+{
+    vec3 a;
+    float b;
+};
+
+struct S2_1
+{
+    vec3 a[1];
+    float b;
+};
+
+struct S3_1
+{
+    vec2 a;
+    float b;
+};
+
+struct S4_1
+{
+    vec2 c;
+};
+
+struct Content_1
+{
+    S0_1 m0s[1];
+    S1_1 m1s[1];
+    S2_1 m2s[1];
+    S0_1 m0;
+    S1_1 m1;
+    S2_1 m2;
+    S3_1 m3;
+    float m4;
+    S4_1 m3s[8];
+};
+
+layout(set = 0, binding = 1, scalar) restrict buffer SSBO1
+{
+    Content content;
+    Content content1[2];
+    Content content2;
+    mat2 m0;
+    mat2 m1;
+    mat2x3 m2[4];
+    mat3x2 m3;
+    layout(row_major) mat2 m4;
+    layout(row_major) mat2 m5[9];
+    layout(row_major) mat2x3 m6[4][2];
+    layout(row_major) mat3x2 m7;
+    float array[];
+} ssbo_430;
+
+layout(set = 0, binding = 0, std140) restrict buffer SSBO0
+{
+    Content_1 content;
+    Content_1 content1[2];
+    Content_1 content2;
+    mat2 m0;
+    mat2 m1;
+    mat2x3 m2[4];
+    mat3x2 m3;
+    layout(row_major) mat2 m4;
+    layout(row_major) mat2 m5[9];
+    layout(row_major) mat2x3 m6[4][2];
+    layout(row_major) mat3x2 m7;
+    float array[];
+} ssbo_140;
+
+void main()
+{
+    ssbo_430.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
+    ssbo_430.content.m0s[0].b = ssbo_140.content.m0s[0].b;
+    ssbo_430.content.m1s[0].a = ssbo_140.content.m1s[0].a;
+    ssbo_430.content.m1s[0].b = ssbo_140.content.m1s[0].b;
+    ssbo_430.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
+    ssbo_430.content.m2s[0].b = ssbo_140.content.m2s[0].b;
+    ssbo_430.content.m0.a[0] = ssbo_140.content.m0.a[0];
+    ssbo_430.content.m0.b = ssbo_140.content.m0.b;
+    ssbo_430.content.m1.a = ssbo_140.content.m1.a;
+    ssbo_430.content.m1.b = ssbo_140.content.m1.b;
+    ssbo_430.content.m2.a[0] = ssbo_140.content.m2.a[0];
+    ssbo_430.content.m2.b = ssbo_140.content.m2.b;
+    ssbo_430.content.m3.a = ssbo_140.content.m3.a;
+    ssbo_430.content.m3.b = ssbo_140.content.m3.b;
+    ssbo_430.content.m4 = ssbo_140.content.m4;
+    ssbo_430.content.m3s[0].c = ssbo_140.content.m3s[0].c;
+    ssbo_430.content.m3s[1].c = ssbo_140.content.m3s[1].c;
+    ssbo_430.content.m3s[2].c = ssbo_140.content.m3s[2].c;
+    ssbo_430.content.m3s[3].c = ssbo_140.content.m3s[3].c;
+    ssbo_430.content.m3s[4].c = ssbo_140.content.m3s[4].c;
+    ssbo_430.content.m3s[5].c = ssbo_140.content.m3s[5].c;
+    ssbo_430.content.m3s[6].c = ssbo_140.content.m3s[6].c;
+    ssbo_430.content.m3s[7].c = ssbo_140.content.m3s[7].c;
+}
+

3rdparty/spirv-cross/shaders-hlsl-no-opt/asm/comp/atomic-result-temporary.asm.comp

@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 35
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "count"
+               OpMemberName %SSBO 1 "data"
+               OpName %_ ""
+               OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpMemberDecorate %SSBO 1 Offset 4
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+               OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+       %SSBO = OpTypeStruct %uint %_runtimearr_uint
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+  %uint_1024 = OpConstant %uint 1024
+       %bool = OpTypeBool
+      %int_1 = OpConstant %int 1
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Input_uint = OpTypePointer Input %uint
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %16 = OpAccessChain %_ptr_Uniform_uint %_ %int_0
+         %19 = OpAtomicIAdd %uint %16 %uint_1 %uint_0 %uint_1
+         %23 = OpULessThan %bool %19 %uint_1024
+               OpSelectionMerge %25 None
+               OpBranchConditional %23 %24 %25
+         %24 = OpLabel
+         %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+         %33 = OpLoad %uint %32
+         %34 = OpAccessChain %_ptr_Uniform_uint %_ %int_1 %19
+               OpStore %34 %33
+               OpBranch %25
+         %25 = OpLabel
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/shaders-msl-no-opt/asm/comp/atomic-result-temporary.asm.comp

@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 35
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "count"
+               OpMemberName %SSBO 1 "data"
+               OpName %_ ""
+               OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpMemberDecorate %SSBO 1 Offset 4
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+               OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+       %SSBO = OpTypeStruct %uint %_runtimearr_uint
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+  %uint_1024 = OpConstant %uint 1024
+       %bool = OpTypeBool
+      %int_1 = OpConstant %int 1
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Input_uint = OpTypePointer Input %uint
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %16 = OpAccessChain %_ptr_Uniform_uint %_ %int_0
+         %19 = OpAtomicIAdd %uint %16 %uint_1 %uint_0 %uint_1
+         %23 = OpULessThan %bool %19 %uint_1024
+               OpSelectionMerge %25 None
+               OpBranchConditional %23 %24 %25
+         %24 = OpLabel
+         %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+         %33 = OpLoad %uint %32
+         %34 = OpAccessChain %_ptr_Uniform_uint %_ %int_1 %19
+               OpStore %34 %33
+               OpBranch %25
+         %25 = OpLabel
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/shaders-msl/asm/vert/copy-memory-interface.asm.vert

@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Wine VKD3D Shader Compiler; 1
+; Bound: 13
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %1 "main" %8 %9 %11 %12
+               OpName %1 "main"
+               OpName %8 "v0"
+               OpName %9 "v1"
+               OpName %11 "o0"
+               OpName %12 "o1"
+               OpDecorate %8 Location 0
+               OpDecorate %9 Location 1
+               OpDecorate %11 BuiltIn Position
+               OpDecorate %12 Location 1
+          %2 = OpTypeVoid
+          %3 = OpTypeFunction %2
+          %5 = OpTypeFloat 32
+          %6 = OpTypeVector %5 4
+          %7 = OpTypePointer Input %6
+          %8 = OpVariable %7 Input
+          %9 = OpVariable %7 Input
+         %10 = OpTypePointer Output %6
+         %11 = OpVariable %10 Output
+         %12 = OpVariable %10 Output
+          %1 = OpFunction %2 None %3
+          %4 = OpLabel
+               OpCopyMemory %11 %8
+               OpCopyMemory %12 %9
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/shaders-msl/frag/buffer-read.frag → 3rdparty/spirv-cross/shaders-msl/frag/buffer-read-write.frag

@@ -1,10 +1,12 @@
 #version 450
 
 layout(rgba8, binding = 0) uniform readonly imageBuffer buf;
+layout(rgba8, binding = 1) uniform writeonly imageBuffer bufOut;
 
 layout(location = 0) out vec4 FragColor;
 
 void main()
 {
 	FragColor = imageLoad(buf, 0);
+	imageStore(bufOut, int(gl_FragCoord.x), FragColor);
 }

3rdparty/spirv-cross/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag

@@ -0,0 +1,12 @@
+#version 450
+
+layout(rgba8, binding = 0) uniform readonly imageBuffer buf;
+layout(rgba8, binding = 1) uniform writeonly imageBuffer bufOut;
+
+layout(location = 0) out vec4 FragColor;
+
+void main()
+{
+	FragColor = imageLoad(buf, 0);
+	imageStore(bufOut, int(gl_FragCoord.x), FragColor);
+}

3rdparty/spirv-cross/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.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);
+}

3rdparty/spirv-cross/shaders-no-opt/asm/comp/atomic-result-temporary.asm.comp

@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 35
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpName %main "main"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "count"
+               OpMemberName %SSBO 1 "data"
+               OpName %_ ""
+               OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpMemberDecorate %SSBO 1 Offset 4
+               OpDecorate %SSBO BufferBlock
+               OpDecorate %_ DescriptorSet 0
+               OpDecorate %_ Binding 0
+               OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+       %SSBO = OpTypeStruct %uint %_runtimearr_uint
+%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
+          %_ = OpVariable %_ptr_Uniform_SSBO Uniform
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+  %uint_1024 = OpConstant %uint 1024
+       %bool = OpTypeBool
+      %int_1 = OpConstant %int 1
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Input_uint = OpTypePointer Input %uint
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %16 = OpAccessChain %_ptr_Uniform_uint %_ %int_0
+         %19 = OpAtomicIAdd %uint %16 %uint_1 %uint_0 %uint_1
+         %23 = OpULessThan %bool %19 %uint_1024
+               OpSelectionMerge %25 None
+               OpBranchConditional %23 %24 %25
+         %24 = OpLabel
+         %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+         %33 = OpLoad %uint %32
+         %34 = OpAccessChain %_ptr_Uniform_uint %_ %int_1 %19
+               OpStore %34 %33
+               OpBranch %25
+         %25 = OpLabel
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer-2.asm.nocompat.vk.comp

@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 27
+; Schema: 0
+               OpCapability Shader
+               OpCapability Int64
+               OpCapability PhysicalStorageBufferAddressesEXT
+               OpExtension "SPV_EXT_physical_storage_buffer"
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel PhysicalStorageBuffer64EXT GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpSourceExtension "GL_ARB_gpu_shader_int64"
+               OpSourceExtension "GL_EXT_buffer_reference"
+               OpDecorate %ptr AliasedPointerEXT
+               OpMemberDecorate %Registers 0 Offset 0
+               OpDecorate %Registers Block
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+%_ptr_PhysicalStorageBufferEXT_uint = OpTypePointer PhysicalStorageBufferEXT %uint
+%_ptr_Function__ptr_PhysicalStorageBufferEXT_uint = OpTypePointer Function %_ptr_PhysicalStorageBufferEXT_uint
+      %ulong = OpTypeInt 64 0
+  %Registers = OpTypeStruct %ulong
+%_ptr_PushConstant_Registers = OpTypePointer PushConstant %Registers
+  %registers = OpVariable %_ptr_PushConstant_Registers PushConstant
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_PushConstant_ulong = OpTypePointer PushConstant %ulong
+     %int_10 = OpConstant %int 10
+    %uint_20 = OpConstant %uint 20
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+        %ptr = OpVariable %_ptr_Function__ptr_PhysicalStorageBufferEXT_uint Function
+         %19 = OpAccessChain %_ptr_PushConstant_ulong %registers %int_0
+         %20 = OpLoad %ulong %19
+         %21 = OpConvertUToPtr %_ptr_PhysicalStorageBufferEXT_uint %20
+               OpStore %ptr %21
+         %22 = OpLoad %_ptr_PhysicalStorageBufferEXT_uint %ptr
+               OpStore %22 %uint_20 Aligned 4
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer.asm.nocompat.vk.comp

@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 27
+; Schema: 0
+               OpCapability Shader
+               OpCapability Int64
+               OpCapability PhysicalStorageBufferAddressesEXT
+               OpExtension "SPV_EXT_physical_storage_buffer"
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel PhysicalStorageBuffer64EXT GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpSource GLSL 450
+               OpSourceExtension "GL_ARB_gpu_shader_int64"
+               OpSourceExtension "GL_EXT_buffer_reference"
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %uintPtr 0 Offset 0
+               OpDecorate %uintPtr Block
+               OpDecorate %ptr AliasedPointerEXT
+               OpMemberDecorate %Registers 0 Offset 0
+               OpDecorate %Registers Block
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+    %uintPtr = OpTypeStruct %_runtimearr_uint
+%_ptr_PhysicalStorageBufferEXT_uint_array = OpTypePointer PhysicalStorageBufferEXT %_runtimearr_uint
+%_ptr_Function__ptr_PhysicalStorageBufferEXT_uint_array = OpTypePointer Function %_ptr_PhysicalStorageBufferEXT_uint_array
+      %ulong = OpTypeInt 64 0
+  %Registers = OpTypeStruct %ulong
+%_ptr_PushConstant_Registers = OpTypePointer PushConstant %Registers
+  %registers = OpVariable %_ptr_PushConstant_Registers PushConstant
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_PushConstant_ulong = OpTypePointer PushConstant %ulong
+     %int_10 = OpConstant %int 10
+    %uint_20 = OpConstant %uint 20
+%_ptr_PhysicalStorageBufferEXT_uint = OpTypePointer PhysicalStorageBufferEXT %uint
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+        %ptr = OpVariable %_ptr_Function__ptr_PhysicalStorageBufferEXT_uint_array Function
+         %19 = OpAccessChain %_ptr_PushConstant_ulong %registers %int_0
+         %20 = OpLoad %ulong %19
+         %21 = OpConvertUToPtr %_ptr_PhysicalStorageBufferEXT_uint_array %20
+               OpStore %ptr %21
+         %22 = OpLoad %_ptr_PhysicalStorageBufferEXT_uint_array %ptr
+         %26 = OpAccessChain %_ptr_PhysicalStorageBufferEXT_uint %22 %int_10
+               OpStore %26 %uint_20 Aligned 4
+               OpReturn
+               OpFunctionEnd

3rdparty/spirv-cross/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp

@@ -0,0 +1,23 @@
+#version 450
+#extension GL_EXT_buffer_reference : require
+layout(local_size_x = 1) in;
+
+layout(buffer_reference) buffer Block
+{
+	float v;
+};
+
+layout(std140, set = 0, binding = 0) uniform UBO
+{
+	Block blocks[4];
+} ubo;
+
+void main()
+{
+	Block blocks[4];
+	blocks[0] = ubo.blocks[0];
+	blocks[1] = ubo.blocks[1];
+	blocks[2] = ubo.blocks[2];
+	blocks[3] = ubo.blocks[3];
+	blocks[gl_WorkGroupID.x].v = 20.0;
+}

3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp

@@ -0,0 +1,40 @@
+#version 450
+#extension GL_EXT_buffer_reference : require
+#extension GL_ARB_gpu_shader_int64 : require
+
+layout(buffer_reference) buffer Node;
+layout(buffer_reference) buffer Node
+{
+	int value;
+	layout(offset = 16) Node next;
+	layout(offset = 32) Node prev;
+};
+
+layout(std430, set = 0, binding = 0) buffer LinkedList
+{
+	restrict Node head1;
+	restrict Node head2;
+};
+
+void copy_node(restrict Node dst, restrict Node a, restrict Node b)
+{
+	dst.value = a.value + b.value;
+}
+
+void overwrite_node(out Node dst, Node src)
+{
+	dst = src;
+}
+
+void main()
+{
+	restrict Node n = gl_WorkGroupID.x < 4u ? head1 : head2;
+	copy_node(n.next, head1, head2);
+	overwrite_node(n, head1);
+	int v = head2.value;
+	n.value = 20;
+	n.value = v * 10;
+
+	uint64_t uptr = uint64_t(head2.next);
+	Node unode = Node(uptr);
+}

3rdparty/spirv-cross/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp

@@ -0,0 +1,88 @@
+#version 310 es
+#extension GL_EXT_scalar_block_layout : require
+
+layout(local_size_x = 1) in;
+
+struct S0
+{
+    vec2 a[1];
+    float b;
+};
+
+struct S1
+{
+    vec3 a;
+    float b;
+};
+
+struct S2
+{
+    vec3 a[1];
+    float b;
+};
+
+struct S3
+{
+    vec2 a;
+    float b;
+};
+
+struct S4
+{
+	vec2 c;
+};
+
+struct Content
+{
+    S0 m0s[1];
+    S1 m1s[1];
+    S2 m2s[1];
+    S0 m0;
+    S1 m1;
+    S2 m2;
+    S3 m3;
+    float m4;
+
+	S4 m3s[8];
+};
+
+layout(binding = 1, scalar) restrict buffer SSBO1
+{
+    Content content;
+    Content content1[2];
+    Content content2;
+
+    layout(column_major) mat2 m0;
+    layout(column_major) mat2 m1;
+    layout(column_major) mat2x3 m2[4];
+    layout(column_major) mat3x2 m3;
+    layout(row_major) mat2 m4;
+    layout(row_major) mat2 m5[9];
+    layout(row_major) mat2x3 m6[4][2];
+    layout(row_major) mat3x2 m7;
+    float array[];
+} ssbo_430;
+
+layout(binding = 0, std140) restrict buffer SSBO0
+{
+    Content content;
+    Content content1[2];
+    Content content2;
+
+    layout(column_major) mat2 m0;
+    layout(column_major) mat2 m1;
+    layout(column_major) mat2x3 m2[4];
+    layout(column_major) mat3x2 m3;
+    layout(row_major) mat2 m4;
+    layout(row_major) mat2 m5[9];
+    layout(row_major) mat2x3 m6[4][2];
+    layout(row_major) mat3x2 m7;
+
+    float array[];
+} ssbo_140;
+
+void main()
+{
+    ssbo_430.content = ssbo_140.content;
+}
+

3rdparty/spirv-cross/spirv_cross.cpp

@@ -74,6 +74,7 @@ bool Compiler::variable_storage_is_aliased(const SPIRVariable &v)
 	            ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock);
 	bool image = type.basetype == SPIRType::Image;
 	bool counter = type.basetype == SPIRType::AtomicCounter;
+	bool buffer_reference = type.storage == StorageClassPhysicalStorageBufferEXT;
 
 	bool is_restrict;
 	if (ssbo)
@@ -81,7 +82,7 @@ bool Compiler::variable_storage_is_aliased(const SPIRVariable &v)
 	else
 		is_restrict = has_decoration(v.self, DecorationRestrict);
 
-	return !is_restrict && (ssbo || image || counter);
+	return !is_restrict && (ssbo || image || counter || buffer_reference);
 }
 
 bool Compiler::block_is_pure(const SPIRBlock &block)
@@ -300,18 +301,41 @@ void Compiler::register_write(uint32_t chain)
 
 	if (var)
 	{
+		bool check_argument_storage_qualifier = true;
+		auto &type = expression_type(chain);
+
 		// If our variable is in a storage class which can alias with other buffers,
 		// invalidate all variables which depend on aliased variables. And if this is a
 		// variable pointer, then invalidate all variables regardless.
 		if (get_variable_data_type(*var).pointer)
+		{
 			flush_all_active_variables();
-		if (variable_storage_is_aliased(*var))
+
+			if (type.pointer_depth == 1)
+			{
+				// We have a backing variable which is a pointer-to-pointer type.
+				// We are storing some data through a pointer acquired through that variable,
+				// but we are not writing to the value of the variable itself,
+				// i.e., we are not modifying the pointer directly.
+				// If we are storing a non-pointer type (pointer_depth == 1),
+				// we know that we are storing some unrelated data.
+				// A case here would be
+				// void foo(Foo * const *arg) {
+				//   Foo *bar = *arg;
+				//   bar->unrelated = 42;
+				// }
+				// arg, the argument is constant.
+				check_argument_storage_qualifier = false;
+			}
+		}
+
+		if (type.storage == StorageClassPhysicalStorageBufferEXT || variable_storage_is_aliased(*var))
 			flush_all_aliased_variables();
 		else if (var)
 			flush_dependees(*var);
 
 		// We tried to write to a parameter which is not marked with out qualifier, force a recompile.
-		if (var->parameter && var->parameter->write_count == 0)
+		if (check_argument_storage_qualifier && var->parameter && var->parameter->write_count == 0)
 		{
 			var->parameter->write_count++;
 			force_recompile();
@@ -624,11 +648,11 @@ bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t
 
 		auto *var = compiler.maybe_get<SPIRVariable>(args[0]);
 		if (var && storage_class_is_interface(var->storage))
-			variables.insert(variable);
+			variables.insert(args[0]);
 
 		var = compiler.maybe_get<SPIRVariable>(args[1]);
 		if (var && storage_class_is_interface(var->storage))
-			variables.insert(variable);
+			variables.insert(args[1]);
 		break;
 	}
 
@@ -4114,8 +4138,13 @@ Bitset Compiler::combined_decoration_for_member(const SPIRType &type, uint32_t i
 
 		// If our type is a struct, traverse all the members as well recursively.
 		flags.merge_or(dec.decoration_flags);
+
 		for (uint32_t i = 0; i < type.member_types.size(); i++)
-			flags.merge_or(combined_decoration_for_member(get<SPIRType>(type.member_types[i]), i));
+		{
+			auto &memb_type = get<SPIRType>(type.member_types[i]);
+			if (!memb_type.pointer)
+				flags.merge_or(combined_decoration_for_member(memb_type, i));
+		}
 	}
 
 	return flags;
@@ -4180,3 +4209,44 @@ void Compiler::clear_force_recompile()
 {
 	is_force_recompile = false;
 }
+
+Compiler::PhysicalStorageBufferPointerHandler::PhysicalStorageBufferPointerHandler(Compiler &compiler_)
+    : compiler(compiler_)
+{
+}
+
+bool Compiler::PhysicalStorageBufferPointerHandler::handle(Op op, const uint32_t *args, uint32_t)
+{
+	if (op == OpConvertUToPtr)
+	{
+		auto &type = compiler.get<SPIRType>(args[0]);
+		if (type.storage == StorageClassPhysicalStorageBufferEXT && type.pointer && type.pointer_depth == 1)
+		{
+			// If we need to cast to a pointer type which is not a block, we might need to synthesize ourselves
+			// a block type which wraps this POD type.
+			if (type.basetype != SPIRType::Struct)
+				types.insert(args[0]);
+		}
+	}
+
+	return true;
+}
+
+void Compiler::analyze_non_block_pointer_types()
+{
+	PhysicalStorageBufferPointerHandler handler(*this);
+	traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
+	physical_storage_non_block_pointer_types.reserve(handler.types.size());
+	for (auto type : handler.types)
+		physical_storage_non_block_pointer_types.push_back(type);
+	sort(begin(physical_storage_non_block_pointer_types), end(physical_storage_non_block_pointer_types));
+}
+
+bool Compiler::type_is_array_of_pointers(const SPIRType &type) const
+{
+	if (!type.pointer)
+		return false;
+
+	// If parent type has same pointer depth, we must have an array of pointers.
+	return type.pointer_depth == get<SPIRType>(type.parent_type).pointer_depth;
+}

3rdparty/spirv-cross/spirv_cross.hpp

@@ -106,7 +106,9 @@ enum BufferPackingStandard
 	BufferPackingStd140EnhancedLayout,
 	BufferPackingStd430EnhancedLayout,
 	BufferPackingHLSLCbuffer,
-	BufferPackingHLSLCbufferPackOffset
+	BufferPackingHLSLCbufferPackOffset,
+	BufferPackingScalar,
+	BufferPackingScalarEnhancedLayout
 };
 
 struct EntryPoint
@@ -932,6 +934,16 @@ protected:
 		uint32_t write_count = 0;
 	};
 
+	struct PhysicalStorageBufferPointerHandler : OpcodeHandler
+	{
+		PhysicalStorageBufferPointerHandler(Compiler &compiler_);
+		bool handle(spv::Op op, const uint32_t *args, uint32_t length) override;
+		Compiler &compiler;
+		std::unordered_set<uint32_t> types;
+	};
+	void analyze_non_block_pointer_types();
+	SmallVector<uint32_t> physical_storage_non_block_pointer_types;
+
 	void analyze_variable_scope(SPIRFunction &function, AnalyzeVariableScopeAccessHandler &handler);
 	void find_function_local_luts(SPIRFunction &function, const AnalyzeVariableScopeAccessHandler &handler,
 	                              bool single_function);
@@ -959,6 +971,8 @@ protected:
 	bool has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const;
 	void unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration);
 
+	bool type_is_array_of_pointers(const SPIRType &type) const;
+
 private:
 	// Used only to implement the old deprecated get_entry_point() interface.
 	const SPIREntryPoint &get_first_entry_point(const std::string &name) const;

3rdparty/spirv-cross/spirv_cross_c.cpp

@@ -526,6 +526,10 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
 	case SPVC_COMPILER_OPTION_MSL_ARGUMENT_BUFFERS:
 		options->msl.argument_buffers = value != 0;
 		break;
+
+	case SPVC_COMPILER_OPTION_MSL_TEXTURE_BUFFER_NATIVE:
+		options->msl.texture_buffer_native = value != 0;
+		break;
 #endif
 
 	default:

3rdparty/spirv-cross/spirv_cross_c.h

@@ -33,7 +33,7 @@ extern "C" {
 /* Bumped if ABI or API breaks backwards compatibility. */
 #define SPVC_C_API_VERSION_MAJOR 0
 /* Bumped if APIs or enumerations are added in a backwards compatible way. */
-#define SPVC_C_API_VERSION_MINOR 5
+#define SPVC_C_API_VERSION_MINOR 6
 /* Bumped if internal implementation details change. */
 #define SPVC_C_API_VERSION_PATCH 0
 
@@ -424,6 +424,8 @@ typedef enum spvc_compiler_option
 
 	SPVC_COMPILER_OPTION_GLSL_EMIT_PUSH_CONSTANT_AS_UNIFORM_BUFFER = 33 | SPVC_COMPILER_OPTION_GLSL_BIT,
 
+	SPVC_COMPILER_OPTION_MSL_TEXTURE_BUFFER_NATIVE = 34 | SPVC_COMPILER_OPTION_MSL_BIT,
+
 	SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff
 } spvc_compiler_option;
 

3rdparty/spirv-cross/spirv_cross_containers.hpp

@@ -297,7 +297,10 @@ public:
 
 	void pop_back()
 	{
-		resize(this->buffer_size - 1);
+		// Work around false positive warning on GCC 8.3.
+		// Calling pop_back on empty vector is undefined.
+		if (!this->empty())
+			resize(this->buffer_size - 1);
 	}
 
 	template <typename... Ts>

3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp

@@ -66,6 +66,8 @@ ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
 		continue_block_to_loop_header = move(other.continue_block_to_loop_header);
 		entry_points = move(other.entry_points);
 		ids = move(other.ids);
+		addressing_model = other.addressing_model;
+		memory_model = other.memory_model;
 
 		default_entry_point = other.default_entry_point;
 		source = other.source;
@@ -98,6 +100,8 @@ ParsedIR &ParsedIR::operator=(const ParsedIR &other)
 		default_entry_point = other.default_entry_point;
 		source = other.source;
 		loop_iteration_depth = other.loop_iteration_depth;
+		addressing_model = other.addressing_model;
+		memory_model = other.memory_model;
 
 		// Very deliberate copying of IDs. There is no default copy constructor, nor a simple default constructor.
 		// Construct object first so we have the correct allocator set-up, then we can copy object into our new pool group.
@@ -692,24 +696,27 @@ void ParsedIR::add_typed_id(Types type, uint32_t id)
 	if (loop_iteration_depth)
 		SPIRV_CROSS_THROW("Cannot add typed ID while looping over it.");
 
-	switch (type)
+	if (ids[id].empty() || ids[id].get_type() != type)
 	{
-	case TypeConstant:
-		ids_for_constant_or_variable.push_back(id);
-		ids_for_constant_or_type.push_back(id);
-		break;
+		switch (type)
+		{
+		case TypeConstant:
+			ids_for_constant_or_variable.push_back(id);
+			ids_for_constant_or_type.push_back(id);
+			break;
 
-	case TypeVariable:
-		ids_for_constant_or_variable.push_back(id);
-		break;
+		case TypeVariable:
+			ids_for_constant_or_variable.push_back(id);
+			break;
 
-	case TypeType:
-	case TypeConstantOp:
-		ids_for_constant_or_type.push_back(id);
-		break;
+		case TypeType:
+		case TypeConstantOp:
+			ids_for_constant_or_type.push_back(id);
+			break;
 
-	default:
-		break;
+		default:
+			break;
+		}
 	}
 
 	if (ids[id].empty())

3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp

@@ -107,6 +107,9 @@ public:
 
 	Source source;
 
+	spv::AddressingModel addressing_model = spv::AddressingModelMax;
+	spv::MemoryModel memory_model = spv::MemoryModelMax;
+
 	// Decoration handling methods.
 	// Can be useful for simple "raw" reflection.
 	// However, most members are here because the Parser needs most of these,

3rdparty/spirv-cross/spirv_glsl.cpp

@@ -106,6 +106,7 @@ static bool packing_has_flexible_offset(BufferPackingStandard packing)
 	{
 	case BufferPackingStd140:
 	case BufferPackingStd430:
+	case BufferPackingScalar:
 	case BufferPackingHLSLCbuffer:
 		return false;
 
@@ -114,6 +115,19 @@ static bool packing_has_flexible_offset(BufferPackingStandard packing)
 	}
 }
 
+static bool packing_is_scalar(BufferPackingStandard packing)
+{
+	switch (packing)
+	{
+	case BufferPackingScalar:
+	case BufferPackingScalarEnhancedLayout:
+		return true;
+
+	default:
+		return false;
+	}
+}
+
 static BufferPackingStandard packing_to_substruct_packing(BufferPackingStandard packing)
 {
 	switch (packing)
@@ -124,6 +138,8 @@ static BufferPackingStandard packing_to_substruct_packing(BufferPackingStandard
 		return BufferPackingStd430;
 	case BufferPackingHLSLCbufferPackOffset:
 		return BufferPackingHLSLCbuffer;
+	case BufferPackingScalarEnhancedLayout:
+		return BufferPackingScalar;
 	default:
 		return packing;
 	}
@@ -430,6 +446,21 @@ void CompilerGLSL::find_static_extensions()
 
 	if (options.separate_shader_objects && !options.es && options.version < 410)
 		require_extension_internal("GL_ARB_separate_shader_objects");
+
+	if (ir.addressing_model == AddressingModelPhysicalStorageBuffer64EXT)
+	{
+		if (!options.vulkan_semantics)
+			SPIRV_CROSS_THROW("GL_EXT_buffer_reference is only supported in Vulkan GLSL.");
+		if (options.es && options.version < 320)
+			SPIRV_CROSS_THROW("GL_EXT_buffer_reference requires ESSL 320.");
+		else if (!options.es && options.version < 450)
+			SPIRV_CROSS_THROW("GL_EXT_buffer_reference requires GLSL 450.");
+		require_extension_internal("GL_EXT_buffer_reference");
+	}
+	else if (ir.addressing_model != AddressingModelLogical)
+	{
+		SPIRV_CROSS_THROW("Only Logical and PhysicalStorageBuffer64EXT addressing models are supported.");
+	}
 }
 
 string CompilerGLSL::compile()
@@ -446,6 +477,11 @@ string CompilerGLSL::compile()
 	update_active_builtins();
 	analyze_image_and_sampler_usage();
 
+	// Shaders might cast unrelated data to pointers of non-block types.
+	// Find all such instances and make sure we can cast the pointers to a synthesized block type.
+	if (ir.addressing_model == AddressingModelPhysicalStorageBuffer64EXT)
+		analyze_non_block_pointer_types();
+
 	uint32_t pass_count = 0;
 	do
 	{
@@ -972,6 +1008,24 @@ uint32_t CompilerGLSL::type_to_packed_base_size(const SPIRType &type, BufferPack
 uint32_t CompilerGLSL::type_to_packed_alignment(const SPIRType &type, const Bitset &flags,
                                                 BufferPackingStandard packing)
 {
+	// If using PhysicalStorageBufferEXT storage class, this is a pointer,
+	// and is 64-bit.
+	if (type.storage == StorageClassPhysicalStorageBufferEXT)
+	{
+		if (!type.pointer)
+			SPIRV_CROSS_THROW("Types in PhysicalStorageBufferEXT must be pointers.");
+
+		if (ir.addressing_model == AddressingModelPhysicalStorageBuffer64EXT)
+		{
+			if (packing_is_vec4_padded(packing) && type_is_array_of_pointers(type))
+				return 16;
+			else
+				return 8;
+		}
+		else
+			SPIRV_CROSS_THROW("AddressingModelPhysicalStorageBuffer64EXT must be used for PhysicalStorageBufferEXT.");
+	}
+
 	if (!type.array.empty())
 	{
 		uint32_t minimum_alignment = 1;
@@ -1007,6 +1061,10 @@ uint32_t CompilerGLSL::type_to_packed_alignment(const SPIRType &type, const Bits
 	{
 		const uint32_t base_alignment = type_to_packed_base_size(type, packing);
 
+		// Alignment requirement for scalar block layout is always the alignment for the most basic component.
+		if (packing_is_scalar(packing))
+			return base_alignment;
+
 		// Vectors are *not* aligned in HLSL, but there's an extra rule where vectors cannot straddle
 		// a vec4, this is handled outside since that part knows our current offset.
 		if (type.columns == 1 && packing_is_hlsl(packing))
@@ -1088,6 +1146,19 @@ uint32_t CompilerGLSL::type_to_packed_size(const SPIRType &type, const Bitset &f
 		return to_array_size_literal(type) * type_to_packed_array_stride(type, flags, packing);
 	}
 
+	// If using PhysicalStorageBufferEXT storage class, this is a pointer,
+	// and is 64-bit.
+	if (type.storage == StorageClassPhysicalStorageBufferEXT)
+	{
+		if (!type.pointer)
+			SPIRV_CROSS_THROW("Types in PhysicalStorageBufferEXT must be pointers.");
+
+		if (ir.addressing_model == AddressingModelPhysicalStorageBuffer64EXT)
+			return 8;
+		else
+			SPIRV_CROSS_THROW("AddressingModelPhysicalStorageBuffer64EXT must be used for PhysicalStorageBufferEXT.");
+	}
+
 	uint32_t size = 0;
 
 	if (type.basetype == SPIRType::Struct)
@@ -1117,27 +1188,34 @@ uint32_t CompilerGLSL::type_to_packed_size(const SPIRType &type, const Bitset &f
 	{
 		const uint32_t base_alignment = type_to_packed_base_size(type, packing);
 
-		if (type.columns == 1)
-			size = type.vecsize * base_alignment;
-
-		if (flags.get(DecorationColMajor) && type.columns > 1)
+		if (packing_is_scalar(packing))
 		{
-			if (packing_is_vec4_padded(packing))
-				size = type.columns * 4 * base_alignment;
-			else if (type.vecsize == 3)
-				size = type.columns * 4 * base_alignment;
-			else
-				size = type.columns * type.vecsize * base_alignment;
+			size = type.vecsize * type.columns * base_alignment;
 		}
-
-		if (flags.get(DecorationRowMajor) && type.vecsize > 1)
+		else
 		{
-			if (packing_is_vec4_padded(packing))
-				size = type.vecsize * 4 * base_alignment;
-			else if (type.columns == 3)
-				size = type.vecsize * 4 * base_alignment;
-			else
-				size = type.vecsize * type.columns * base_alignment;
+			if (type.columns == 1)
+				size = type.vecsize * base_alignment;
+
+			if (flags.get(DecorationColMajor) && type.columns > 1)
+			{
+				if (packing_is_vec4_padded(packing))
+					size = type.columns * 4 * base_alignment;
+				else if (type.vecsize == 3)
+					size = type.columns * 4 * base_alignment;
+				else
+					size = type.columns * type.vecsize * base_alignment;
+			}
+
+			if (flags.get(DecorationRowMajor) && type.vecsize > 1)
+			{
+				if (packing_is_vec4_padded(packing))
+					size = type.vecsize * 4 * base_alignment;
+				else if (type.columns == 3)
+					size = type.vecsize * 4 * base_alignment;
+				else
+					size = type.vecsize * type.columns * base_alignment;
+			}
 		}
 	}
 
@@ -1211,7 +1289,7 @@ bool CompilerGLSL::buffer_is_packing_standard(const SPIRType &type, BufferPackin
 
 		// The next member following a struct member is aligned to the base alignment of the struct that came before.
 		// GL 4.5 spec, 7.6.2.2.
-		if (memb_type.basetype == SPIRType::Struct)
+		if (memb_type.basetype == SPIRType::Struct && !memb_type.pointer)
 			pad_alignment = packed_alignment;
 		else
 			pad_alignment = 1;
@@ -1237,8 +1315,11 @@ bool CompilerGLSL::buffer_is_packing_standard(const SPIRType &type, BufferPackin
 			// We cannot use enhanced layouts on substructs, so they better be up to spec.
 			auto substruct_packing = packing_to_substruct_packing(packing);
 
-			if (!memb_type.member_types.empty() && !buffer_is_packing_standard(memb_type, substruct_packing))
+			if (!memb_type.pointer && !memb_type.member_types.empty() &&
+			    !buffer_is_packing_standard(memb_type, substruct_packing))
+			{
 				return false;
+			}
 		}
 
 		// Bump size.
@@ -1382,71 +1463,11 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var)
 	// If SPIR-V does not comply with either layout, we cannot really work around it.
 	if (can_use_buffer_blocks && (ubo_block || emulated_ubo))
 	{
-		if (buffer_is_packing_standard(type, BufferPackingStd140))
-			attr.push_back("std140");
-		else if (buffer_is_packing_standard(type, BufferPackingStd140EnhancedLayout))
-		{
-			attr.push_back("std140");
-			// Fallback time. We might be able to use the ARB_enhanced_layouts to deal with this difference,
-			// however, we can only use layout(offset) on the block itself, not any substructs, so the substructs better be the appropriate layout.
-			// Enhanced layouts seem to always work in Vulkan GLSL, so no need for extensions there.
-			if (options.es && !options.vulkan_semantics)
-				SPIRV_CROSS_THROW("Uniform buffer block cannot be expressed as std140. ES-targets do "
-				                  "not support GL_ARB_enhanced_layouts.");
-			if (!options.es && !options.vulkan_semantics && options.version < 440)
-				require_extension_internal("GL_ARB_enhanced_layouts");
-
-			// This is a very last minute to check for this, but use this unused decoration to mark that we should emit
-			// explicit offsets for this block type.
-			// layout_for_variable() will be called before the actual buffer emit.
-			// The alternative is a full pass before codegen where we deduce this decoration,
-			// but then we are just doing the exact same work twice, and more complexity.
-			set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
-		}
-		else
-		{
-			SPIRV_CROSS_THROW("Uniform buffer cannot be expressed as std140, even with enhanced layouts. You can try "
-			                  "flattening this block to "
-			                  "support a more flexible layout.");
-		}
+		attr.push_back(buffer_to_packing_standard(type, false));
 	}
 	else if (can_use_buffer_blocks && (push_constant_block || ssbo_block))
 	{
-		if (buffer_is_packing_standard(type, BufferPackingStd430))
-			attr.push_back("std430");
-		else if (buffer_is_packing_standard(type, BufferPackingStd140))
-			attr.push_back("std140");
-		else if (buffer_is_packing_standard(type, BufferPackingStd140EnhancedLayout))
-		{
-			attr.push_back("std140");
-
-			// Fallback time. We might be able to use the ARB_enhanced_layouts to deal with this difference,
-			// however, we can only use layout(offset) on the block itself, not any substructs, so the substructs better be the appropriate layout.
-			// Enhanced layouts seem to always work in Vulkan GLSL, so no need for extensions there.
-			if (options.es && !options.vulkan_semantics)
-				SPIRV_CROSS_THROW("Push constant block cannot be expressed as neither std430 nor std140. ES-targets do "
-				                  "not support GL_ARB_enhanced_layouts.");
-			if (!options.es && !options.vulkan_semantics && options.version < 440)
-				require_extension_internal("GL_ARB_enhanced_layouts");
-
-			set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
-		}
-		else if (buffer_is_packing_standard(type, BufferPackingStd430EnhancedLayout))
-		{
-			attr.push_back("std430");
-			if (options.es && !options.vulkan_semantics)
-				SPIRV_CROSS_THROW("Push constant block cannot be expressed as neither std430 nor std140. ES-targets do "
-				                  "not support GL_ARB_enhanced_layouts.");
-			if (!options.es && !options.vulkan_semantics && options.version < 440)
-				require_extension_internal("GL_ARB_enhanced_layouts");
-
-			set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
-		}
-		else
-		{
-			SPIRV_CROSS_THROW("Buffer block cannot be expressed as neither std430 nor std140, even with enhanced "
-			                  "layouts. You can try flattening this block to support a more flexible layout.");
-		}
+		attr.push_back(buffer_to_packing_standard(type, true));
 	}
 
 	// For images, the type itself adds a layout qualifer.
@@ -1467,6 +1488,55 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var)
 	return res;
 }
 
+string CompilerGLSL::buffer_to_packing_standard(const SPIRType &type, bool check_std430)
+{
+	if (check_std430 && buffer_is_packing_standard(type, BufferPackingStd430))
+		return "std430";
+	else if (buffer_is_packing_standard(type, BufferPackingStd140))
+		return "std140";
+	else if (options.vulkan_semantics && buffer_is_packing_standard(type, BufferPackingScalar))
+	{
+		require_extension_internal("GL_EXT_scalar_block_layout");
+		return "scalar";
+	}
+	else if (check_std430 && buffer_is_packing_standard(type, BufferPackingStd430EnhancedLayout))
+	{
+		if (options.es && !options.vulkan_semantics)
+			SPIRV_CROSS_THROW("Push constant block cannot be expressed as neither std430 nor std140. ES-targets do "
+			                  "not support GL_ARB_enhanced_layouts.");
+		if (!options.es && !options.vulkan_semantics && options.version < 440)
+			require_extension_internal("GL_ARB_enhanced_layouts");
+
+		set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
+		return "std430";
+	}
+	else if (buffer_is_packing_standard(type, BufferPackingStd140EnhancedLayout))
+	{
+		// Fallback time. We might be able to use the ARB_enhanced_layouts to deal with this difference,
+		// however, we can only use layout(offset) on the block itself, not any substructs, so the substructs better be the appropriate layout.
+		// Enhanced layouts seem to always work in Vulkan GLSL, so no need for extensions there.
+		if (options.es && !options.vulkan_semantics)
+			SPIRV_CROSS_THROW("Push constant block cannot be expressed as neither std430 nor std140. ES-targets do "
+			                  "not support GL_ARB_enhanced_layouts.");
+		if (!options.es && !options.vulkan_semantics && options.version < 440)
+			require_extension_internal("GL_ARB_enhanced_layouts");
+
+		set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
+		return "std140";
+	}
+	else if (options.vulkan_semantics && buffer_is_packing_standard(type, BufferPackingScalarEnhancedLayout))
+	{
+		set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
+		require_extension_internal("GL_EXT_scalar_block_layout");
+		return "scalar";
+	}
+	else
+	{
+		SPIRV_CROSS_THROW("Buffer block cannot be expressed as any of std430, std140, scalar, even with enhanced "
+		                  "layouts. You can try flattening this block to support a more flexible layout.");
+	}
+}
+
 void CompilerGLSL::emit_push_constant_block(const SPIRVariable &var)
 {
 	if (flattened_buffer_blocks.count(var.self))
@@ -1544,6 +1614,81 @@ void CompilerGLSL::emit_buffer_block_legacy(const SPIRVariable &var)
 	statement("");
 }
 
+void CompilerGLSL::emit_buffer_reference_block(SPIRType &type, bool forward_declaration)
+{
+	string buffer_name;
+
+	if (forward_declaration)
+	{
+		// Block names should never alias, but from HLSL input they kind of can because block types are reused for UAVs ...
+		// Allow aliased name since we might be declaring the block twice. Once with buffer reference (forward declared) and one proper declaration.
+		// The names must match up.
+		buffer_name = to_name(type.self, false);
+
+		// Shaders never use the block by interface name, so we don't
+		// have to track this other than updating name caches.
+		// If we have a collision for any reason, just fallback immediately.
+		if (ir.meta[type.self].decoration.alias.empty() ||
+		    block_ssbo_names.find(buffer_name) != end(block_ssbo_names) ||
+		    resource_names.find(buffer_name) != end(resource_names))
+		{
+			buffer_name = join("_", type.self);
+		}
+
+		// Make sure we get something unique for both global name scope and block name scope.
+		// See GLSL 4.5 spec: section 4.3.9 for details.
+		add_variable(block_ssbo_names, resource_names, buffer_name);
+
+		// If for some reason buffer_name is an illegal name, make a final fallback to a workaround name.
+		// This cannot conflict with anything else, so we're safe now.
+		// We cannot reuse this fallback name in neither global scope (blocked by block_names) nor block name scope.
+		if (buffer_name.empty())
+			buffer_name = join("_", type.self);
+
+		block_names.insert(buffer_name);
+		block_ssbo_names.insert(buffer_name);
+	}
+	else if (type.basetype != SPIRType::Struct)
+		buffer_name = type_to_glsl(type);
+	else
+		buffer_name = to_name(type.self, false);
+
+	if (!forward_declaration)
+	{
+		if (type.basetype == SPIRType::Struct)
+			statement("layout(buffer_reference, ", buffer_to_packing_standard(type, true), ") buffer ", buffer_name);
+		else
+			statement("layout(buffer_reference) buffer ", buffer_name);
+
+		begin_scope();
+
+		if (type.basetype == SPIRType::Struct)
+		{
+			type.member_name_cache.clear();
+
+			uint32_t i = 0;
+			for (auto &member : type.member_types)
+			{
+				add_member_name(type, i);
+				emit_struct_member(type, member, i);
+				i++;
+			}
+		}
+		else
+		{
+			auto &pointee_type = get_pointee_type(type);
+			statement(type_to_glsl(pointee_type), " value", type_to_array_glsl(pointee_type), ";");
+		}
+
+		end_scope_decl();
+		statement("");
+	}
+	else
+	{
+		statement("layout(buffer_reference) buffer ", buffer_name, ";");
+	}
+}
+
 void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var)
 {
 	auto &type = get<SPIRType>(var.basetype);
@@ -1629,7 +1774,7 @@ void CompilerGLSL::emit_buffer_block_flattened(const SPIRVariable &var)
 			SPIRV_CROSS_THROW("Basic types in a flattened UBO must be float, int or uint.");
 
 		auto flags = ir.get_buffer_block_flags(var);
-		statement("uniform ", flags_to_precision_qualifiers_glsl(tmp, flags), type_to_glsl(tmp), " ", buffer_name, "[",
+		statement("uniform ", flags_to_qualifiers_glsl(tmp, flags), type_to_glsl(tmp), " ", buffer_name, "[",
 		          buffer_size, "];");
 	}
 	else
@@ -2333,6 +2478,36 @@ void CompilerGLSL::emit_resources()
 
 	emitted = false;
 
+	if (ir.addressing_model == AddressingModelPhysicalStorageBuffer64EXT)
+	{
+		for (auto type : physical_storage_non_block_pointer_types)
+		{
+			emit_buffer_reference_block(get<SPIRType>(type), false);
+		}
+
+		// Output buffer reference blocks.
+		// Do this in two stages, one with forward declaration,
+		// and one without. Buffer reference blocks can reference themselves
+		// to support things like linked lists.
+		ir.for_each_typed_id<SPIRType>([&](uint32_t, SPIRType &type) {
+			bool has_block_flags = has_decoration(type.self, DecorationBlock);
+			if (has_block_flags && type.pointer && type.pointer_depth == 1 && !type_is_array_of_pointers(type) &&
+			    type.storage == StorageClassPhysicalStorageBufferEXT)
+			{
+				emit_buffer_reference_block(type, true);
+			}
+		});
+
+		ir.for_each_typed_id<SPIRType>([&](uint32_t, SPIRType &type) {
+			bool has_block_flags = has_decoration(type.self, DecorationBlock);
+			if (has_block_flags && type.pointer && type.pointer_depth == 1 && !type_is_array_of_pointers(type) &&
+			    type.storage == StorageClassPhysicalStorageBufferEXT)
+			{
+				emit_buffer_reference_block(type, false);
+			}
+		});
+	}
+
 	// Output UBOs and SSBOs
 	ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
 		auto &type = this->get<SPIRType>(var.basetype);
@@ -2534,15 +2709,22 @@ string CompilerGLSL::enclose_expression(const string &expr)
 		return expr;
 }
 
-string CompilerGLSL::dereference_expression(const std::string &expr)
+string CompilerGLSL::dereference_expression(const SPIRType &expr_type, const std::string &expr)
 {
 	// If this expression starts with an address-of operator ('&'), then
 	// just return the part after the operator.
 	// TODO: Strip parens if unnecessary?
 	if (expr.front() == '&')
 		return expr.substr(1);
-	else
+	else if (backend.native_pointers)
 		return join('*', expr);
+	else if (expr_type.storage == StorageClassPhysicalStorageBufferEXT && expr_type.basetype != SPIRType::Struct &&
+	         expr_type.pointer_depth == 1)
+	{
+		return join(enclose_expression(expr), ".value");
+	}
+	else
+		return expr;
 }
 
 string CompilerGLSL::address_of_expression(const std::string &expr)
@@ -2590,7 +2772,7 @@ string CompilerGLSL::to_dereferenced_expression(uint32_t id, bool register_expre
 {
 	auto &type = expression_type(id);
 	if (type.pointer && should_dereference(id))
-		return dereference_expression(to_enclosed_expression(id, register_expression_read));
+		return dereference_expression(type, to_enclosed_expression(id, register_expression_read));
 	else
 		return to_expression(id, register_expression_read);
 }
@@ -3629,7 +3811,7 @@ void CompilerGLSL::emit_uninitialized_temporary(uint32_t result_type, uint32_t r
 
 		// The result_id has not been made into an expression yet, so use flags interface.
 		add_local_variable_name(result_id);
-		statement(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(result_id)), ";");
+		statement(flags_to_qualifiers_glsl(type, flags), variable_decl(type, to_name(result_id)), ";");
 	}
 }
 
@@ -3664,7 +3846,7 @@ string CompilerGLSL::declare_temporary(uint32_t result_type, uint32_t result_id)
 	{
 		// The result_id has not been made into an expression yet, so use flags interface.
 		add_local_variable_name(result_id);
-		return join(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(result_id)), " = ");
+		return join(flags_to_qualifiers_glsl(type, flags), variable_decl(type, to_name(result_id)), " = ");
 	}
 }
 
@@ -5933,6 +6115,21 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
 	// Start traversing type hierarchy at the proper non-pointer types,
 	// but keep type_id referencing the original pointer for use below.
 	uint32_t type_id = expression_type_id(base);
+
+	if (!backend.native_pointers)
+	{
+		if (ptr_chain)
+			SPIRV_CROSS_THROW("Backend does not support native pointers and does not support OpPtrAccessChain.");
+
+		// Wrapped buffer reference pointer types will need to poke into the internal "value" member before
+		// continuing the access chain.
+		if (should_dereference(base))
+		{
+			auto &type = get<SPIRType>(type_id);
+			expr = dereference_expression(type, expr);
+		}
+	}
+
 	const auto *type = &get_pointee_type(type_id);
 
 	bool access_chain_is_arrayed = expr.find_first_of('[') != string::npos;
@@ -6780,8 +6977,7 @@ void CompilerGLSL::flush_variable_declaration(uint32_t id)
 		{
 			auto &type = get<SPIRType>(var->basetype);
 			auto &flags = ir.meta[id].decoration.decoration_flags;
-			statement(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, join("_", id, "_copy")),
-			          ";");
+			statement(flags_to_qualifiers_glsl(type, flags), variable_decl(type, join("_", id, "_copy")), ";");
 		}
 		var->deferred_declaration = false;
 	}
@@ -8418,8 +8614,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		flush_all_atomic_capable_variables();
 		// FIXME: Image?
 		// OpAtomicLoad seems to only be relevant for atomic counters.
+		forced_temporaries.insert(ops[1]);
 		GLSL_UFOP(atomicCounter);
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 
 	case OpAtomicStore:
@@ -8459,7 +8655,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		}
 
 		flush_all_atomic_capable_variables();
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 	}
 
@@ -8469,7 +8664,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		forced_temporaries.insert(ops[1]);
 		emit_binary_func_op(ops[0], ops[1], ops[2], ops[5], op);
 		flush_all_atomic_capable_variables();
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 	}
 
@@ -8480,7 +8674,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		auto expr = join(op, "(", to_expression(ops[2]), ", -", to_enclosed_expression(ops[5]), ")");
 		emit_op(ops[0], ops[1], expr, should_forward(ops[2]) && should_forward(ops[5]));
 		flush_all_atomic_capable_variables();
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 	}
 
@@ -8491,7 +8684,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		forced_temporaries.insert(ops[1]);
 		emit_binary_func_op(ops[0], ops[1], ops[2], ops[5], op);
 		flush_all_atomic_capable_variables();
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 	}
 
@@ -8502,7 +8694,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		forced_temporaries.insert(ops[1]);
 		emit_binary_func_op(ops[0], ops[1], ops[2], ops[5], op);
 		flush_all_atomic_capable_variables();
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 	}
 
@@ -8512,7 +8703,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		forced_temporaries.insert(ops[1]);
 		emit_binary_func_op(ops[0], ops[1], ops[2], ops[5], op);
 		flush_all_atomic_capable_variables();
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 	}
 
@@ -8522,7 +8712,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		forced_temporaries.insert(ops[1]);
 		emit_binary_func_op(ops[0], ops[1], ops[2], ops[5], op);
 		flush_all_atomic_capable_variables();
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 	}
 
@@ -8532,7 +8721,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		forced_temporaries.insert(ops[1]);
 		emit_binary_func_op(ops[0], ops[1], ops[2], ops[5], op);
 		flush_all_atomic_capable_variables();
-		register_read(ops[1], ops[2], should_forward(ops[2]));
 		break;
 	}
 
@@ -9293,6 +9481,29 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		statement("executeCallableNV(", to_expression(ops[0]), ", ", to_expression(ops[1]), ");");
 		break;
 
+	case OpConvertUToPtr:
+	{
+		auto &type = get<SPIRType>(ops[0]);
+		if (type.storage != StorageClassPhysicalStorageBufferEXT)
+			SPIRV_CROSS_THROW("Only StorageClassPhysicalStorageBufferEXT is supported by OpConvertUToPtr.");
+
+		auto op = type_to_glsl(type);
+		emit_unary_func_op(ops[0], ops[1], ops[2], op.c_str());
+		break;
+	}
+
+	case OpConvertPtrToU:
+	{
+		auto &type = get<SPIRType>(ops[0]);
+		auto &ptr_type = expression_type(ops[2]);
+		if (ptr_type.storage != StorageClassPhysicalStorageBufferEXT)
+			SPIRV_CROSS_THROW("Only StorageClassPhysicalStorageBufferEXT is supported by OpConvertPtrToU.");
+
+		auto op = type_to_glsl(type);
+		emit_unary_func_op(ops[0], ops[1], ops[2], op.c_str());
+		break;
+	}
+
 	case OpUndef:
 		// Undefined value has been declared.
 		break;
@@ -9450,13 +9661,16 @@ void CompilerGLSL::emit_struct_member(const SPIRType &type, uint32_t member_type
 	if (is_block)
 		qualifiers = to_interpolation_qualifiers(memberflags);
 
-	statement(layout_for_member(type, index), qualifiers, qualifier,
-	          flags_to_precision_qualifiers_glsl(membertype, memberflags),
+	statement(layout_for_member(type, index), qualifiers, qualifier, flags_to_qualifiers_glsl(membertype, memberflags),
 	          variable_decl(membertype, to_member_name(type, index)), ";");
 }
 
-const char *CompilerGLSL::flags_to_precision_qualifiers_glsl(const SPIRType &type, const Bitset &flags)
+const char *CompilerGLSL::flags_to_qualifiers_glsl(const SPIRType &type, const Bitset &flags)
 {
+	// GL_EXT_buffer_reference variables can be marked as restrict.
+	if (flags.get(DecorationRestrictPointerEXT))
+		return "restrict ";
+
 	// Structs do not have precision qualifiers, neither do doubles (desktop only anyways, so no mediump/highp).
 	if (type.basetype != SPIRType::Float && type.basetype != SPIRType::Int && type.basetype != SPIRType::UInt &&
 	    type.basetype != SPIRType::Image && type.basetype != SPIRType::SampledImage &&
@@ -9509,7 +9723,7 @@ const char *CompilerGLSL::flags_to_precision_qualifiers_glsl(const SPIRType &typ
 
 const char *CompilerGLSL::to_precision_qualifiers_glsl(uint32_t id)
 {
-	return flags_to_precision_qualifiers_glsl(expression_type(id), ir.meta[id].decoration.decoration_flags);
+	return flags_to_qualifiers_glsl(expression_type(id), ir.meta[id].decoration.decoration_flags);
 }
 
 string CompilerGLSL::to_qualifiers_glsl(uint32_t id)
@@ -9672,6 +9886,12 @@ string CompilerGLSL::to_array_size(const SPIRType &type, uint32_t index)
 
 string CompilerGLSL::type_to_array_glsl(const SPIRType &type)
 {
+	if (type.pointer && type.storage == StorageClassPhysicalStorageBufferEXT && type.basetype != SPIRType::Struct)
+	{
+		// We are using a wrapped pointer type, and we should not emit any array declarations here.
+		return "";
+	}
+
 	if (type.array.empty())
 		return "";
 
@@ -9825,7 +10045,20 @@ string CompilerGLSL::type_to_glsl_constructor(const SPIRType &type)
 // depend on a specific object's use of that type.
 string CompilerGLSL::type_to_glsl(const SPIRType &type, uint32_t id)
 {
-	// Ignore the pointer type since GLSL doesn't have pointers.
+	if (type.pointer && type.storage == StorageClassPhysicalStorageBufferEXT && type.basetype != SPIRType::Struct)
+	{
+		// Need to create a magic type name which compacts the entire type information.
+		string name = type_to_glsl(get_pointee_type(type));
+		for (size_t i = 0; i < type.array.size(); i++)
+		{
+			if (type.array_size_literal[i])
+				name += join(type.array[i], "_");
+			else
+				name += join("id", type.array[i], "_");
+		}
+		name += "Pointer";
+		return name;
+	}
 
 	switch (type.basetype)
 	{
@@ -10132,7 +10365,7 @@ void CompilerGLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret
 	string decl;
 
 	auto &type = get<SPIRType>(func.return_type);
-	decl += flags_to_precision_qualifiers_glsl(type, return_flags);
+	decl += flags_to_qualifiers_glsl(type, return_flags);
 	decl += type_to_glsl(type);
 	decl += type_to_array_glsl(type);
 	decl += " ";
@@ -10939,7 +11172,7 @@ void CompilerGLSL::emit_hoisted_temporaries(SmallVector<pair<uint32_t, uint32_t>
 		add_local_variable_name(tmp.second);
 		auto &flags = ir.meta[tmp.second].decoration.decoration_flags;
 		auto &type = get<SPIRType>(tmp.first);
-		statement(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(tmp.second)), ";");
+		statement(flags_to_qualifiers_glsl(type, flags), variable_decl(type, to_name(tmp.second)), ";");
 
 		hoisted_temporaries.insert(tmp.second);
 		forced_temporaries.insert(tmp.second);

3rdparty/spirv-cross/spirv_glsl.hpp

@@ -392,11 +392,13 @@ protected:
 		bool supports_empty_struct = false;
 		bool array_is_value_type = true;
 		bool comparison_image_samples_scalar = false;
+		bool native_pointers = false;
 	} backend;
 
 	void emit_struct(SPIRType &type);
 	void emit_resources();
 	void emit_buffer_block_native(const SPIRVariable &var);
+	void emit_buffer_reference_block(SPIRType &type, bool forward_declaration);
 	void emit_buffer_block_legacy(const SPIRVariable &var);
 	void emit_buffer_block_flattened(const SPIRVariable &type);
 	void emit_declared_builtin_block(spv::StorageClass storage, spv::ExecutionModel model);
@@ -495,7 +497,7 @@ protected:
 	std::string to_enclosed_pointer_expression(uint32_t id, bool register_expression_read = true);
 	std::string to_extract_component_expression(uint32_t id, uint32_t index);
 	std::string enclose_expression(const std::string &expr);
-	std::string dereference_expression(const std::string &expr);
+	std::string dereference_expression(const SPIRType &expression_type, const std::string &expr);
 	std::string address_of_expression(const std::string &expr);
 	void strip_enclosed_expression(std::string &expr);
 	std::string to_member_name(const SPIRType &type, uint32_t index);
@@ -505,7 +507,7 @@ protected:
 	virtual std::string to_qualifiers_glsl(uint32_t id);
 	const char *to_precision_qualifiers_glsl(uint32_t id);
 	virtual const char *to_storage_qualifiers_glsl(const SPIRVariable &var);
-	const char *flags_to_precision_qualifiers_glsl(const SPIRType &type, const Bitset &flags);
+	const char *flags_to_qualifiers_glsl(const SPIRType &type, const Bitset &flags);
 	const char *format_to_glsl(spv::ImageFormat format);
 	virtual std::string layout_for_member(const SPIRType &type, uint32_t index);
 	virtual std::string to_interpolation_qualifiers(const Bitset &flags);
@@ -518,6 +520,8 @@ protected:
 
 	bool buffer_is_packing_standard(const SPIRType &type, BufferPackingStandard packing, uint32_t start_offset = 0,
 	                                uint32_t end_offset = ~(0u));
+	std::string buffer_to_packing_standard(const SPIRType &type, bool enable_std430);
+
 	uint32_t type_to_packed_base_size(const SPIRType &type, BufferPackingStandard packing);
 	uint32_t type_to_packed_alignment(const SPIRType &type, const Bitset &flags, BufferPackingStandard packing);
 	uint32_t type_to_packed_array_stride(const SPIRType &type, const Bitset &flags, BufferPackingStandard packing);

3rdparty/spirv-cross/spirv_hlsl.cpp

@@ -2006,7 +2006,7 @@ void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret
 	auto &type = get<SPIRType>(func.return_type);
 	if (type.array.empty())
 	{
-		decl += flags_to_precision_qualifiers_glsl(type, return_flags);
+		decl += flags_to_qualifiers_glsl(type, return_flags);
 		decl += type_to_glsl(type);
 		decl += " ";
 	}
@@ -3713,7 +3713,6 @@ void CompilerHLSL::emit_atomic(const uint32_t *ops, uint32_t length, spv::Op op)
 	auto expr = bitcast_expression(type, expr_type, to_name(id));
 	set<SPIRExpression>(id, expr, result_type, true);
 	flush_all_atomic_capable_variables();
-	register_read(ops[1], ops[2], should_forward(ops[2]));
 }
 
 void CompilerHLSL::emit_subgroup_op(const Instruction &i)

3rdparty/spirv-cross/spirv_msl.cpp

@@ -584,6 +584,7 @@ string CompilerMSL::compile()
 	backend.allow_truncated_access_chain = true;
 	backend.array_is_value_type = false;
 	backend.comparison_image_samples_scalar = true;
+	backend.native_pointers = true;
 
 	capture_output_to_buffer = msl_options.capture_output_to_buffer;
 	is_rasterization_disabled = msl_options.disable_rasterization || capture_output_to_buffer;
@@ -3729,6 +3730,9 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		break;
 	}
 
+	case OpImageTexelPointer:
+		SPIRV_CROSS_THROW("MSL does not support atomic operations on images or texel buffers.");
+
 	// Casting
 	case OpQuantizeToF16:
 	{
@@ -4483,9 +4487,16 @@ string CompilerMSL::to_function_args(uint32_t img, const SPIRType &imgtype, bool
 		if (coord_type.vecsize > 1)
 			tex_coords = enclose_expression(tex_coords) + ".x";
 
-		// Metal texel buffer textures are 2D, so convert 1D coord to 2D.
-		if (is_fetch)
-			tex_coords = "spvTexelBufferCoord(" + round_fp_tex_coords(tex_coords, coord_is_fp) + ")";
+		if (msl_options.texture_buffer_native)
+		{
+			tex_coords = "uint(" + round_fp_tex_coords(tex_coords, coord_is_fp) + ")";
+		}
+		else
+		{
+			// Metal texel buffer textures are 2D, so convert 1D coord to 2D.
+			if (is_fetch)
+				tex_coords = "spvTexelBufferCoord(" + round_fp_tex_coords(tex_coords, coord_is_fp) + ")";
+		}
 
 		alt_coord_component = 1;
 		break;
@@ -5613,14 +5624,14 @@ void CompilerMSL::entry_point_args_builtin(string &ep_args)
 			if (!ep_args.empty())
 				ep_args += ", ";
 			ep_args +=
-				join("constant uint* spvIndirectParams [[buffer(", msl_options.indirect_params_buffer_index, ")]]");
+			    join("constant uint* spvIndirectParams [[buffer(", msl_options.indirect_params_buffer_index, ")]]");
 		}
 		else if (stage_out_var_id)
 		{
 			if (!ep_args.empty())
 				ep_args += ", ";
 			ep_args +=
-				join("device uint* spvIndirectParams [[buffer(", msl_options.indirect_params_buffer_index, ")]]");
+			    join("device uint* spvIndirectParams [[buffer(", msl_options.indirect_params_buffer_index, ")]]");
 		}
 
 		// Tessellation control shaders get three additional parameters:
@@ -6643,6 +6654,18 @@ string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id)
 			img_type_name += (img_type.arrayed ? "texture1d_array" : "texture1d");
 			break;
 		case DimBuffer:
+			if (img_type.ms || img_type.arrayed)
+				SPIRV_CROSS_THROW("Cannot use texel buffers with multisampling or array layers.");
+
+			if (msl_options.texture_buffer_native)
+			{
+				if (!msl_options.supports_msl_version(2, 1))
+					SPIRV_CROSS_THROW("Native texture_buffer type is only supported in MSL 2.1.");
+				img_type_name = "texture_buffer";
+			}
+			else
+				img_type_name += "texture2d";
+			break;
 		case Dim2D:
 		case DimSubpassData:
 			if (img_type.ms && img_type.arrayed)
@@ -7328,7 +7351,7 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
 	{
 		// Retrieve the image type, and if it's a Buffer, emit a texel coordinate function
 		uint32_t tid = result_types[args[opcode == OpImageWrite ? 0 : 2]];
-		if (tid && compiler.get<SPIRType>(tid).image.dim == DimBuffer)
+		if (tid && compiler.get<SPIRType>(tid).image.dim == DimBuffer && !compiler.msl_options.texture_buffer_native)
 			return SPVFuncImplTexelBufferCoords;
 
 		if (opcode == OpImageFetch && compiler.msl_options.swizzle_texture_samples)

3rdparty/spirv-cross/spirv_msl.hpp

@@ -194,6 +194,9 @@ public:
 		// Add support to explicit pad out components.
 		bool pad_fragment_output_components = false;
 
+		// Requires MSL 2.1, use the native support for texel buffers.
+		bool texture_buffer_native = false;
+
 		bool is_ios()
 		{
 			return platform == iOS;

3rdparty/spirv-cross/spirv_parser.cpp

@@ -158,7 +158,6 @@ void Parser::parse(const Instruction &instruction)
 
 	switch (op)
 	{
-	case OpMemoryModel:
 	case OpSourceContinued:
 	case OpSourceExtension:
 	case OpNop:
@@ -168,6 +167,11 @@ void Parser::parse(const Instruction &instruction)
 	case OpModuleProcessed:
 		break;
 
+	case OpMemoryModel:
+		ir.addressing_model = static_cast<AddressingModel>(ops[0]);
+		ir.memory_model = static_cast<MemoryModel>(ops[1]);
+		break;
+
 	case OpSource:
 	{
 		auto lang = static_cast<SourceLanguage>(ops[0]);
@@ -598,6 +602,20 @@ void Parser::parse(const Instruction &instruction)
 		break;
 	}
 
+	case OpTypeForwardPointer:
+	{
+		uint32_t id = ops[0];
+		auto &ptrbase = set<SPIRType>(id);
+		ptrbase.pointer = true;
+		ptrbase.pointer_depth++;
+		ptrbase.storage = static_cast<StorageClass>(ops[1]);
+
+		if (ptrbase.storage == StorageClassAtomicCounter)
+			ptrbase.basetype = SPIRType::AtomicCounter;
+
+		break;
+	}
+
 	case OpTypeStruct:
 	{
 		uint32_t id = ops[0];

3rdparty/spirv-cross/test_shaders.py

@@ -172,6 +172,8 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths):
         msl_args.append('--msl-domain-lower-left')
     if '.argument.' in shader:
         msl_args.append('--msl-argument-buffers')
+    if '.texture-buffer-native.' in shader:
+        msl_args.append('--msl-texture-buffer-native')
     if '.discrete.' in shader:
         # Arbitrary for testing purposes.
         msl_args.append('--msl-discrete-descriptor-set')