diff --git a/reference/opt/shaders-hlsl/comp/rayquery.nofxc.fxconly.comp b/reference/opt/shaders-hlsl/comp/rayquery.nofxc.fxconly.comp index a18be3bb0..80394bef7 100644 --- a/reference/opt/shaders-hlsl/comp/rayquery.nofxc.fxconly.comp +++ b/reference/opt/shaders-hlsl/comp/rayquery.nofxc.fxconly.comp @@ -1,11 +1,11 @@ +static float3x4 _362; +static float4x3 _364; + RWByteAddressBuffer _17 : register(u0); uniform RaytracingAccelerationStructure rtas : register(t1); static RayQuery rayQuery; -static float3x4 _362; -static float4x3 _364; - void comp_main() { RayDesc _1ident = {0.0f.xxx, 0.0f, float3(1.0f, 0.0f, 0.0f), 9999.0f}; diff --git a/reference/opt/shaders-hlsl/frag/bvec-operations.frag b/reference/opt/shaders-hlsl/frag/bvec-operations.frag index 2770e97c0..4813cc55a 100644 --- a/reference/opt/shaders-hlsl/frag/bvec-operations.frag +++ b/reference/opt/shaders-hlsl/frag/bvec-operations.frag @@ -1,3 +1,5 @@ +static bool _47; + static float2 value; static float4 FragColor; @@ -11,8 +13,6 @@ struct SPIRV_Cross_Output float4 FragColor : SV_Target0; }; -static bool _47; - void frag_main() { bool2 _25 = bool2(value.x == 0.0f, _47); diff --git a/reference/opt/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp b/reference/opt/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp new file mode 100644 index 000000000..986e90966 --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp @@ -0,0 +1,23 @@ +#include +#include + +using namespace metal; + +struct _19 +{ +}; +struct _5 +{ + int _m0; + _19 _m1; + char _m2_pad[4]; + _19 _m2; + char _m3_pad[4]; + int _m3; +}; + +kernel void main0(device _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]]) +{ + _4 = _3; +} + diff --git a/reference/opt/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp b/reference/opt/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp new file mode 100644 index 000000000..4bcfeb21a --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp @@ -0,0 +1,24 @@ +#include +#include + +using namespace metal; + +struct _19 +{ +}; +struct _5 +{ + int _m0; + char _m1_pad[12]; + _19 _m1; + char _m2_pad[16]; + _19 _m2; + char _m3_pad[16]; + int _m3; +}; + +kernel void main0(constant _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]]) +{ + _4 = _3; +} + diff --git a/reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp b/reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp index 7c3eba06b..d643379aa 100644 --- a/reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp +++ b/reference/opt/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp @@ -54,6 +54,7 @@ struct _7 int _m0[1]; }; +constant int3 _32 = {}; constant int _3_tmp [[function_constant(0)]]; constant int _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 0; constant int _4_tmp [[function_constant(1)]]; @@ -102,8 +103,6 @@ constant int _74 = (_73 * _72); constant spvUnsafeArray _33 = spvUnsafeArray({ 0, 0, 0 }); constant spvUnsafeArray, 3> _34 = spvUnsafeArray, 3>({ spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }) }); -constant int3 _32 = {}; - kernel void main0(device _7& _8 [[buffer(0)]], device _7& _9 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { _9._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + ((((1 - _59) * _60) * (_61 - 1)) * _74); diff --git a/reference/opt/shaders-msl/asm/comp/uint_smulextended.asm.comp b/reference/opt/shaders-msl/asm/comp/uint_smulextended.asm.comp new file mode 100644 index 000000000..6996f7fd2 --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/uint_smulextended.asm.comp @@ -0,0 +1,25 @@ +#include +#include + +using namespace metal; + +struct _4 +{ + uint _m0[1]; +}; + +struct _20 +{ + uint _m0; + uint _m1; +}; + +kernel void main0(device _4& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]], device _4& _7 [[buffer(2)]], device _4& _8 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _20 _28; + _28._m0 = uint(int(_5._m0[gl_GlobalInvocationID.x]) * int(_6._m0[gl_GlobalInvocationID.x])); + _28._m1 = uint(mulhi(int(_5._m0[gl_GlobalInvocationID.x]), int(_6._m0[gl_GlobalInvocationID.x]))); + _7._m0[gl_GlobalInvocationID.x] = _28._m0; + _8._m0[gl_GlobalInvocationID.x] = _28._m1; +} + diff --git a/reference/opt/shaders-msl/asm/comp/undefined-constant-composite.asm.comp b/reference/opt/shaders-msl/asm/comp/undefined-constant-composite.asm.comp new file mode 100644 index 000000000..359e8913f --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/undefined-constant-composite.asm.comp @@ -0,0 +1,28 @@ +#include +#include + +using namespace metal; + +struct _20 +{ + int _m0; + int _m1; +}; + +struct _5 +{ + int _m0[10]; +}; + +struct _7 +{ + int _m0[10]; +}; + +constant int _28 = {}; + +kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _6._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + (_20{ _28, 200 })._m1; +} + diff --git a/reference/opt/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp b/reference/opt/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp new file mode 100644 index 000000000..0cb22e176 --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp @@ -0,0 +1,31 @@ +#include +#include + +using namespace metal; + +struct _21 +{ + int _m0; + int _m1; +}; + +struct _5 +{ + int _m0[10]; +}; + +struct _7 +{ + int _m0[10]; +}; + +constant int _29 = {}; +constant int _9_tmp [[function_constant(0)]]; +constant int _9 = is_function_constant_defined(_9_tmp) ? _9_tmp : 0; +constant _21 _30 = _21{ _9, _29 }; + +kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _6._m0[gl_GlobalInvocationID.x] = (_8._m0[gl_GlobalInvocationID.x] + _30._m0) + (_21{ _29, 200 })._m1; +} + diff --git a/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag b/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag index ac6b15d73..3c9be2985 100644 --- a/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag +++ b/reference/opt/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag @@ -16,6 +16,8 @@ struct type_Globals uint2 ShadowTileListGroupSize; }; +constant float3 _70 = {}; + struct spvDescriptorSetBuffer0 { const device type_StructuredBuffer_v4float* CulledObjectBoxBounds [[id(0)]]; @@ -24,8 +26,6 @@ struct spvDescriptorSetBuffer0 device atomic_uint* RWShadowTileNumCulledObjects_atomic [[id(3)]]; }; -constant float3 _70 = {}; - struct main0_out { float4 out_var_SV_Target0 [[color(0)]]; diff --git a/reference/opt/shaders/asm/comp/undefined-constant-composite.asm.comp b/reference/opt/shaders/asm/comp/undefined-constant-composite.asm.comp new file mode 100644 index 000000000..eb94756d7 --- /dev/null +++ b/reference/opt/shaders/asm/comp/undefined-constant-composite.asm.comp @@ -0,0 +1,26 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct _20 +{ + int _m0; + int _m1; +}; + +int _28; + +layout(binding = 1, std430) buffer _5_6 +{ + int _m0[10]; +} _6; + +layout(binding = 0, std430) buffer _7_8 +{ + int _m0[10]; +} _8; + +void main() +{ + _6._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + _20(_28, 200)._m1; +} + diff --git a/reference/opt/shaders/asm/frag/default-member-names.asm.frag b/reference/opt/shaders/asm/frag/default-member-names.asm.frag index 13f81b11a..ad64761ed 100644 --- a/reference/opt/shaders/asm/frag/default-member-names.asm.frag +++ b/reference/opt/shaders/asm/frag/default-member-names.asm.frag @@ -1,9 +1,9 @@ #version 450 -layout(location = 0) out vec4 _3; - float _49; +layout(location = 0) out vec4 _3; + void main() { _3 = vec4(_49); diff --git a/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag b/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag index 554bb4677..eb16828e6 100644 --- a/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag +++ b/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag @@ -12,6 +12,8 @@ #define SPIRV_CROSS_LOOP #endif +int _231; + layout(binding = 0, std140) uniform Foo { layout(row_major) mat4 lightVP[64]; @@ -22,8 +24,6 @@ layout(binding = 0, std140) uniform Foo layout(location = 0) in vec3 fragWorld; layout(location = 0) out int _entryPointOutput; -int _231; - mat4 spvWorkaroundRowMajor(mat4 wrap) { return wrap; } void main() diff --git a/reference/opt/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag b/reference/opt/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag index b2473f4d0..7930ca3b4 100644 --- a/reference/opt/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag +++ b/reference/opt/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag @@ -8,12 +8,12 @@ struct Foo float var2; }; +Foo _22; + layout(binding = 0) uniform mediump sampler2D uSampler; layout(location = 0) out vec4 FragColor; -Foo _22; - void main() { FragColor = texture(uSampler, vec2(_22.var1, _22.var2)); diff --git a/reference/opt/shaders/comp/cfg.comp b/reference/opt/shaders/comp/cfg.comp index af2073784..f6e02a855 100644 --- a/reference/opt/shaders/comp/cfg.comp +++ b/reference/opt/shaders/comp/cfg.comp @@ -1,13 +1,13 @@ #version 310 es layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +float _188; + layout(binding = 0, std430) buffer SSBO { float data; } _11; -float _188; - void main() { if (_11.data != 0.0) diff --git a/reference/opt/shaders/frag/loop-dominator-and-switch-default.frag b/reference/opt/shaders/frag/loop-dominator-and-switch-default.frag index 2c193483d..ee64d8335 100644 --- a/reference/opt/shaders/frag/loop-dominator-and-switch-default.frag +++ b/reference/opt/shaders/frag/loop-dominator-and-switch-default.frag @@ -2,10 +2,10 @@ precision mediump float; precision highp int; -layout(location = 0) out vec4 fragColor; - vec4 _80; +layout(location = 0) out vec4 fragColor; + void main() { mediump int _18 = int(_80.x); diff --git a/reference/shaders-hlsl-no-opt/asm/comp/constant-composite-undef.asm.comp b/reference/shaders-hlsl-no-opt/asm/comp/constant-composite-undef.asm.comp index e05b2f1d2..a9eab1cca 100644 --- a/reference/shaders-hlsl-no-opt/asm/comp/constant-composite-undef.asm.comp +++ b/reference/shaders-hlsl-no-opt/asm/comp/constant-composite-undef.asm.comp @@ -1,7 +1,7 @@ -RWByteAddressBuffer block : register(u0); - static float _15; +RWByteAddressBuffer block : register(u0); + void comp_main() { block.Store4(0, asuint(float4(0.100000001490116119384765625f, 0.20000000298023223876953125f, 0.300000011920928955078125f, 0.0f))); diff --git a/reference/shaders-hlsl-no-opt/asm/frag/composite-insert-inheritance.asm.frag b/reference/shaders-hlsl-no-opt/asm/frag/composite-insert-inheritance.asm.frag index 69f6a1915..8c61e61b5 100644 --- a/reference/shaders-hlsl-no-opt/asm/frag/composite-insert-inheritance.asm.frag +++ b/reference/shaders-hlsl-no-opt/asm/frag/composite-insert-inheritance.asm.frag @@ -1,3 +1,5 @@ +static float4 _32; + static const float4 _34[2] = { 0.0f.xxxx, 0.0f.xxxx }; static float4 vInput; @@ -13,8 +15,6 @@ struct SPIRV_Cross_Output float4 FragColor : SV_Target0; }; -static float4 _32; - void frag_main() { float4 _37 = vInput; diff --git a/reference/shaders-hlsl-no-opt/asm/frag/phi.zero-initialize.asm.frag b/reference/shaders-hlsl-no-opt/asm/frag/phi.zero-initialize.asm.frag index c9c152bef..2ce5fd41c 100644 --- a/reference/shaders-hlsl-no-opt/asm/frag/phi.zero-initialize.asm.frag +++ b/reference/shaders-hlsl-no-opt/asm/frag/phi.zero-initialize.asm.frag @@ -1,8 +1,14 @@ +static int uninit_int = 0; +static int4 uninit_vector = int4(0, 0, 0, 0); +static float4x4 uninit_matrix = float4x4(0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx); + struct Foo { int a; }; +static Foo uninit_foo = { 0 }; + static float4 vColor; static float4 FragColor; @@ -16,11 +22,6 @@ struct SPIRV_Cross_Output float4 FragColor : SV_Target0; }; -static int uninit_int = 0; -static int4 uninit_vector = int4(0, 0, 0, 0); -static float4x4 uninit_matrix = float4x4(0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx); -static Foo uninit_foo = { 0 }; - void frag_main() { int _39 = 0; diff --git a/reference/shaders-hlsl-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag b/reference/shaders-hlsl-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag index fa41b13c5..95a48835d 100644 --- a/reference/shaders-hlsl-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag +++ b/reference/shaders-hlsl-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag @@ -1,3 +1,5 @@ +static float4 undef; + static float4 FragColor; static float4 vFloat; @@ -11,8 +13,6 @@ struct SPIRV_Cross_Output float4 FragColor : SV_Target0; }; -static float4 undef; - void frag_main() { FragColor = float4(undef.x, vFloat.y, 0.0f, vFloat.w) + float4(vFloat.z, vFloat.y, 0.0f, vFloat.w); diff --git a/reference/shaders-hlsl/asm/frag/unreachable.asm.frag b/reference/shaders-hlsl/asm/frag/unreachable.asm.frag index 86c1daa24..5eb70adf8 100644 --- a/reference/shaders-hlsl/asm/frag/unreachable.asm.frag +++ b/reference/shaders-hlsl/asm/frag/unreachable.asm.frag @@ -1,3 +1,5 @@ +static float4 _21; + static int counter; static float4 FragColor; @@ -11,8 +13,6 @@ struct SPIRV_Cross_Output float4 FragColor : SV_Target0; }; -static float4 _21; - void frag_main() { float4 _24; diff --git a/reference/shaders-msl-no-opt/asm/frag/composite-insert-inheritance.asm.frag b/reference/shaders-msl-no-opt/asm/frag/composite-insert-inheritance.asm.frag index 8f3153b70..a4bb56283 100644 --- a/reference/shaders-msl-no-opt/asm/frag/composite-insert-inheritance.asm.frag +++ b/reference/shaders-msl-no-opt/asm/frag/composite-insert-inheritance.asm.frag @@ -44,10 +44,10 @@ struct spvUnsafeArray } }; -constant spvUnsafeArray _34 = spvUnsafeArray({ float4(0.0), float4(0.0) }); - constant float4 _32 = {}; +constant spvUnsafeArray _34 = spvUnsafeArray({ float4(0.0), float4(0.0) }); + struct main0_out { float4 FragColor [[color(0)]]; diff --git a/reference/shaders-msl-no-opt/asm/frag/phi.zero-initialize.asm.frag b/reference/shaders-msl-no-opt/asm/frag/phi.zero-initialize.asm.frag index 32188581c..cffd0bd1a 100644 --- a/reference/shaders-msl-no-opt/asm/frag/phi.zero-initialize.asm.frag +++ b/reference/shaders-msl-no-opt/asm/frag/phi.zero-initialize.asm.frag @@ -3,14 +3,15 @@ using namespace metal; +constant int uninit_int = {}; +constant int4 uninit_vector = {}; +constant float4x4 uninit_matrix = {}; + struct Foo { int a; }; -constant int uninit_int = {}; -constant int4 uninit_vector = {}; -constant float4x4 uninit_matrix = {}; constant Foo uninit_foo = {}; struct main0_out diff --git a/reference/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp b/reference/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp new file mode 100644 index 000000000..986e90966 --- /dev/null +++ b/reference/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp @@ -0,0 +1,23 @@ +#include +#include + +using namespace metal; + +struct _19 +{ +}; +struct _5 +{ + int _m0; + _19 _m1; + char _m2_pad[4]; + _19 _m2; + char _m3_pad[4]; + int _m3; +}; + +kernel void main0(device _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]]) +{ + _4 = _3; +} + diff --git a/reference/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp b/reference/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp new file mode 100644 index 000000000..4bcfeb21a --- /dev/null +++ b/reference/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp @@ -0,0 +1,24 @@ +#include +#include + +using namespace metal; + +struct _19 +{ +}; +struct _5 +{ + int _m0; + char _m1_pad[12]; + _19 _m1; + char _m2_pad[16]; + _19 _m2; + char _m3_pad[16]; + int _m3; +}; + +kernel void main0(constant _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]]) +{ + _4 = _3; +} + diff --git a/reference/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp b/reference/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp index 050dca4ab..42c13f7a4 100644 --- a/reference/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp +++ b/reference/shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp @@ -54,6 +54,7 @@ struct _7 int _m0[1]; }; +constant int3 _32 = {}; constant int _3_tmp [[function_constant(0)]]; constant int _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 0; constant int _4_tmp [[function_constant(1)]]; @@ -103,8 +104,6 @@ constant int _74 = (_73 * _72); constant spvUnsafeArray _33 = spvUnsafeArray({ 0, 0, 0 }); constant spvUnsafeArray, 3> _34 = spvUnsafeArray, 3>({ spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }), spvUnsafeArray({ 0, 0, 0 }) }); -constant int3 _32 = {}; - kernel void main0(device _7& _8 [[buffer(0)]], device _7& _9 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { _9._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + ((((1 - _59) * _60) * (_61 - 1)) * _74); diff --git a/reference/shaders-msl/asm/comp/uint_smulextended.asm.comp b/reference/shaders-msl/asm/comp/uint_smulextended.asm.comp new file mode 100644 index 000000000..6996f7fd2 --- /dev/null +++ b/reference/shaders-msl/asm/comp/uint_smulextended.asm.comp @@ -0,0 +1,25 @@ +#include +#include + +using namespace metal; + +struct _4 +{ + uint _m0[1]; +}; + +struct _20 +{ + uint _m0; + uint _m1; +}; + +kernel void main0(device _4& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]], device _4& _7 [[buffer(2)]], device _4& _8 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _20 _28; + _28._m0 = uint(int(_5._m0[gl_GlobalInvocationID.x]) * int(_6._m0[gl_GlobalInvocationID.x])); + _28._m1 = uint(mulhi(int(_5._m0[gl_GlobalInvocationID.x]), int(_6._m0[gl_GlobalInvocationID.x]))); + _7._m0[gl_GlobalInvocationID.x] = _28._m0; + _8._m0[gl_GlobalInvocationID.x] = _28._m1; +} + diff --git a/reference/shaders-msl/asm/comp/undefined-constant-composite.asm.comp b/reference/shaders-msl/asm/comp/undefined-constant-composite.asm.comp new file mode 100644 index 000000000..e3ded1f20 --- /dev/null +++ b/reference/shaders-msl/asm/comp/undefined-constant-composite.asm.comp @@ -0,0 +1,38 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct _20 +{ + int _m0; + int _m1; +}; + +struct _5 +{ + int _m0[10]; +}; + +struct _7 +{ + int _m0[10]; +}; + +constant int _28 = {}; + +static inline __attribute__((always_inline)) +int _39(thread const int& _41, thread const _20& _42) +{ + return _41 + _42._m1; +} + +kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + int _32 = _8._m0[gl_GlobalInvocationID.x]; + _20 _33 = _20{ _28, 200 }; + _6._m0[gl_GlobalInvocationID.x] = _39(_32, _33); +} + diff --git a/reference/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp b/reference/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp new file mode 100644 index 000000000..d0f2790d8 --- /dev/null +++ b/reference/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp @@ -0,0 +1,42 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct _21 +{ + int _m0; + int _m1; +}; + +struct _5 +{ + int _m0[10]; +}; + +struct _7 +{ + int _m0[10]; +}; + +constant int _29 = {}; +constant int _9_tmp [[function_constant(0)]]; +constant int _9 = is_function_constant_defined(_9_tmp) ? _9_tmp : 0; +constant _21 _30 = _21{ _9, _29 }; + +static inline __attribute__((always_inline)) +int _42(thread const int& _44, thread const _21& _45, thread const _21& _46) +{ + return (_44 + _45._m0) + _46._m1; +} + +kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + int _34 = _8._m0[gl_GlobalInvocationID.x]; + _21 _35 = _30; + _21 _36 = _21{ _29, 200 }; + _6._m0[gl_GlobalInvocationID.x] = _42(_34, _35, _36); +} + diff --git a/reference/shaders-no-opt/asm/comp/constant-composite-undef.asm.comp b/reference/shaders-no-opt/asm/comp/constant-composite-undef.asm.comp index 279dede11..d8f1f19b1 100644 --- a/reference/shaders-no-opt/asm/comp/constant-composite-undef.asm.comp +++ b/reference/shaders-no-opt/asm/comp/constant-composite-undef.asm.comp @@ -1,13 +1,13 @@ #version 450 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +float _15; + layout(binding = 0, std430) buffer Block { vec4 f; } block; -float _15; - void main() { block.f = vec4(0.100000001490116119384765625, 0.20000000298023223876953125, 0.300000011920928955078125, 0.0); diff --git a/reference/shaders-no-opt/asm/comp/ray-query-function-object.spv14.asm.vk.nocompat.comp.vk b/reference/shaders-no-opt/asm/comp/ray-query-function-object.spv14.asm.vk.nocompat.comp.vk index 04b9eca6f..ccbbc02ff 100644 --- a/reference/shaders-no-opt/asm/comp/ray-query-function-object.spv14.asm.vk.nocompat.comp.vk +++ b/reference/shaders-no-opt/asm/comp/ray-query-function-object.spv14.asm.vk.nocompat.comp.vk @@ -2,11 +2,11 @@ #extension GL_EXT_ray_query : require layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in; -layout(set = 0, binding = 0) uniform accelerationStructureEXT RTAS; - float _16; vec3 _17; +layout(set = 0, binding = 0) uniform accelerationStructureEXT RTAS; + void main() { rayQueryEXT _19; diff --git a/reference/shaders-no-opt/asm/frag/composite-insert-inheritance.asm.frag b/reference/shaders-no-opt/asm/frag/composite-insert-inheritance.asm.frag index f7e9f6f27..11c1f4ca3 100644 --- a/reference/shaders-no-opt/asm/frag/composite-insert-inheritance.asm.frag +++ b/reference/shaders-no-opt/asm/frag/composite-insert-inheritance.asm.frag @@ -2,11 +2,11 @@ precision mediump float; precision highp int; +vec4 _32; + layout(location = 0) in vec4 vInput; layout(location = 0) out vec4 FragColor; -vec4 _32; - void main() { vec4 _37 = vInput; diff --git a/reference/shaders-no-opt/asm/frag/early-conditional-return-switch.asm.frag b/reference/shaders-no-opt/asm/frag/early-conditional-return-switch.asm.frag index 34fffb85a..b03d5a4d7 100644 --- a/reference/shaders-no-opt/asm/frag/early-conditional-return-switch.asm.frag +++ b/reference/shaders-no-opt/asm/frag/early-conditional-return-switch.asm.frag @@ -1,5 +1,7 @@ #version 450 +vec4 _32; + layout(binding = 0, std140) uniform type_gCBuffarrayIndex { uint gArrayIndex; @@ -12,8 +14,6 @@ uniform sampler2D SPIRV_Cross_Combinedg_textureArray3SPIRV_Cross_DummySampler; layout(location = 0) out vec4 out_var_SV_TARGET; -vec4 _32; - void main() { vec4 _80; diff --git a/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag b/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag index 88b94d8c5..874bc6de1 100644 --- a/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag +++ b/reference/shaders-no-opt/asm/frag/inliner-dominator-inside-loop.asm.frag @@ -61,6 +61,15 @@ struct Params vec4 LqmatFarTilingFactor; }; +VertexOutput _121; +SurfaceInput _122; +vec2 _123; +vec4 _124; +Surface _125; +vec4 _192; +vec4 _219; +vec4 _297; + layout(binding = 0, std140) uniform CB0 { Globals CB0; @@ -86,15 +95,6 @@ layout(location = 7) in vec4 IN_PosLightSpace_Reflectance; layout(location = 8) in float IN_studIndex; layout(location = 0) out vec4 _entryPointOutput; -VertexOutput _121; -SurfaceInput _122; -vec2 _123; -vec4 _124; -Surface _125; -vec4 _192; -vec4 _219; -vec4 _297; - void main() { VertexOutput _128; diff --git a/reference/shaders-no-opt/asm/frag/phi.zero-initialize.asm.frag b/reference/shaders-no-opt/asm/frag/phi.zero-initialize.asm.frag index ce14c4616..59bac9945 100644 --- a/reference/shaders-no-opt/asm/frag/phi.zero-initialize.asm.frag +++ b/reference/shaders-no-opt/asm/frag/phi.zero-initialize.asm.frag @@ -1,18 +1,19 @@ #version 450 +int uninit_int = 0; +ivec4 uninit_vector = ivec4(0); +mat4 uninit_matrix = mat4(vec4(0.0), vec4(0.0), vec4(0.0), vec4(0.0)); + struct Foo { int a; }; +Foo uninit_foo = Foo(0); + layout(location = 0) in vec4 vColor; layout(location = 0) out vec4 FragColor; -int uninit_int = 0; -ivec4 uninit_vector = ivec4(0); -mat4 uninit_matrix = mat4(vec4(0.0), vec4(0.0), vec4(0.0), vec4(0.0)); -Foo uninit_foo = Foo(0); - void main() { int _39 = 0; diff --git a/reference/shaders-no-opt/asm/frag/switch-non-default-fallthrough-no-phi.asm.frag b/reference/shaders-no-opt/asm/frag/switch-non-default-fallthrough-no-phi.asm.frag index 34fce5648..331518096 100644 --- a/reference/shaders-no-opt/asm/frag/switch-non-default-fallthrough-no-phi.asm.frag +++ b/reference/shaders-no-opt/asm/frag/switch-non-default-fallthrough-no-phi.asm.frag @@ -12,12 +12,12 @@ struct _5 int _m1; }; -layout(location = 0) flat in int _2; -layout(location = 0) out int _3; - _4 _16; int _21; +layout(location = 0) flat in int _2; +layout(location = 0) out int _3; + void main() { bool _25 = false; diff --git a/reference/shaders-no-opt/asm/frag/switch-single-case-multiple-exit-cfg.asm.frag b/reference/shaders-no-opt/asm/frag/switch-single-case-multiple-exit-cfg.asm.frag index 2ec47169d..c9ddbe689 100644 --- a/reference/shaders-no-opt/asm/frag/switch-single-case-multiple-exit-cfg.asm.frag +++ b/reference/shaders-no-opt/asm/frag/switch-single-case-multiple-exit-cfg.asm.frag @@ -2,10 +2,10 @@ precision mediump float; precision highp int; -layout(location = 0) out highp vec4 _GLF_color; - vec2 _19; +layout(location = 0) out highp vec4 _GLF_color; + void main() { highp vec2 _30; diff --git a/reference/shaders-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag b/reference/shaders-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag index b32d18748..b6d3bc849 100644 --- a/reference/shaders-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag +++ b/reference/shaders-no-opt/asm/frag/vector-shuffle-undef-index.asm.frag @@ -1,10 +1,10 @@ #version 450 +vec4 undef; + layout(location = 0) out vec4 FragColor; layout(location = 0) in vec4 vFloat; -vec4 undef; - void main() { FragColor = vec4(undef.x, vFloat.y, 0.0, vFloat.w) + vec4(vFloat.z, vFloat.y, 0.0, vFloat.w); diff --git a/reference/shaders-ue4/asm/frag/global-constant-arrays.asm.frag b/reference/shaders-ue4/asm/frag/global-constant-arrays.asm.frag index a741c4342..88618a851 100644 --- a/reference/shaders-ue4/asm/frag/global-constant-arrays.asm.frag +++ b/reference/shaders-ue4/asm/frag/global-constant-arrays.asm.frag @@ -93,6 +93,8 @@ struct type_Globals float ExpandGamut; }; +constant float3 _391 = {}; + constant spvUnsafeArray _475 = spvUnsafeArray({ -4.0, -4.0, -3.1573765277862548828125, -0.485249996185302734375, 1.84773242473602294921875, 1.84773242473602294921875 }); constant spvUnsafeArray _476 = spvUnsafeArray({ -0.718548238277435302734375, 2.0810306072235107421875, 3.66812419891357421875, 4.0, 4.0, 4.0 }); constant spvUnsafeArray _479 = spvUnsafeArray({ -4.97062206268310546875, -3.0293781757354736328125, -2.1261999607086181640625, -1.5104999542236328125, -1.0578000545501708984375, -0.4668000042438507080078125, 0.11937999725341796875, 0.7088134288787841796875, 1.2911865711212158203125, 1.2911865711212158203125 }); @@ -100,8 +102,6 @@ constant spvUnsafeArray _480 = spvUnsafeArray({ 0.80891323 constant spvUnsafeArray _482 = spvUnsafeArray({ -2.3010299205780029296875, -2.3010299205780029296875, -1.9312000274658203125, -1.5204999446868896484375, -1.0578000545501708984375, -0.4668000042438507080078125, 0.11937999725341796875, 0.7088134288787841796875, 1.2911865711212158203125, 1.2911865711212158203125 }); constant spvUnsafeArray _483 = spvUnsafeArray({ 0.801995217800140380859375, 1.19800484180450439453125, 1.5943000316619873046875, 1.99730002880096435546875, 2.3782999515533447265625, 2.7683999538421630859375, 3.0515000820159912109375, 3.2746293544769287109375, 3.32743072509765625, 3.32743072509765625 }); -constant float3 _391 = {}; - struct main0_out { float4 out_var_SV_Target0 [[color(0)]]; diff --git a/reference/shaders-ue4/asm/frag/padded-float-array-member-defef.asm.frag b/reference/shaders-ue4/asm/frag/padded-float-array-member-defef.asm.frag index 08bf76098..f405fbb68 100644 --- a/reference/shaders-ue4/asm/frag/padded-float-array-member-defef.asm.frag +++ b/reference/shaders-ue4/asm/frag/padded-float-array-member-defef.asm.frag @@ -94,6 +94,9 @@ struct type_Globals float ExpandGamut; }; +constant float3 _523 = {}; +constant float3 _525 = {}; + constant spvUnsafeArray _499 = spvUnsafeArray({ -4.0, -4.0, -3.1573765277862548828125, -0.485249996185302734375, 1.84773242473602294921875, 1.84773242473602294921875 }); constant spvUnsafeArray _500 = spvUnsafeArray({ -0.718548238277435302734375, 2.0810306072235107421875, 3.66812419891357421875, 4.0, 4.0, 4.0 }); constant spvUnsafeArray _503 = spvUnsafeArray({ -4.97062206268310546875, -3.0293781757354736328125, -2.1261999607086181640625, -1.5104999542236328125, -1.0578000545501708984375, -0.4668000042438507080078125, 0.11937999725341796875, 0.7088134288787841796875, 1.2911865711212158203125, 1.2911865711212158203125 }); @@ -101,9 +104,6 @@ constant spvUnsafeArray _504 = spvUnsafeArray({ 0.80891323 constant spvUnsafeArray _506 = spvUnsafeArray({ -2.3010299205780029296875, -2.3010299205780029296875, -1.9312000274658203125, -1.5204999446868896484375, -1.0578000545501708984375, -0.4668000042438507080078125, 0.11937999725341796875, 0.7088134288787841796875, 1.2911865711212158203125, 1.2911865711212158203125 }); constant spvUnsafeArray _507 = spvUnsafeArray({ 0.801995217800140380859375, 1.19800484180450439453125, 1.5943000316619873046875, 1.99730002880096435546875, 2.3782999515533447265625, 2.7683999538421630859375, 3.0515000820159912109375, 3.2746293544769287109375, 3.32743072509765625, 3.32743072509765625 }); -constant float3 _523 = {}; -constant float3 _525 = {}; - struct main0_out { float4 out_var_SV_Target0 [[color(0)]]; diff --git a/reference/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag b/reference/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag index 505165071..575c9ddb5 100644 --- a/reference/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag +++ b/reference/shaders-ue4/asm/frag/texture-atomics.asm.argument.msl2.frag @@ -56,6 +56,8 @@ struct type_Globals uint2 ShadowTileListGroupSize; }; +constant float3 _70 = {}; + struct spvDescriptorSetBuffer0 { const device type_StructuredBuffer_v4float* CulledObjectBoxBounds [[id(0)]]; @@ -64,8 +66,6 @@ struct spvDescriptorSetBuffer0 device atomic_uint* RWShadowTileNumCulledObjects_atomic [[id(3)]]; }; -constant float3 _70 = {}; - struct main0_out { float4 out_var_SV_Target0 [[color(0)]]; diff --git a/reference/shaders/asm/comp/undefined-constant-composite.asm.comp b/reference/shaders/asm/comp/undefined-constant-composite.asm.comp new file mode 100644 index 000000000..dae558dc8 --- /dev/null +++ b/reference/shaders/asm/comp/undefined-constant-composite.asm.comp @@ -0,0 +1,33 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct _20 +{ + int _m0; + int _m1; +}; + +int _28; + +layout(binding = 1, std430) buffer _5_6 +{ + int _m0[10]; +} _6; + +layout(binding = 0, std430) buffer _7_8 +{ + int _m0[10]; +} _8; + +int _39(int _41, _20 _42) +{ + return _41 + _42._m1; +} + +void main() +{ + int _32 = _8._m0[gl_GlobalInvocationID.x]; + _20 _33 = _20(_28, 200); + _6._m0[gl_GlobalInvocationID.x] = _39(_32, _33); +} + diff --git a/reference/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag b/reference/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag index b2473f4d0..7930ca3b4 100644 --- a/reference/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag +++ b/reference/shaders/asm/frag/struct-composite-extract-swizzle.asm.frag @@ -8,12 +8,12 @@ struct Foo float var2; }; +Foo _22; + layout(binding = 0) uniform mediump sampler2D uSampler; layout(location = 0) out vec4 FragColor; -Foo _22; - void main() { FragColor = texture(uSampler, vec2(_22.var1, _22.var2)); diff --git a/reference/shaders/asm/frag/undef-variable-store.asm.frag b/reference/shaders/asm/frag/undef-variable-store.asm.frag index 26ad568ad..a3b1290de 100644 --- a/reference/shaders/asm/frag/undef-variable-store.asm.frag +++ b/reference/shaders/asm/frag/undef-variable-store.asm.frag @@ -1,10 +1,10 @@ #version 450 -layout(location = 0) out vec4 _entryPointOutput; - vec4 _38; vec4 _47; +layout(location = 0) out vec4 _entryPointOutput; + void main() { vec4 _27; diff --git a/reference/shaders/asm/frag/unreachable.asm.frag b/reference/shaders/asm/frag/unreachable.asm.frag index 8bc88b9f0..beb8708e1 100644 --- a/reference/shaders/asm/frag/unreachable.asm.frag +++ b/reference/shaders/asm/frag/unreachable.asm.frag @@ -1,10 +1,10 @@ #version 450 +vec4 _21; + layout(location = 0) flat in int counter; layout(location = 0) out vec4 FragColor; -vec4 _21; - void main() { vec4 _24; diff --git a/reference/shaders/asm/frag/vector-shuffle-oom.asm.frag b/reference/shaders/asm/frag/vector-shuffle-oom.asm.frag index e84fb894f..97c9a2eb9 100644 --- a/reference/shaders/asm/frag/vector-shuffle-oom.asm.frag +++ b/reference/shaders/asm/frag/vector-shuffle-oom.asm.frag @@ -17,6 +17,8 @@ struct _28 vec4 _m0; }; +_28 _74; + layout(binding = 0, std140) uniform _6_7 { vec4 _m0; @@ -102,8 +104,6 @@ uniform sampler2D SPIRV_Cross_Combined_2; layout(location = 0) out vec4 _5; -_28 _74; - void main() { _28 _77; diff --git a/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp b/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp new file mode 100644 index 000000000..b01262f5b --- /dev/null +++ b/shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp @@ -0,0 +1,43 @@ +OpCapability Shader +OpExtension "SPV_KHR_storage_buffer_storage_class" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %main "main" %var_id +OpExecutionMode %main LocalSize 1 1 1 +OpDecorate %var_id BuiltIn GlobalInvocationId +OpDecorate %var_input Binding 0 +OpDecorate %var_input DescriptorSet 0 +OpDecorate %var_outdata Binding 1 +OpDecorate %var_outdata DescriptorSet 0 +OpMemberDecorate %type_container_struct 0 Offset 0 +OpMemberDecorate %type_container_struct 1 Offset 4 +OpMemberDecorate %type_container_struct 2 Offset 8 +OpMemberDecorate %type_container_struct 3 Offset 12 +OpDecorate %type_container_struct Block +%bool = OpTypeBool +%void = OpTypeVoid +%voidf = OpTypeFunction %void +%u32 = OpTypeInt 32 0 +%i32 = OpTypeInt 32 1 +%f32 = OpTypeFloat 32 +%uvec3 = OpTypeVector %u32 3 +%fvec3 = OpTypeVector %f32 3 +%uvec3ptr = OpTypePointer Input %uvec3 +%i32ptr = OpTypePointer Uniform %i32 +%f32ptr = OpTypePointer Uniform %f32 +%i32arr = OpTypeRuntimeArray %i32 +%f32arr = OpTypeRuntimeArray %f32 +%type_empty_struct = OpTypeStruct +%type_container_struct = OpTypeStruct %i32 %type_empty_struct %type_empty_struct %i32 +%type_container_struct_ubo_ptr = OpTypePointer Uniform %type_container_struct +%type_container_struct_ssbo_ptr = OpTypePointer StorageBuffer %type_container_struct +%var_id = OpVariable %uvec3ptr Input +%var_input = OpVariable %type_container_struct_ssbo_ptr StorageBuffer +%var_outdata = OpVariable %type_container_struct_ssbo_ptr StorageBuffer + +%main = OpFunction %void None %voidf +%label = OpLabel +%input_copy = OpCopyObject %type_container_struct_ssbo_ptr %var_input +%result = OpLoad %type_container_struct %input_copy +OpStore %var_outdata %result +OpReturn +OpFunctionEnd diff --git a/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp b/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp new file mode 100644 index 000000000..63df59ac3 --- /dev/null +++ b/shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp @@ -0,0 +1,43 @@ +OpCapability Shader +OpExtension "SPV_KHR_storage_buffer_storage_class" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %main "main" %var_id +OpExecutionMode %main LocalSize 1 1 1 +OpDecorate %var_id BuiltIn GlobalInvocationId +OpDecorate %var_input Binding 0 +OpDecorate %var_input DescriptorSet 0 +OpDecorate %var_outdata Binding 1 +OpDecorate %var_outdata DescriptorSet 0 +OpMemberDecorate %type_container_struct 0 Offset 0 +OpMemberDecorate %type_container_struct 1 Offset 16 +OpMemberDecorate %type_container_struct 2 Offset 32 +OpMemberDecorate %type_container_struct 3 Offset 48 +OpDecorate %type_container_struct Block +%bool = OpTypeBool +%void = OpTypeVoid +%voidf = OpTypeFunction %void +%u32 = OpTypeInt 32 0 +%i32 = OpTypeInt 32 1 +%f32 = OpTypeFloat 32 +%uvec3 = OpTypeVector %u32 3 +%fvec3 = OpTypeVector %f32 3 +%uvec3ptr = OpTypePointer Input %uvec3 +%i32ptr = OpTypePointer Uniform %i32 +%f32ptr = OpTypePointer Uniform %f32 +%i32arr = OpTypeRuntimeArray %i32 +%f32arr = OpTypeRuntimeArray %f32 +%type_empty_struct = OpTypeStruct +%type_container_struct = OpTypeStruct %i32 %type_empty_struct %type_empty_struct %i32 +%type_container_struct_ubo_ptr = OpTypePointer Uniform %type_container_struct +%type_container_struct_ssbo_ptr = OpTypePointer StorageBuffer %type_container_struct +%var_id = OpVariable %uvec3ptr Input +%var_input = OpVariable %type_container_struct_ubo_ptr Uniform +%var_outdata = OpVariable %type_container_struct_ssbo_ptr StorageBuffer + +%main = OpFunction %void None %voidf +%label = OpLabel +%input_copy = OpCopyObject %type_container_struct_ubo_ptr %var_input +%result = OpLoad %type_container_struct %input_copy +OpStore %var_outdata %result +OpReturn +OpFunctionEnd diff --git a/shaders-msl/asm/comp/uint_smulextended.asm.comp b/shaders-msl/asm/comp/uint_smulextended.asm.comp new file mode 100644 index 000000000..32d483636 --- /dev/null +++ b/shaders-msl/asm/comp/uint_smulextended.asm.comp @@ -0,0 +1,61 @@ + OpCapability Shader + + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationId + OpExecutionMode %main LocalSize 1 1 1 + + OpDecorate %gl_GlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %ra_uint ArrayStride 4 + OpDecorate %struct_uint4 BufferBlock + OpMemberDecorate %struct_uint4 0 Offset 0 + OpDecorate %input0 DescriptorSet 0 + OpDecorate %input0 Binding 0 + OpDecorate %input1 DescriptorSet 0 + OpDecorate %input1 Binding 1 + OpDecorate %output0 DescriptorSet 0 + OpDecorate %output0 Binding 2 + OpDecorate %output1 DescriptorSet 0 + OpDecorate %output1 Binding 3 + + %uint = OpTypeInt 32 0 + %ptr_uint = OpTypePointer Uniform %uint + %ptr_input_uint = OpTypePointer Input %uint + %uint3 = OpTypeVector %uint 3 + %ptr_input_uint3 = OpTypePointer Input %uint3 + %void = OpTypeVoid + %voidFn = OpTypeFunction %void + + %uint_0 = OpConstant %uint 0 + %uint_1 = OpConstant %uint 1 + %ra_uint = OpTypeRuntimeArray %uint + %uint4 = OpTypeVector %uint 4 + %struct_uint4 = OpTypeStruct %ra_uint + %ptr_struct_uint4 = OpTypePointer Uniform %struct_uint4 + %resulttype = OpTypeStruct %uint %uint +%gl_GlobalInvocationId = OpVariable %ptr_input_uint3 Input + %input0 = OpVariable %ptr_struct_uint4 Uniform + %input1 = OpVariable %ptr_struct_uint4 Uniform + + %output0 = OpVariable %ptr_struct_uint4 Uniform + %output1 = OpVariable %ptr_struct_uint4 Uniform + + %main = OpFunction %void None %voidFn + %mainStart = OpLabel + %index_ptr = OpAccessChain %ptr_input_uint %gl_GlobalInvocationId %uint_0 + %index = OpLoad %uint %index_ptr + %in_ptr0 = OpAccessChain %ptr_uint %input0 %uint_0 %index + %invalue0 = OpLoad %uint %in_ptr0 + %in_ptr1 = OpAccessChain %ptr_uint %input1 %uint_0 %index + %invalue1 = OpLoad %uint %in_ptr1 + + %outvalue = OpSMulExtended %resulttype %invalue0 %invalue1 + %outvalue0 = OpCompositeExtract %uint %outvalue 0 + %out_ptr0 = OpAccessChain %ptr_uint %output0 %uint_0 %index + OpStore %out_ptr0 %outvalue0 + %outvalue1 = OpCompositeExtract %uint %outvalue 1 + %out_ptr1 = OpAccessChain %ptr_uint %output1 %uint_0 %index + OpStore %out_ptr1 %outvalue1 + + + OpReturn + OpFunctionEnd diff --git a/shaders-msl/asm/comp/undefined-constant-composite.asm.comp b/shaders-msl/asm/comp/undefined-constant-composite.asm.comp new file mode 100644 index 000000000..9de0501fe --- /dev/null +++ b/shaders-msl/asm/comp/undefined-constant-composite.asm.comp @@ -0,0 +1,102 @@ +; +; The shader below is based on the following GLSL shader: +; +; #version 450 +; +; struct Pair { +; int first; +; int second; +; }; +; +; const Pair constant_pair = { 100, 200 }; +; +; layout(set=0, binding=0, std430) buffer InputBlock { +; int array[10]; +; } inputValues; +; +; layout(set=0, binding=1, std430) buffer OutputBlock { +; int array[10]; +; } outputValues; +; +; int add_second (int value, Pair pair) { +; return value + pair.second; +; } +; +; void main() { +; uint idx = gl_GlobalInvocationID.x; +; outputValues.array[idx] = add_second(inputValues.array[idx], constant_pair); +; } +; +; However, the first element of constant_pair has been modified to be undefined. +; + OpCapability Shader + %std450 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID + OpExecutionMode %main LocalSize 1 1 1 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %_arr_int_uint_10 ArrayStride 4 + OpMemberDecorate %OutputBlock 0 Offset 0 + OpDecorate %OutputBlock BufferBlock + OpDecorate %outputValues DescriptorSet 0 + OpDecorate %outputValues Binding 1 + OpMemberDecorate %InputBlock 0 Offset 0 + OpDecorate %InputBlock BufferBlock + OpDecorate %inputValues DescriptorSet 0 + OpDecorate %inputValues Binding 0 + %void = OpTypeVoid + %void_func = OpTypeFunction %void + %int = OpTypeInt 32 1 + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %int_0 = OpConstant %int 0 + %int_1 = OpConstant %int 1 + %int_200 = OpConstant %int 200 + %uint_0 = OpConstant %uint 0 + %uint_10 = OpConstant %uint 10 + %_ptr_Function_int = OpTypePointer Function %int + %Pair = OpTypeStruct %int %int + %_ptr_Function_Pair = OpTypePointer Function %Pair + %add_second_func_type = OpTypeFunction %int %_ptr_Function_int %_ptr_Function_Pair + %_ptr_Function_uint = OpTypePointer Function %uint + %_ptr_Input_v3uint = OpTypePointer Input %v3uint + %_ptr_Input_uint = OpTypePointer Input %uint + %_arr_int_uint_10 = OpTypeArray %int %uint_10 + %OutputBlock = OpTypeStruct %_arr_int_uint_10 +%_ptr_Uniform_OutputBlock = OpTypePointer Uniform %OutputBlock + %outputValues = OpVariable %_ptr_Uniform_OutputBlock Uniform + %InputBlock = OpTypeStruct %_arr_int_uint_10 + %_ptr_Uniform_InputBlock = OpTypePointer Uniform %InputBlock + %inputValues = OpVariable %_ptr_Uniform_InputBlock Uniform + ; Replaced %int_100 with an undefined int. + %undef_int = OpUndef %int + ; Composed a constant Pair with the undefined int in the first member. + %const_Pair = OpConstantComposite %Pair %undef_int %int_200 + %_ptr_Uniform_int = OpTypePointer Uniform %int + %gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %main = OpFunction %void None %void_func + %main_label = OpLabel + %param_1 = OpVariable %_ptr_Function_int Function + %param_2 = OpVariable %_ptr_Function_Pair Function + %gidx_ptr = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %gidx = OpLoad %uint %gidx_ptr + %input_value_ptr = OpAccessChain %_ptr_Uniform_int %inputValues %int_0 %gidx + %input_value = OpLoad %int %input_value_ptr + OpStore %param_1 %input_value + OpStore %param_2 %const_Pair + %retval = OpFunctionCall %int %add_second %param_1 %param_2 + %output_value_ptr = OpAccessChain %_ptr_Uniform_int %outputValues %int_0 %gidx + OpStore %output_value_ptr %retval + OpReturn + OpFunctionEnd + %add_second = OpFunction %int None %add_second_func_type + %value_ptr = OpFunctionParameter %_ptr_Function_int + %pair = OpFunctionParameter %_ptr_Function_Pair + %add_second_label = OpLabel + %value = OpLoad %int %value_ptr + ; Access the second struct member, which is defined. + %pair_second_ptr = OpAccessChain %_ptr_Function_int %pair %int_1 + %pair_second = OpLoad %int %pair_second_ptr + %add_result = OpIAdd %int %value %pair_second + OpReturnValue %add_result + OpFunctionEnd diff --git a/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp b/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp new file mode 100644 index 000000000..d89a402bf --- /dev/null +++ b/shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp @@ -0,0 +1,122 @@ +; +; The shader below is based on the following GLSL shader: +; +; #version 450 +; +; struct Pair { +; int first; +; int second; +; }; +; +; const Pair constant_pair = { 100, 200 }; +; +; layout (constant_id=0) const int constantFirst = 0; +; +; Pair spec_constant_pair = { constantFirst, 200 }; +; +; layout(set=0, binding=0, std430) buffer InputBlock { +; int array[10]; +; } inputValues; +; +; layout(set=0, binding=1, std430) buffer OutputBlock { +; int array[10]; +; } outputValues; +; +; int add_first_and_second (int value, Pair p1, Pair p2) { +; return value + p1.first + p2.second; +; } +; +; void main() { +; uint idx = gl_GlobalInvocationID.x; +; outputValues.array[idx] = add_first_and_second(inputValues.array[idx], spec_constant_pair, constant_pair); +; } +; +; However, both the constant_pair and the spec_constant_pair have one of their members replaced by undefined values. +; + OpCapability Shader + %std450 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID + OpExecutionMode %main LocalSize 1 1 1 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %_arr_int_uint_10 ArrayStride 4 + OpMemberDecorate %OutputBlock 0 Offset 0 + OpDecorate %OutputBlock BufferBlock + OpDecorate %outputValues DescriptorSet 0 + OpDecorate %outputValues Binding 1 + OpMemberDecorate %InputBlock 0 Offset 0 + OpDecorate %InputBlock BufferBlock + OpDecorate %inputValues DescriptorSet 0 + OpDecorate %inputValues Binding 0 + OpDecorate %spec_constant SpecId 0 + %void = OpTypeVoid + %void_func = OpTypeFunction %void + %int = OpTypeInt 32 1 + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %int_0 = OpConstant %int 0 + %int_1 = OpConstant %int 1 + %int_200 = OpConstant %int 200 + %uint_0 = OpConstant %uint 0 + %uint_10 = OpConstant %uint 10 + %_ptr_Function_int = OpTypePointer Function %int + %Pair = OpTypeStruct %int %int + %_ptr_Function_Pair = OpTypePointer Function %Pair +%add_pair_members_func_type = OpTypeFunction %int %_ptr_Function_int %_ptr_Function_Pair %_ptr_Function_Pair + %_ptr_Function_uint = OpTypePointer Function %uint + %_ptr_Input_v3uint = OpTypePointer Input %v3uint + %_ptr_Input_uint = OpTypePointer Input %uint + %_arr_int_uint_10 = OpTypeArray %int %uint_10 + %OutputBlock = OpTypeStruct %_arr_int_uint_10 + %_ptr_Uniform_OutputBlock = OpTypePointer Uniform %OutputBlock + %outputValues = OpVariable %_ptr_Uniform_OutputBlock Uniform + %InputBlock = OpTypeStruct %_arr_int_uint_10 + %_ptr_Uniform_InputBlock = OpTypePointer Uniform %InputBlock + %inputValues = OpVariable %_ptr_Uniform_InputBlock Uniform + ; Replaced %int_100 with an undefined int. + %undef_int = OpUndef %int + ; Composed a spec constant Pair with an undefined int in the second member. + %spec_constant = OpSpecConstant %int 0 + %spec_const_Pair = OpSpecConstantComposite %Pair %spec_constant %undef_int + ; Composed a constant Pair with the undefined int in the first member. + %const_Pair = OpConstantComposite %Pair %undef_int %int_200 + %_ptr_Uniform_int = OpTypePointer Uniform %int + %gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %main = OpFunction %void None %void_func + %main_label = OpLabel + %param_1 = OpVariable %_ptr_Function_int Function + %param_2 = OpVariable %_ptr_Function_Pair Function + %param_3 = OpVariable %_ptr_Function_Pair Function + %gidx_ptr = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %gidx = OpLoad %uint %gidx_ptr + %input_value_ptr = OpAccessChain %_ptr_Uniform_int %inputValues %int_0 %gidx + %input_value = OpLoad %int %input_value_ptr + OpStore %param_1 %input_value + OpStore %param_2 %spec_const_Pair + OpStore %param_3 %const_Pair + ; Pass the input value as the first argument. + ; Pass the specialization constant Pair as the second argument. + ; Pass the constant Pair as the third argument. + %retval = OpFunctionCall %int %add_pair_members %param_1 %param_2 %param_3 + %output_value_ptr = OpAccessChain %_ptr_Uniform_int %outputValues %int_0 %gidx + OpStore %output_value_ptr %retval + OpReturn + OpFunctionEnd + %add_pair_members = OpFunction %int None %add_pair_members_func_type + %value_ptr = OpFunctionParameter %_ptr_Function_int + %pair_1 = OpFunctionParameter %_ptr_Function_Pair + %pair_2 = OpFunctionParameter %_ptr_Function_Pair + %add_pair_members_label = OpLabel + %value = OpLoad %int %value_ptr + ; Access the first struct member from the first pair. + ; Access the second struct member from the second pair. + ; Both should be defined according to the function call above. + %pair_1_first_ptr = OpAccessChain %_ptr_Function_int %pair_1 %int_0 + %pair_2_second_ptr = OpAccessChain %_ptr_Function_int %pair_2 %int_1 + %pair_1_first = OpLoad %int %pair_1_first_ptr + %pair_2_second = OpLoad %int %pair_2_second_ptr + %partial_result = OpIAdd %int %value %pair_1_first + %final_result = OpIAdd %int %partial_result %pair_2_second + OpReturnValue %final_result + OpFunctionEnd + diff --git a/shaders/asm/comp/undefined-constant-composite.asm.comp b/shaders/asm/comp/undefined-constant-composite.asm.comp new file mode 100644 index 000000000..9de0501fe --- /dev/null +++ b/shaders/asm/comp/undefined-constant-composite.asm.comp @@ -0,0 +1,102 @@ +; +; The shader below is based on the following GLSL shader: +; +; #version 450 +; +; struct Pair { +; int first; +; int second; +; }; +; +; const Pair constant_pair = { 100, 200 }; +; +; layout(set=0, binding=0, std430) buffer InputBlock { +; int array[10]; +; } inputValues; +; +; layout(set=0, binding=1, std430) buffer OutputBlock { +; int array[10]; +; } outputValues; +; +; int add_second (int value, Pair pair) { +; return value + pair.second; +; } +; +; void main() { +; uint idx = gl_GlobalInvocationID.x; +; outputValues.array[idx] = add_second(inputValues.array[idx], constant_pair); +; } +; +; However, the first element of constant_pair has been modified to be undefined. +; + OpCapability Shader + %std450 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID + OpExecutionMode %main LocalSize 1 1 1 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %_arr_int_uint_10 ArrayStride 4 + OpMemberDecorate %OutputBlock 0 Offset 0 + OpDecorate %OutputBlock BufferBlock + OpDecorate %outputValues DescriptorSet 0 + OpDecorate %outputValues Binding 1 + OpMemberDecorate %InputBlock 0 Offset 0 + OpDecorate %InputBlock BufferBlock + OpDecorate %inputValues DescriptorSet 0 + OpDecorate %inputValues Binding 0 + %void = OpTypeVoid + %void_func = OpTypeFunction %void + %int = OpTypeInt 32 1 + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 + %int_0 = OpConstant %int 0 + %int_1 = OpConstant %int 1 + %int_200 = OpConstant %int 200 + %uint_0 = OpConstant %uint 0 + %uint_10 = OpConstant %uint 10 + %_ptr_Function_int = OpTypePointer Function %int + %Pair = OpTypeStruct %int %int + %_ptr_Function_Pair = OpTypePointer Function %Pair + %add_second_func_type = OpTypeFunction %int %_ptr_Function_int %_ptr_Function_Pair + %_ptr_Function_uint = OpTypePointer Function %uint + %_ptr_Input_v3uint = OpTypePointer Input %v3uint + %_ptr_Input_uint = OpTypePointer Input %uint + %_arr_int_uint_10 = OpTypeArray %int %uint_10 + %OutputBlock = OpTypeStruct %_arr_int_uint_10 +%_ptr_Uniform_OutputBlock = OpTypePointer Uniform %OutputBlock + %outputValues = OpVariable %_ptr_Uniform_OutputBlock Uniform + %InputBlock = OpTypeStruct %_arr_int_uint_10 + %_ptr_Uniform_InputBlock = OpTypePointer Uniform %InputBlock + %inputValues = OpVariable %_ptr_Uniform_InputBlock Uniform + ; Replaced %int_100 with an undefined int. + %undef_int = OpUndef %int + ; Composed a constant Pair with the undefined int in the first member. + %const_Pair = OpConstantComposite %Pair %undef_int %int_200 + %_ptr_Uniform_int = OpTypePointer Uniform %int + %gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %main = OpFunction %void None %void_func + %main_label = OpLabel + %param_1 = OpVariable %_ptr_Function_int Function + %param_2 = OpVariable %_ptr_Function_Pair Function + %gidx_ptr = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %gidx = OpLoad %uint %gidx_ptr + %input_value_ptr = OpAccessChain %_ptr_Uniform_int %inputValues %int_0 %gidx + %input_value = OpLoad %int %input_value_ptr + OpStore %param_1 %input_value + OpStore %param_2 %const_Pair + %retval = OpFunctionCall %int %add_second %param_1 %param_2 + %output_value_ptr = OpAccessChain %_ptr_Uniform_int %outputValues %int_0 %gidx + OpStore %output_value_ptr %retval + OpReturn + OpFunctionEnd + %add_second = OpFunction %int None %add_second_func_type + %value_ptr = OpFunctionParameter %_ptr_Function_int + %pair = OpFunctionParameter %_ptr_Function_Pair + %add_second_label = OpLabel + %value = OpLoad %int %value_ptr + ; Access the second struct member, which is defined. + %pair_second_ptr = OpAccessChain %_ptr_Function_int %pair %int_1 + %pair_second = OpLoad %int %pair_second_ptr + %add_result = OpIAdd %int %value %pair_second + OpReturnValue %add_result + OpFunctionEnd diff --git a/spirv_cpp.cpp b/spirv_cpp.cpp index 8d934d2c5..dd0a84c83 100644 --- a/spirv_cpp.cpp +++ b/spirv_cpp.cpp @@ -274,8 +274,6 @@ void CompilerCPP::emit_resources() if (emitted) statement(""); - declare_undefined_values(); - statement("inline void init(spirv_cross_shader& s)"); begin_scope(); statement(resource_type, "::init(s);"); diff --git a/spirv_cross_parsed_ir.cpp b/spirv_cross_parsed_ir.cpp index 3b4d6e5ca..8d1acf69f 100644 --- a/spirv_cross_parsed_ir.cpp +++ b/spirv_cross_parsed_ir.cpp @@ -66,7 +66,7 @@ ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT meta = std::move(other.meta); for (int i = 0; i < TypeCount; i++) ids_for_type[i] = std::move(other.ids_for_type[i]); - ids_for_constant_or_type = std::move(other.ids_for_constant_or_type); + ids_for_constant_undef_or_type = std::move(other.ids_for_constant_undef_or_type); ids_for_constant_or_variable = std::move(other.ids_for_constant_or_variable); declared_capabilities = std::move(other.declared_capabilities); declared_extensions = std::move(other.declared_extensions); @@ -102,7 +102,7 @@ ParsedIR &ParsedIR::operator=(const ParsedIR &other) meta = other.meta; for (int i = 0; i < TypeCount; i++) ids_for_type[i] = other.ids_for_type[i]; - ids_for_constant_or_type = other.ids_for_constant_or_type; + ids_for_constant_undef_or_type = other.ids_for_constant_undef_or_type; ids_for_constant_or_variable = other.ids_for_constant_or_variable; declared_capabilities = other.declared_capabilities; declared_extensions = other.declared_extensions; @@ -934,7 +934,7 @@ void ParsedIR::add_typed_id(Types type, ID id) { case TypeConstant: ids_for_constant_or_variable.push_back(id); - ids_for_constant_or_type.push_back(id); + ids_for_constant_undef_or_type.push_back(id); break; case TypeVariable: @@ -943,7 +943,8 @@ void ParsedIR::add_typed_id(Types type, ID id) case TypeType: case TypeConstantOp: - ids_for_constant_or_type.push_back(id); + case TypeUndef: + ids_for_constant_undef_or_type.push_back(id); break; default: diff --git a/spirv_cross_parsed_ir.hpp b/spirv_cross_parsed_ir.hpp index 138d9dd43..7f35c3815 100644 --- a/spirv_cross_parsed_ir.hpp +++ b/spirv_cross_parsed_ir.hpp @@ -74,8 +74,8 @@ class ParsedIR // Special purpose lists which contain a union of types. // This is needed so we can declare specialization constants and structs in an interleaved fashion, // among other things. - // Constants can be of struct type, and struct array sizes can use specialization constants. - SmallVector ids_for_constant_or_type; + // Constants can be undef or of struct type, and struct array sizes can use specialization constants. + SmallVector ids_for_constant_undef_or_type; SmallVector ids_for_constant_or_variable; // We need to keep track of the width the Ops that contains a type for the diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 4d2254f95..ea56a439c 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -3416,27 +3416,6 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo statement(""); } -void CompilerGLSL::declare_undefined_values() -{ - bool emitted = false; - ir.for_each_typed_id([&](uint32_t, const SPIRUndef &undef) { - auto &type = this->get(undef.basetype); - // OpUndef can be void for some reason ... - if (type.basetype == SPIRType::Void) - return; - - string initializer; - if (options.force_zero_initialized_variables && type_can_zero_initialize(type)) - initializer = join(" = ", to_zero_initialized_expression(undef.basetype)); - - statement(variable_decl(type, to_name(undef.self), undef.self), initializer, ";"); - emitted = true; - }); - - if (emitted) - statement(""); -} - bool CompilerGLSL::variable_is_lut(const SPIRVariable &var) const { bool statically_assigned = var.statically_assigned && var.static_expression != ID(0) && var.remapped_variable; @@ -3537,7 +3516,7 @@ void CompilerGLSL::emit_resources() // { auto loop_lock = ir.create_loop_hard_lock(); - for (auto &id_ : ir.ids_for_constant_or_type) + for (auto &id_ : ir.ids_for_constant_undef_or_type) { auto &id = ir.ids[id_]; @@ -3590,6 +3569,22 @@ void CompilerGLSL::emit_resources() emit_struct(*type); } } + else if (id.get_type() == TypeUndef) + { + auto &undef = id.get(); + auto &type = this->get(undef.basetype); + // OpUndef can be void for some reason ... + if (type.basetype == SPIRType::Void) + return; + + string initializer; + if (options.force_zero_initialized_variables && type_can_zero_initialize(type)) + initializer = join(" = ", to_zero_initialized_expression(undef.basetype)); + + // FIXME: If used in a constant, we must declare it as one. + statement(variable_decl(type, to_name(undef.self), undef.self), initializer, ";"); + emitted = true; + } } } @@ -3806,8 +3801,6 @@ void CompilerGLSL::emit_resources() if (emitted) statement(""); - - declare_undefined_values(); } void CompilerGLSL::emit_output_variable_initializer(const SPIRVariable &var) @@ -5342,6 +5335,10 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c, bool inside_bloc { res += constant_op_expression(*op); } + else if (maybe_get(elem) != nullptr) + { + res += to_name(elem); + } else { auto &subc = get(elem); @@ -10064,9 +10061,32 @@ bool CompilerGLSL::should_dereference(uint32_t id) if (auto *var = maybe_get(id)) return var->phi_variable; - // If id is an access chain, we should not dereference it. if (auto *expr = maybe_get(id)) - return !expr->access_chain; + { + // If id is an access chain, we should not dereference it. + if (expr->access_chain) + return false; + + // If id is a forwarded copy of a variable pointer, we should not dereference it. + SPIRVariable *var = nullptr; + while (expr->loaded_from && expression_is_forwarded(expr->self)) + { + auto &src_type = expression_type(expr->loaded_from); + // To be a copy, the pointer and its source expression must be the + // same type. Can't check type.self, because for some reason that's + // usually the base type with pointers stripped off. This check is + // complex enough that I've hoisted it out of the while condition. + if (src_type.pointer != type.pointer || src_type.pointer_depth != type.pointer || + src_type.parent_type != type.parent_type) + break; + if ((var = maybe_get(expr->loaded_from))) + break; + if (!(expr = maybe_get(expr->loaded_from))) + break; + } + + return !var || var->phi_variable; + } // Otherwise, we should dereference this pointer expression. return true; @@ -11663,7 +11683,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // RHS expression is immutable, so just forward it. // Copying these things really make no sense, but // seems to be allowed anyways. - auto &e = set(id, to_expression(rhs), result_type, true); + auto &e = emit_op(result_type, id, to_expression(rhs), true, true); if (pointer) { auto *var = maybe_get_backing_variable(rhs); @@ -16900,7 +16920,7 @@ void CompilerGLSL::reorder_type_alias() if (alias_itr < master_itr) { // Must also swap the type order for the constant-type joined array. - auto &joined_types = ir.ids_for_constant_or_type; + auto &joined_types = ir.ids_for_constant_undef_or_type; auto alt_alias_itr = find(begin(joined_types), end(joined_types), *alias_itr); auto alt_master_itr = find(begin(joined_types), end(joined_types), *master_itr); assert(alt_alias_itr != end(joined_types)); diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index 2ccb36cb6..7320daeda 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -937,8 +937,6 @@ class CompilerGLSL : public Compiler bool type_is_empty(const SPIRType &type); - virtual void declare_undefined_values(); - bool can_use_io_location(spv::StorageClass storage, bool block); const Instruction *get_next_instruction_in_block(const Instruction &instr); static uint32_t mask_relevant_memory_semantics(uint32_t semantics); diff --git a/spirv_hlsl.cpp b/spirv_hlsl.cpp index 48aabef44..b3ba58041 100644 --- a/spirv_hlsl.cpp +++ b/spirv_hlsl.cpp @@ -1409,7 +1409,7 @@ void CompilerHLSL::emit_specialization_constants_and_structs() }); auto loop_lock = ir.create_loop_hard_lock(); - for (auto &id_ : ir.ids_for_constant_or_type) + for (auto &id_ : ir.ids_for_constant_undef_or_type) { auto &id = ir.ids[id_]; @@ -1471,6 +1471,21 @@ void CompilerHLSL::emit_specialization_constants_and_structs() emit_struct(type); } } + else if (id.get_type() == TypeUndef) + { + auto &undef = id.get(); + auto &type = this->get(undef.basetype); + // OpUndef can be void for some reason ... + if (type.basetype == SPIRType::Void) + return; + + string initializer; + if (options.force_zero_initialized_variables && type_can_zero_initialize(type)) + initializer = join(" = ", to_zero_initialized_expression(undef.basetype)); + + statement("static ", variable_decl(type, to_name(undef.self), undef.self), initializer, ";"); + emitted = true; + } } if (emitted) @@ -1514,27 +1529,6 @@ void CompilerHLSL::replace_illegal_names() CompilerGLSL::replace_illegal_names(); } -void CompilerHLSL::declare_undefined_values() -{ - bool emitted = false; - ir.for_each_typed_id([&](uint32_t, const SPIRUndef &undef) { - auto &type = this->get(undef.basetype); - // OpUndef can be void for some reason ... - if (type.basetype == SPIRType::Void) - return; - - string initializer; - if (options.force_zero_initialized_variables && type_can_zero_initialize(type)) - initializer = join(" = ", to_zero_initialized_expression(undef.basetype)); - - statement("static ", variable_decl(type, to_name(undef.self), undef.self), initializer, ";"); - emitted = true; - }); - - if (emitted) - statement(""); -} - void CompilerHLSL::emit_resources() { auto &execution = get_entry_point(); @@ -1840,8 +1834,6 @@ void CompilerHLSL::emit_resources() if (emitted) statement(""); - declare_undefined_values(); - if (requires_op_fmod) { static const char *types[] = { diff --git a/spirv_hlsl.hpp b/spirv_hlsl.hpp index 773823630..57d1c2cdc 100644 --- a/spirv_hlsl.hpp +++ b/spirv_hlsl.hpp @@ -230,7 +230,6 @@ class CompilerHLSL : public CompilerGLSL void emit_hlsl_entry_point(); void emit_header() override; void emit_resources(); - void declare_undefined_values() override; void emit_interface_block_globally(const SPIRVariable &type); void emit_interface_block_in_struct(const SPIRVariable &var, std::unordered_set &active_locations); void emit_interface_block_member_in_struct(const SPIRVariable &var, uint32_t member_index, uint32_t location, diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 58090ebb5..ba7736901 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -7093,28 +7093,6 @@ static string inject_top_level_storage_qualifier(const string &expr, const strin } } -// Undefined global memory is not allowed in MSL. -// Declare constant and init to zeros. Use {}, as global constructors can break Metal. -void CompilerMSL::declare_undefined_values() -{ - bool emitted = false; - ir.for_each_typed_id([&](uint32_t, SPIRUndef &undef) { - auto &type = this->get(undef.basetype); - // OpUndef can be void for some reason ... - if (type.basetype == SPIRType::Void) - return; - - statement(inject_top_level_storage_qualifier( - variable_decl(type, to_name(undef.self), undef.self), - "constant"), - " = {};"); - emitted = true; - }); - - if (emitted) - statement(""); -} - void CompilerMSL::declare_constant_arrays() { bool fully_inlined = ir.ids_for_type[TypeFunction].size() == 1; @@ -7180,7 +7158,6 @@ void CompilerMSL::declare_complex_constant_arrays() void CompilerMSL::emit_resources() { declare_constant_arrays(); - declare_undefined_values(); // Emit the special [[stage_in]] and [[stage_out]] interface blocks which we created. emit_interface_block(stage_out_var_id); @@ -7243,7 +7220,7 @@ void CompilerMSL::emit_specialization_constants_and_structs() emitted = false; declared_structs.clear(); - for (auto &id_ : ir.ids_for_constant_or_type) + for (auto &id_ : ir.ids_for_constant_undef_or_type) { auto &id = ir.ids[id_]; @@ -7356,6 +7333,21 @@ void CompilerMSL::emit_specialization_constants_and_structs() emit_struct(get(type_id)); } } + else if (id.get_type() == TypeUndef) + { + auto &undef = id.get(); + auto &type = get(undef.basetype); + // OpUndef can be void for some reason ... + if (type.basetype == SPIRType::Void) + return; + + // Undefined global memory is not allowed in MSL. + // Declare constant and init to zeros. Use {}, as global constructors can break Metal. + statement( + inject_top_level_storage_qualifier(variable_decl(type, to_name(undef.self), undef.self), "constant"), + " = {};"); + emitted = true; + } } if (emitted) @@ -8937,12 +8929,33 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) uint32_t op0 = ops[2]; uint32_t op1 = ops[3]; auto &type = get(result_type); + auto input_type = opcode == OpSMulExtended ? int_type : uint_type; + auto &output_type = get_type(result_type); + string cast_op0, cast_op1; + + auto expected_type = binary_op_bitcast_helper(cast_op0, cast_op1, input_type, op0, op1, false); + emit_uninitialized_temporary_expression(result_type, result_id); - statement(to_expression(result_id), ".", to_member_name(type, 0), " = ", - to_enclosed_unpacked_expression(op0), " * ", to_enclosed_unpacked_expression(op1), ";"); - statement(to_expression(result_id), ".", to_member_name(type, 1), " = mulhi(", - to_unpacked_expression(op0), ", ", to_unpacked_expression(op1), ");"); + string mullo_expr, mulhi_expr; + mullo_expr = join(cast_op0, " * ", cast_op1); + mulhi_expr = join("mulhi(", cast_op0, ", ", cast_op1, ")"); + + auto &low_type = get_type(output_type.member_types[0]); + auto &high_type = get_type(output_type.member_types[1]); + if (low_type.basetype != input_type) + { + expected_type.basetype = input_type; + mullo_expr = join(bitcast_glsl_op(low_type, expected_type), "(", mullo_expr, ")"); + } + if (high_type.basetype != input_type) + { + expected_type.basetype = input_type; + mulhi_expr = join(bitcast_glsl_op(high_type, expected_type), "(", mulhi_expr, ")"); + } + + statement(to_expression(result_id), ".", to_member_name(type, 0), " = ", mullo_expr, ";"); + statement(to_expression(result_id), ".", to_member_name(type, 1), " = ", mulhi_expr, ";"); break; } diff --git a/spirv_msl.hpp b/spirv_msl.hpp index c7701b20a..2a165be57 100644 --- a/spirv_msl.hpp +++ b/spirv_msl.hpp @@ -820,7 +820,6 @@ class CompilerMSL : public CompilerGLSL std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain) override; std::string to_qualifiers_glsl(uint32_t id) override; void replace_illegal_names() override; - void declare_undefined_values() override; void declare_constant_arrays(); void replace_illegal_entry_point_names();