mirror of https://github.com/bkaradzic/bgfx
Updated spirv-cross.
This commit is contained in:
Бранимир Караџић
parent
65001ee075
commit
e12a9c48b1
25
3rdparty/spirv-cross/reference/opt/shaders-hlsl/frag/for-loop-continue-control-flow.frag
vendored
Normal file
25
3rdparty/spirv-cross/reference/opt/shaders-hlsl/frag/for-loop-continue-control-flow.frag
vendored
Normal file
|
|
@ -0,0 +1,25 @@
|
|||
static float4 FragColor;
|
||||
|
||||
struct SPIRV_Cross_Output
|
||||
{
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
FragColor = 0.0f.xxxx;
|
||||
for (int _43 = 0; _43 < 3; )
|
||||
{
|
||||
FragColor[_43] += float(_43);
|
||||
_43++;
|
||||
continue;
|
||||
}
|
||||
}
|
||||
|
||||
SPIRV_Cross_Output main()
|
||||
{
|
||||
frag_main();
|
||||
SPIRV_Cross_Output stage_output;
|
||||
stage_output.FragColor = FragColor;
|
||||
return stage_output;
|
||||
}
|
||||
|
|
@ -3,7 +3,7 @@
|
|||
|
||||
using namespace metal;
|
||||
|
||||
typedef float3x2 packed_float2x3;
|
||||
typedef packed_float2 packed_rm_float2x3[3];
|
||||
|
||||
struct S0
|
||||
{
|
||||
|
|
@ -58,7 +58,7 @@ struct SSBO1
|
|||
float3x2 m3;
|
||||
float2x2 m4;
|
||||
float2x2 m5[9];
|
||||
packed_float2x3 m6[4][2];
|
||||
packed_rm_float2x3 m6[4][2];
|
||||
char _m10_pad[8];
|
||||
float3x2 m7;
|
||||
char _m11_pad[8];
|
||||
|
|
@ -142,6 +142,6 @@ kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [
|
|||
ssbo_430.content.m3s[5].c = _60.m3s[5].c;
|
||||
ssbo_430.content.m3s[6].c = _60.m3s[6].c;
|
||||
ssbo_430.content.m3s[7].c = _60.m3s[7].c;
|
||||
ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
|
||||
ssbo_430.content.m1.a = ssbo_430.content.m3.a * float3x2(float2(ssbo_430.m6[1][1][0]), float2(ssbo_430.m6[1][1][1]), float2(ssbo_430.m6[1][1][2]));
|
||||
}
|
||||
|
||||
|
|
|
|||
23
3rdparty/spirv-cross/reference/opt/shaders-msl/frag/for-loop-continue-control-flow.frag
vendored
Normal file
23
3rdparty/spirv-cross/reference/opt/shaders-msl/frag/for-loop-continue-control-flow.frag
vendored
Normal file
|
|
@ -0,0 +1,23 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0()
|
||||
{
|
||||
main0_out out = {};
|
||||
out.FragColor = float4(0.0);
|
||||
for (int _43 = 0; _43 < 3; )
|
||||
{
|
||||
out.FragColor[_43] += float(_43);
|
||||
_43++;
|
||||
continue;
|
||||
}
|
||||
return out;
|
||||
}
|
||||
|
||||
|
|
@ -3,12 +3,12 @@
|
|||
|
||||
using namespace metal;
|
||||
|
||||
typedef float3x4 packed_float4x3;
|
||||
typedef packed_float4 packed_rm_float4x3[3];
|
||||
|
||||
struct _15
|
||||
{
|
||||
packed_float4x3 _m0;
|
||||
packed_float4x3 _m1;
|
||||
packed_rm_float4x3 _m0;
|
||||
packed_rm_float4x3 _m1;
|
||||
};
|
||||
|
||||
struct _42
|
||||
|
|
@ -41,7 +41,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant _15& _17 [[buffer(0)]]
|
|||
{
|
||||
main0_out out = {};
|
||||
float4 _70 = _44._m0 * float4(float3(_44._m3) + (in.m_25.xyz * (_44._m6 + _44._m7)), 1.0);
|
||||
out.m_72 = normalize(float4(in.m_25.xyz, 0.0) * _17._m1);
|
||||
out.m_72 = normalize(float4(in.m_25.xyz, 0.0) * float3x4(float4(_17._m1[0]), float4(_17._m1[1]), float4(_17._m1[2])));
|
||||
float4 _94 = _70;
|
||||
_94.y = -_70.y;
|
||||
out.gl_Position = _94;
|
||||
|
|
|
|||
|
|
@ -0,0 +1,171 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
typedef packed_float2 packed_float2x2[2];
|
||||
typedef packed_float2 packed_rm_float2x3[3];
|
||||
typedef packed_float3 packed_float2x3[2];
|
||||
typedef packed_float3 packed_rm_float3x2[2];
|
||||
|
||||
struct S0
|
||||
{
|
||||
packed_float2 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S1
|
||||
{
|
||||
packed_float3 a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S2
|
||||
{
|
||||
packed_float3 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S3
|
||||
{
|
||||
packed_float2 a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S4
|
||||
{
|
||||
float2 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 SSBO1
|
||||
{
|
||||
Content content;
|
||||
Content content1[2];
|
||||
Content content2;
|
||||
float2x2 m0;
|
||||
float2x2 m1;
|
||||
packed_float2x3 m2[4];
|
||||
float3x2 m3;
|
||||
float2x2 m4;
|
||||
float2x2 m5[9];
|
||||
packed_rm_float2x3 m6[4][2];
|
||||
float3x2 m7;
|
||||
float array[1];
|
||||
};
|
||||
|
||||
struct S0_1
|
||||
{
|
||||
float4 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S1_1
|
||||
{
|
||||
packed_float3 a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S2_1
|
||||
{
|
||||
float3 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S3_1
|
||||
{
|
||||
float2 a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S4_1
|
||||
{
|
||||
float2 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;
|
||||
char _m8_pad[12];
|
||||
/* FIXME: A padded struct is needed here. If you see this message, file a bug! */ S4_1 m3s[8];
|
||||
};
|
||||
|
||||
struct SSBO0
|
||||
{
|
||||
Content_1 content;
|
||||
Content_1 content1[2];
|
||||
Content_1 content2;
|
||||
float2x2 m0;
|
||||
char _m4_pad[16];
|
||||
float2x2 m1;
|
||||
char _m5_pad[16];
|
||||
float2x3 m2[4];
|
||||
float3x2 m3;
|
||||
char _m7_pad[24];
|
||||
float2x2 m4;
|
||||
char _m8_pad[16];
|
||||
float2x2 m5[9];
|
||||
float2x3 m6[4][2];
|
||||
float3x2 m7;
|
||||
float4 array[1];
|
||||
};
|
||||
|
||||
struct SSBO2
|
||||
{
|
||||
float m0;
|
||||
packed_float2x2 m1;
|
||||
packed_rm_float3x2 m2;
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
|
||||
{
|
||||
ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0].xy;
|
||||
ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].b;
|
||||
ssbo_scalar.content.m1s[0].a = float3(ssbo_140.content.m1s[0].a);
|
||||
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
||||
ssbo_scalar.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
|
||||
ssbo_scalar.content.m2s[0].b = ssbo_140.content.m2s[0].b;
|
||||
ssbo_scalar.content.m0.a[0] = ssbo_140.content.m0.a[0].xy;
|
||||
ssbo_scalar.content.m0.b = ssbo_140.content.m0.b;
|
||||
ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
|
||||
ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;
|
||||
ssbo_scalar.content.m2.a[0] = ssbo_140.content.m2.a[0];
|
||||
ssbo_scalar.content.m2.b = ssbo_140.content.m2.b;
|
||||
ssbo_scalar.content.m3.a = ssbo_140.content.m3.a;
|
||||
ssbo_scalar.content.m3.b = ssbo_140.content.m3.b;
|
||||
ssbo_scalar.content.m4 = ssbo_140.content.m4;
|
||||
ssbo_scalar.content.m3s[0].c = ssbo_140.content.m3s[0].c;
|
||||
ssbo_scalar.content.m3s[1].c = ssbo_140.content.m3s[1].c;
|
||||
ssbo_scalar.content.m3s[2].c = ssbo_140.content.m3s[2].c;
|
||||
ssbo_scalar.content.m3s[3].c = ssbo_140.content.m3s[3].c;
|
||||
ssbo_scalar.content.m3s[4].c = ssbo_140.content.m3s[4].c;
|
||||
ssbo_scalar.content.m3s[5].c = ssbo_140.content.m3s[5].c;
|
||||
ssbo_scalar.content.m3s[6].c = ssbo_140.content.m3s[6].c;
|
||||
ssbo_scalar.content.m3s[7].c = ssbo_140.content.m3s[7].c;
|
||||
ssbo_scalar.content.m1.a = float2x3(float3(ssbo_scalar.m2[1][0]), float3(ssbo_scalar.m2[1][1])) * float2(ssbo_scalar.content.m0.a[0]);
|
||||
ssbo_scalar.m0 = float2x2(float2(ssbo_scalar2.m1[0]), float2(ssbo_scalar2.m1[1]));
|
||||
ssbo_scalar2.m1[0] = transpose(ssbo_scalar.m4)[0];
|
||||
ssbo_scalar2.m1[1] = transpose(ssbo_scalar.m4)[1];
|
||||
ssbo_scalar2.m2[0] = spvConvertFromRowMajor3x2(ssbo_scalar.m3)[0];
|
||||
ssbo_scalar2.m2[1] = spvConvertFromRowMajor3x2(ssbo_scalar.m3)[1];
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,36 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct UBO
|
||||
{
|
||||
float a[1];
|
||||
float2 b[2];
|
||||
};
|
||||
|
||||
struct UBOEnhancedLayout
|
||||
{
|
||||
float c[1];
|
||||
float2 d[2];
|
||||
char _m2_pad[9976];
|
||||
float e;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
int vIndex [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], constant UBO& _17 [[buffer(0)]], constant UBOEnhancedLayout& _30 [[buffer(1)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
out.FragColor = (_17.a[in.vIndex] + _30.c[in.vIndex]) + _30.e;
|
||||
return out;
|
||||
}
|
||||
|
||||
18
3rdparty/spirv-cross/reference/opt/shaders/asm/extended-debug-extinst.invalid.asm.comp
vendored
Normal file
18
3rdparty/spirv-cross/reference/opt/shaders/asm/extended-debug-extinst.invalid.asm.comp
vendored
Normal file
|
|
@ -0,0 +1,18 @@
|
|||
#version 430
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
layout(binding = 0, std430) buffer _8_9
|
||||
{
|
||||
float _m0[];
|
||||
} _9;
|
||||
|
||||
layout(binding = 1, std430) buffer _8_10
|
||||
{
|
||||
float _m0[];
|
||||
} _10;
|
||||
|
||||
void main()
|
||||
{
|
||||
_10._m0[gl_GlobalInvocationID.x] = -_9._m0[gl_GlobalInvocationID.x];
|
||||
}
|
||||
|
||||
15
3rdparty/spirv-cross/reference/opt/shaders/frag/for-loop-continue-control-flow.frag
vendored
Normal file
15
3rdparty/spirv-cross/reference/opt/shaders/frag/for-loop-continue-control-flow.frag
vendored
Normal file
|
|
@ -0,0 +1,15 @@
|
|||
#version 450
|
||||
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = vec4(0.0);
|
||||
for (int _43 = 0; _43 < 3; )
|
||||
{
|
||||
FragColor[_43] += float(_43);
|
||||
_43++;
|
||||
continue;
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,52 @@
|
|||
struct UBO_1_1
|
||||
{
|
||||
float4 v[64];
|
||||
};
|
||||
|
||||
ConstantBuffer<UBO_1_1> ubos[] : register(b0, space2);
|
||||
ByteAddressBuffer ssbos[] : register(t0, space3);
|
||||
Texture2D<float4> uSamplers[] : register(t0, space0);
|
||||
SamplerState uSamps[] : register(s0, space1);
|
||||
Texture2D<float4> uCombinedSamplers[] : register(t4, space0);
|
||||
SamplerState _uCombinedSamplers_sampler[] : register(s4, space0);
|
||||
|
||||
static int vIndex;
|
||||
static float4 FragColor;
|
||||
static float2 vUV;
|
||||
|
||||
struct SPIRV_Cross_Input
|
||||
{
|
||||
nointerpolation int vIndex : TEXCOORD0;
|
||||
float2 vUV : TEXCOORD1;
|
||||
};
|
||||
|
||||
struct SPIRV_Cross_Output
|
||||
{
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
int i = vIndex;
|
||||
int _59 = i + 10;
|
||||
int _64 = i + 40;
|
||||
FragColor = uSamplers[NonUniformResourceIndex(_59)].Sample(uSamps[NonUniformResourceIndex(_64)], vUV);
|
||||
int _71 = i + 10;
|
||||
FragColor = uCombinedSamplers[NonUniformResourceIndex(_71)].Sample(_uCombinedSamplers_sampler[NonUniformResourceIndex(_71)], vUV);
|
||||
int _77 = i + 20;
|
||||
int _80 = i + 40;
|
||||
FragColor += ubos[NonUniformResourceIndex(_77)].v[_80];
|
||||
int _87 = i + 50;
|
||||
int _90 = i + 60;
|
||||
FragColor += asfloat(ssbos[NonUniformResourceIndex(_87)].Load4(_90 * 16 + 0));
|
||||
}
|
||||
|
||||
SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)
|
||||
{
|
||||
vIndex = stage_input.vIndex;
|
||||
vUV = stage_input.vUV;
|
||||
frag_main();
|
||||
SPIRV_Cross_Output stage_output;
|
||||
stage_output.FragColor = FragColor;
|
||||
return stage_output;
|
||||
}
|
||||
44
3rdparty/spirv-cross/reference/shaders-hlsl/frag/for-loop-continue-control-flow.frag
vendored
Normal file
44
3rdparty/spirv-cross/reference/shaders-hlsl/frag/for-loop-continue-control-flow.frag
vendored
Normal file
|
|
@ -0,0 +1,44 @@
|
|||
static float4 FragColor;
|
||||
|
||||
struct SPIRV_Cross_Output
|
||||
{
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
FragColor = 0.0f.xxxx;
|
||||
int i = 0;
|
||||
int _36;
|
||||
for (;;)
|
||||
{
|
||||
if (i < 3)
|
||||
{
|
||||
int a = i;
|
||||
FragColor[a] += float(i);
|
||||
if (false)
|
||||
{
|
||||
_36 = 1;
|
||||
}
|
||||
else
|
||||
{
|
||||
int _41 = i;
|
||||
i = _41 + 1;
|
||||
_36 = _41;
|
||||
}
|
||||
continue;
|
||||
}
|
||||
else
|
||||
{
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
SPIRV_Cross_Output main()
|
||||
{
|
||||
frag_main();
|
||||
SPIRV_Cross_Output stage_output;
|
||||
stage_output.FragColor = FragColor;
|
||||
return stage_output;
|
||||
}
|
||||
|
|
@ -73,17 +73,11 @@ kernel void main0(device BUF& o [[buffer(0)]])
|
|||
o.a = int(c[1][1][1]);
|
||||
float _43[2] = { o.b, o.c };
|
||||
float _48[2] = { o.b, o.b };
|
||||
float _49[2][2];
|
||||
spvArrayCopyFromStack1(_49[0], _43);
|
||||
spvArrayCopyFromStack1(_49[1], _48);
|
||||
float _49[2][2] = { { _43[0], _43[1] }, { _48[0], _48[1] } };
|
||||
float _54[2] = { o.c, o.c };
|
||||
float _59[2] = { o.c, o.b };
|
||||
float _60[2][2];
|
||||
spvArrayCopyFromStack1(_60[0], _54);
|
||||
spvArrayCopyFromStack1(_60[1], _59);
|
||||
float _61[2][2][2];
|
||||
spvArrayCopyFromStack2(_61[0], _49);
|
||||
spvArrayCopyFromStack2(_61[1], _60);
|
||||
float _60[2][2] = { { _54[0], _54[1] }, { _59[0], _59[1] } };
|
||||
float _61[2][2][2] = { { { _49[0][0], _49[0][1] }, { _49[1][0], _49[1][1] } }, { { _60[0][0], _60[0][1] }, { _60[1][0], _60[1][1] } } };
|
||||
float d[2][2][2];
|
||||
spvArrayCopyFromStack3(d, _61);
|
||||
float e[2][2][2];
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
|
||||
using namespace metal;
|
||||
|
||||
typedef float3x2 packed_float2x3;
|
||||
typedef packed_float2 packed_rm_float2x3[3];
|
||||
|
||||
struct S0
|
||||
{
|
||||
|
|
@ -58,7 +58,7 @@ struct SSBO1
|
|||
float3x2 m3;
|
||||
float2x2 m4;
|
||||
float2x2 m5[9];
|
||||
packed_float2x3 m6[4][2];
|
||||
packed_rm_float2x3 m6[4][2];
|
||||
char _m10_pad[8];
|
||||
float3x2 m7;
|
||||
char _m11_pad[8];
|
||||
|
|
@ -142,6 +142,6 @@ kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [
|
|||
ssbo_430.content.m3s[5].c = _60.m3s[5].c;
|
||||
ssbo_430.content.m3s[6].c = _60.m3s[6].c;
|
||||
ssbo_430.content.m3s[7].c = _60.m3s[7].c;
|
||||
ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
|
||||
ssbo_430.content.m1.a = ssbo_430.content.m3.a * float3x2(float2(ssbo_430.m6[1][1][0]), float2(ssbo_430.m6[1][1][1]), float2(ssbo_430.m6[1][1][2]));
|
||||
}
|
||||
|
||||
|
|
|
|||
42
3rdparty/spirv-cross/reference/shaders-msl/frag/for-loop-continue-control-flow.frag
vendored
Normal file
42
3rdparty/spirv-cross/reference/shaders-msl/frag/for-loop-continue-control-flow.frag
vendored
Normal file
|
|
@ -0,0 +1,42 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0()
|
||||
{
|
||||
main0_out out = {};
|
||||
out.FragColor = float4(0.0);
|
||||
int i = 0;
|
||||
int _36;
|
||||
for (;;)
|
||||
{
|
||||
if (i < 3)
|
||||
{
|
||||
int a = i;
|
||||
out.FragColor[a] += float(i);
|
||||
if (false)
|
||||
{
|
||||
_36 = 1;
|
||||
}
|
||||
else
|
||||
{
|
||||
int _41 = i;
|
||||
i = _41 + 1;
|
||||
_36 = _41;
|
||||
}
|
||||
continue;
|
||||
}
|
||||
else
|
||||
{
|
||||
break;
|
||||
}
|
||||
}
|
||||
return out;
|
||||
}
|
||||
|
||||
|
|
@ -3,12 +3,12 @@
|
|||
|
||||
using namespace metal;
|
||||
|
||||
typedef float3x4 packed_float4x3;
|
||||
typedef packed_float4 packed_rm_float4x3[3];
|
||||
|
||||
struct _15
|
||||
{
|
||||
packed_float4x3 _m0;
|
||||
packed_float4x3 _m1;
|
||||
packed_rm_float4x3 _m0;
|
||||
packed_rm_float4x3 _m1;
|
||||
};
|
||||
|
||||
struct _42
|
||||
|
|
@ -44,7 +44,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant _15& _17 [[buffer(0)]]
|
|||
float3 _13;
|
||||
do
|
||||
{
|
||||
_13 = normalize(float4(in.m_25.xyz, 0.0) * _17._m1);
|
||||
_13 = normalize(float4(in.m_25.xyz, 0.0) * float3x4(float4(_17._m1[0]), float4(_17._m1[1]), float4(_17._m1[2])));
|
||||
break;
|
||||
} while (false);
|
||||
float4 _39 = _44._m0 * float4(float3(_44._m3) + (in.m_25.xyz * (_44._m6 + _44._m7)), 1.0);
|
||||
|
|
|
|||
|
|
@ -0,0 +1,171 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
typedef packed_float2 packed_float2x2[2];
|
||||
typedef packed_float2 packed_rm_float2x3[3];
|
||||
typedef packed_float3 packed_float2x3[2];
|
||||
typedef packed_float3 packed_rm_float3x2[2];
|
||||
|
||||
struct S0
|
||||
{
|
||||
packed_float2 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S1
|
||||
{
|
||||
packed_float3 a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S2
|
||||
{
|
||||
packed_float3 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S3
|
||||
{
|
||||
packed_float2 a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S4
|
||||
{
|
||||
float2 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 SSBO1
|
||||
{
|
||||
Content content;
|
||||
Content content1[2];
|
||||
Content content2;
|
||||
float2x2 m0;
|
||||
float2x2 m1;
|
||||
packed_float2x3 m2[4];
|
||||
float3x2 m3;
|
||||
float2x2 m4;
|
||||
float2x2 m5[9];
|
||||
packed_rm_float2x3 m6[4][2];
|
||||
float3x2 m7;
|
||||
float array[1];
|
||||
};
|
||||
|
||||
struct S0_1
|
||||
{
|
||||
float4 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S1_1
|
||||
{
|
||||
packed_float3 a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S2_1
|
||||
{
|
||||
float3 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S3_1
|
||||
{
|
||||
float2 a;
|
||||
float b;
|
||||
};
|
||||
|
||||
struct S4_1
|
||||
{
|
||||
float2 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;
|
||||
char _m8_pad[12];
|
||||
/* FIXME: A padded struct is needed here. If you see this message, file a bug! */ S4_1 m3s[8];
|
||||
};
|
||||
|
||||
struct SSBO0
|
||||
{
|
||||
Content_1 content;
|
||||
Content_1 content1[2];
|
||||
Content_1 content2;
|
||||
float2x2 m0;
|
||||
char _m4_pad[16];
|
||||
float2x2 m1;
|
||||
char _m5_pad[16];
|
||||
float2x3 m2[4];
|
||||
float3x2 m3;
|
||||
char _m7_pad[24];
|
||||
float2x2 m4;
|
||||
char _m8_pad[16];
|
||||
float2x2 m5[9];
|
||||
float2x3 m6[4][2];
|
||||
float3x2 m7;
|
||||
float4 array[1];
|
||||
};
|
||||
|
||||
struct SSBO2
|
||||
{
|
||||
float m0;
|
||||
packed_float2x2 m1;
|
||||
packed_rm_float3x2 m2;
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
|
||||
{
|
||||
ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0].xy;
|
||||
ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].b;
|
||||
ssbo_scalar.content.m1s[0].a = float3(ssbo_140.content.m1s[0].a);
|
||||
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
||||
ssbo_scalar.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
|
||||
ssbo_scalar.content.m2s[0].b = ssbo_140.content.m2s[0].b;
|
||||
ssbo_scalar.content.m0.a[0] = ssbo_140.content.m0.a[0].xy;
|
||||
ssbo_scalar.content.m0.b = ssbo_140.content.m0.b;
|
||||
ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
|
||||
ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;
|
||||
ssbo_scalar.content.m2.a[0] = ssbo_140.content.m2.a[0];
|
||||
ssbo_scalar.content.m2.b = ssbo_140.content.m2.b;
|
||||
ssbo_scalar.content.m3.a = ssbo_140.content.m3.a;
|
||||
ssbo_scalar.content.m3.b = ssbo_140.content.m3.b;
|
||||
ssbo_scalar.content.m4 = ssbo_140.content.m4;
|
||||
ssbo_scalar.content.m3s[0].c = ssbo_140.content.m3s[0].c;
|
||||
ssbo_scalar.content.m3s[1].c = ssbo_140.content.m3s[1].c;
|
||||
ssbo_scalar.content.m3s[2].c = ssbo_140.content.m3s[2].c;
|
||||
ssbo_scalar.content.m3s[3].c = ssbo_140.content.m3s[3].c;
|
||||
ssbo_scalar.content.m3s[4].c = ssbo_140.content.m3s[4].c;
|
||||
ssbo_scalar.content.m3s[5].c = ssbo_140.content.m3s[5].c;
|
||||
ssbo_scalar.content.m3s[6].c = ssbo_140.content.m3s[6].c;
|
||||
ssbo_scalar.content.m3s[7].c = ssbo_140.content.m3s[7].c;
|
||||
ssbo_scalar.content.m1.a = float2x3(float3(ssbo_scalar.m2[1][0]), float3(ssbo_scalar.m2[1][1])) * float2(ssbo_scalar.content.m0.a[0]);
|
||||
ssbo_scalar.m0 = float2x2(float2(ssbo_scalar2.m1[0]), float2(ssbo_scalar2.m1[1]));
|
||||
ssbo_scalar2.m1[0] = transpose(ssbo_scalar.m4)[0];
|
||||
ssbo_scalar2.m1[1] = transpose(ssbo_scalar.m4)[1];
|
||||
ssbo_scalar2.m2[0] = spvConvertFromRowMajor3x2(ssbo_scalar.m3)[0];
|
||||
ssbo_scalar2.m2[1] = spvConvertFromRowMajor3x2(ssbo_scalar.m3)[1];
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,36 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct UBO
|
||||
{
|
||||
float a[1];
|
||||
float2 b[2];
|
||||
};
|
||||
|
||||
struct UBOEnhancedLayout
|
||||
{
|
||||
float c[1];
|
||||
float2 d[2];
|
||||
char _m2_pad[9976];
|
||||
float e;
|
||||
};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float FragColor [[color(0)]];
|
||||
};
|
||||
|
||||
struct main0_in
|
||||
{
|
||||
int vIndex [[user(locn0)]];
|
||||
};
|
||||
|
||||
fragment main0_out main0(main0_in in [[stage_in]], constant UBO& _17 [[buffer(0)]], constant UBOEnhancedLayout& _30 [[buffer(1)]])
|
||||
{
|
||||
main0_out out = {};
|
||||
out.FragColor = (_17.a[in.vIndex] + _30.c[in.vIndex]) + _30.e;
|
||||
return out;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,37 @@
|
|||
#version 450
|
||||
#extension GL_EXT_nonuniform_qualifier : require
|
||||
|
||||
layout(set = 0, binding = 2, std140) uniform UBO
|
||||
{
|
||||
vec4 v[64];
|
||||
} ubos[];
|
||||
|
||||
layout(set = 0, binding = 3, std430) readonly buffer SSBO
|
||||
{
|
||||
vec4 v[];
|
||||
} ssbos[];
|
||||
|
||||
layout(set = 0, binding = 0) uniform texture2D uSamplers[];
|
||||
layout(set = 0, binding = 1) uniform sampler uSamps[];
|
||||
layout(set = 0, binding = 4) uniform sampler2D uCombinedSamplers[];
|
||||
|
||||
layout(location = 0) flat in int vIndex;
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
layout(location = 1) in vec2 vUV;
|
||||
|
||||
void main()
|
||||
{
|
||||
int i = vIndex;
|
||||
int _59 = i + 10;
|
||||
int _64 = i + 40;
|
||||
FragColor = texture(sampler2D(uSamplers[nonuniformEXT(_59)], uSamps[nonuniformEXT(_64)]), vUV);
|
||||
int _71 = i + 10;
|
||||
FragColor = texture(uCombinedSamplers[nonuniformEXT(_71)], vUV);
|
||||
int _77 = i + 20;
|
||||
int _80 = i + 40;
|
||||
FragColor += ubos[nonuniformEXT(_77)].v[_80];
|
||||
int _87 = i + 50;
|
||||
int _90 = i + 60;
|
||||
FragColor += ssbos[nonuniformEXT(_87)].v[_90];
|
||||
}
|
||||
|
||||
18
3rdparty/spirv-cross/reference/shaders/asm/extended-debug-extinst.invalid.asm.comp
vendored
Normal file
18
3rdparty/spirv-cross/reference/shaders/asm/extended-debug-extinst.invalid.asm.comp
vendored
Normal file
|
|
@ -0,0 +1,18 @@
|
|||
#version 430
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
layout(binding = 0, std430) buffer _8_9
|
||||
{
|
||||
float _m0[];
|
||||
} _9;
|
||||
|
||||
layout(binding = 1, std430) buffer _8_10
|
||||
{
|
||||
float _m0[];
|
||||
} _10;
|
||||
|
||||
void main()
|
||||
{
|
||||
_10._m0[gl_GlobalInvocationID.x] = -_9._m0[gl_GlobalInvocationID.x];
|
||||
}
|
||||
|
||||
34
3rdparty/spirv-cross/reference/shaders/frag/for-loop-continue-control-flow.frag
vendored
Normal file
34
3rdparty/spirv-cross/reference/shaders/frag/for-loop-continue-control-flow.frag
vendored
Normal file
|
|
@ -0,0 +1,34 @@
|
|||
#version 450
|
||||
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = vec4(0.0);
|
||||
int i = 0;
|
||||
int _36;
|
||||
for (;;)
|
||||
{
|
||||
if (i < 3)
|
||||
{
|
||||
int a = i;
|
||||
FragColor[a] += float(i);
|
||||
if (false)
|
||||
{
|
||||
_36 = 1;
|
||||
}
|
||||
else
|
||||
{
|
||||
int _41 = i;
|
||||
i = _41 + 1;
|
||||
_36 = _41;
|
||||
}
|
||||
continue;
|
||||
}
|
||||
else
|
||||
{
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,159 @@
|
|||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 93
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability ShaderNonUniformEXT
|
||||
OpCapability RuntimeDescriptorArrayEXT
|
||||
OpCapability UniformBufferArrayNonUniformIndexingEXT
|
||||
OpCapability SampledImageArrayNonUniformIndexingEXT
|
||||
OpCapability StorageBufferArrayNonUniformIndexingEXT
|
||||
OpExtension "SPV_EXT_descriptor_indexing"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main" %vIndex %FragColor %vUV
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpSource GLSL 450
|
||||
OpSourceExtension "GL_EXT_nonuniform_qualifier"
|
||||
OpName %main "main"
|
||||
OpName %i "i"
|
||||
OpName %vIndex "vIndex"
|
||||
OpName %FragColor "FragColor"
|
||||
OpName %uSamplers "uSamplers"
|
||||
OpName %uSamps "uSamps"
|
||||
OpName %vUV "vUV"
|
||||
OpName %uCombinedSamplers "uCombinedSamplers"
|
||||
OpName %UBO "UBO"
|
||||
OpMemberName %UBO 0 "v"
|
||||
OpName %ubos "ubos"
|
||||
OpName %SSBO "SSBO"
|
||||
OpMemberName %SSBO 0 "v"
|
||||
OpName %ssbos "ssbos"
|
||||
OpDecorate %vIndex Flat
|
||||
OpDecorate %vIndex Location 0
|
||||
OpDecorate %FragColor Location 0
|
||||
OpDecorate %uSamplers DescriptorSet 0
|
||||
OpDecorate %uSamplers Binding 0
|
||||
|
||||
OpDecorate %sampled_image NonUniformEXT
|
||||
OpDecorate %combined_sampler NonUniformEXT
|
||||
OpDecorate %ubo_ptr_copy NonUniformEXT
|
||||
OpDecorate %ssbo_ptr_copy NonUniformEXT
|
||||
|
||||
OpDecorate %uSamps DescriptorSet 1
|
||||
OpDecorate %uSamps Binding 0
|
||||
OpDecorate %vUV Location 1
|
||||
OpDecorate %uCombinedSamplers DescriptorSet 0
|
||||
OpDecorate %uCombinedSamplers Binding 4
|
||||
OpDecorate %_arr_v4float_uint_64 ArrayStride 16
|
||||
OpMemberDecorate %UBO 0 Offset 0
|
||||
OpDecorate %UBO Block
|
||||
OpDecorate %ubos DescriptorSet 2
|
||||
OpDecorate %ubos Binding 0
|
||||
OpDecorate %_runtimearr_v4float ArrayStride 16
|
||||
OpMemberDecorate %SSBO 0 NonWritable
|
||||
OpMemberDecorate %SSBO 0 Offset 0
|
||||
OpDecorate %SSBO BufferBlock
|
||||
OpDecorate %ssbos DescriptorSet 3
|
||||
OpDecorate %ssbos Binding 0
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%_ptr_Input_int = OpTypePointer Input %int
|
||||
%vIndex = OpVariable %_ptr_Input_int Input
|
||||
%float = OpTypeFloat 32
|
||||
%v4float = OpTypeVector %float 4
|
||||
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||
%FragColor = OpVariable %_ptr_Output_v4float Output
|
||||
%16 = OpTypeImage %float 2D 0 0 0 1 Unknown
|
||||
%_runtimearr_16 = OpTypeRuntimeArray %16
|
||||
%_ptr_UniformConstant__runtimearr_16 = OpTypePointer UniformConstant %_runtimearr_16
|
||||
%uSamplers = OpVariable %_ptr_UniformConstant__runtimearr_16 UniformConstant
|
||||
%int_10 = OpConstant %int 10
|
||||
%_ptr_UniformConstant_16 = OpTypePointer UniformConstant %16
|
||||
%27 = OpTypeSampler
|
||||
%_runtimearr_27 = OpTypeRuntimeArray %27
|
||||
%_ptr_UniformConstant__runtimearr_27 = OpTypePointer UniformConstant %_runtimearr_27
|
||||
%uSamps = OpVariable %_ptr_UniformConstant__runtimearr_27 UniformConstant
|
||||
%int_40 = OpConstant %int 40
|
||||
%_ptr_UniformConstant_27 = OpTypePointer UniformConstant %27
|
||||
%38 = OpTypeSampledImage %16
|
||||
%v2float = OpTypeVector %float 2
|
||||
%_ptr_Input_v2float = OpTypePointer Input %v2float
|
||||
%vUV = OpVariable %_ptr_Input_v2float Input
|
||||
%_runtimearr_38 = OpTypeRuntimeArray %38
|
||||
%_ptr_UniformConstant__runtimearr_38 = OpTypePointer UniformConstant %_runtimearr_38
|
||||
%uCombinedSamplers = OpVariable %_ptr_UniformConstant__runtimearr_38 UniformConstant
|
||||
%_ptr_UniformConstant_38 = OpTypePointer UniformConstant %38
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_64 = OpConstant %uint 64
|
||||
%_arr_v4float_uint_64 = OpTypeArray %v4float %uint_64
|
||||
%UBO = OpTypeStruct %_arr_v4float_uint_64
|
||||
%_runtimearr_UBO = OpTypeRuntimeArray %UBO
|
||||
%_ptr_Uniform__runtimearr_UBO = OpTypePointer Uniform %_runtimearr_UBO
|
||||
%ubos = OpVariable %_ptr_Uniform__runtimearr_UBO Uniform
|
||||
%int_20 = OpConstant %int 20
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
|
||||
%_runtimearr_v4float = OpTypeRuntimeArray %v4float
|
||||
%SSBO = OpTypeStruct %_runtimearr_v4float
|
||||
%_runtimearr_SSBO = OpTypeRuntimeArray %SSBO
|
||||
%_ptr_Uniform__runtimearr_SSBO = OpTypePointer Uniform %_runtimearr_SSBO
|
||||
%ssbos = OpVariable %_ptr_Uniform__runtimearr_SSBO Uniform
|
||||
%int_50 = OpConstant %int 50
|
||||
%int_60 = OpConstant %int 60
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%i = OpVariable %_ptr_Function_int Function
|
||||
%11 = OpLoad %int %vIndex
|
||||
OpStore %i %11
|
||||
%20 = OpLoad %int %i
|
||||
%22 = OpIAdd %int %20 %int_10
|
||||
%23 = OpCopyObject %int %22
|
||||
%25 = OpAccessChain %_ptr_UniformConstant_16 %uSamplers %23
|
||||
%26 = OpLoad %16 %25
|
||||
%31 = OpLoad %int %i
|
||||
%33 = OpIAdd %int %31 %int_40
|
||||
%34 = OpCopyObject %int %33
|
||||
%36 = OpAccessChain %_ptr_UniformConstant_27 %uSamps %34
|
||||
%37 = OpLoad %27 %36
|
||||
%sampled_image = OpSampledImage %38 %26 %37
|
||||
%43 = OpLoad %v2float %vUV
|
||||
%44 = OpImageSampleImplicitLod %v4float %sampled_image %43
|
||||
OpStore %FragColor %44
|
||||
%48 = OpLoad %int %i
|
||||
%49 = OpIAdd %int %48 %int_10
|
||||
%50 = OpCopyObject %int %49
|
||||
%52 = OpAccessChain %_ptr_UniformConstant_38 %uCombinedSamplers %50
|
||||
%combined_sampler = OpLoad %38 %52
|
||||
%54 = OpLoad %v2float %vUV
|
||||
%55 = OpImageSampleImplicitLod %v4float %combined_sampler %54
|
||||
OpStore %FragColor %55
|
||||
%63 = OpLoad %int %i
|
||||
%65 = OpIAdd %int %63 %int_20
|
||||
%66 = OpCopyObject %int %65
|
||||
%68 = OpLoad %int %i
|
||||
%69 = OpIAdd %int %68 %int_40
|
||||
%70 = OpCopyObject %int %69
|
||||
%ubo_ptr = OpAccessChain %_ptr_Uniform_v4float %ubos %66 %int_0 %70
|
||||
%ubo_ptr_copy = OpCopyObject %_ptr_Uniform_v4float %ubo_ptr
|
||||
%73 = OpLoad %v4float %ubo_ptr_copy
|
||||
%74 = OpLoad %v4float %FragColor
|
||||
%75 = OpFAdd %v4float %74 %73
|
||||
OpStore %FragColor %75
|
||||
%81 = OpLoad %int %i
|
||||
%83 = OpIAdd %int %81 %int_50
|
||||
%84 = OpCopyObject %int %83
|
||||
%85 = OpLoad %int %i
|
||||
%87 = OpIAdd %int %85 %int_60
|
||||
%88 = OpCopyObject %int %87
|
||||
%ssbo_ptr = OpAccessChain %_ptr_Uniform_v4float %ssbos %84 %int_0 %88
|
||||
%ssbo_ptr_copy = OpCopyObject %_ptr_Uniform_v4float %ssbo_ptr
|
||||
%90 = OpLoad %v4float %ssbo_ptr_copy
|
||||
%91 = OpLoad %v4float %FragColor
|
||||
%92 = OpFAdd %v4float %91 %90
|
||||
OpStore %FragColor %92
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,11 @@
|
|||
#version 450
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
void main()
|
||||
{
|
||||
FragColor = vec4(0.0);
|
||||
for (int i = 0; i < 3; (0 > 1) ? 1 : i ++)
|
||||
{
|
||||
int a = i;
|
||||
FragColor[a] += float(i);
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,11 @@
|
|||
#version 450
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
void main()
|
||||
{
|
||||
FragColor = vec4(0.0);
|
||||
for (int i = 0; i < 3; (0 > 1) ? 1 : i ++)
|
||||
{
|
||||
int a = i;
|
||||
FragColor[a] += float(i);
|
||||
}
|
||||
}
|
||||
99
3rdparty/spirv-cross/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp
vendored
Normal file
99
3rdparty/spirv-cross/shaders-msl/vulkan/comp/struct-packing-scalar.nocompat.invalid.vk.comp
vendored
Normal file
|
|
@ -0,0 +1,99 @@
|
|||
#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 = 2, scalar) restrict buffer SSBO2
|
||||
{
|
||||
float m0;
|
||||
mat2 m1;
|
||||
layout(row_major) mat3x2 m2;
|
||||
} ssbo_scalar2;
|
||||
|
||||
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_scalar;
|
||||
|
||||
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_scalar.content = ssbo_140.content;
|
||||
ssbo_scalar.content.m1.a = ssbo_scalar.m2[1] * ssbo_scalar.content.m0.a[0]; // test packed matrix access
|
||||
ssbo_scalar.m0 = ssbo_scalar2.m1;
|
||||
ssbo_scalar2.m1 = ssbo_scalar.m4;
|
||||
ssbo_scalar2.m2 = ssbo_scalar.m3;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,23 @@
|
|||
#version 450
|
||||
#extension GL_EXT_scalar_block_layout : require
|
||||
|
||||
layout(std430, binding = 0) uniform UBO
|
||||
{
|
||||
float a[1];
|
||||
vec2 b[2];
|
||||
};
|
||||
|
||||
layout(std430, binding = 1) uniform UBOEnhancedLayout
|
||||
{
|
||||
float c[1];
|
||||
vec2 d[2];
|
||||
layout(offset = 10000) float e;
|
||||
};
|
||||
|
||||
layout(location = 0) flat in int vIndex;
|
||||
layout(location = 0) out float FragColor;
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = a[vIndex] + c[vIndex] + e;
|
||||
}
|
||||
159
3rdparty/spirv-cross/shaders-no-opt/asm/frag/nonuniform-qualifier-propagation.vk.nocompat.asm.frag
vendored
Normal file
159
3rdparty/spirv-cross/shaders-no-opt/asm/frag/nonuniform-qualifier-propagation.vk.nocompat.asm.frag
vendored
Normal file
|
|
@ -0,0 +1,159 @@
|
|||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 93
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability ShaderNonUniformEXT
|
||||
OpCapability RuntimeDescriptorArrayEXT
|
||||
OpCapability UniformBufferArrayNonUniformIndexingEXT
|
||||
OpCapability SampledImageArrayNonUniformIndexingEXT
|
||||
OpCapability StorageBufferArrayNonUniformIndexingEXT
|
||||
OpExtension "SPV_EXT_descriptor_indexing"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %main "main" %vIndex %FragColor %vUV
|
||||
OpExecutionMode %main OriginUpperLeft
|
||||
OpSource GLSL 450
|
||||
OpSourceExtension "GL_EXT_nonuniform_qualifier"
|
||||
OpName %main "main"
|
||||
OpName %i "i"
|
||||
OpName %vIndex "vIndex"
|
||||
OpName %FragColor "FragColor"
|
||||
OpName %uSamplers "uSamplers"
|
||||
OpName %uSamps "uSamps"
|
||||
OpName %vUV "vUV"
|
||||
OpName %uCombinedSamplers "uCombinedSamplers"
|
||||
OpName %UBO "UBO"
|
||||
OpMemberName %UBO 0 "v"
|
||||
OpName %ubos "ubos"
|
||||
OpName %SSBO "SSBO"
|
||||
OpMemberName %SSBO 0 "v"
|
||||
OpName %ssbos "ssbos"
|
||||
OpDecorate %vIndex Flat
|
||||
OpDecorate %vIndex Location 0
|
||||
OpDecorate %FragColor Location 0
|
||||
OpDecorate %uSamplers DescriptorSet 0
|
||||
OpDecorate %uSamplers Binding 0
|
||||
|
||||
OpDecorate %sampled_image NonUniformEXT
|
||||
OpDecorate %combined_sampler NonUniformEXT
|
||||
OpDecorate %ubo_ptr_copy NonUniformEXT
|
||||
OpDecorate %ssbo_ptr_copy NonUniformEXT
|
||||
|
||||
OpDecorate %uSamps DescriptorSet 0
|
||||
OpDecorate %uSamps Binding 1
|
||||
OpDecorate %vUV Location 1
|
||||
OpDecorate %uCombinedSamplers DescriptorSet 0
|
||||
OpDecorate %uCombinedSamplers Binding 4
|
||||
OpDecorate %_arr_v4float_uint_64 ArrayStride 16
|
||||
OpMemberDecorate %UBO 0 Offset 0
|
||||
OpDecorate %UBO Block
|
||||
OpDecorate %ubos DescriptorSet 0
|
||||
OpDecorate %ubos Binding 2
|
||||
OpDecorate %_runtimearr_v4float ArrayStride 16
|
||||
OpMemberDecorate %SSBO 0 NonWritable
|
||||
OpMemberDecorate %SSBO 0 Offset 0
|
||||
OpDecorate %SSBO BufferBlock
|
||||
OpDecorate %ssbos DescriptorSet 0
|
||||
OpDecorate %ssbos Binding 3
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%_ptr_Input_int = OpTypePointer Input %int
|
||||
%vIndex = OpVariable %_ptr_Input_int Input
|
||||
%float = OpTypeFloat 32
|
||||
%v4float = OpTypeVector %float 4
|
||||
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||
%FragColor = OpVariable %_ptr_Output_v4float Output
|
||||
%16 = OpTypeImage %float 2D 0 0 0 1 Unknown
|
||||
%_runtimearr_16 = OpTypeRuntimeArray %16
|
||||
%_ptr_UniformConstant__runtimearr_16 = OpTypePointer UniformConstant %_runtimearr_16
|
||||
%uSamplers = OpVariable %_ptr_UniformConstant__runtimearr_16 UniformConstant
|
||||
%int_10 = OpConstant %int 10
|
||||
%_ptr_UniformConstant_16 = OpTypePointer UniformConstant %16
|
||||
%27 = OpTypeSampler
|
||||
%_runtimearr_27 = OpTypeRuntimeArray %27
|
||||
%_ptr_UniformConstant__runtimearr_27 = OpTypePointer UniformConstant %_runtimearr_27
|
||||
%uSamps = OpVariable %_ptr_UniformConstant__runtimearr_27 UniformConstant
|
||||
%int_40 = OpConstant %int 40
|
||||
%_ptr_UniformConstant_27 = OpTypePointer UniformConstant %27
|
||||
%38 = OpTypeSampledImage %16
|
||||
%v2float = OpTypeVector %float 2
|
||||
%_ptr_Input_v2float = OpTypePointer Input %v2float
|
||||
%vUV = OpVariable %_ptr_Input_v2float Input
|
||||
%_runtimearr_38 = OpTypeRuntimeArray %38
|
||||
%_ptr_UniformConstant__runtimearr_38 = OpTypePointer UniformConstant %_runtimearr_38
|
||||
%uCombinedSamplers = OpVariable %_ptr_UniformConstant__runtimearr_38 UniformConstant
|
||||
%_ptr_UniformConstant_38 = OpTypePointer UniformConstant %38
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_64 = OpConstant %uint 64
|
||||
%_arr_v4float_uint_64 = OpTypeArray %v4float %uint_64
|
||||
%UBO = OpTypeStruct %_arr_v4float_uint_64
|
||||
%_runtimearr_UBO = OpTypeRuntimeArray %UBO
|
||||
%_ptr_Uniform__runtimearr_UBO = OpTypePointer Uniform %_runtimearr_UBO
|
||||
%ubos = OpVariable %_ptr_Uniform__runtimearr_UBO Uniform
|
||||
%int_20 = OpConstant %int 20
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
|
||||
%_runtimearr_v4float = OpTypeRuntimeArray %v4float
|
||||
%SSBO = OpTypeStruct %_runtimearr_v4float
|
||||
%_runtimearr_SSBO = OpTypeRuntimeArray %SSBO
|
||||
%_ptr_Uniform__runtimearr_SSBO = OpTypePointer Uniform %_runtimearr_SSBO
|
||||
%ssbos = OpVariable %_ptr_Uniform__runtimearr_SSBO Uniform
|
||||
%int_50 = OpConstant %int 50
|
||||
%int_60 = OpConstant %int 60
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%i = OpVariable %_ptr_Function_int Function
|
||||
%11 = OpLoad %int %vIndex
|
||||
OpStore %i %11
|
||||
%20 = OpLoad %int %i
|
||||
%22 = OpIAdd %int %20 %int_10
|
||||
%23 = OpCopyObject %int %22
|
||||
%25 = OpAccessChain %_ptr_UniformConstant_16 %uSamplers %23
|
||||
%26 = OpLoad %16 %25
|
||||
%31 = OpLoad %int %i
|
||||
%33 = OpIAdd %int %31 %int_40
|
||||
%34 = OpCopyObject %int %33
|
||||
%36 = OpAccessChain %_ptr_UniformConstant_27 %uSamps %34
|
||||
%37 = OpLoad %27 %36
|
||||
%sampled_image = OpSampledImage %38 %26 %37
|
||||
%43 = OpLoad %v2float %vUV
|
||||
%44 = OpImageSampleImplicitLod %v4float %sampled_image %43
|
||||
OpStore %FragColor %44
|
||||
%48 = OpLoad %int %i
|
||||
%49 = OpIAdd %int %48 %int_10
|
||||
%50 = OpCopyObject %int %49
|
||||
%52 = OpAccessChain %_ptr_UniformConstant_38 %uCombinedSamplers %50
|
||||
%combined_sampler = OpLoad %38 %52
|
||||
%54 = OpLoad %v2float %vUV
|
||||
%55 = OpImageSampleImplicitLod %v4float %combined_sampler %54
|
||||
OpStore %FragColor %55
|
||||
%63 = OpLoad %int %i
|
||||
%65 = OpIAdd %int %63 %int_20
|
||||
%66 = OpCopyObject %int %65
|
||||
%68 = OpLoad %int %i
|
||||
%69 = OpIAdd %int %68 %int_40
|
||||
%70 = OpCopyObject %int %69
|
||||
%ubo_ptr = OpAccessChain %_ptr_Uniform_v4float %ubos %66 %int_0 %70
|
||||
%ubo_ptr_copy = OpCopyObject %_ptr_Uniform_v4float %ubo_ptr
|
||||
%73 = OpLoad %v4float %ubo_ptr_copy
|
||||
%74 = OpLoad %v4float %FragColor
|
||||
%75 = OpFAdd %v4float %74 %73
|
||||
OpStore %FragColor %75
|
||||
%81 = OpLoad %int %i
|
||||
%83 = OpIAdd %int %81 %int_50
|
||||
%84 = OpCopyObject %int %83
|
||||
%85 = OpLoad %int %i
|
||||
%87 = OpIAdd %int %85 %int_60
|
||||
%88 = OpCopyObject %int %87
|
||||
%ssbo_ptr = OpAccessChain %_ptr_Uniform_v4float %ssbos %84 %int_0 %88
|
||||
%ssbo_ptr_copy = OpCopyObject %_ptr_Uniform_v4float %ssbo_ptr
|
||||
%90 = OpLoad %v4float %ssbo_ptr_copy
|
||||
%91 = OpLoad %v4float %FragColor
|
||||
%92 = OpFAdd %v4float %91 %90
|
||||
OpStore %FragColor %92
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,67 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Khronos SPIR-V Tools Assembler; 0
|
||||
; Bound: 37
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "DebugInfo"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %2 "main" %3
|
||||
OpExecutionMode %2 LocalSize 1 1 1
|
||||
%4 = OpString "negateInputs.comp"
|
||||
%5 = OpString "negateInputs"
|
||||
%6 = OpString "main"
|
||||
%7 = OpString ""
|
||||
OpSource GLSL 430
|
||||
OpName %2 "main"
|
||||
OpName %3 "gl_GlobalInvocationID"
|
||||
OpDecorate %3 BuiltIn GlobalInvocationId
|
||||
OpDecorate %8 BufferBlock
|
||||
OpDecorate %9 DescriptorSet 0
|
||||
OpDecorate %9 Binding 0
|
||||
OpDecorate %10 DescriptorSet 0
|
||||
OpDecorate %10 Binding 1
|
||||
OpDecorate %11 ArrayStride 4
|
||||
OpMemberDecorate %8 0 Offset 0
|
||||
OpLine %4 0 0
|
||||
%12 = OpTypeBool
|
||||
%13 = OpTypeVoid
|
||||
%14 = OpTypeFunction %13
|
||||
%15 = OpTypeInt 32 0
|
||||
%16 = OpTypeInt 32 1
|
||||
%17 = OpTypeFloat 32
|
||||
%18 = OpTypeVector %15 3
|
||||
%19 = OpTypeVector %17 3
|
||||
%20 = OpTypePointer Input %18
|
||||
%21 = OpTypePointer Uniform %16
|
||||
%22 = OpTypePointer Uniform %17
|
||||
%23 = OpTypeRuntimeArray %16
|
||||
%11 = OpTypeRuntimeArray %17
|
||||
%8 = OpTypeStruct %11
|
||||
%24 = OpTypePointer Uniform %8
|
||||
%9 = OpVariable %24 Uniform
|
||||
%10 = OpVariable %24 Uniform
|
||||
OpLine %4 0 1
|
||||
OpLine %5 1 0
|
||||
OpLine %4 1000 100000
|
||||
%3 = OpVariable %20 Input
|
||||
%25 = OpConstant %16 0
|
||||
OpNoLine
|
||||
OpLine %4 1 1
|
||||
%26 = OpExtInst %13 %1 DebugInfoNone
|
||||
%27 = OpExtInst %13 %1 DebugTypeFunction %13
|
||||
%28 = OpExtInst %13 %1 DebugFunction %6 %27 %4 1 1 %4 %7 FlagIsDefinition|FlagPrototyped|FlagIsOptimized 1 %26 %26
|
||||
%2 = OpFunction %13 None %14
|
||||
%29 = OpLabel
|
||||
%30 = OpExtInst %13 %1 DebugScope %28
|
||||
OpLine %4 1 1
|
||||
%31 = OpLoad %18 %3
|
||||
%32 = OpCompositeExtract %15 %31 0
|
||||
%33 = OpAccessChain %22 %9 %25 %32
|
||||
%34 = OpLoad %17 %33
|
||||
%35 = OpFNegate %17 %34
|
||||
%36 = OpAccessChain %22 %10 %25 %32
|
||||
OpStore %36 %35
|
||||
OpNoLine
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,11 @@
|
|||
#version 450
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
void main()
|
||||
{
|
||||
FragColor = vec4(0.0);
|
||||
for (int i = 0; i < 3; (0 > 1) ? 1 : i ++)
|
||||
{
|
||||
int a = i;
|
||||
FragColor[a] += float(i);
|
||||
}
|
||||
}
|
||||
|
|
@ -476,6 +476,7 @@ struct SPIRExtension : IVariant
|
|||
{
|
||||
Unsupported,
|
||||
GLSL,
|
||||
SPV_debug_info,
|
||||
SPV_AMD_shader_ballot,
|
||||
SPV_AMD_shader_explicit_vertex_parameter,
|
||||
SPV_AMD_shader_trinary_minmax,
|
||||
|
|
|
|||
|
|
@ -3079,6 +3079,7 @@ bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint3
|
|||
|
||||
case OpArrayLength:
|
||||
case OpLine:
|
||||
case OpNoLine:
|
||||
// Uses literals, but cannot be a phi variable or temporary, so ignore.
|
||||
break;
|
||||
|
||||
|
|
@ -4266,6 +4267,7 @@ bool Compiler::instruction_to_result_type(uint32_t &result_type, uint32_t &resul
|
|||
case OpGroupCommitReadPipe:
|
||||
case OpGroupCommitWritePipe:
|
||||
case OpLine:
|
||||
case OpNoLine:
|
||||
return false;
|
||||
|
||||
default:
|
||||
|
|
|
|||
|
|
@ -832,8 +832,6 @@ void CompilerGLSL::emit_struct(SPIRType &type)
|
|||
string CompilerGLSL::to_interpolation_qualifiers(const Bitset &flags)
|
||||
{
|
||||
string res;
|
||||
if (flags.get(DecorationNonUniformEXT))
|
||||
res += "nonuniformEXT ";
|
||||
//if (flags & (1ull << DecorationSmooth))
|
||||
// res += "smooth ";
|
||||
if (flags.get(DecorationFlat))
|
||||
|
|
@ -2888,6 +2886,49 @@ string CompilerGLSL::to_extract_component_expression(uint32_t id, uint32_t index
|
|||
return join(expr, ".", index_to_swizzle(index));
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_rerolled_array_expression(const string &base_expr, const SPIRType &type)
|
||||
{
|
||||
uint32_t size = to_array_size_literal(type);
|
||||
auto &parent = get<SPIRType>(type.parent_type);
|
||||
string expr = "{ ";
|
||||
|
||||
for (uint32_t i = 0; i < size; i++)
|
||||
{
|
||||
auto subexpr = join(base_expr, "[", convert_to_string(i), "]");
|
||||
if (parent.array.empty())
|
||||
expr += subexpr;
|
||||
else
|
||||
expr += to_rerolled_array_expression(subexpr, parent);
|
||||
|
||||
if (i + 1 < size)
|
||||
expr += ", ";
|
||||
}
|
||||
|
||||
expr += " }";
|
||||
return expr;
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_composite_constructor_expression(uint32_t id)
|
||||
{
|
||||
auto &type = expression_type(id);
|
||||
if (!backend.array_is_value_type && !type.array.empty())
|
||||
{
|
||||
// For this case, we need to "re-roll" an array initializer from a temporary.
|
||||
// We cannot simply pass the array directly, since it decays to a pointer and it cannot
|
||||
// participate in a struct initializer. E.g.
|
||||
// float arr[2] = { 1.0, 2.0 };
|
||||
// Foo foo = { arr }; must be transformed to
|
||||
// Foo foo = { { arr[0], arr[1] } };
|
||||
// The array sizes cannot be deduced from specialization constants since we cannot use any loops.
|
||||
|
||||
// We're only triggering one read of the array expression, but this is fine since arrays have to be declared
|
||||
// as temporaries anyways.
|
||||
return to_rerolled_array_expression(to_enclosed_expression(id), type);
|
||||
}
|
||||
else
|
||||
return to_expression(id);
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read)
|
||||
{
|
||||
auto itr = invalid_expressions.find(id);
|
||||
|
|
@ -4617,6 +4658,10 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
|
|||
|
||||
inherited_expressions.push_back(coord);
|
||||
|
||||
// Make sure non-uniform decoration is back-propagated to where it needs to be.
|
||||
if (has_decoration(img, DecorationNonUniformEXT))
|
||||
propagate_nonuniform_qualifier(img);
|
||||
|
||||
switch (op)
|
||||
{
|
||||
case OpImageSampleDrefImplicitLod:
|
||||
|
|
@ -5874,13 +5919,38 @@ string CompilerGLSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &i
|
|||
// Floating <-> Integer special casts. Just have to enumerate all cases. :(
|
||||
// 16-bit, 32-bit and 64-bit floats.
|
||||
if (out_type.basetype == SPIRType::UInt && in_type.basetype == SPIRType::Float)
|
||||
{
|
||||
if (is_legacy_es())
|
||||
SPIRV_CROSS_THROW("Float -> Uint bitcast not supported on legacy ESSL.");
|
||||
else if (!options.es && options.version < 330)
|
||||
require_extension_internal("GL_ARB_shader_bit_encoding");
|
||||
return "floatBitsToUint";
|
||||
}
|
||||
else if (out_type.basetype == SPIRType::Int && in_type.basetype == SPIRType::Float)
|
||||
{
|
||||
if (is_legacy_es())
|
||||
SPIRV_CROSS_THROW("Float -> Int bitcast not supported on legacy ESSL.");
|
||||
else if (!options.es && options.version < 330)
|
||||
require_extension_internal("GL_ARB_shader_bit_encoding");
|
||||
return "floatBitsToInt";
|
||||
}
|
||||
else if (out_type.basetype == SPIRType::Float && in_type.basetype == SPIRType::UInt)
|
||||
{
|
||||
if (is_legacy_es())
|
||||
SPIRV_CROSS_THROW("Uint -> Float bitcast not supported on legacy ESSL.");
|
||||
else if (!options.es && options.version < 330)
|
||||
require_extension_internal("GL_ARB_shader_bit_encoding");
|
||||
return "uintBitsToFloat";
|
||||
}
|
||||
else if (out_type.basetype == SPIRType::Float && in_type.basetype == SPIRType::Int)
|
||||
{
|
||||
if (is_legacy_es())
|
||||
SPIRV_CROSS_THROW("Int -> Float bitcast not supported on legacy ESSL.");
|
||||
else if (!options.es && options.version < 330)
|
||||
require_extension_internal("GL_ARB_shader_bit_encoding");
|
||||
return "intBitsToFloat";
|
||||
}
|
||||
|
||||
else if (out_type.basetype == SPIRType::Int64 && in_type.basetype == SPIRType::Double)
|
||||
return "doubleBitsToInt64";
|
||||
else if (out_type.basetype == SPIRType::UInt64 && in_type.basetype == SPIRType::Double)
|
||||
|
|
@ -7300,7 +7370,7 @@ string CompilerGLSL::build_composite_combiner(uint32_t return_type, const uint32
|
|||
|
||||
if (i)
|
||||
op += ", ";
|
||||
subop = to_expression(elems[i]);
|
||||
subop = to_composite_constructor_expression(elems[i]);
|
||||
}
|
||||
|
||||
base = e ? e->base_expression : 0;
|
||||
|
|
@ -7567,8 +7637,13 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
unroll_array_from_complex_load(id, ptr, expr);
|
||||
|
||||
auto &type = get<SPIRType>(result_type);
|
||||
if (has_decoration(id, DecorationNonUniformEXT))
|
||||
// Shouldn't need to check for ID, but current glslang codegen requires it in some cases
|
||||
// when loading Image/Sampler descriptors. It does not hurt to check ID as well.
|
||||
if (has_decoration(id, DecorationNonUniformEXT) || has_decoration(ptr, DecorationNonUniformEXT))
|
||||
{
|
||||
propagate_nonuniform_qualifier(ptr);
|
||||
convert_non_uniform_expression(type, expr);
|
||||
}
|
||||
|
||||
if (ptr_expression)
|
||||
ptr_expression->need_transpose = old_need_transpose;
|
||||
|
|
@ -7650,6 +7725,9 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
{
|
||||
auto *var = maybe_get<SPIRVariable>(ops[0]);
|
||||
|
||||
if (has_decoration(ops[0], DecorationNonUniformEXT))
|
||||
propagate_nonuniform_qualifier(ops[0]);
|
||||
|
||||
if (var && var->statically_assigned)
|
||||
var->static_expression = ops[1];
|
||||
else if (var && var->loop_variable && !var->loop_variable_enable)
|
||||
|
|
@ -7848,15 +7926,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
forward = false;
|
||||
|
||||
string constructor_op;
|
||||
if (!backend.array_is_value_type && out_type.array.size() > 1)
|
||||
{
|
||||
// We cannot construct array of arrays because we cannot treat the inputs
|
||||
// as value types. Need to declare the array-of-arrays, and copy in elements one by one.
|
||||
emit_uninitialized_temporary_expression(result_type, id);
|
||||
for (uint32_t i = 0; i < length; i++)
|
||||
emit_array_copy(join(to_expression(id), "[", i, "]"), elems[i]);
|
||||
}
|
||||
else if (backend.use_initializer_list && composite)
|
||||
if (backend.use_initializer_list && composite)
|
||||
{
|
||||
// Only use this path if we are building composites.
|
||||
// This path cannot be used for arithmetic.
|
||||
|
|
@ -8032,7 +8102,14 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
uint32_t rhs = ops[2];
|
||||
bool pointer = get<SPIRType>(result_type).pointer;
|
||||
|
||||
if (expression_is_lvalue(rhs) && !pointer)
|
||||
auto *chain = maybe_get<SPIRAccessChain>(rhs);
|
||||
if (chain)
|
||||
{
|
||||
// Cannot lower to a SPIRExpression, just copy the object.
|
||||
auto &e = set<SPIRAccessChain>(id, *chain);
|
||||
e.self = id;
|
||||
}
|
||||
else if (expression_is_lvalue(rhs) && !pointer)
|
||||
{
|
||||
// Need a copy.
|
||||
// For pointer types, we copy the pointer itself.
|
||||
|
|
@ -8051,6 +8128,11 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
auto *var = maybe_get_backing_variable(rhs);
|
||||
e.loaded_from = var ? var->self : 0;
|
||||
}
|
||||
|
||||
// If we're copying an access chain, need to inherit the read expressions.
|
||||
auto *rhs_expr = maybe_get<SPIRExpression>(rhs);
|
||||
if (rhs_expr)
|
||||
e.implied_read_expressions = rhs_expr->implied_read_expressions;
|
||||
}
|
||||
break;
|
||||
}
|
||||
|
|
@ -8985,6 +9067,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
uint32_t result_type = ops[0];
|
||||
uint32_t id = ops[1];
|
||||
emit_sampled_image_op(result_type, id, ops[2], ops[3]);
|
||||
inherit_expression_dependencies(id, ops[2]);
|
||||
inherit_expression_dependencies(id, ops[3]);
|
||||
break;
|
||||
}
|
||||
|
||||
|
|
@ -9406,6 +9490,10 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
{
|
||||
emit_spv_amd_gcn_shader_op(ops[0], ops[1], ops[3], &ops[4], length - 4);
|
||||
}
|
||||
else if (get<SPIRExtension>(extension_set).ext == SPIRExtension::SPV_debug_info)
|
||||
{
|
||||
break; // Ignore SPIR-V debug information extended instructions.
|
||||
}
|
||||
else
|
||||
{
|
||||
statement("// unimplemented ext op ", instruction.op);
|
||||
|
|
@ -9672,6 +9760,9 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
break;
|
||||
}
|
||||
|
||||
case OpNoLine:
|
||||
break;
|
||||
|
||||
default:
|
||||
statement("// unimplemented op ", instruction.op);
|
||||
break;
|
||||
|
|
@ -11219,7 +11310,10 @@ bool CompilerGLSL::attempt_emit_loop_header(SPIRBlock &block, SPIRBlock::Method
|
|||
}
|
||||
|
||||
default:
|
||||
SPIRV_CROSS_THROW("For/while loop detected, but need while/for loop semantics.");
|
||||
block.disable_block_optimization = true;
|
||||
force_recompile();
|
||||
begin_scope(); // We'll see an end_scope() later.
|
||||
return false;
|
||||
}
|
||||
|
||||
begin_scope();
|
||||
|
|
@ -11293,7 +11387,10 @@ bool CompilerGLSL::attempt_emit_loop_header(SPIRBlock &block, SPIRBlock::Method
|
|||
}
|
||||
|
||||
default:
|
||||
SPIRV_CROSS_THROW("For/while loop detected, but need while/for loop semantics.");
|
||||
block.disable_block_optimization = true;
|
||||
force_recompile();
|
||||
begin_scope(); // We'll see an end_scope() later.
|
||||
return false;
|
||||
}
|
||||
|
||||
begin_scope();
|
||||
|
|
@ -11885,6 +11982,11 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block)
|
|||
uint32_t var = block.dominated_variables[i];
|
||||
get<SPIRVariable>(var).deferred_declaration = rearm_dominated_variables[i];
|
||||
}
|
||||
|
||||
// Just like for deferred declaration, we need to forget about loop variable enable
|
||||
// if our block chain is reinstantiated later.
|
||||
for (auto &var_id : block.loop_variables)
|
||||
get<SPIRVariable>(var_id).loop_variable_enable = false;
|
||||
}
|
||||
|
||||
void CompilerGLSL::begin_scope()
|
||||
|
|
@ -12225,3 +12327,37 @@ void CompilerGLSL::emit_line_directive(uint32_t file_id, uint32_t line_literal)
|
|||
statement_no_indent("#line ", line_literal, " \"", get<SPIRString>(file_id).str, "\"");
|
||||
}
|
||||
}
|
||||
|
||||
void CompilerGLSL::propagate_nonuniform_qualifier(uint32_t id)
|
||||
{
|
||||
// SPIR-V might only tag the very last ID with NonUniformEXT, but for codegen,
|
||||
// we need to know NonUniformEXT a little earlier, when the resource is actually loaded.
|
||||
// Back-propagate the qualifier based on the expression dependency chain.
|
||||
|
||||
if (!has_decoration(id, DecorationNonUniformEXT))
|
||||
{
|
||||
set_decoration(id, DecorationNonUniformEXT);
|
||||
force_recompile();
|
||||
}
|
||||
|
||||
auto *e = maybe_get<SPIRExpression>(id);
|
||||
auto *combined = maybe_get<SPIRCombinedImageSampler>(id);
|
||||
auto *chain = maybe_get<SPIRAccessChain>(id);
|
||||
if (e)
|
||||
{
|
||||
for (auto &expr : e->expression_dependencies)
|
||||
propagate_nonuniform_qualifier(expr);
|
||||
for (auto &expr : e->implied_read_expressions)
|
||||
propagate_nonuniform_qualifier(expr);
|
||||
}
|
||||
else if (combined)
|
||||
{
|
||||
propagate_nonuniform_qualifier(combined->image);
|
||||
propagate_nonuniform_qualifier(combined->sampler);
|
||||
}
|
||||
else if (chain)
|
||||
{
|
||||
for (auto &expr : chain->implied_read_expressions)
|
||||
propagate_nonuniform_qualifier(expr);
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -502,6 +502,8 @@ protected:
|
|||
SPIRExpression &emit_uninitialized_temporary_expression(uint32_t type, uint32_t id);
|
||||
void append_global_func_args(const SPIRFunction &func, uint32_t index, SmallVector<std::string> &arglist);
|
||||
std::string to_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_composite_constructor_expression(uint32_t id);
|
||||
std::string to_rerolled_array_expression(const std::string &expr, const SPIRType &type);
|
||||
std::string to_enclosed_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_unpacked_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_enclosed_unpacked_expression(uint32_t id, bool register_expression_read = true);
|
||||
|
|
@ -672,6 +674,8 @@ protected:
|
|||
void fixup_type_alias();
|
||||
void reorder_type_alias();
|
||||
|
||||
void propagate_nonuniform_qualifier(uint32_t id);
|
||||
|
||||
private:
|
||||
void init();
|
||||
};
|
||||
|
|
|
|||
|
|
@ -2479,6 +2479,10 @@ void CompilerHLSL::emit_texture_op(const Instruction &i)
|
|||
|
||||
inherited_expressions.push_back(coord);
|
||||
|
||||
// Make sure non-uniform decoration is back-propagated to where it needs to be.
|
||||
if (has_decoration(img, DecorationNonUniformEXT))
|
||||
propagate_nonuniform_qualifier(img);
|
||||
|
||||
switch (op)
|
||||
{
|
||||
case OpImageSampleDrefImplicitLod:
|
||||
|
|
@ -3458,6 +3462,9 @@ void CompilerHLSL::emit_load(const Instruction &instruction)
|
|||
uint32_t id = ops[1];
|
||||
uint32_t ptr = ops[2];
|
||||
|
||||
if (has_decoration(ptr, DecorationNonUniformEXT))
|
||||
propagate_nonuniform_qualifier(ptr);
|
||||
|
||||
auto load_expr = read_access_chain(*chain);
|
||||
|
||||
bool forward = should_forward(ptr) && forced_temporaries.find(id) == end(forced_temporaries);
|
||||
|
|
@ -3491,6 +3498,9 @@ void CompilerHLSL::write_access_chain(const SPIRAccessChain &chain, uint32_t val
|
|||
// Make sure we trigger a read of the constituents in the access chain.
|
||||
track_expression_read(chain.self);
|
||||
|
||||
if (has_decoration(chain.self, DecorationNonUniformEXT))
|
||||
propagate_nonuniform_qualifier(chain.self);
|
||||
|
||||
SPIRType target_type;
|
||||
target_type.basetype = SPIRType::UInt;
|
||||
target_type.vecsize = type.vecsize;
|
||||
|
|
|
|||
|
|
@ -776,6 +776,10 @@ string CompilerMSL::compile()
|
|||
capture_output_to_buffer = msl_options.capture_output_to_buffer;
|
||||
is_rasterization_disabled = msl_options.disable_rasterization || capture_output_to_buffer;
|
||||
|
||||
// Initialize array here rather than constructor, MSVC 2013 workaround.
|
||||
for (auto &id : next_metal_resource_ids)
|
||||
id = 0;
|
||||
|
||||
fixup_type_alias();
|
||||
replace_illegal_names();
|
||||
|
||||
|
|
@ -2416,7 +2420,7 @@ void CompilerMSL::align_struct(SPIRType &ib_type)
|
|||
{
|
||||
set_extended_member_decoration(ib_type_id, mbr_idx, SPIRVCrossDecorationPacked);
|
||||
set_extended_member_decoration(ib_type_id, mbr_idx, SPIRVCrossDecorationPackedType,
|
||||
ib_type.member_types[mbr_idx]);
|
||||
get_member_packed_type(ib_type, mbr_idx));
|
||||
}
|
||||
|
||||
// Align current offset to the current member's default alignment.
|
||||
|
|
@ -2443,7 +2447,7 @@ void CompilerMSL::align_struct(SPIRType &ib_type)
|
|||
|
||||
// Returns whether the specified struct member supports a packable type
|
||||
// variation that is smaller than the unpacked variation of that type.
|
||||
bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index)
|
||||
bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index, uint32_t base_offset)
|
||||
{
|
||||
// We've already marked it as packable
|
||||
if (has_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPacked))
|
||||
|
|
@ -2466,30 +2470,81 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index)
|
|||
return true;
|
||||
}
|
||||
|
||||
// Check for array of struct, where the SPIR-V declares an array stride which is larger than the struct itself.
|
||||
// This can happen for struct A { float a }; A a[]; in std140 layout.
|
||||
// TODO: Emit a padded struct which can be used for this purpose.
|
||||
if (is_array(mbr_type) && mbr_type.basetype == SPIRType::Struct)
|
||||
uint32_t mbr_offset_curr = base_offset + get_member_decoration(ib_type.self, index, DecorationOffset);
|
||||
if (mbr_type.basetype == SPIRType::Struct)
|
||||
{
|
||||
// If this is a struct type, check if any of its members need packing.
|
||||
for (uint32_t i = 0; i < mbr_type.member_types.size(); i++)
|
||||
{
|
||||
if (is_member_packable(mbr_type, i, mbr_offset_curr))
|
||||
{
|
||||
set_extended_member_decoration(mbr_type.self, i, SPIRVCrossDecorationPacked);
|
||||
set_extended_member_decoration(mbr_type.self, i, SPIRVCrossDecorationPackedType,
|
||||
get_member_packed_type(mbr_type, i));
|
||||
}
|
||||
}
|
||||
size_t declared_struct_size = get_declared_struct_size(mbr_type);
|
||||
size_t alignment = get_declared_struct_member_alignment(ib_type, index);
|
||||
declared_struct_size = (declared_struct_size + alignment - 1) & ~(alignment - 1);
|
||||
if (type_struct_member_array_stride(ib_type, index) > declared_struct_size)
|
||||
return true;
|
||||
// Check for array of struct, where the SPIR-V declares an array stride which is larger than the struct itself.
|
||||
// This can happen for struct A { float a }; A a[]; in std140 layout.
|
||||
// TODO: Emit a padded struct which can be used for this purpose.
|
||||
if (is_array(mbr_type))
|
||||
{
|
||||
size_t array_stride = type_struct_member_array_stride(ib_type, index);
|
||||
if (array_stride > declared_struct_size)
|
||||
return true;
|
||||
if (array_stride < declared_struct_size)
|
||||
{
|
||||
// If the stride is *less* (i.e. more tightly packed), then
|
||||
// we need to pack the members of the struct itself.
|
||||
for (uint32_t i = 0; i < mbr_type.member_types.size(); i++)
|
||||
{
|
||||
if (is_member_packable(mbr_type, i, mbr_offset_curr + array_stride))
|
||||
{
|
||||
set_extended_member_decoration(mbr_type.self, i, SPIRVCrossDecorationPacked);
|
||||
set_extended_member_decoration(mbr_type.self, i, SPIRVCrossDecorationPackedType,
|
||||
get_member_packed_type(mbr_type, i));
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
// Pack if there is not enough space between this member and next.
|
||||
if (index < ib_type.member_types.size() - 1)
|
||||
{
|
||||
uint32_t mbr_offset_next =
|
||||
base_offset + get_member_decoration(ib_type.self, index + 1, DecorationOffset);
|
||||
if (declared_struct_size > mbr_offset_next - mbr_offset_curr)
|
||||
{
|
||||
for (uint32_t i = 0; i < mbr_type.member_types.size(); i++)
|
||||
{
|
||||
if (is_member_packable(mbr_type, i, mbr_offset_next))
|
||||
{
|
||||
set_extended_member_decoration(mbr_type.self, i, SPIRVCrossDecorationPacked);
|
||||
set_extended_member_decoration(mbr_type.self, i, SPIRVCrossDecorationPackedType,
|
||||
get_member_packed_type(mbr_type, i));
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// TODO: Another sanity check for matrices. We currently do not support std140 matrices which need to be padded out per column.
|
||||
//if (is_matrix(mbr_type) && mbr_type.vecsize <= 2 && type_struct_member_matrix_stride(ib_type, index) == 16)
|
||||
// SPIRV_CROSS_THROW("Currently cannot support matrices with small vector size in std140 layout.");
|
||||
|
||||
// Pack if the member's offset doesn't conform to the type's usual
|
||||
// alignment. For example, a float3 at offset 4.
|
||||
if (mbr_offset_curr % get_declared_struct_member_alignment(ib_type, index))
|
||||
return true;
|
||||
|
||||
// Only vectors or 3-row matrices need to be packed.
|
||||
if (mbr_type.vecsize == 1 || (is_matrix(mbr_type) && mbr_type.vecsize != 3))
|
||||
return false;
|
||||
|
||||
// Only row-major matrices need to be packed.
|
||||
if (is_matrix(mbr_type) && !has_member_decoration(ib_type.self, index, DecorationRowMajor))
|
||||
return false;
|
||||
|
||||
if (is_array(mbr_type))
|
||||
{
|
||||
// If member is an array, and the array stride is larger than the type needs, don't pack it.
|
||||
|
|
@ -2505,16 +2560,11 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index)
|
|||
}
|
||||
else
|
||||
{
|
||||
uint32_t mbr_offset_curr = get_member_decoration(ib_type.self, index, DecorationOffset);
|
||||
// For vectors, pack if the member's offset doesn't conform to the
|
||||
// type's usual alignment. For example, a float3 at offset 4.
|
||||
if (!is_matrix(mbr_type) && (mbr_offset_curr % unpacked_mbr_size))
|
||||
return true;
|
||||
// Pack if there is not enough space between this member and next.
|
||||
// If last member, only pack if it's a row-major matrix.
|
||||
if (index < ib_type.member_types.size() - 1)
|
||||
{
|
||||
uint32_t mbr_offset_next = get_member_decoration(ib_type.self, index + 1, DecorationOffset);
|
||||
uint32_t mbr_offset_next = base_offset + get_member_decoration(ib_type.self, index + 1, DecorationOffset);
|
||||
return unpacked_mbr_size > mbr_offset_next - mbr_offset_curr;
|
||||
}
|
||||
else
|
||||
|
|
@ -2522,6 +2572,25 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index)
|
|||
}
|
||||
}
|
||||
|
||||
uint32_t CompilerMSL::get_member_packed_type(SPIRType &type, uint32_t index)
|
||||
{
|
||||
auto &mbr_type = get<SPIRType>(type.member_types[index]);
|
||||
if (is_matrix(mbr_type) && has_member_decoration(type.self, index, DecorationRowMajor))
|
||||
{
|
||||
// Packed row-major matrices are stored transposed. But, we don't know if
|
||||
// we're dealing with a row-major matrix at the time we need to load it.
|
||||
// So, we'll set a packed type with the columns and rows transposed, so we'll
|
||||
// know to use the correct constructor.
|
||||
uint32_t new_type_id = ir.increase_bound_by(1);
|
||||
auto &transpose_type = set<SPIRType>(new_type_id);
|
||||
transpose_type = mbr_type;
|
||||
transpose_type.vecsize = mbr_type.columns;
|
||||
transpose_type.columns = mbr_type.vecsize;
|
||||
return new_type_id;
|
||||
}
|
||||
return type.member_types[index];
|
||||
}
|
||||
|
||||
// Returns a combination of type ID and member index for use as hash key
|
||||
MSLStructMemberKey CompilerMSL::get_struct_member_key(uint32_t type_id, uint32_t index)
|
||||
{
|
||||
|
|
@ -2542,19 +2611,51 @@ void CompilerMSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_exp
|
|||
{
|
||||
// Special handling when storing to a float[] or float2[] in std140 layout.
|
||||
|
||||
auto &type = get<SPIRType>(get_extended_decoration(lhs_expression, SPIRVCrossDecorationPackedType));
|
||||
uint32_t type_id = get_extended_decoration(lhs_expression, SPIRVCrossDecorationPackedType);
|
||||
auto &type = get<SPIRType>(type_id);
|
||||
string lhs = to_dereferenced_expression(lhs_expression);
|
||||
string rhs = to_pointer_expression(rhs_expression);
|
||||
uint32_t stride = get_decoration(type_id, DecorationArrayStride);
|
||||
|
||||
// Unpack the expression so we can store to it with a float or float2.
|
||||
// It's still an l-value, so it's fine. Most other unpacking of expressions turn them into r-values instead.
|
||||
if (is_scalar(type) && is_array(type))
|
||||
lhs = enclose_expression(lhs) + ".x";
|
||||
else if (is_vector(type) && type.vecsize == 2 && is_array(type))
|
||||
lhs = enclose_expression(lhs) + ".xy";
|
||||
if (is_matrix(type))
|
||||
{
|
||||
// Packed matrices are stored as arrays of packed vectors, so we need
|
||||
// to assign the vectors one at a time.
|
||||
// For row-major matrices, we need to transpose the *right-hand* side,
|
||||
// not the left-hand side. Otherwise, the changes will be lost.
|
||||
auto *lhs_e = maybe_get<SPIRExpression>(lhs_expression);
|
||||
auto *rhs_e = maybe_get<SPIRExpression>(rhs_expression);
|
||||
bool transpose = lhs_e && lhs_e->need_transpose;
|
||||
if (transpose)
|
||||
{
|
||||
lhs_e->need_transpose = false;
|
||||
if (rhs_e) rhs_e->need_transpose = !rhs_e->need_transpose;
|
||||
lhs = to_dereferenced_expression(lhs_expression);
|
||||
rhs = to_pointer_expression(rhs_expression);
|
||||
}
|
||||
for (uint32_t i = 0; i < type.columns; i++)
|
||||
statement(enclose_expression(lhs), "[", i, "] = ", enclose_expression(rhs), "[", i, "];");
|
||||
if (transpose)
|
||||
{
|
||||
lhs_e->need_transpose = true;
|
||||
if (rhs_e) rhs_e->need_transpose = !rhs_e->need_transpose;
|
||||
}
|
||||
}
|
||||
else if (is_array(type) && stride == 4 * type.width / 8)
|
||||
{
|
||||
// Unpack the expression so we can store to it with a float or float2.
|
||||
// It's still an l-value, so it's fine. Most other unpacking of expressions turn them into r-values instead.
|
||||
if (is_scalar(type))
|
||||
lhs = enclose_expression(lhs) + ".x";
|
||||
else if (is_vector(type) && type.vecsize == 2)
|
||||
lhs = enclose_expression(lhs) + ".xy";
|
||||
}
|
||||
|
||||
if (!optimize_read_modify_write(expression_type(rhs_expression), lhs, rhs))
|
||||
statement(lhs, " = ", rhs, ";");
|
||||
if (!is_matrix(type))
|
||||
{
|
||||
if (!optimize_read_modify_write(expression_type(rhs_expression), lhs, rhs))
|
||||
statement(lhs, " = ", rhs, ";");
|
||||
}
|
||||
register_write(lhs_expression);
|
||||
}
|
||||
}
|
||||
|
|
@ -2564,14 +2665,37 @@ void CompilerMSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_exp
|
|||
string CompilerMSL::unpack_expression_type(string expr_str, const SPIRType &type, uint32_t packed_type_id)
|
||||
{
|
||||
const SPIRType *packed_type = nullptr;
|
||||
uint32_t stride = 0;
|
||||
if (packed_type_id)
|
||||
{
|
||||
packed_type = &get<SPIRType>(packed_type_id);
|
||||
stride = get_decoration(packed_type_id, DecorationArrayStride);
|
||||
}
|
||||
|
||||
// float[] and float2[] cases are really just padding, so directly swizzle from the backing float4 instead.
|
||||
if (packed_type && is_array(*packed_type) && is_scalar(*packed_type))
|
||||
if (packed_type && is_array(*packed_type) && is_scalar(*packed_type) && stride == 4 * packed_type->width / 8)
|
||||
return enclose_expression(expr_str) + ".x";
|
||||
else if (packed_type && is_array(*packed_type) && is_vector(*packed_type) && packed_type->vecsize == 2)
|
||||
else if (packed_type && is_array(*packed_type) && is_vector(*packed_type) && packed_type->vecsize == 2 &&
|
||||
stride == 4 * packed_type->width / 8)
|
||||
return enclose_expression(expr_str) + ".xy";
|
||||
else if (is_matrix(type))
|
||||
{
|
||||
// Packed matrices are stored as arrays of packed vectors. Unfortunately,
|
||||
// we can't just pass the array straight to the matrix constructor. We have to
|
||||
// pass each vector individually, so that they can be unpacked to normal vectors.
|
||||
if (!packed_type)
|
||||
packed_type = &type;
|
||||
const char *base_type = packed_type->width == 16 ? "half" : "float";
|
||||
string unpack_expr = join(type_to_glsl(*packed_type), "(");
|
||||
for (uint32_t i = 0; i < packed_type->columns; i++)
|
||||
{
|
||||
if (i > 0)
|
||||
unpack_expr += ", ";
|
||||
unpack_expr += join(base_type, packed_type->vecsize, "(", expr_str, "[", i, "])");
|
||||
}
|
||||
unpack_expr += ")";
|
||||
return unpack_expr;
|
||||
}
|
||||
else
|
||||
return join(type_to_glsl(type), "(", expr_str, ")");
|
||||
}
|
||||
|
|
@ -4176,16 +4300,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
|
|||
if (e && e->need_transpose && (t.columns == t.vecsize || is_packed))
|
||||
{
|
||||
e->need_transpose = false;
|
||||
|
||||
// This is important for matrices. Packed matrices
|
||||
// are generally transposed, so unpacking using a constructor argument
|
||||
// will result in an error.
|
||||
// The simplest solution for now is to just avoid unpacking the matrix in this operation.
|
||||
unset_extended_decoration(mtx_id, SPIRVCrossDecorationPacked);
|
||||
|
||||
emit_binary_op(ops[0], ops[1], ops[3], ops[2], "*");
|
||||
if (is_packed)
|
||||
set_extended_decoration(mtx_id, SPIRVCrossDecorationPacked);
|
||||
e->need_transpose = true;
|
||||
}
|
||||
else
|
||||
|
|
@ -5640,16 +5755,28 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
|
|||
}
|
||||
else if (membertype.vecsize > 1 && membertype.columns > 1)
|
||||
{
|
||||
uint32_t rows = membertype.vecsize;
|
||||
uint32_t cols = membertype.columns;
|
||||
pack_pfx = "packed_";
|
||||
if (has_member_decoration(type.self, index, DecorationRowMajor))
|
||||
{
|
||||
// These are stored transposed.
|
||||
rows = membertype.columns;
|
||||
cols = membertype.vecsize;
|
||||
pack_pfx = "packed_rm_";
|
||||
}
|
||||
string base_type = membertype.width == 16 ? "half" : "float";
|
||||
string td_line = "typedef ";
|
||||
td_line += base_type + to_string(membertype.vecsize) + "x" + to_string(membertype.columns);
|
||||
td_line += "packed_" + base_type + to_string(rows);
|
||||
td_line += " " + pack_pfx;
|
||||
// Use the actual matrix size here.
|
||||
td_line += base_type + to_string(membertype.columns) + "x" + to_string(membertype.vecsize);
|
||||
td_line += "[" + to_string(cols) + "]";
|
||||
td_line += ";";
|
||||
add_typedef_line(td_line);
|
||||
}
|
||||
else if (is_array(membertype) && membertype.vecsize <= 2 && membertype.basetype != SPIRType::Struct)
|
||||
else if (is_array(membertype) && membertype.vecsize <= 2 && membertype.basetype != SPIRType::Struct &&
|
||||
type_struct_member_array_stride(type, index) == 4 * membertype.width / 8)
|
||||
{
|
||||
// A "packed" float array, but we pad here instead to 4-vector.
|
||||
override_type = membertype;
|
||||
|
|
@ -8428,9 +8555,15 @@ size_t CompilerMSL::get_declared_struct_member_alignment(const SPIRType &struct_
|
|||
const SPIRType *packed_type = packed_type_id != 0 ? &get<SPIRType>(packed_type_id) : nullptr;
|
||||
if (packed_type && is_array(*packed_type) && !is_matrix(*packed_type) &&
|
||||
packed_type->basetype != SPIRType::Struct)
|
||||
return (packed_type->width / 8) * 4;
|
||||
{
|
||||
uint32_t stride = type_struct_member_array_stride(struct_type, index);
|
||||
if (stride == (packed_type->width / 8) * 4)
|
||||
return stride;
|
||||
else
|
||||
return packed_type->width / 8;
|
||||
}
|
||||
else
|
||||
return (type.width / 8) * (type.columns == 3 ? 4 : type.columns);
|
||||
return type.width / 8;
|
||||
}
|
||||
else
|
||||
return (type.width / 8) * (type.vecsize == 3 ? 4 : type.vecsize);
|
||||
|
|
@ -8691,14 +8824,6 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
|
|||
return SPVFuncImplTextureSwizzle;
|
||||
break;
|
||||
|
||||
case OpCompositeConstruct:
|
||||
{
|
||||
auto &type = compiler.get<SPIRType>(args[0]);
|
||||
if (type.array.size() > 1) // We need to use copies to build the composite.
|
||||
return static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + type.array.size() - 1);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpExtInst:
|
||||
{
|
||||
uint32_t extension_set = args[2];
|
||||
|
|
|
|||
|
|
@ -511,7 +511,8 @@ protected:
|
|||
size_t get_declared_struct_member_alignment(const SPIRType &struct_type, uint32_t index) const;
|
||||
std::string to_component_argument(uint32_t id);
|
||||
void align_struct(SPIRType &ib_type);
|
||||
bool is_member_packable(SPIRType &ib_type, uint32_t index);
|
||||
bool is_member_packable(SPIRType &ib_type, uint32_t index, uint32_t base_offset = 0);
|
||||
uint32_t get_member_packed_type(SPIRType &ib_type, uint32_t index);
|
||||
MSLStructMemberKey get_struct_member_key(uint32_t type_id, uint32_t index);
|
||||
std::string get_argument_address_space(const SPIRVariable &argument);
|
||||
std::string get_type_address_space(const SPIRType &type, uint32_t id);
|
||||
|
|
@ -593,10 +594,12 @@ protected:
|
|||
};
|
||||
|
||||
std::unordered_map<StageSetBinding, std::pair<MSLResourceBinding, bool>, InternalHasher> resource_bindings;
|
||||
|
||||
uint32_t next_metal_resource_index_buffer = 0;
|
||||
uint32_t next_metal_resource_index_texture = 0;
|
||||
uint32_t next_metal_resource_index_sampler = 0;
|
||||
uint32_t next_metal_resource_ids[kMaxArgumentBuffers] = {};
|
||||
// Intentionally uninitialized, works around MSVC 2013 bug.
|
||||
uint32_t next_metal_resource_ids[kMaxArgumentBuffers];
|
||||
|
||||
uint32_t stage_in_var_id = 0;
|
||||
uint32_t stage_out_var_id = 0;
|
||||
|
|
|
|||
|
|
@ -162,7 +162,6 @@ void Parser::parse(const Instruction &instruction)
|
|||
case OpSourceContinued:
|
||||
case OpSourceExtension:
|
||||
case OpNop:
|
||||
case OpNoLine:
|
||||
case OpModuleProcessed:
|
||||
break;
|
||||
|
||||
|
|
@ -244,6 +243,8 @@ void Parser::parse(const Instruction &instruction)
|
|||
auto ext = extract_string(ir.spirv, instruction.offset + 1);
|
||||
if (ext == "GLSL.std.450")
|
||||
set<SPIRExtension>(id, SPIRExtension::GLSL);
|
||||
else if (ext == "DebugInfo")
|
||||
set<SPIRExtension>(id, SPIRExtension::SPV_debug_info);
|
||||
else if (ext == "SPV_AMD_shader_ballot")
|
||||
set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_ballot);
|
||||
else if (ext == "SPV_AMD_shader_explicit_vertex_parameter")
|
||||
|
|
@ -260,6 +261,14 @@ void Parser::parse(const Instruction &instruction)
|
|||
break;
|
||||
}
|
||||
|
||||
case OpExtInst:
|
||||
{
|
||||
// The SPIR-V debug information extended instructions might come at global scope.
|
||||
if (current_block)
|
||||
current_block->ops.push_back(instruction);
|
||||
break;
|
||||
}
|
||||
|
||||
case OpEntryPoint:
|
||||
{
|
||||
auto itr =
|
||||
|
|
@ -560,10 +569,6 @@ void Parser::parse(const Instruction &instruction)
|
|||
type.image.sampled = ops[6];
|
||||
type.image.format = static_cast<ImageFormat>(ops[7]);
|
||||
type.image.access = (length >= 9) ? static_cast<AccessQualifier>(ops[8]) : AccessQualifierMax;
|
||||
|
||||
if (type.image.sampled == 0)
|
||||
SPIRV_CROSS_THROW("OpTypeImage Sampled parameter must not be zero.");
|
||||
|
||||
break;
|
||||
}
|
||||
|
||||
|
|
@ -1057,6 +1062,14 @@ void Parser::parse(const Instruction &instruction)
|
|||
break;
|
||||
}
|
||||
|
||||
case OpNoLine:
|
||||
{
|
||||
// OpNoLine might come at global scope.
|
||||
if (current_block)
|
||||
current_block->ops.push_back(instruction);
|
||||
break;
|
||||
}
|
||||
|
||||
// Actual opcodes.
|
||||
default:
|
||||
{
|
||||
|
|
|
|||
Loading…
Reference in New Issue