Update tests.

This commit is contained in:
Hans-Kristian Arntzen 2018-11-01 11:22:14 +01:00
parent fd6ff3617a
commit 62db535b3f
36 changed files with 278 additions and 125 deletions

View File

@ -1,5 +1,11 @@
static const uint _3 = 1u;
static const uint _4 = 3u;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 1u
#endif
static const uint _3 = SPIRV_CROSS_CONSTANT_ID_0;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 3u
#endif
static const uint _4 = SPIRV_CROSS_CONSTANT_ID_2;
static const uint3 gl_WorkGroupSize = uint3(_3, 2u, _4);
RWByteAddressBuffer _8 : register(u0);
@ -18,7 +24,7 @@ void comp_main()
_8.Store(gl_WorkGroupID.x * 4 + 0, asuint(asfloat(_9.Load(gl_WorkGroupID.x * 4 + 0)) + asfloat(_8.Load(gl_WorkGroupID.x * 4 + 0))));
}
[numthreads(1, 2, 3)]
[numthreads(SPIRV_CROSS_CONSTANT_ID_0, 2, SPIRV_CROSS_CONSTANT_ID_2)]
void main(SPIRV_Cross_Input stage_input)
{
gl_WorkGroupID = stage_input.gl_WorkGroupID;

View File

@ -1,8 +1,17 @@
static const int _7 = -10;
static const uint _8 = 100u;
#ifndef SPIRV_CROSS_CONSTANT_ID_201
#define SPIRV_CROSS_CONSTANT_ID_201 -10
#endif
static const int _7 = SPIRV_CROSS_CONSTANT_ID_201;
#ifndef SPIRV_CROSS_CONSTANT_ID_202
#define SPIRV_CROSS_CONSTANT_ID_202 100u
#endif
static const uint _8 = SPIRV_CROSS_CONSTANT_ID_202;
static const int _20 = (_7 + 2);
static const uint _25 = (_8 % 5u);
static const int4 _30 = int4(20, 30, _20, _20);
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 int4(20, 30, _20, _20)
#endif
static const int4 _30 = SPIRV_CROSS_CONSTANT_ID_0;
static const int2 _32 = int2(_30.y, _30.x);
static const int _33 = _30.y;

View File

@ -1,4 +1,7 @@
static const float X = 4.0f;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 4.0f
#endif
static const float X = SPIRV_CROSS_CONSTANT_ID_0;
static const uint3 gl_WorkGroupSize = uint3(2u, 1u, 1u);
struct Data

View File

@ -1,4 +1,7 @@
static const int Value = 2;
#ifndef SPIRV_CROSS_CONSTANT_ID_10
#define SPIRV_CROSS_CONSTANT_ID_10 2
#endif
static const int Value = SPIRV_CROSS_CONSTANT_ID_10;
cbuffer SpecConstArray : register(b0)
{

View File

@ -1,4 +1,7 @@
static const uint s = 10u;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 10u
#endif
static const uint s = SPIRV_CROSS_CONSTANT_ID_0;
static const bool _13 = (s > 20u);
static const uint _16 = _13 ? 30u : 50u;

View File

@ -3,9 +3,14 @@
using namespace metal;
#ifndef SPIRV_CROSS_CONSTANT_ID_10
#define SPIRV_CROSS_CONSTANT_ID_10 2
#endif
constant int Value = SPIRV_CROSS_CONSTANT_ID_10;
struct SpecConstArray
{
float4 samples[2];
float4 samples[Value];
};
struct main0_out

View File

@ -3,6 +3,11 @@
using namespace metal;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 3
#endif
constant int arraySize = SPIRV_CROSS_CONSTANT_ID_0;
struct storage_block
{
uint4 baz;

View File

@ -3,6 +3,11 @@
using namespace metal;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 3
#endif
constant int arraySize = SPIRV_CROSS_CONSTANT_ID_0;
struct storage_block
{
uint4 baz;

View File

@ -3,8 +3,14 @@
using namespace metal;
constant float a = 1.0;
constant float b = 2.0;
#ifndef SPIRV_CROSS_CONSTANT_ID_1
#define SPIRV_CROSS_CONSTANT_ID_1 1.0
#endif
constant float a = SPIRV_CROSS_CONSTANT_ID_1;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 2.0
#endif
constant float b = SPIRV_CROSS_CONSTANT_ID_2;
struct main0_out
{

View File

@ -1,6 +1,13 @@
#version 310 es
layout(local_size_x = 9, local_size_y = 20, local_size_z = 4) in;
#ifndef SPIRV_CROSS_CONSTANT_ID_10
#define SPIRV_CROSS_CONSTANT_ID_10 9u
#endif
#ifndef SPIRV_CROSS_CONSTANT_ID_12
#define SPIRV_CROSS_CONSTANT_ID_12 4u
#endif
layout(local_size_x = SPIRV_CROSS_CONSTANT_ID_10, local_size_y = 20, local_size_z = SPIRV_CROSS_CONSTANT_ID_12) in;
layout(binding = 0, std430) buffer SSBO
{

View File

@ -1,6 +1,13 @@
#version 450
layout(local_size_x = 1, local_size_y = 2, local_size_z = 3) in;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 1u
#endif
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 3u
#endif
layout(local_size_x = SPIRV_CROSS_CONSTANT_ID_0, local_size_y = 2, local_size_z = SPIRV_CROSS_CONSTANT_ID_2) in;
layout(binding = 0, std430) buffer _6_8
{

View File

@ -4,12 +4,10 @@
#define SPIRV_CROSS_CONSTANT_ID_201 -10
#endif
const int _7 = SPIRV_CROSS_CONSTANT_ID_201;
#ifndef SPIRV_CROSS_CONSTANT_ID_202
#define SPIRV_CROSS_CONSTANT_ID_202 100u
#endif
const uint _8 = SPIRV_CROSS_CONSTANT_ID_202;
const int _20 = (_7 + 2);
const uint _25 = (_8 % 5u);
const ivec4 _30 = ivec4(20, 30, _20, _20);

View File

@ -6,7 +6,6 @@ layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;
#endif
const float X = SPIRV_CROSS_CONSTANT_ID_0;
struct Data
{
float a;

View File

@ -7,7 +7,6 @@ precision highp int;
#endif
const int Value = SPIRV_CROSS_CONSTANT_ID_10;
layout(binding = 0, std140) uniform SpecConstArray
{
vec4 samples[Value];

View File

@ -4,7 +4,6 @@
#define SPIRV_CROSS_CONSTANT_ID_0 10u
#endif
const uint s = SPIRV_CROSS_CONSTANT_ID_0;
const bool _13 = (s > 20u);
const uint _16 = _13 ? 30u : 50u;

View File

@ -1,5 +1,11 @@
static const uint _5 = 9u;
static const uint _6 = 4u;
#ifndef SPIRV_CROSS_CONSTANT_ID_10
#define SPIRV_CROSS_CONSTANT_ID_10 9u
#endif
static const uint _5 = SPIRV_CROSS_CONSTANT_ID_10;
#ifndef SPIRV_CROSS_CONSTANT_ID_12
#define SPIRV_CROSS_CONSTANT_ID_12 4u
#endif
static const uint _6 = SPIRV_CROSS_CONSTANT_ID_12;
static const uint3 gl_WorkGroupSize = uint3(_5, 20u, _6);
RWByteAddressBuffer _4 : register(u0);
@ -9,7 +15,7 @@ void comp_main()
_4.Store(0, asuint(asfloat(_4.Load(0)) + 1.0f));
}
[numthreads(9, 20, 4)]
[numthreads(SPIRV_CROSS_CONSTANT_ID_10, 20, SPIRV_CROSS_CONSTANT_ID_12)]
void main()
{
comp_main();

View File

@ -1,16 +1,34 @@
static const float a = 1.0f;
static const float b = 2.0f;
static const int c = 3;
#ifndef SPIRV_CROSS_CONSTANT_ID_1
#define SPIRV_CROSS_CONSTANT_ID_1 1.0f
#endif
static const float a = SPIRV_CROSS_CONSTANT_ID_1;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 2.0f
#endif
static const float b = SPIRV_CROSS_CONSTANT_ID_2;
#ifndef SPIRV_CROSS_CONSTANT_ID_3
#define SPIRV_CROSS_CONSTANT_ID_3 3
#endif
static const int c = SPIRV_CROSS_CONSTANT_ID_3;
static const uint _18 = (uint(c) + 0u);
static const int _21 = (-c);
static const int _23 = (~c);
static const int d = 4;
#ifndef SPIRV_CROSS_CONSTANT_ID_4
#define SPIRV_CROSS_CONSTANT_ID_4 4
#endif
static const int d = SPIRV_CROSS_CONSTANT_ID_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;
#ifndef SPIRV_CROSS_CONSTANT_ID_5
#define SPIRV_CROSS_CONSTANT_ID_5 5u
#endif
static const uint e = SPIRV_CROSS_CONSTANT_ID_5;
#ifndef SPIRV_CROSS_CONSTANT_ID_6
#define SPIRV_CROSS_CONSTANT_ID_6 6u
#endif
static const uint f = SPIRV_CROSS_CONSTANT_ID_6;
static const uint _36 = (e / f);
static const int _38 = (c % d);
static const uint _40 = (e % f);
@ -20,8 +38,14 @@ 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;
#ifndef SPIRV_CROSS_CONSTANT_ID_7
#define SPIRV_CROSS_CONSTANT_ID_7 false
#endif
static const bool g = SPIRV_CROSS_CONSTANT_ID_7;
#ifndef SPIRV_CROSS_CONSTANT_ID_8
#define SPIRV_CROSS_CONSTANT_ID_8 true
#endif
static const bool h = SPIRV_CROSS_CONSTANT_ID_8;
static const bool _58 = (g || h);
static const bool _60 = (g && h);
static const bool _62 = (!g);

View File

@ -1,5 +1,11 @@
static const uint _3 = 1u;
static const uint _4 = 3u;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 1u
#endif
static const uint _3 = SPIRV_CROSS_CONSTANT_ID_0;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 3u
#endif
static const uint _4 = SPIRV_CROSS_CONSTANT_ID_2;
static const uint3 gl_WorkGroupSize = uint3(_3, 2u, _4);
RWByteAddressBuffer _8 : register(u0);
@ -18,7 +24,7 @@ void comp_main()
_8.Store(gl_WorkGroupID.x * 4 + 0, asuint(asfloat(_9.Load(gl_WorkGroupID.x * 4 + 0)) + asfloat(_8.Load(gl_WorkGroupID.x * 4 + 0))));
}
[numthreads(1, 2, 3)]
[numthreads(SPIRV_CROSS_CONSTANT_ID_0, 2, SPIRV_CROSS_CONSTANT_ID_2)]
void main(SPIRV_Cross_Input stage_input)
{
gl_WorkGroupID = stage_input.gl_WorkGroupID;

View File

@ -1,9 +1,21 @@
static const int _7 = -10;
static const uint _8 = 100u;
static const float _9 = 3.141590118408203125f;
#ifndef SPIRV_CROSS_CONSTANT_ID_201
#define SPIRV_CROSS_CONSTANT_ID_201 -10
#endif
static const int _7 = SPIRV_CROSS_CONSTANT_ID_201;
#ifndef SPIRV_CROSS_CONSTANT_ID_202
#define SPIRV_CROSS_CONSTANT_ID_202 100u
#endif
static const uint _8 = SPIRV_CROSS_CONSTANT_ID_202;
#ifndef SPIRV_CROSS_CONSTANT_ID_200
#define SPIRV_CROSS_CONSTANT_ID_200 3.141590118408203125f
#endif
static const float _9 = SPIRV_CROSS_CONSTANT_ID_200;
static const int _20 = (_7 + 2);
static const uint _25 = (_8 % 5u);
static const int4 _30 = int4(20, 30, _20, _20);
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 int4(20, 30, _20, _20)
#endif
static const int4 _30 = SPIRV_CROSS_CONSTANT_ID_0;
static const int2 _32 = int2(_30.y, _30.x);
static const int _33 = _30.y;

View File

@ -1,4 +1,7 @@
static const float X = 4.0f;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 4.0f
#endif
static const float X = SPIRV_CROSS_CONSTANT_ID_0;
static const uint3 gl_WorkGroupSize = uint3(2u, 1u, 1u);
struct Data

View File

@ -1,4 +1,7 @@
static const int Value = 2;
#ifndef SPIRV_CROSS_CONSTANT_ID_10
#define SPIRV_CROSS_CONSTANT_ID_10 2
#endif
static const int Value = SPIRV_CROSS_CONSTANT_ID_10;
cbuffer SpecConstArray : register(b0)
{

View File

@ -1,4 +1,7 @@
static const uint s = 10u;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 10u
#endif
static const uint s = SPIRV_CROSS_CONSTANT_ID_0;
static const bool _13 = (s > 20u);
static const uint _16 = _13 ? 30u : 50u;

View File

@ -3,9 +3,14 @@
using namespace metal;
#ifndef SPIRV_CROSS_CONSTANT_ID_10
#define SPIRV_CROSS_CONSTANT_ID_10 2
#endif
constant int Value = SPIRV_CROSS_CONSTANT_ID_10;
struct SpecConstArray
{
float4 samples[2];
float4 samples[Value];
};
struct main0_out

View File

@ -5,6 +5,11 @@
using namespace metal;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 3
#endif
constant int arraySize = SPIRV_CROSS_CONSTANT_ID_0;
struct storage_block
{
uint4 baz;

View File

@ -3,6 +3,11 @@
using namespace metal;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 3
#endif
constant int arraySize = SPIRV_CROSS_CONSTANT_ID_0;
struct storage_block
{
uint4 baz;

View File

@ -3,19 +3,37 @@
using namespace metal;
constant float a = 1.0;
constant float b = 2.0;
constant int c = 3;
#ifndef SPIRV_CROSS_CONSTANT_ID_1
#define SPIRV_CROSS_CONSTANT_ID_1 1.0
#endif
constant float a = SPIRV_CROSS_CONSTANT_ID_1;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 2.0
#endif
constant float b = SPIRV_CROSS_CONSTANT_ID_2;
#ifndef SPIRV_CROSS_CONSTANT_ID_3
#define SPIRV_CROSS_CONSTANT_ID_3 3
#endif
constant int c = SPIRV_CROSS_CONSTANT_ID_3;
constant uint _18 = (uint(c) + 0u);
constant int _21 = (-c);
constant int _23 = (~c);
constant int d = 4;
#ifndef SPIRV_CROSS_CONSTANT_ID_4
#define SPIRV_CROSS_CONSTANT_ID_4 4
#endif
constant int d = SPIRV_CROSS_CONSTANT_ID_4;
constant int _26 = (c + d);
constant int _28 = (c - d);
constant int _30 = (c * d);
constant int _32 = (c / d);
constant uint e = 5u;
constant uint f = 6u;
#ifndef SPIRV_CROSS_CONSTANT_ID_5
#define SPIRV_CROSS_CONSTANT_ID_5 5u
#endif
constant uint e = SPIRV_CROSS_CONSTANT_ID_5;
#ifndef SPIRV_CROSS_CONSTANT_ID_6
#define SPIRV_CROSS_CONSTANT_ID_6 6u
#endif
constant uint f = SPIRV_CROSS_CONSTANT_ID_6;
constant uint _36 = (e / f);
constant int _38 = (c % d);
constant uint _40 = (e % f);
@ -25,8 +43,14 @@ constant int _46 = (c << d);
constant int _48 = (c | d);
constant int _50 = (c ^ d);
constant int _52 = (c & d);
constant bool g = false;
constant bool h = true;
#ifndef SPIRV_CROSS_CONSTANT_ID_7
#define SPIRV_CROSS_CONSTANT_ID_7 false
#endif
constant bool g = SPIRV_CROSS_CONSTANT_ID_7;
#ifndef SPIRV_CROSS_CONSTANT_ID_8
#define SPIRV_CROSS_CONSTANT_ID_8 true
#endif
constant bool h = SPIRV_CROSS_CONSTANT_ID_8;
constant bool _58 = (g || h);
constant bool _60 = (g && h);
constant bool _62 = (!g);

View File

@ -6,17 +6,14 @@ precision highp int;
#define SPIRV_CROSS_CONSTANT_ID_1 1.0
#endif
const float a = SPIRV_CROSS_CONSTANT_ID_1;
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 2.0
#endif
const float b = SPIRV_CROSS_CONSTANT_ID_2;
#ifndef SPIRV_CROSS_CONSTANT_ID_3
#define SPIRV_CROSS_CONSTANT_ID_3 3
#endif
const int c = SPIRV_CROSS_CONSTANT_ID_3;
const uint _18 = (uint(c) + 0u);
const int _21 = (-c);
const int _23 = (~c);
@ -24,7 +21,6 @@ const int _23 = (~c);
#define SPIRV_CROSS_CONSTANT_ID_4 4
#endif
const int d = SPIRV_CROSS_CONSTANT_ID_4;
const int _26 = (c + d);
const int _28 = (c - d);
const int _30 = (c * d);
@ -33,12 +29,10 @@ const int _32 = (c / d);
#define SPIRV_CROSS_CONSTANT_ID_5 5u
#endif
const uint e = SPIRV_CROSS_CONSTANT_ID_5;
#ifndef SPIRV_CROSS_CONSTANT_ID_6
#define SPIRV_CROSS_CONSTANT_ID_6 6u
#endif
const uint f = SPIRV_CROSS_CONSTANT_ID_6;
const uint _36 = (e / f);
const int _38 = (c % d);
const uint _40 = (e % f);
@ -52,12 +46,10 @@ const int _52 = (c & d);
#define SPIRV_CROSS_CONSTANT_ID_7 false
#endif
const bool g = SPIRV_CROSS_CONSTANT_ID_7;
#ifndef SPIRV_CROSS_CONSTANT_ID_8
#define SPIRV_CROSS_CONSTANT_ID_8 true
#endif
const bool h = SPIRV_CROSS_CONSTANT_ID_8;
const bool _58 = (g || h);
const bool _60 = (g && h);
const bool _62 = (!g);

View File

@ -1,6 +1,13 @@
#version 310 es
layout(local_size_x = 9, local_size_y = 20, local_size_z = 4) in;
#ifndef SPIRV_CROSS_CONSTANT_ID_10
#define SPIRV_CROSS_CONSTANT_ID_10 9u
#endif
#ifndef SPIRV_CROSS_CONSTANT_ID_12
#define SPIRV_CROSS_CONSTANT_ID_12 4u
#endif
layout(local_size_x = SPIRV_CROSS_CONSTANT_ID_10, local_size_y = 20, local_size_z = SPIRV_CROSS_CONSTANT_ID_12) in;
layout(binding = 0, std430) buffer SSBO
{

View File

@ -1,6 +1,13 @@
#version 450
layout(local_size_x = 1, local_size_y = 2, local_size_z = 3) in;
#ifndef SPIRV_CROSS_CONSTANT_ID_0
#define SPIRV_CROSS_CONSTANT_ID_0 1u
#endif
#ifndef SPIRV_CROSS_CONSTANT_ID_2
#define SPIRV_CROSS_CONSTANT_ID_2 3u
#endif
layout(local_size_x = SPIRV_CROSS_CONSTANT_ID_0, local_size_y = 2, local_size_z = SPIRV_CROSS_CONSTANT_ID_2) in;
layout(binding = 0, std430) buffer _6_8
{

View File

@ -4,17 +4,14 @@
#define SPIRV_CROSS_CONSTANT_ID_201 -10
#endif
const int _7 = SPIRV_CROSS_CONSTANT_ID_201;
#ifndef SPIRV_CROSS_CONSTANT_ID_202
#define SPIRV_CROSS_CONSTANT_ID_202 100u
#endif
const uint _8 = SPIRV_CROSS_CONSTANT_ID_202;
#ifndef SPIRV_CROSS_CONSTANT_ID_200
#define SPIRV_CROSS_CONSTANT_ID_200 3.141590118408203125
#endif
const float _9 = SPIRV_CROSS_CONSTANT_ID_200;
const int _20 = (_7 + 2);
const uint _25 = (_8 % 5u);
const ivec4 _30 = ivec4(20, 30, _20, _20);

View File

@ -6,7 +6,6 @@ layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;
#endif
const float X = SPIRV_CROSS_CONSTANT_ID_0;
struct Data
{
float a;

View File

@ -7,7 +7,6 @@ precision highp int;
#endif
const int Value = SPIRV_CROSS_CONSTANT_ID_10;
layout(binding = 0, std140) uniform SpecConstArray
{
vec4 samples[Value];

View File

@ -4,7 +4,6 @@
#define SPIRV_CROSS_CONSTANT_ID_0 10u
#endif
const uint s = SPIRV_CROSS_CONSTANT_ID_0;
const bool _13 = (s > 20u);
const uint _16 = _13 ? 30u : 50u;

View File

@ -1025,23 +1025,29 @@ void CompilerHLSL::emit_specialization_constants()
if (id.get_type() == TypeConstant)
{
auto &c = id.get<SPIRConstant>();
if (!c.specialization)
continue;
if (c.self == workgroup_size_id)
continue;
{
statement("static const uint3 gl_WorkGroupSize = ",
constant_expression(get<SPIRConstant>(workgroup_size_id)),
";");
emitted = true;
}
else if (c.specialization)
{
auto &type = get<SPIRType>(c.constant_type);
auto name = to_name(c.self);
auto &type = get<SPIRType>(c.constant_type);
auto name = to_name(c.self);
// HLSL does not support specialization constants, so fallback to macros.
c.specialization_constant_macro_name =
constant_value_macro_name(get_decoration(c.self, DecorationSpecId));
// HLSL does not support specialization constants, so fallback to macros.
c.specialization_constant_macro_name =
constant_value_macro_name(get_decoration(c.self, DecorationSpecId));
statement("#ifndef ", c.specialization_constant_macro_name);
statement("#define ", c.specialization_constant_macro_name, " ", constant_expression(c));
statement("#endif");
statement("static const ", variable_decl(type, name), " = ", c.specialization_constant_macro_name, ";");
emitted = true;
statement("#ifndef ", c.specialization_constant_macro_name);
statement("#define ", c.specialization_constant_macro_name, " ", constant_expression(c));
statement("#endif");
statement("static const ", variable_decl(type, name), " = ", c.specialization_constant_macro_name, ";");
emitted = true;
}
}
else if (id.get_type() == TypeConstantOp)
{
@ -1052,13 +1058,6 @@ void CompilerHLSL::emit_specialization_constants()
}
}
if (workgroup_size_id)
{
statement("static const uint3 gl_WorkGroupSize = ", constant_expression(get<SPIRConstant>(workgroup_size_id)),
";");
emitted = true;
}
if (emitted)
statement("");
}

View File

@ -443,11 +443,6 @@ string CompilerMSL::compile()
// Mark any non-stage-in structs to be tightly packed.
mark_packable_structs();
// Metal does not allow dynamic array lengths.
// Resolve any specialization constants that are used for array lengths.
if (msl_options.resolve_specialized_array_lengths)
resolve_specialized_array_lengths();
uint32_t pass_count = 0;
do
{
@ -1927,35 +1922,55 @@ void CompilerMSL::emit_specialization_constants()
if (id.get_type() == TypeConstant)
{
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";
// Function constants are only supported in MSL 1.2 and later.
// If we don't support it just declare the "default" directly.
// This "default" value can be overridden to the true specialization constant by the API user.
if (msl_options.supports_msl_version(1, 2) && 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), ";");
// TODO: This can be expressed as a [[threads_per_threadgroup]] input semantic, but we need to know
// 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.
statement("constant uint3 ", builtin_to_glsl(BuiltInWorkgroupSize, StorageClassWorkgroup), " = ",
constant_expression(get<SPIRConstant>(workgroup_size_id)), ";");
emitted = true;
}
else
else if (c.specialization)
{
// Composite specialization constants must be built from other specialization constants.
statement("constant ", sc_type_name, " ", sc_name, " = ", constant_expression(c), ";");
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";
// Function constants are only supported in MSL 1.2 and later.
// If we don't support it just declare the "default" directly.
// This "default" value can be overridden to the true specialization constant by the API user.
// Specialization constants which are used as array length expressions cannot be function constants in MSL,
// so just fall back to macros.
if (msl_options.supports_msl_version(1, 2) && has_decoration(c.self, DecorationSpecId) && !c.is_used_as_array_length)
{
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 if (has_decoration(c.self, DecorationSpecId))
{
// Fallback to macro overrides.
c.specialization_constant_macro_name =
constant_value_macro_name(get_decoration(c.self, DecorationSpecId));
statement("#ifndef ", c.specialization_constant_macro_name);
statement("#define ", c.specialization_constant_macro_name, " ", constant_expression(c));
statement("#endif");
statement("constant ", sc_type_name, " ", sc_name, " = ", c.specialization_constant_macro_name, ";");
}
else
{
// Composite specialization constants must be built from other specialization constants.
statement("constant ", sc_type_name, " ", sc_name, " = ", constant_expression(c), ";");
}
emitted = true;
}
emitted = true;
}
else if (id.get_type() == TypeConstantOp)
{
@ -1967,16 +1982,6 @@ void CompilerMSL::emit_specialization_constants()
}
}
// TODO: This can be expressed as a [[threads_per_threadgroup]] input semantic, but we need to know
// 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 (emitted)
statement("");
}

View File

@ -155,7 +155,6 @@ public:
uint32_t aux_buffer_index = 0;
bool enable_point_size_builtin = true;
bool disable_rasterization = false;
bool resolve_specialized_array_lengths = true;
bool swizzle_texture_samples = false;
bool is_ios()