Declare OpSpecConstantOp up-front on relevant targets.
Required, since spec constants can include results from constant ops.
This commit is contained in:
Hans-Kristian Arntzen
parent
0617b98613
commit
991b655c72
@ -0,0 +1,37 @@
|
||||
static const int _7 = -10;
|
||||
static const uint _8 = 100u;
|
||||
static const int _20 = (_7 + 2);
|
||||
static const uint _25 = (_8 % 5u);
|
||||
static const int4 _30 = int4(20, 30, _20, _20);
|
||||
static const int2 _32 = int2(_30.y, _30.x);
|
||||
static const int _33 = _30.y;
|
||||
|
||||
static float4 gl_Position;
|
||||
static int _4;
|
||||
|
||||
struct SPIRV_Cross_Output
|
||||
{
|
||||
nointerpolation int _4 : TEXCOORD0;
|
||||
float4 gl_Position : SV_Position;
|
||||
};
|
||||
|
||||
void vert_main()
|
||||
{
|
||||
float4 _64 = 0.0f.xxxx;
|
||||
_64.y = float(_20);
|
||||
float4 _68 = _64;
|
||||
_68.z = float(_25);
|
||||
float4 _52 = _68 + float4(_30);
|
||||
float2 _56 = _52.xy + float2(_32);
|
||||
gl_Position = float4(_56.x, _56.y, _52.z, _52.w);
|
||||
_4 = _33;
|
||||
}
|
||||
|
||||
SPIRV_Cross_Output main()
|
||||
{
|
||||
vert_main();
|
||||
SPIRV_Cross_Output stage_output;
|
||||
stage_output.gl_Position = gl_Position;
|
||||
stage_output._4 = _4;
|
||||
return stage_output;
|
||||
}
|
||||
@ -0,0 +1,35 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
constant int _7_tmp [[function_constant(201)]];
|
||||
constant int _7 = is_function_constant_defined(_7_tmp) ? _7_tmp : -10;
|
||||
constant uint _8_tmp [[function_constant(202)]];
|
||||
constant uint _8 = is_function_constant_defined(_8_tmp) ? _8_tmp : 100u;
|
||||
constant int _20 = (_7 + 2);
|
||||
constant uint _25 = (_8 % 5u);
|
||||
constant int4 _30 = int4(20, 30, _20, _20);
|
||||
constant int2 _32 = int2(_30.y, _30.x);
|
||||
constant int _33 = _30.y;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
int m_4 [[user(locn0)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
vertex main0_out main0()
|
||||
{
|
||||
main0_out out = {};
|
||||
float4 _64 = float4(0.0);
|
||||
_64.y = float(_20);
|
||||
float4 _68 = _64;
|
||||
_68.z = float(_25);
|
||||
float4 _52 = _68 + float4(_30);
|
||||
float2 _56 = _52.xy + float2(_32);
|
||||
out.gl_Position = float4(_56.x, _56.y, _52.z, _52.w);
|
||||
out.m_4 = _33;
|
||||
return out;
|
||||
}
|
||||
|
||||
@ -2,19 +2,23 @@
|
||||
|
||||
layout(constant_id = 201) const int _7 = -10;
|
||||
layout(constant_id = 202) const uint _8 = 100u;
|
||||
const ivec4 _30 = ivec4(20, 30, 0, 0);
|
||||
const int _20 = (_7 + 2);
|
||||
const uint _25 = (_8 % 5u);
|
||||
const ivec4 _30 = ivec4(20, 30, _20, _20);
|
||||
const ivec2 _32 = ivec2(_30.y, _30.x);
|
||||
const int _33 = _30.y;
|
||||
|
||||
layout(location = 0) flat out int _4;
|
||||
|
||||
void main()
|
||||
{
|
||||
vec4 _64 = vec4(0.0);
|
||||
_64.y = float((_7 + 2));
|
||||
_64.y = float(_20);
|
||||
vec4 _68 = _64;
|
||||
_68.z = float((_8 % 5u));
|
||||
_68.z = float(_25);
|
||||
vec4 _52 = _68 + vec4(_30);
|
||||
vec2 _56 = _52.xy + vec2(ivec2(_30.y, _30.x));
|
||||
vec2 _56 = _52.xy + vec2(_32);
|
||||
gl_Position = vec4(_56.x, _56.y, _52.z, _52.w);
|
||||
_4 = _30.y;
|
||||
_4 = _33;
|
||||
}
|
||||
|
||||
|
||||
@ -1,15 +1,54 @@
|
||||
static const float a = 1.0f;
|
||||
static const float b = 2.0f;
|
||||
static const int c = 3;
|
||||
static const uint _18 = (uint(c) + 0u);
|
||||
static const int _21 = (-c);
|
||||
static const int _23 = (~c);
|
||||
static const int d = 4;
|
||||
static const int _26 = (c + d);
|
||||
static const int _28 = (c - d);
|
||||
static const int _30 = (c * d);
|
||||
static const int _32 = (c / d);
|
||||
static const uint e = 5u;
|
||||
static const uint f = 6u;
|
||||
static const uint _36 = (e / f);
|
||||
static const int _38 = (c % d);
|
||||
static const uint _40 = (e % f);
|
||||
static const int _42 = (c >> d);
|
||||
static const uint _44 = (e >> f);
|
||||
static const int _46 = (c << d);
|
||||
static const int _48 = (c | d);
|
||||
static const int _50 = (c ^ d);
|
||||
static const int _52 = (c & d);
|
||||
static const bool g = false;
|
||||
static const bool h = true;
|
||||
static const bool _58 = (g || h);
|
||||
static const bool _60 = (g && h);
|
||||
static const bool _62 = (!g);
|
||||
static const bool _64 = (g == h);
|
||||
static const bool _66 = (g != h);
|
||||
static const bool _68 = (c == d);
|
||||
static const bool _70 = (c != d);
|
||||
static const bool _72 = (c < d);
|
||||
static const bool _74 = (e < f);
|
||||
static const bool _76 = (c > d);
|
||||
static const bool _78 = (e > f);
|
||||
static const bool _80 = (c <= d);
|
||||
static const bool _82 = (e <= f);
|
||||
static const bool _84 = (c >= d);
|
||||
static const bool _86 = (e >= f);
|
||||
static const int _92 = int(e + 0u);
|
||||
static const bool _94 = (c != int(0u));
|
||||
static const bool _96 = (e != 0u);
|
||||
static const int _100 = int(g);
|
||||
static const uint _103 = uint(g);
|
||||
static const int _111 = (c + 3);
|
||||
static const int _118 = (c + 2);
|
||||
static const int _124 = (d + 2);
|
||||
|
||||
struct Foo
|
||||
{
|
||||
float elems[(d + 2)];
|
||||
float elems[_124];
|
||||
};
|
||||
|
||||
static float4 FragColor;
|
||||
@ -23,47 +62,47 @@ void frag_main()
|
||||
{
|
||||
float t0 = a;
|
||||
float t1 = b;
|
||||
uint c0 = (uint(c) + 0u);
|
||||
int c1 = (-c);
|
||||
int c2 = (~c);
|
||||
int c3 = (c + d);
|
||||
int c4 = (c - d);
|
||||
int c5 = (c * d);
|
||||
int c6 = (c / d);
|
||||
uint c7 = (e / f);
|
||||
int c8 = (c % d);
|
||||
uint c9 = (e % f);
|
||||
int c10 = (c >> d);
|
||||
uint c11 = (e >> f);
|
||||
int c12 = (c << d);
|
||||
int c13 = (c | d);
|
||||
int c14 = (c ^ d);
|
||||
int c15 = (c & d);
|
||||
bool c16 = (g || h);
|
||||
bool c17 = (g && h);
|
||||
bool c18 = (!g);
|
||||
bool c19 = (g == h);
|
||||
bool c20 = (g != h);
|
||||
bool c21 = (c == d);
|
||||
bool c22 = (c != d);
|
||||
bool c23 = (c < d);
|
||||
bool c24 = (e < f);
|
||||
bool c25 = (c > d);
|
||||
bool c26 = (e > f);
|
||||
bool c27 = (c <= d);
|
||||
bool c28 = (e <= f);
|
||||
bool c29 = (c >= d);
|
||||
bool c30 = (e >= f);
|
||||
uint c0 = _18;
|
||||
int c1 = _21;
|
||||
int c2 = _23;
|
||||
int c3 = _26;
|
||||
int c4 = _28;
|
||||
int c5 = _30;
|
||||
int c6 = _32;
|
||||
uint c7 = _36;
|
||||
int c8 = _38;
|
||||
uint c9 = _40;
|
||||
int c10 = _42;
|
||||
uint c11 = _44;
|
||||
int c12 = _46;
|
||||
int c13 = _48;
|
||||
int c14 = _50;
|
||||
int c15 = _52;
|
||||
bool c16 = _58;
|
||||
bool c17 = _60;
|
||||
bool c18 = _62;
|
||||
bool c19 = _64;
|
||||
bool c20 = _66;
|
||||
bool c21 = _68;
|
||||
bool c22 = _70;
|
||||
bool c23 = _72;
|
||||
bool c24 = _74;
|
||||
bool c25 = _76;
|
||||
bool c26 = _78;
|
||||
bool c27 = _80;
|
||||
bool c28 = _82;
|
||||
bool c29 = _84;
|
||||
bool c30 = _86;
|
||||
int c31 = c8 + c3;
|
||||
int c32 = int(e + 0u);
|
||||
bool c33 = (c != int(0u));
|
||||
bool c34 = (e != 0u);
|
||||
int c35 = int(g);
|
||||
uint c36 = uint(g);
|
||||
int c32 = _92;
|
||||
bool c33 = _94;
|
||||
bool c34 = _96;
|
||||
int c35 = _100;
|
||||
uint c36 = _103;
|
||||
float c37 = float(g);
|
||||
float vec0[(c + 3)][8];
|
||||
float vec0[_111][8];
|
||||
vec0[0][0] = 10.0f;
|
||||
float vec1[(c + 2)];
|
||||
float vec1[_118];
|
||||
vec1[0] = 20.0f;
|
||||
Foo foo;
|
||||
foo.elems[c] = 10.0f;
|
||||
|
||||
@ -0,0 +1,38 @@
|
||||
static const int _7 = -10;
|
||||
static const uint _8 = 100u;
|
||||
static const float _9 = 3.141590118408203125f;
|
||||
static const int _20 = (_7 + 2);
|
||||
static const uint _25 = (_8 % 5u);
|
||||
static const int4 _30 = int4(20, 30, _20, _20);
|
||||
static const int2 _32 = int2(_30.y, _30.x);
|
||||
static const int _33 = _30.y;
|
||||
|
||||
static float4 gl_Position;
|
||||
static int _4;
|
||||
|
||||
struct SPIRV_Cross_Output
|
||||
{
|
||||
nointerpolation int _4 : TEXCOORD0;
|
||||
float4 gl_Position : SV_Position;
|
||||
};
|
||||
|
||||
void vert_main()
|
||||
{
|
||||
float4 pos = 0.0f.xxxx;
|
||||
pos.y += float(_20);
|
||||
pos.z += float(_25);
|
||||
pos += float4(_30);
|
||||
float2 _56 = pos.xy + float2(_32);
|
||||
pos = float4(_56.x, _56.y, pos.z, pos.w);
|
||||
gl_Position = pos;
|
||||
_4 = _33;
|
||||
}
|
||||
|
||||
SPIRV_Cross_Output main()
|
||||
{
|
||||
vert_main();
|
||||
SPIRV_Cross_Output stage_output;
|
||||
stage_output.gl_Position = gl_Position;
|
||||
stage_output._4 = _4;
|
||||
return stage_output;
|
||||
}
|
||||
@ -0,0 +1,37 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
constant int _7_tmp [[function_constant(201)]];
|
||||
constant int _7 = is_function_constant_defined(_7_tmp) ? _7_tmp : -10;
|
||||
constant uint _8_tmp [[function_constant(202)]];
|
||||
constant uint _8 = is_function_constant_defined(_8_tmp) ? _8_tmp : 100u;
|
||||
constant float _9_tmp [[function_constant(200)]];
|
||||
constant float _9 = is_function_constant_defined(_9_tmp) ? _9_tmp : 3.141590118408203125;
|
||||
constant int _20 = (_7 + 2);
|
||||
constant uint _25 = (_8 % 5u);
|
||||
constant int4 _30 = int4(20, 30, _20, _20);
|
||||
constant int2 _32 = int2(_30.y, _30.x);
|
||||
constant int _33 = _30.y;
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
int m_4 [[user(locn0)]];
|
||||
float4 gl_Position [[position]];
|
||||
};
|
||||
|
||||
vertex main0_out main0()
|
||||
{
|
||||
main0_out out = {};
|
||||
float4 pos = float4(0.0);
|
||||
pos.y += float(_20);
|
||||
pos.z += float(_25);
|
||||
pos += float4(_30);
|
||||
float2 _56 = pos.xy + float2(_32);
|
||||
pos = float4(_56.x, _56.y, pos.z, pos.w);
|
||||
out.gl_Position = pos;
|
||||
out.m_4 = _33;
|
||||
return out;
|
||||
}
|
||||
|
||||
@ -9,16 +9,52 @@ constant float b_tmp [[function_constant(2)]];
|
||||
constant float b = is_function_constant_defined(b_tmp) ? b_tmp : 2.0;
|
||||
constant int c_tmp [[function_constant(3)]];
|
||||
constant int c = is_function_constant_defined(c_tmp) ? c_tmp : 3;
|
||||
constant uint _18 = (uint(c) + 0u);
|
||||
constant int _21 = (-c);
|
||||
constant int _23 = (~c);
|
||||
constant int d_tmp [[function_constant(4)]];
|
||||
constant int d = is_function_constant_defined(d_tmp) ? d_tmp : 4;
|
||||
constant int _26 = (c + d);
|
||||
constant int _28 = (c - d);
|
||||
constant int _30 = (c * d);
|
||||
constant int _32 = (c / d);
|
||||
constant uint e_tmp [[function_constant(5)]];
|
||||
constant uint e = is_function_constant_defined(e_tmp) ? e_tmp : 5u;
|
||||
constant uint f_tmp [[function_constant(6)]];
|
||||
constant uint f = is_function_constant_defined(f_tmp) ? f_tmp : 6u;
|
||||
constant uint _36 = (e / f);
|
||||
constant int _38 = (c % d);
|
||||
constant uint _40 = (e % f);
|
||||
constant int _42 = (c >> d);
|
||||
constant uint _44 = (e >> f);
|
||||
constant int _46 = (c << d);
|
||||
constant int _48 = (c | d);
|
||||
constant int _50 = (c ^ d);
|
||||
constant int _52 = (c & d);
|
||||
constant bool g_tmp [[function_constant(7)]];
|
||||
constant bool g = is_function_constant_defined(g_tmp) ? g_tmp : false;
|
||||
constant bool h_tmp [[function_constant(8)]];
|
||||
constant bool h = is_function_constant_defined(h_tmp) ? h_tmp : true;
|
||||
constant bool _58 = (g || h);
|
||||
constant bool _60 = (g && h);
|
||||
constant bool _62 = (!g);
|
||||
constant bool _64 = (g == h);
|
||||
constant bool _66 = (g != h);
|
||||
constant bool _68 = (c == d);
|
||||
constant bool _70 = (c != d);
|
||||
constant bool _72 = (c < d);
|
||||
constant bool _74 = (e < f);
|
||||
constant bool _76 = (c > d);
|
||||
constant bool _78 = (e > f);
|
||||
constant bool _80 = (c <= d);
|
||||
constant bool _82 = (e <= f);
|
||||
constant bool _84 = (c >= d);
|
||||
constant bool _86 = (e >= f);
|
||||
constant int _92 = int(e + 0u);
|
||||
constant bool _94 = (c != int(0u));
|
||||
constant bool _96 = (e != 0u);
|
||||
constant int _100 = int(g);
|
||||
constant uint _103 = uint(g);
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
@ -30,43 +66,43 @@ fragment main0_out main0()
|
||||
main0_out out = {};
|
||||
float t0 = a;
|
||||
float t1 = b;
|
||||
uint c0 = (uint(c) + 0u);
|
||||
int c1 = (-c);
|
||||
int c2 = (~c);
|
||||
int c3 = (c + d);
|
||||
int c4 = (c - d);
|
||||
int c5 = (c * d);
|
||||
int c6 = (c / d);
|
||||
uint c7 = (e / f);
|
||||
int c8 = (c % d);
|
||||
uint c9 = (e % f);
|
||||
int c10 = (c >> d);
|
||||
uint c11 = (e >> f);
|
||||
int c12 = (c << d);
|
||||
int c13 = (c | d);
|
||||
int c14 = (c ^ d);
|
||||
int c15 = (c & d);
|
||||
bool c16 = (g || h);
|
||||
bool c17 = (g && h);
|
||||
bool c18 = (!g);
|
||||
bool c19 = (g == h);
|
||||
bool c20 = (g != h);
|
||||
bool c21 = (c == d);
|
||||
bool c22 = (c != d);
|
||||
bool c23 = (c < d);
|
||||
bool c24 = (e < f);
|
||||
bool c25 = (c > d);
|
||||
bool c26 = (e > f);
|
||||
bool c27 = (c <= d);
|
||||
bool c28 = (e <= f);
|
||||
bool c29 = (c >= d);
|
||||
bool c30 = (e >= f);
|
||||
uint c0 = _18;
|
||||
int c1 = _21;
|
||||
int c2 = _23;
|
||||
int c3 = _26;
|
||||
int c4 = _28;
|
||||
int c5 = _30;
|
||||
int c6 = _32;
|
||||
uint c7 = _36;
|
||||
int c8 = _38;
|
||||
uint c9 = _40;
|
||||
int c10 = _42;
|
||||
uint c11 = _44;
|
||||
int c12 = _46;
|
||||
int c13 = _48;
|
||||
int c14 = _50;
|
||||
int c15 = _52;
|
||||
bool c16 = _58;
|
||||
bool c17 = _60;
|
||||
bool c18 = _62;
|
||||
bool c19 = _64;
|
||||
bool c20 = _66;
|
||||
bool c21 = _68;
|
||||
bool c22 = _70;
|
||||
bool c23 = _72;
|
||||
bool c24 = _74;
|
||||
bool c25 = _76;
|
||||
bool c26 = _78;
|
||||
bool c27 = _80;
|
||||
bool c28 = _82;
|
||||
bool c29 = _84;
|
||||
bool c30 = _86;
|
||||
int c31 = c8 + c3;
|
||||
int c32 = int(e + 0u);
|
||||
bool c33 = (c != int(0u));
|
||||
bool c34 = (e != 0u);
|
||||
int c35 = int(g);
|
||||
uint c36 = uint(g);
|
||||
int c32 = _92;
|
||||
bool c33 = _94;
|
||||
bool c34 = _96;
|
||||
int c35 = _100;
|
||||
uint c36 = _103;
|
||||
float c37 = float(g);
|
||||
out.FragColor = float4(t0 + t1);
|
||||
return out;
|
||||
|
||||
@ -5,15 +5,54 @@ precision highp int;
|
||||
layout(constant_id = 1) const float a = 1.0;
|
||||
layout(constant_id = 2) const float b = 2.0;
|
||||
layout(constant_id = 3) const int c = 3;
|
||||
const uint _18 = (uint(c) + 0u);
|
||||
const int _21 = (-c);
|
||||
const int _23 = (~c);
|
||||
layout(constant_id = 4) const int d = 4;
|
||||
const int _26 = (c + d);
|
||||
const int _28 = (c - d);
|
||||
const int _30 = (c * d);
|
||||
const int _32 = (c / d);
|
||||
layout(constant_id = 5) const uint e = 5u;
|
||||
layout(constant_id = 6) const uint f = 6u;
|
||||
const uint _36 = (e / f);
|
||||
const int _38 = (c % d);
|
||||
const uint _40 = (e % f);
|
||||
const int _42 = (c >> d);
|
||||
const uint _44 = (e >> f);
|
||||
const int _46 = (c << d);
|
||||
const int _48 = (c | d);
|
||||
const int _50 = (c ^ d);
|
||||
const int _52 = (c & d);
|
||||
layout(constant_id = 7) const bool g = false;
|
||||
layout(constant_id = 8) const bool h = true;
|
||||
const bool _58 = (g || h);
|
||||
const bool _60 = (g && h);
|
||||
const bool _62 = (!g);
|
||||
const bool _64 = (g == h);
|
||||
const bool _66 = (g != h);
|
||||
const bool _68 = (c == d);
|
||||
const bool _70 = (c != d);
|
||||
const bool _72 = (c < d);
|
||||
const bool _74 = (e < f);
|
||||
const bool _76 = (c > d);
|
||||
const bool _78 = (e > f);
|
||||
const bool _80 = (c <= d);
|
||||
const bool _82 = (e <= f);
|
||||
const bool _84 = (c >= d);
|
||||
const bool _86 = (e >= f);
|
||||
const int _92 = int(e + 0u);
|
||||
const bool _94 = (c != int(0u));
|
||||
const bool _96 = (e != 0u);
|
||||
const int _100 = int(g);
|
||||
const uint _103 = uint(g);
|
||||
const int _118 = (c + 3);
|
||||
const int _127 = (c + 2);
|
||||
const int _135 = (d + 2);
|
||||
|
||||
struct Foo
|
||||
{
|
||||
float elems[(d + 2)];
|
||||
float elems[_135];
|
||||
};
|
||||
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
@ -22,46 +61,46 @@ void main()
|
||||
{
|
||||
float t0 = a;
|
||||
float t1 = b;
|
||||
mediump uint c0 = (uint(c) + 0u);
|
||||
mediump int c1 = (-c);
|
||||
mediump int c2 = (~c);
|
||||
mediump int c3 = (c + d);
|
||||
mediump int c4 = (c - d);
|
||||
mediump int c5 = (c * d);
|
||||
mediump int c6 = (c / d);
|
||||
mediump uint c7 = (e / f);
|
||||
mediump int c8 = (c % d);
|
||||
mediump uint c9 = (e % f);
|
||||
mediump int c10 = (c >> d);
|
||||
mediump uint c11 = (e >> f);
|
||||
mediump int c12 = (c << d);
|
||||
mediump int c13 = (c | d);
|
||||
mediump int c14 = (c ^ d);
|
||||
mediump int c15 = (c & d);
|
||||
bool c16 = (g || h);
|
||||
bool c17 = (g && h);
|
||||
bool c18 = (!g);
|
||||
bool c19 = (g == h);
|
||||
bool c20 = (g != h);
|
||||
bool c21 = (c == d);
|
||||
bool c22 = (c != d);
|
||||
bool c23 = (c < d);
|
||||
bool c24 = (e < f);
|
||||
bool c25 = (c > d);
|
||||
bool c26 = (e > f);
|
||||
bool c27 = (c <= d);
|
||||
bool c28 = (e <= f);
|
||||
bool c29 = (c >= d);
|
||||
bool c30 = (e >= f);
|
||||
mediump uint c0 = _18;
|
||||
mediump int c1 = _21;
|
||||
mediump int c2 = _23;
|
||||
mediump int c3 = _26;
|
||||
mediump int c4 = _28;
|
||||
mediump int c5 = _30;
|
||||
mediump int c6 = _32;
|
||||
mediump uint c7 = _36;
|
||||
mediump int c8 = _38;
|
||||
mediump uint c9 = _40;
|
||||
mediump int c10 = _42;
|
||||
mediump uint c11 = _44;
|
||||
mediump int c12 = _46;
|
||||
mediump int c13 = _48;
|
||||
mediump int c14 = _50;
|
||||
mediump int c15 = _52;
|
||||
bool c16 = _58;
|
||||
bool c17 = _60;
|
||||
bool c18 = _62;
|
||||
bool c19 = _64;
|
||||
bool c20 = _66;
|
||||
bool c21 = _68;
|
||||
bool c22 = _70;
|
||||
bool c23 = _72;
|
||||
bool c24 = _74;
|
||||
bool c25 = _76;
|
||||
bool c26 = _78;
|
||||
bool c27 = _80;
|
||||
bool c28 = _82;
|
||||
bool c29 = _84;
|
||||
bool c30 = _86;
|
||||
mediump int c31 = c8 + c3;
|
||||
mediump int c32 = int(e + 0u);
|
||||
bool c33 = (c != int(0u));
|
||||
bool c34 = (e != 0u);
|
||||
mediump int c35 = int(g);
|
||||
mediump uint c36 = uint(g);
|
||||
mediump int c32 = _92;
|
||||
bool c33 = _94;
|
||||
bool c34 = _96;
|
||||
mediump int c35 = _100;
|
||||
mediump uint c36 = _103;
|
||||
float c37 = float(g);
|
||||
float vec0[(c + 3)][8];
|
||||
float vec1[(c + 2)];
|
||||
float vec0[_118][8];
|
||||
float vec1[_127];
|
||||
Foo foo;
|
||||
FragColor = ((vec4(t0 + t1) + vec4(vec0[0][0])) + vec4(vec1[0])) + vec4(foo.elems[c]);
|
||||
}
|
||||
|
||||
@ -3,19 +3,23 @@
|
||||
layout(constant_id = 201) const int _7 = -10;
|
||||
layout(constant_id = 202) const uint _8 = 100u;
|
||||
layout(constant_id = 200) const float _9 = 3.141590118408203125;
|
||||
const ivec4 _30 = ivec4(20, 30, 0, 0);
|
||||
const int _20 = (_7 + 2);
|
||||
const uint _25 = (_8 % 5u);
|
||||
const ivec4 _30 = ivec4(20, 30, _20, _20);
|
||||
const ivec2 _32 = ivec2(_30.y, _30.x);
|
||||
const int _33 = _30.y;
|
||||
|
||||
layout(location = 0) flat out int _4;
|
||||
|
||||
void main()
|
||||
{
|
||||
vec4 pos = vec4(0.0);
|
||||
pos.y += float((_7 + 2));
|
||||
pos.z += float((_8 % 5u));
|
||||
pos.y += float(_20);
|
||||
pos.z += float(_25);
|
||||
pos += vec4(_30);
|
||||
vec2 _56 = pos.xy + vec2(ivec2(_30.y, _30.x));
|
||||
vec2 _56 = pos.xy + vec2(_32);
|
||||
pos = vec4(_56.x, _56.y, pos.z, pos.w);
|
||||
gl_Position = pos;
|
||||
_4 = _30.y;
|
||||
_4 = _33;
|
||||
}
|
||||
|
||||
|
||||
98
shaders-hlsl/asm/vert/spec-constant-op-composite.asm.vert
Normal file
98
shaders-hlsl/asm/vert/spec-constant-op-composite.asm.vert
Normal file
@ -0,0 +1,98 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 1
|
||||
; Bound: 58
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability ClipDistance
|
||||
OpCapability CullDistance
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %4 "main" %52 %output
|
||||
OpSource GLSL 450
|
||||
OpName %4 "main"
|
||||
OpName %9 "pos"
|
||||
OpName %50 "gl_PerVertex"
|
||||
OpMemberName %50 0 "gl_Position"
|
||||
OpMemberName %50 1 "gl_PointSize"
|
||||
OpMemberName %50 2 "gl_ClipDistance"
|
||||
OpMemberName %50 3 "gl_CullDistance"
|
||||
OpName %52 ""
|
||||
OpDecorate %13 SpecId 201
|
||||
OpDecorate %24 SpecId 202
|
||||
OpMemberDecorate %50 0 BuiltIn Position
|
||||
OpMemberDecorate %50 1 BuiltIn PointSize
|
||||
OpMemberDecorate %50 2 BuiltIn ClipDistance
|
||||
OpMemberDecorate %50 3 BuiltIn CullDistance
|
||||
OpDecorate %50 Block
|
||||
OpDecorate %57 SpecId 200
|
||||
OpDecorate %output Flat
|
||||
OpDecorate %output Location 0
|
||||
%2 = OpTypeVoid
|
||||
%3 = OpTypeFunction %2
|
||||
%6 = OpTypeFloat 32
|
||||
%7 = OpTypeVector %6 4
|
||||
%8 = OpTypePointer Function %7
|
||||
%10 = OpConstant %6 0
|
||||
%11 = OpConstantComposite %7 %10 %10 %10 %10
|
||||
%12 = OpTypeInt 32 1
|
||||
%int_ptr = OpTypePointer Output %12
|
||||
%13 = OpSpecConstant %12 -10
|
||||
%14 = OpConstant %12 2
|
||||
%15 = OpSpecConstantOp %12 IAdd %13 %14
|
||||
%17 = OpTypeInt 32 0
|
||||
%18 = OpConstant %17 1
|
||||
%19 = OpTypePointer Function %6
|
||||
%24 = OpSpecConstant %17 100
|
||||
%25 = OpConstant %17 5
|
||||
%26 = OpSpecConstantOp %17 UMod %24 %25
|
||||
%28 = OpConstant %17 2
|
||||
%33 = OpConstant %12 20
|
||||
%34 = OpConstant %12 30
|
||||
%35 = OpTypeVector %12 4
|
||||
%36 = OpSpecConstantComposite %35 %33 %34 %15 %15
|
||||
%40 = OpTypeVector %12 2
|
||||
%41 = OpSpecConstantOp %40 VectorShuffle %36 %36 1 0
|
||||
%foo = OpSpecConstantOp %12 CompositeExtract %36 1
|
||||
%42 = OpTypeVector %6 2
|
||||
%49 = OpTypeArray %6 %18
|
||||
%50 = OpTypeStruct %7 %6 %49 %49
|
||||
%51 = OpTypePointer Output %50
|
||||
%52 = OpVariable %51 Output
|
||||
%output = OpVariable %int_ptr Output
|
||||
%53 = OpConstant %12 0
|
||||
%55 = OpTypePointer Output %7
|
||||
%57 = OpSpecConstant %6 3.14159
|
||||
%4 = OpFunction %2 None %3
|
||||
%5 = OpLabel
|
||||
%9 = OpVariable %8 Function
|
||||
OpStore %9 %11
|
||||
%16 = OpConvertSToF %6 %15
|
||||
%20 = OpAccessChain %19 %9 %18
|
||||
%21 = OpLoad %6 %20
|
||||
%22 = OpFAdd %6 %21 %16
|
||||
%23 = OpAccessChain %19 %9 %18
|
||||
OpStore %23 %22
|
||||
%27 = OpConvertUToF %6 %26
|
||||
%29 = OpAccessChain %19 %9 %28
|
||||
%30 = OpLoad %6 %29
|
||||
%31 = OpFAdd %6 %30 %27
|
||||
%32 = OpAccessChain %19 %9 %28
|
||||
OpStore %32 %31
|
||||
%37 = OpConvertSToF %7 %36
|
||||
%38 = OpLoad %7 %9
|
||||
%39 = OpFAdd %7 %38 %37
|
||||
OpStore %9 %39
|
||||
%43 = OpConvertSToF %42 %41
|
||||
%44 = OpLoad %7 %9
|
||||
%45 = OpVectorShuffle %42 %44 %44 0 1
|
||||
%46 = OpFAdd %42 %45 %43
|
||||
%47 = OpLoad %7 %9
|
||||
%48 = OpVectorShuffle %7 %47 %46 4 5 2 3
|
||||
OpStore %9 %48
|
||||
%54 = OpLoad %7 %9
|
||||
%56 = OpAccessChain %55 %52 %53
|
||||
OpStore %56 %54
|
||||
OpStore %output %foo
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
98
shaders-msl/asm/vert/spec-constant-op-composite.asm.vert
Normal file
98
shaders-msl/asm/vert/spec-constant-op-composite.asm.vert
Normal file
@ -0,0 +1,98 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 1
|
||||
; Bound: 58
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability ClipDistance
|
||||
OpCapability CullDistance
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %4 "main" %52 %output
|
||||
OpSource GLSL 450
|
||||
OpName %4 "main"
|
||||
OpName %9 "pos"
|
||||
OpName %50 "gl_PerVertex"
|
||||
OpMemberName %50 0 "gl_Position"
|
||||
OpMemberName %50 1 "gl_PointSize"
|
||||
OpMemberName %50 2 "gl_ClipDistance"
|
||||
OpMemberName %50 3 "gl_CullDistance"
|
||||
OpName %52 ""
|
||||
OpDecorate %13 SpecId 201
|
||||
OpDecorate %24 SpecId 202
|
||||
OpMemberDecorate %50 0 BuiltIn Position
|
||||
OpMemberDecorate %50 1 BuiltIn PointSize
|
||||
OpMemberDecorate %50 2 BuiltIn ClipDistance
|
||||
OpMemberDecorate %50 3 BuiltIn CullDistance
|
||||
OpDecorate %50 Block
|
||||
OpDecorate %57 SpecId 200
|
||||
OpDecorate %output Flat
|
||||
OpDecorate %output Location 0
|
||||
%2 = OpTypeVoid
|
||||
%3 = OpTypeFunction %2
|
||||
%6 = OpTypeFloat 32
|
||||
%7 = OpTypeVector %6 4
|
||||
%8 = OpTypePointer Function %7
|
||||
%10 = OpConstant %6 0
|
||||
%11 = OpConstantComposite %7 %10 %10 %10 %10
|
||||
%12 = OpTypeInt 32 1
|
||||
%int_ptr = OpTypePointer Output %12
|
||||
%13 = OpSpecConstant %12 -10
|
||||
%14 = OpConstant %12 2
|
||||
%15 = OpSpecConstantOp %12 IAdd %13 %14
|
||||
%17 = OpTypeInt 32 0
|
||||
%18 = OpConstant %17 1
|
||||
%19 = OpTypePointer Function %6
|
||||
%24 = OpSpecConstant %17 100
|
||||
%25 = OpConstant %17 5
|
||||
%26 = OpSpecConstantOp %17 UMod %24 %25
|
||||
%28 = OpConstant %17 2
|
||||
%33 = OpConstant %12 20
|
||||
%34 = OpConstant %12 30
|
||||
%35 = OpTypeVector %12 4
|
||||
%36 = OpSpecConstantComposite %35 %33 %34 %15 %15
|
||||
%40 = OpTypeVector %12 2
|
||||
%41 = OpSpecConstantOp %40 VectorShuffle %36 %36 1 0
|
||||
%foo = OpSpecConstantOp %12 CompositeExtract %36 1
|
||||
%42 = OpTypeVector %6 2
|
||||
%49 = OpTypeArray %6 %18
|
||||
%50 = OpTypeStruct %7 %6 %49 %49
|
||||
%51 = OpTypePointer Output %50
|
||||
%52 = OpVariable %51 Output
|
||||
%output = OpVariable %int_ptr Output
|
||||
%53 = OpConstant %12 0
|
||||
%55 = OpTypePointer Output %7
|
||||
%57 = OpSpecConstant %6 3.14159
|
||||
%4 = OpFunction %2 None %3
|
||||
%5 = OpLabel
|
||||
%9 = OpVariable %8 Function
|
||||
OpStore %9 %11
|
||||
%16 = OpConvertSToF %6 %15
|
||||
%20 = OpAccessChain %19 %9 %18
|
||||
%21 = OpLoad %6 %20
|
||||
%22 = OpFAdd %6 %21 %16
|
||||
%23 = OpAccessChain %19 %9 %18
|
||||
OpStore %23 %22
|
||||
%27 = OpConvertUToF %6 %26
|
||||
%29 = OpAccessChain %19 %9 %28
|
||||
%30 = OpLoad %6 %29
|
||||
%31 = OpFAdd %6 %30 %27
|
||||
%32 = OpAccessChain %19 %9 %28
|
||||
OpStore %32 %31
|
||||
%37 = OpConvertSToF %7 %36
|
||||
%38 = OpLoad %7 %9
|
||||
%39 = OpFAdd %7 %38 %37
|
||||
OpStore %9 %39
|
||||
%43 = OpConvertSToF %42 %41
|
||||
%44 = OpLoad %7 %9
|
||||
%45 = OpVectorShuffle %42 %44 %44 0 1
|
||||
%46 = OpFAdd %42 %45 %43
|
||||
%47 = OpLoad %7 %9
|
||||
%48 = OpVectorShuffle %7 %47 %46 4 5 2 3
|
||||
OpStore %9 %48
|
||||
%54 = OpLoad %7 %9
|
||||
%56 = OpAccessChain %55 %52 %53
|
||||
OpStore %56 %54
|
||||
OpStore %output %foo
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -2039,6 +2039,7 @@ void Compiler::parse(const Instruction &instruction)
|
||||
remapped_constant_ops[i].make_null(get<SPIRType>(constant_op->basetype));
|
||||
remapped_constant_ops[i].self = constant_op->self;
|
||||
remapped_constant_ops[i].constant_type = constant_op->basetype;
|
||||
remapped_constant_ops[i].specialization = true;
|
||||
c[i] = &remapped_constant_ops[i];
|
||||
}
|
||||
else
|
||||
|
||||
@ -1688,6 +1688,13 @@ void CompilerGLSL::emit_uniform(const SPIRVariable &var)
|
||||
statement(layout_for_variable(var), variable_decl(var), ";");
|
||||
}
|
||||
|
||||
void CompilerGLSL::emit_specialization_constant_op(const SPIRConstantOp &constant)
|
||||
{
|
||||
auto &type = get<SPIRType>(constant.basetype);
|
||||
auto name = to_name(constant.self);
|
||||
statement("const ", variable_decl(type, name), " = ", constant_op_expression(constant), ";");
|
||||
}
|
||||
|
||||
void CompilerGLSL::emit_specialization_constant(const SPIRConstant &constant)
|
||||
{
|
||||
auto &type = get<SPIRType>(constant.constant_type);
|
||||
@ -2121,6 +2128,11 @@ void CompilerGLSL::emit_resources()
|
||||
emit_specialization_constant(c);
|
||||
emitted = true;
|
||||
}
|
||||
else if (id.get_type() == TypeConstantOp)
|
||||
{
|
||||
emit_specialization_constant_op(id.get<SPIRConstantOp>());
|
||||
emitted = true;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -2433,7 +2445,10 @@ string CompilerGLSL::to_expression(uint32_t id)
|
||||
}
|
||||
|
||||
case TypeConstantOp:
|
||||
return constant_op_expression(get<SPIRConstantOp>(id));
|
||||
if (options.vulkan_semantics)
|
||||
return to_name(id);
|
||||
else
|
||||
return constant_op_expression(get<SPIRConstantOp>(id));
|
||||
|
||||
case TypeVariable:
|
||||
{
|
||||
|
||||
@ -371,6 +371,7 @@ protected:
|
||||
void emit_block_chain(SPIRBlock &block);
|
||||
void emit_hoisted_temporaries(std::vector<std::pair<uint32_t, uint32_t>> &temporaries);
|
||||
void emit_specialization_constant(const SPIRConstant &constant);
|
||||
void emit_specialization_constant_op(const SPIRConstantOp &constant);
|
||||
std::string emit_continue_block(uint32_t continue_block);
|
||||
bool attempt_emit_loop_header(SPIRBlock &block, SPIRBlock::Method method);
|
||||
void propagate_loop_dominators(const SPIRBlock &block);
|
||||
|
||||
@ -1034,6 +1034,13 @@ void CompilerHLSL::emit_specialization_constants()
|
||||
statement("static const ", variable_decl(type, name), " = ", constant_expression(c), ";");
|
||||
emitted = true;
|
||||
}
|
||||
else if (id.get_type() == TypeConstantOp)
|
||||
{
|
||||
auto &c = id.get<SPIRConstantOp>();
|
||||
auto &type = get<SPIRType>(c.basetype);
|
||||
auto name = to_name(c.self);
|
||||
statement("static const ", variable_decl(type, name), " = ", constant_op_expression(c), ";");
|
||||
}
|
||||
}
|
||||
|
||||
if (workgroup_size_id)
|
||||
|
||||
@ -1559,35 +1559,51 @@ void CompilerMSL::emit_resources()
|
||||
// Emit declarations for the specialization Metal function constants
|
||||
void CompilerMSL::emit_specialization_constants()
|
||||
{
|
||||
const vector<SpecializationConstant> spec_consts = get_specialization_constants();
|
||||
|
||||
SpecializationConstant wg_x, wg_y, wg_z;
|
||||
uint32_t workgroup_size_id = get_work_group_size_specialization_constants(wg_x, wg_y, wg_z);
|
||||
bool emitted = false;
|
||||
|
||||
for (auto &sc : spec_consts)
|
||||
for (auto &id : ids)
|
||||
{
|
||||
// If WorkGroupSize is a specialization constant, it will be declared explicitly below.
|
||||
if (sc.id == workgroup_size_id)
|
||||
continue;
|
||||
|
||||
auto &type = expression_type(sc.id);
|
||||
string sc_type_name = type_to_glsl(type);
|
||||
string sc_name = to_name(sc.id);
|
||||
string sc_tmp_name = to_name(sc.id) + "_tmp";
|
||||
|
||||
if (type.vecsize == 1 && type.columns == 1 && type.basetype != SPIRType::Struct && type.array.empty())
|
||||
if (id.get_type() == TypeConstant)
|
||||
{
|
||||
// Only scalar, non-composite values can be function constants.
|
||||
statement("constant ", sc_type_name, " ", sc_tmp_name, " [[function_constant(",
|
||||
convert_to_string(sc.constant_id), ")]];");
|
||||
statement("constant ", sc_type_name, " ", sc_name, " = is_function_constant_defined(", sc_tmp_name, ") ? ",
|
||||
sc_tmp_name, " : ", constant_expression(get<SPIRConstant>(sc.id)), ";");
|
||||
auto &c = id.get<SPIRConstant>();
|
||||
if (!c.specialization)
|
||||
continue;
|
||||
|
||||
// If WorkGroupSize is a specialization constant, it will be declared explicitly below.
|
||||
if (c.self == workgroup_size_id)
|
||||
continue;
|
||||
|
||||
auto &type = get<SPIRType>(c.constant_type);
|
||||
string sc_type_name = type_to_glsl(type);
|
||||
string sc_name = to_name(c.self);
|
||||
string sc_tmp_name = sc_name + "_tmp";
|
||||
|
||||
if (has_decoration(c.self, DecorationSpecId))
|
||||
{
|
||||
uint32_t constant_id = get_decoration(c.self, DecorationSpecId);
|
||||
// Only scalar, non-composite values can be function constants.
|
||||
statement("constant ", sc_type_name, " ", sc_tmp_name, " [[function_constant(",
|
||||
constant_id, ")]];");
|
||||
statement("constant ", sc_type_name, " ", sc_name, " = is_function_constant_defined(", sc_tmp_name,
|
||||
") ? ",
|
||||
sc_tmp_name, " : ", constant_expression(c), ";");
|
||||
}
|
||||
else
|
||||
{
|
||||
// Composite specialization constants must be built from other specialization constants.
|
||||
statement("constant ", sc_type_name, " ", sc_name, " = ", constant_expression(c), ";");
|
||||
}
|
||||
emitted = true;
|
||||
}
|
||||
else
|
||||
else if (id.get_type() == TypeConstantOp)
|
||||
{
|
||||
// Composite specialization constants must be built from other specialization constants.
|
||||
statement("constant ", sc_type_name, " ", sc_name, " = ", constant_expression(get<SPIRConstant>(sc.id)),
|
||||
";");
|
||||
auto &c = id.get<SPIRConstantOp>();
|
||||
auto &type = get<SPIRType>(c.basetype);
|
||||
auto name = to_name(c.self);
|
||||
statement("constant ", variable_decl(type, name), " = ", constant_op_expression(c), ";");
|
||||
emitted = true;
|
||||
}
|
||||
}
|
||||
|
||||
@ -1595,10 +1611,13 @@ void CompilerMSL::emit_specialization_constants()
|
||||
// the work group size at compile time in SPIR-V, and [[threads_per_threadgroup]] would need to be passed around as a global.
|
||||
// The work group size may be a specialization constant.
|
||||
if (workgroup_size_id)
|
||||
{
|
||||
statement("constant uint3 ", builtin_to_glsl(BuiltInWorkgroupSize, StorageClassWorkgroup), " = ",
|
||||
constant_expression(get<SPIRConstant>(workgroup_size_id)), ";");
|
||||
emitted = true;
|
||||
}
|
||||
|
||||
if (!spec_consts.empty() || workgroup_size_id)
|
||||
if (emitted)
|
||||
statement("");
|
||||
}
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user