From e108f742474497b4a331ef6a8400cb58bfca489b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=91=D1=80=D0=B0=D0=BD=D0=B8=D0=BC=D0=B8=D1=80=20=D0=9A?= =?UTF-8?q?=D0=B0=D1=80=D0=B0=D1=9F=D0=B8=D1=9B?= Date: Sat, 27 Apr 2019 14:14:44 -0700 Subject: [PATCH] Updated spirv-cross. --- 3rdparty/spirv-cross/CMakeLists.txt | 2 +- 3rdparty/spirv-cross/main.cpp | 4 + .../asm/vert/copy-memory-interface.asm.vert | 25 + ...uffer-read.frag => buffer-read-write.frag} | 3 +- ...ead-write.texture-buffer-native.msl21.frag | 18 + ...re_buffer.texture-buffer-native.msl21.vert | 17 + ...ay-of-buffer-reference.nocompat.vk.comp.vk | 25 + .../comp/buffer-reference.nocompat.vk.comp.vk | 45 ++ ...packing-scalar.nocompat.invalid.vk.comp.vk | 147 ++++++ .../asm/comp/atomic-result-temporary.asm.comp | 24 + .../asm/comp/atomic-result-temporary.asm.comp | 23 + .../asm/vert/copy-memory-interface.asm.vert | 25 + ...uffer-read.frag => buffer-read-write.frag} | 3 +- ...ead-write.texture-buffer-native.msl21.frag | 18 + ...re_buffer.texture-buffer-native.msl21.vert | 17 + .../asm/comp/atomic-result-temporary.asm.comp | 18 + ...thesized-pointer-2.asm.nocompat.vk.comp.vk | 21 + ...ynthesized-pointer.asm.nocompat.vk.comp.vk | 21 + ...ay-of-buffer-reference.nocompat.vk.comp.vk | 25 + .../comp/buffer-reference.nocompat.vk.comp.vk | 56 +++ ...packing-scalar.nocompat.invalid.vk.comp.vk | 147 ++++++ .../asm/comp/atomic-result-temporary.asm.comp | 59 +++ .../asm/comp/atomic-result-temporary.asm.comp | 59 +++ .../asm/vert/copy-memory-interface.asm.vert | 33 ++ ...uffer-read.frag => buffer-read-write.frag} | 2 + ...ead-write.texture-buffer-native.msl21.frag | 12 + ...re_buffer.texture-buffer-native.msl21.vert | 10 + .../asm/comp/atomic-result-temporary.asm.comp | 59 +++ ...synthesized-pointer-2.asm.nocompat.vk.comp | 44 ++ ...e-synthesized-pointer.asm.nocompat.vk.comp | 51 ++ ...array-of-buffer-reference.nocompat.vk.comp | 23 + .../comp/buffer-reference.nocompat.vk.comp | 40 ++ ...ct-packing-scalar.nocompat.invalid.vk.comp | 88 ++++ 3rdparty/spirv-cross/spirv_cross.cpp | 82 +++- 3rdparty/spirv-cross/spirv_cross.hpp | 16 +- 3rdparty/spirv-cross/spirv_cross_c.cpp | 4 + 3rdparty/spirv-cross/spirv_cross_c.h | 4 +- .../spirv-cross/spirv_cross_containers.hpp | 5 +- .../spirv-cross/spirv_cross_parsed_ir.cpp | 35 +- .../spirv-cross/spirv_cross_parsed_ir.hpp | 3 + 3rdparty/spirv-cross/spirv_glsl.cpp | 445 +++++++++++++----- 3rdparty/spirv-cross/spirv_glsl.hpp | 8 +- 3rdparty/spirv-cross/spirv_hlsl.cpp | 3 +- 3rdparty/spirv-cross/spirv_msl.cpp | 35 +- 3rdparty/spirv-cross/spirv_msl.hpp | 3 + 3rdparty/spirv-cross/spirv_parser.cpp | 20 +- 3rdparty/spirv-cross/test_shaders.py | 2 + 47 files changed, 1686 insertions(+), 143 deletions(-) create mode 100644 3rdparty/spirv-cross/reference/opt/shaders-msl/asm/vert/copy-memory-interface.asm.vert rename 3rdparty/spirv-cross/reference/opt/shaders-msl/frag/{buffer-read.frag => buffer-read-write.frag} (65%) create mode 100644 3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag create mode 100644 3rdparty/spirv-cross/reference/opt/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert create mode 100644 3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp.vk create mode 100644 3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk create mode 100644 3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp.vk create mode 100644 3rdparty/spirv-cross/reference/shaders-hlsl-no-opt/asm/comp/atomic-result-temporary.asm.comp create mode 100644 3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/comp/atomic-result-temporary.asm.comp create mode 100644 3rdparty/spirv-cross/reference/shaders-msl/asm/vert/copy-memory-interface.asm.vert rename 3rdparty/spirv-cross/reference/shaders-msl/frag/{buffer-read.frag => buffer-read-write.frag} (65%) create mode 100644 3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag create mode 100644 3rdparty/spirv-cross/reference/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert create mode 100644 3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/atomic-result-temporary.asm.comp create mode 100644 3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer-2.asm.nocompat.vk.comp.vk create mode 100644 3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer.asm.nocompat.vk.comp.vk create mode 100644 3rdparty/spirv-cross/reference/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp.vk create mode 100644 3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk create mode 100644 3rdparty/spirv-cross/reference/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp.vk create mode 100644 3rdparty/spirv-cross/shaders-hlsl-no-opt/asm/comp/atomic-result-temporary.asm.comp create mode 100644 3rdparty/spirv-cross/shaders-msl-no-opt/asm/comp/atomic-result-temporary.asm.comp create mode 100644 3rdparty/spirv-cross/shaders-msl/asm/vert/copy-memory-interface.asm.vert rename 3rdparty/spirv-cross/shaders-msl/frag/{buffer-read.frag => buffer-read-write.frag} (58%) create mode 100644 3rdparty/spirv-cross/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag create mode 100644 3rdparty/spirv-cross/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert create mode 100644 3rdparty/spirv-cross/shaders-no-opt/asm/comp/atomic-result-temporary.asm.comp create mode 100644 3rdparty/spirv-cross/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer-2.asm.nocompat.vk.comp create mode 100644 3rdparty/spirv-cross/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer.asm.nocompat.vk.comp create mode 100644 3rdparty/spirv-cross/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp create mode 100644 3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp create mode 100644 3rdparty/spirv-cross/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp diff --git a/3rdparty/spirv-cross/CMakeLists.txt b/3rdparty/spirv-cross/CMakeLists.txt index 24cd32066..664ba4211 100644 --- a/3rdparty/spirv-cross/CMakeLists.txt +++ b/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) diff --git a/3rdparty/spirv-cross/main.cpp b/3rdparty/spirv-cross/main.cpp index ed395f6c1..2cfb99dcf 100644 --- a/3rdparty/spirv-cross/main.cpp +++ b/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 msl_discrete_descriptor_sets; SmallVector 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 ]\n" "\t[--hlsl]\n" "\t[--reflect]\n" @@ -727,6 +729,7 @@ static string compile_iteration(const CLIArguments &args, std::vector 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(); diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/vert/copy-memory-interface.asm.vert b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/vert/copy-memory-interface.asm.vert new file mode 100644 index 000000000..63ab796e8 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/asm/vert/copy-memory-interface.asm.vert @@ -0,0 +1,25 @@ +#include +#include + +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; +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read.frag b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read-write.frag similarity index 65% rename from 3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read.frag rename to 3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read-write.frag index fdd88b568..2b2ac7f06 100644 --- a/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read.frag +++ b/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 buf [[texture(0)]]) +fragment main0_out main0(texture2d buf [[texture(0)]], texture2d 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; } diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag new file mode 100644 index 000000000..71496a4ef --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag @@ -0,0 +1,18 @@ +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +fragment main0_out main0(texture_buffer buf [[texture(0)]], texture_buffer 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; +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert b/3rdparty/spirv-cross/reference/opt/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert new file mode 100644 index 000000000..75332f484 --- /dev/null +++ b/3rdparty/spirv-cross/reference/opt/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert @@ -0,0 +1,17 @@ +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 gl_Position [[position]]; +}; + +vertex main0_out main0(texture_buffer uSamp [[texture(4)]], texture_buffer uSampo [[texture(5)]]) +{ + main0_out out = {}; + out.gl_Position = uSamp.read(uint(10)) + uSampo.read(uint(100)); + return out; +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp.vk new file mode 100644 index 000000000..82ebb9608 --- /dev/null +++ b/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; +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk new file mode 100644 index 000000000..dfcaac836 --- /dev/null +++ b/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); +} + diff --git a/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp.vk b/3rdparty/spirv-cross/reference/opt/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp.vk new file mode 100644 index 000000000..d67e0beeb --- /dev/null +++ b/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; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-hlsl-no-opt/asm/comp/atomic-result-temporary.asm.comp b/3rdparty/spirv-cross/reference/shaders-hlsl-no-opt/asm/comp/atomic-result-temporary.asm.comp new file mode 100644 index 000000000..3a03fafe6 --- /dev/null +++ b/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(); +} diff --git a/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/comp/atomic-result-temporary.asm.comp b/3rdparty/spirv-cross/reference/shaders-msl-no-opt/asm/comp/atomic-result-temporary.asm.comp new file mode 100644 index 000000000..8b6694288 --- /dev/null +++ b/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 +#include +#include + +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; + } +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/asm/vert/copy-memory-interface.asm.vert b/3rdparty/spirv-cross/reference/shaders-msl/asm/vert/copy-memory-interface.asm.vert new file mode 100644 index 000000000..63ab796e8 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/asm/vert/copy-memory-interface.asm.vert @@ -0,0 +1,25 @@ +#include +#include + +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; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read.frag b/3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read-write.frag similarity index 65% rename from 3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read.frag rename to 3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read-write.frag index fdd88b568..2b2ac7f06 100644 --- a/3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read.frag +++ b/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 buf [[texture(0)]]) +fragment main0_out main0(texture2d buf [[texture(0)]], texture2d 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; } diff --git a/3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag b/3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag new file mode 100644 index 000000000..71496a4ef --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag @@ -0,0 +1,18 @@ +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 FragColor [[color(0)]]; +}; + +fragment main0_out main0(texture_buffer buf [[texture(0)]], texture_buffer 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; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert b/3rdparty/spirv-cross/reference/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert new file mode 100644 index 000000000..75332f484 --- /dev/null +++ b/3rdparty/spirv-cross/reference/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert @@ -0,0 +1,17 @@ +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 gl_Position [[position]]; +}; + +vertex main0_out main0(texture_buffer uSamp [[texture(4)]], texture_buffer uSampo [[texture(5)]]) +{ + main0_out out = {}; + out.gl_Position = uSamp.read(uint(10)) + uSampo.read(uint(100)); + return out; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/atomic-result-temporary.asm.comp b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/atomic-result-temporary.asm.comp new file mode 100644 index 000000000..b51c6c58d --- /dev/null +++ b/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; + } +} + diff --git a/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer-2.asm.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer-2.asm.nocompat.vk.comp.vk new file mode 100644 index 000000000..028893191 --- /dev/null +++ b/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; +} + diff --git a/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer.asm.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer.asm.nocompat.vk.comp.vk new file mode 100644 index 000000000..9553199b4 --- /dev/null +++ b/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; +} + diff --git a/3rdparty/spirv-cross/reference/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp.vk new file mode 100644 index 000000000..82ebb9608 --- /dev/null +++ b/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; +} + diff --git a/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk b/3rdparty/spirv-cross/reference/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp.vk new file mode 100644 index 000000000..610d60cb4 --- /dev/null +++ b/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); +} + diff --git a/3rdparty/spirv-cross/reference/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp.vk b/3rdparty/spirv-cross/reference/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp.vk new file mode 100644 index 000000000..d67e0beeb --- /dev/null +++ b/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; +} + diff --git a/3rdparty/spirv-cross/shaders-hlsl-no-opt/asm/comp/atomic-result-temporary.asm.comp b/3rdparty/spirv-cross/shaders-hlsl-no-opt/asm/comp/atomic-result-temporary.asm.comp new file mode 100644 index 000000000..a32384159 --- /dev/null +++ b/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 diff --git a/3rdparty/spirv-cross/shaders-msl-no-opt/asm/comp/atomic-result-temporary.asm.comp b/3rdparty/spirv-cross/shaders-msl-no-opt/asm/comp/atomic-result-temporary.asm.comp new file mode 100644 index 000000000..a32384159 --- /dev/null +++ b/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 diff --git a/3rdparty/spirv-cross/shaders-msl/asm/vert/copy-memory-interface.asm.vert b/3rdparty/spirv-cross/shaders-msl/asm/vert/copy-memory-interface.asm.vert new file mode 100644 index 000000000..c52c9bf0b --- /dev/null +++ b/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 diff --git a/3rdparty/spirv-cross/shaders-msl/frag/buffer-read.frag b/3rdparty/spirv-cross/shaders-msl/frag/buffer-read-write.frag similarity index 58% rename from 3rdparty/spirv-cross/shaders-msl/frag/buffer-read.frag rename to 3rdparty/spirv-cross/shaders-msl/frag/buffer-read-write.frag index 297f0a71c..70af7d384 100644 --- a/3rdparty/spirv-cross/shaders-msl/frag/buffer-read.frag +++ b/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); } diff --git a/3rdparty/spirv-cross/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag b/3rdparty/spirv-cross/shaders-msl/frag/buffer-read-write.texture-buffer-native.msl21.frag new file mode 100644 index 000000000..70af7d384 --- /dev/null +++ b/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); +} diff --git a/3rdparty/spirv-cross/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert b/3rdparty/spirv-cross/shaders-msl/vert/texture_buffer.texture-buffer-native.msl21.vert new file mode 100644 index 000000000..6bc7ddfae --- /dev/null +++ b/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); +} diff --git a/3rdparty/spirv-cross/shaders-no-opt/asm/comp/atomic-result-temporary.asm.comp b/3rdparty/spirv-cross/shaders-no-opt/asm/comp/atomic-result-temporary.asm.comp new file mode 100644 index 000000000..a32384159 --- /dev/null +++ b/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 diff --git a/3rdparty/spirv-cross/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer-2.asm.nocompat.vk.comp b/3rdparty/spirv-cross/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer-2.asm.nocompat.vk.comp new file mode 100644 index 000000000..76894aa8c --- /dev/null +++ b/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 diff --git a/3rdparty/spirv-cross/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer.asm.nocompat.vk.comp b/3rdparty/spirv-cross/shaders-no-opt/asm/comp/buffer-reference-synthesized-pointer.asm.nocompat.vk.comp new file mode 100644 index 000000000..d1270d4f2 --- /dev/null +++ b/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 diff --git a/3rdparty/spirv-cross/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp b/3rdparty/spirv-cross/shaders/vulkan/comp/array-of-buffer-reference.nocompat.vk.comp new file mode 100644 index 000000000..a1da941fd --- /dev/null +++ b/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; +} diff --git a/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp b/3rdparty/spirv-cross/shaders/vulkan/comp/buffer-reference.nocompat.vk.comp new file mode 100644 index 000000000..624b8c0a0 --- /dev/null +++ b/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); +} diff --git a/3rdparty/spirv-cross/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp b/3rdparty/spirv-cross/shaders/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp new file mode 100644 index 000000000..808403d96 --- /dev/null +++ b/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; +} + diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index cbe9b26ac..97090af4f 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/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(args[0]); if (var && storage_class_is_interface(var->storage)) - variables.insert(variable); + variables.insert(args[0]); var = compiler.maybe_get(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(type.member_types[i]), i)); + { + auto &memb_type = get(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(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(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(type.parent_type).pointer_depth; +} diff --git a/3rdparty/spirv-cross/spirv_cross.hpp b/3rdparty/spirv-cross/spirv_cross.hpp index 4e0b171cf..4129e8166 100644 --- a/3rdparty/spirv-cross/spirv_cross.hpp +++ b/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 types; + }; + void analyze_non_block_pointer_types(); + SmallVector 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; diff --git a/3rdparty/spirv-cross/spirv_cross_c.cpp b/3rdparty/spirv-cross/spirv_cross_c.cpp index 48e48228b..b1f825f39 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.cpp +++ b/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: diff --git a/3rdparty/spirv-cross/spirv_cross_c.h b/3rdparty/spirv-cross/spirv_cross_c.h index 5491a2e6d..ee8d15d9e 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.h +++ b/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; diff --git a/3rdparty/spirv-cross/spirv_cross_containers.hpp b/3rdparty/spirv-cross/spirv_cross_containers.hpp index 11cb37197..393f4614b 100644 --- a/3rdparty/spirv-cross/spirv_cross_containers.hpp +++ b/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 diff --git a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp index 4d92b8477..108000cd1 100644 --- a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp +++ b/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()) diff --git a/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp b/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp index 78ca8813f..79e9e15bb 100644 --- a/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp +++ b/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, diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index 0671e8800..4616cfc49 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/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(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(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([&](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([&](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([&](uint32_t, SPIRVariable &var) { auto &type = this->get(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(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(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(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(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(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 add_local_variable_name(tmp.second); auto &flags = ir.meta[tmp.second].decoration.decoration_flags; auto &type = get(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); diff --git a/3rdparty/spirv-cross/spirv_glsl.hpp b/3rdparty/spirv-cross/spirv_glsl.hpp index 37789a139..37012aa79 100644 --- a/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/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); diff --git a/3rdparty/spirv-cross/spirv_hlsl.cpp b/3rdparty/spirv-cross/spirv_hlsl.cpp index 871d18afd..0cd2fb97a 100644 --- a/3rdparty/spirv-cross/spirv_hlsl.cpp +++ b/3rdparty/spirv-cross/spirv_hlsl.cpp @@ -2006,7 +2006,7 @@ void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret auto &type = get(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(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) diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index a215b91b0..aaa65d4b5 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/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(tid).image.dim == DimBuffer) + if (tid && compiler.get(tid).image.dim == DimBuffer && !compiler.msl_options.texture_buffer_native) return SPVFuncImplTexelBufferCoords; if (opcode == OpImageFetch && compiler.msl_options.swizzle_texture_samples) diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index 5aff0345f..8d3a8ad8d 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/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; diff --git a/3rdparty/spirv-cross/spirv_parser.cpp b/3rdparty/spirv-cross/spirv_parser.cpp index 1372df606..ea356a6ac 100644 --- a/3rdparty/spirv-cross/spirv_parser.cpp +++ b/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(ops[0]); + ir.memory_model = static_cast(ops[1]); + break; + case OpSource: { auto lang = static_cast(ops[0]); @@ -598,6 +602,20 @@ void Parser::parse(const Instruction &instruction) break; } + case OpTypeForwardPointer: + { + uint32_t id = ops[0]; + auto &ptrbase = set(id); + ptrbase.pointer = true; + ptrbase.pointer_depth++; + ptrbase.storage = static_cast(ops[1]); + + if (ptrbase.storage == StorageClassAtomicCounter) + ptrbase.basetype = SPIRType::AtomicCounter; + + break; + } + case OpTypeStruct: { uint32_t id = ops[0]; diff --git a/3rdparty/spirv-cross/test_shaders.py b/3rdparty/spirv-cross/test_shaders.py index f41349b26..d6037a8ea 100755 --- a/3rdparty/spirv-cross/test_shaders.py +++ b/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')