Merge pull request #686 from KhronosGroup/fix-685
Fix composite array declarations in MSL.
This commit is contained in:
commit
6ca16b2a5e
@ -0,0 +1,50 @@
|
|||||||
|
static const float X = 4.0f;
|
||||||
|
static const uint3 gl_WorkGroupSize = uint3(2u, 1u, 1u);
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
static const Data _21 = { 1.0f, 2.0f };
|
||||||
|
static const Data _24 = { 3.0f, 4.0f };
|
||||||
|
static const Data _25[2] = { { 1.0f, 2.0f }, { 3.0f, 4.0f } };
|
||||||
|
static const Data _30 = { 3.0f, 5.0f };
|
||||||
|
|
||||||
|
RWByteAddressBuffer _61 : register(u0);
|
||||||
|
|
||||||
|
static uint3 gl_WorkGroupID;
|
||||||
|
static uint3 gl_LocalInvocationID;
|
||||||
|
static uint gl_LocalInvocationIndex;
|
||||||
|
struct SPIRV_Cross_Input
|
||||||
|
{
|
||||||
|
uint3 gl_WorkGroupID : SV_GroupID;
|
||||||
|
uint3 gl_LocalInvocationID : SV_GroupThreadID;
|
||||||
|
uint gl_LocalInvocationIndex : SV_GroupIndex;
|
||||||
|
};
|
||||||
|
|
||||||
|
static Data data[2];
|
||||||
|
static Data data2[2];
|
||||||
|
|
||||||
|
void comp_main()
|
||||||
|
{
|
||||||
|
data = _25;
|
||||||
|
Data _28 = { X, 2.0f };
|
||||||
|
Data _31[2] = { _28, _30 };
|
||||||
|
data2 = _31;
|
||||||
|
if (gl_LocalInvocationIndex == 0u)
|
||||||
|
{
|
||||||
|
_61.Store(gl_WorkGroupID.x * 8 + 0, asuint(data[gl_LocalInvocationID.x].a + data2[gl_LocalInvocationID.x].a));
|
||||||
|
_61.Store(gl_WorkGroupID.x * 8 + 4, asuint(data[gl_LocalInvocationID.x].b + data2[gl_LocalInvocationID.x].b));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
[numthreads(2, 1, 1)]
|
||||||
|
void main(SPIRV_Cross_Input stage_input)
|
||||||
|
{
|
||||||
|
gl_WorkGroupID = stage_input.gl_WorkGroupID;
|
||||||
|
gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
|
||||||
|
gl_LocalInvocationIndex = stage_input.gl_LocalInvocationIndex;
|
||||||
|
comp_main();
|
||||||
|
}
|
@ -0,0 +1,54 @@
|
|||||||
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
constant float X_tmp [[function_constant(0)]];
|
||||||
|
constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
|
||||||
|
constant uint3 gl_WorkGroupSize = uint3(2u, 1u, 1u);
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct Data_1
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct SSBO
|
||||||
|
{
|
||||||
|
Data_1 outdata[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
|
||||||
|
|
||||||
|
// Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
|
||||||
|
template<typename T, uint N>
|
||||||
|
void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
|
||||||
|
{
|
||||||
|
for (uint i = 0; i < N; dst[i] = src[i], i++);
|
||||||
|
}
|
||||||
|
|
||||||
|
// An overload for constant arrays.
|
||||||
|
template<typename T, uint N>
|
||||||
|
void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
|
||||||
|
{
|
||||||
|
for (uint i = 0; i < N; dst[i] = src[i], i++);
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
||||||
|
{
|
||||||
|
Data data[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
|
||||||
|
Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } };
|
||||||
|
Data data2[2];
|
||||||
|
spvArrayCopy(data2, _31);
|
||||||
|
_53.outdata[gl_WorkGroupID.x].a = data[gl_LocalInvocationID.x].a + data2[gl_LocalInvocationID.x].a;
|
||||||
|
_53.outdata[gl_WorkGroupID.x].b = data[gl_LocalInvocationID.x].b + data2[gl_LocalInvocationID.x].b;
|
||||||
|
}
|
||||||
|
|
@ -15,7 +15,7 @@ constant float4 _37[3] = {float4(1.0), float4(2.0), float4(3.0)};
|
|||||||
constant float4 _49[2] = { float4(1.0), float4(2.0) };
|
constant float4 _49[2] = { float4(1.0), float4(2.0) };
|
||||||
constant float4 _54[2] = { float4(8.0), float4(10.0) };
|
constant float4 _54[2] = { float4(8.0), float4(10.0) };
|
||||||
constant float4 _55[2][2] = { { float4(1.0), float4(2.0) }, { float4(8.0), float4(10.0) } };
|
constant float4 _55[2][2] = { { float4(1.0), float4(2.0) }, { float4(8.0), float4(10.0) } };
|
||||||
constant Foobar _75[2] = {{10.0, 40.0}, {90.0, 70.0}};
|
constant Foobar _75[2] = { Foobar{ 10.0, 40.0 }, Foobar{ 90.0, 70.0 } };
|
||||||
|
|
||||||
struct main0_out
|
struct main0_out
|
||||||
{
|
{
|
||||||
@ -44,7 +44,7 @@ void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
|
|||||||
fragment main0_out main0(main0_in in [[stage_in]])
|
fragment main0_out main0(main0_in in [[stage_in]])
|
||||||
{
|
{
|
||||||
main0_out out = {};
|
main0_out out = {};
|
||||||
Foobar indexable[2] = {{10.0, 40.0}, {90.0, 70.0}};
|
Foobar indexable[2] = { Foobar{ 10.0, 40.0 }, Foobar{ 90.0, 70.0 } };
|
||||||
out.FragColor = ((_37[in.index] + _55[in.index][in.index + 1]) + float4(30.0)) + float4(indexable[in.index].a + indexable[in.index].b);
|
out.FragColor = ((_37[in.index] + _55[in.index][in.index + 1]) + float4(30.0)) + float4(indexable[in.index].a + indexable[in.index].b);
|
||||||
return out;
|
return out;
|
||||||
}
|
}
|
||||||
|
@ -12,7 +12,7 @@ struct Foo
|
|||||||
};
|
};
|
||||||
|
|
||||||
constant float _16[4] = { 1.0, 4.0, 3.0, 2.0 };
|
constant float _16[4] = { 1.0, 4.0, 3.0, 2.0 };
|
||||||
constant Foo _28[2] = {{10.0, 20.0}, {30.0, 40.0}};
|
constant Foo _28[2] = { Foo{ 10.0, 20.0 }, Foo{ 30.0, 40.0 } };
|
||||||
|
|
||||||
struct main0_out
|
struct main0_out
|
||||||
{
|
{
|
||||||
@ -42,7 +42,7 @@ fragment main0_out main0(main0_in in [[stage_in]])
|
|||||||
{
|
{
|
||||||
main0_out out = {};
|
main0_out out = {};
|
||||||
float lut[4] = { 1.0, 4.0, 3.0, 2.0 };
|
float lut[4] = { 1.0, 4.0, 3.0, 2.0 };
|
||||||
Foo foos[2] = {{10.0, 20.0}, {30.0, 40.0}};
|
Foo foos[2] = { Foo{ 10.0, 20.0 }, Foo{ 30.0, 40.0 } };
|
||||||
out.FragColor = float4(lut[in.line]);
|
out.FragColor = float4(lut[in.line]);
|
||||||
out.FragColor += float4(foos[in.line].a * foos[1 - in.line].a);
|
out.FragColor += float4(foos[in.line].a * foos[1 - in.line].a);
|
||||||
return out;
|
return out;
|
||||||
|
@ -0,0 +1,25 @@
|
|||||||
|
#version 310 es
|
||||||
|
layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
layout(binding = 0, std430) buffer SSBO
|
||||||
|
{
|
||||||
|
Data outdata[];
|
||||||
|
} _53;
|
||||||
|
|
||||||
|
Data data[2];
|
||||||
|
Data data2[2];
|
||||||
|
|
||||||
|
void main()
|
||||||
|
{
|
||||||
|
data = Data[](Data(1.0, 2.0), Data(3.0, 4.0));
|
||||||
|
data2 = Data[](Data(4.0, 2.0), Data(3.0, 5.0));
|
||||||
|
_53.outdata[gl_WorkGroupID.x].a = data[gl_LocalInvocationID.x].a + data2[gl_LocalInvocationID.x].a;
|
||||||
|
_53.outdata[gl_WorkGroupID.x].b = data[gl_LocalInvocationID.x].b + data2[gl_LocalInvocationID.x].b;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,59 @@
|
|||||||
|
static const float X = 4.0f;
|
||||||
|
static const uint3 gl_WorkGroupSize = uint3(2u, 1u, 1u);
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
static const Data _21 = { 1.0f, 2.0f };
|
||||||
|
static const Data _24 = { 3.0f, 4.0f };
|
||||||
|
static const Data _25[2] = { { 1.0f, 2.0f }, { 3.0f, 4.0f } };
|
||||||
|
static const Data _30 = { 3.0f, 5.0f };
|
||||||
|
|
||||||
|
RWByteAddressBuffer _61 : register(u0);
|
||||||
|
|
||||||
|
static uint3 gl_WorkGroupID;
|
||||||
|
static uint3 gl_LocalInvocationID;
|
||||||
|
static uint gl_LocalInvocationIndex;
|
||||||
|
struct SPIRV_Cross_Input
|
||||||
|
{
|
||||||
|
uint3 gl_WorkGroupID : SV_GroupID;
|
||||||
|
uint3 gl_LocalInvocationID : SV_GroupThreadID;
|
||||||
|
uint gl_LocalInvocationIndex : SV_GroupIndex;
|
||||||
|
};
|
||||||
|
|
||||||
|
static Data data[2];
|
||||||
|
static Data data2[2];
|
||||||
|
|
||||||
|
Data combine(Data a, Data b)
|
||||||
|
{
|
||||||
|
Data _46 = { a.a + b.a, a.b + b.b };
|
||||||
|
return _46;
|
||||||
|
}
|
||||||
|
|
||||||
|
void comp_main()
|
||||||
|
{
|
||||||
|
data = _25;
|
||||||
|
Data _28 = { X, 2.0f };
|
||||||
|
Data _31[2] = { _28, _30 };
|
||||||
|
data2 = _31;
|
||||||
|
if (gl_LocalInvocationIndex == 0u)
|
||||||
|
{
|
||||||
|
Data param = data[gl_LocalInvocationID.x];
|
||||||
|
Data param_1 = data2[gl_LocalInvocationID.x];
|
||||||
|
Data _79 = combine(param, param_1);
|
||||||
|
_61.Store(gl_WorkGroupID.x * 8 + 0, asuint(_79.a));
|
||||||
|
_61.Store(gl_WorkGroupID.x * 8 + 4, asuint(_79.b));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
[numthreads(2, 1, 1)]
|
||||||
|
void main(SPIRV_Cross_Input stage_input)
|
||||||
|
{
|
||||||
|
gl_WorkGroupID = stage_input.gl_WorkGroupID;
|
||||||
|
gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
|
||||||
|
gl_LocalInvocationIndex = stage_input.gl_LocalInvocationIndex;
|
||||||
|
comp_main();
|
||||||
|
}
|
@ -22,7 +22,7 @@ fragment main0_out main0()
|
|||||||
float a = 0.0;
|
float a = 0.0;
|
||||||
float4 b = float4(0.0);
|
float4 b = float4(0.0);
|
||||||
float2x3 c = float2x3(float3(0.0), float3(0.0));
|
float2x3 c = float2x3(float3(0.0), float3(0.0));
|
||||||
D d = {float4(0.0), 0.0};
|
D d = D{ float4(0.0), 0.0 };
|
||||||
out.FragColor = a;
|
out.FragColor = a;
|
||||||
return out;
|
return out;
|
||||||
}
|
}
|
||||||
|
@ -0,0 +1,62 @@
|
|||||||
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
constant float X_tmp [[function_constant(0)]];
|
||||||
|
constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
|
||||||
|
constant uint3 gl_WorkGroupSize = uint3(2u, 1u, 1u);
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct Data_1
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct SSBO
|
||||||
|
{
|
||||||
|
Data_1 outdata[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
|
||||||
|
|
||||||
|
// Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
|
||||||
|
template<typename T, uint N>
|
||||||
|
void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
|
||||||
|
{
|
||||||
|
for (uint i = 0; i < N; dst[i] = src[i], i++);
|
||||||
|
}
|
||||||
|
|
||||||
|
// An overload for constant arrays.
|
||||||
|
template<typename T, uint N>
|
||||||
|
void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
|
||||||
|
{
|
||||||
|
for (uint i = 0; i < N; dst[i] = src[i], i++);
|
||||||
|
}
|
||||||
|
|
||||||
|
Data combine(thread const Data& a, thread const Data& b)
|
||||||
|
{
|
||||||
|
return Data{ a.a + b.a, a.b + b.b };
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
||||||
|
{
|
||||||
|
Data data[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
|
||||||
|
Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } };
|
||||||
|
Data data2[2];
|
||||||
|
spvArrayCopy(data2, _31);
|
||||||
|
Data param = data[gl_LocalInvocationID.x];
|
||||||
|
Data param_1 = data2[gl_LocalInvocationID.x];
|
||||||
|
Data _73 = combine(param, param_1);
|
||||||
|
_53.outdata[gl_WorkGroupID.x].a = _73.a;
|
||||||
|
_53.outdata[gl_WorkGroupID.x].b = _73.b;
|
||||||
|
}
|
||||||
|
|
@ -15,7 +15,7 @@ constant float4 _37[3] = {float4(1.0), float4(2.0), float4(3.0)};
|
|||||||
constant float4 _49[2] = { float4(1.0), float4(2.0) };
|
constant float4 _49[2] = { float4(1.0), float4(2.0) };
|
||||||
constant float4 _54[2] = { float4(8.0), float4(10.0) };
|
constant float4 _54[2] = { float4(8.0), float4(10.0) };
|
||||||
constant float4 _55[2][2] = { { float4(1.0), float4(2.0) }, { float4(8.0), float4(10.0) } };
|
constant float4 _55[2][2] = { { float4(1.0), float4(2.0) }, { float4(8.0), float4(10.0) } };
|
||||||
constant Foobar _75[2] = {{10.0, 40.0}, {90.0, 70.0}};
|
constant Foobar _75[2] = { Foobar{ 10.0, 40.0 }, Foobar{ 90.0, 70.0 } };
|
||||||
|
|
||||||
struct main0_out
|
struct main0_out
|
||||||
{
|
{
|
||||||
@ -49,8 +49,8 @@ float4 resolve(thread const Foobar& f)
|
|||||||
fragment main0_out main0(main0_in in [[stage_in]])
|
fragment main0_out main0(main0_in in [[stage_in]])
|
||||||
{
|
{
|
||||||
main0_out out = {};
|
main0_out out = {};
|
||||||
Foobar param = {10.0, 20.0};
|
Foobar param = Foobar{ 10.0, 20.0 };
|
||||||
Foobar indexable[2] = {{10.0, 40.0}, {90.0, 70.0}};
|
Foobar indexable[2] = { Foobar{ 10.0, 40.0 }, Foobar{ 90.0, 70.0 } };
|
||||||
Foobar param_1 = indexable[in.index];
|
Foobar param_1 = indexable[in.index];
|
||||||
out.FragColor = ((_37[in.index] + _55[in.index][in.index + 1]) + resolve(param)) + resolve(param_1);
|
out.FragColor = ((_37[in.index] + _55[in.index][in.index + 1]) + resolve(param)) + resolve(param_1);
|
||||||
return out;
|
return out;
|
||||||
|
@ -12,7 +12,7 @@ struct Foo
|
|||||||
};
|
};
|
||||||
|
|
||||||
constant float _16[4] = { 1.0, 4.0, 3.0, 2.0 };
|
constant float _16[4] = { 1.0, 4.0, 3.0, 2.0 };
|
||||||
constant Foo _28[2] = {{10.0, 20.0}, {30.0, 40.0}};
|
constant Foo _28[2] = { Foo{ 10.0, 20.0 }, Foo{ 30.0, 40.0 } };
|
||||||
|
|
||||||
struct main0_out
|
struct main0_out
|
||||||
{
|
{
|
||||||
@ -42,7 +42,7 @@ fragment main0_out main0(main0_in in [[stage_in]])
|
|||||||
{
|
{
|
||||||
main0_out out = {};
|
main0_out out = {};
|
||||||
float lut[4] = { 1.0, 4.0, 3.0, 2.0 };
|
float lut[4] = { 1.0, 4.0, 3.0, 2.0 };
|
||||||
Foo foos[2] = {{10.0, 20.0}, {30.0, 40.0}};
|
Foo foos[2] = { Foo{ 10.0, 20.0 }, Foo{ 30.0, 40.0 } };
|
||||||
out.FragColor = float4(lut[in.line]);
|
out.FragColor = float4(lut[in.line]);
|
||||||
out.FragColor += float4(foos[in.line].a * foos[1 - in.line].a);
|
out.FragColor += float4(foos[in.line].a * foos[1 - in.line].a);
|
||||||
return out;
|
return out;
|
||||||
|
33
reference/shaders/comp/composite-array-initialization.comp
Normal file
33
reference/shaders/comp/composite-array-initialization.comp
Normal file
@ -0,0 +1,33 @@
|
|||||||
|
#version 310 es
|
||||||
|
layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
layout(binding = 0, std430) buffer SSBO
|
||||||
|
{
|
||||||
|
Data outdata[];
|
||||||
|
} _53;
|
||||||
|
|
||||||
|
Data data[2];
|
||||||
|
Data data2[2];
|
||||||
|
|
||||||
|
Data combine(Data a, Data b)
|
||||||
|
{
|
||||||
|
return Data(a.a + b.a, a.b + b.b);
|
||||||
|
}
|
||||||
|
|
||||||
|
void main()
|
||||||
|
{
|
||||||
|
data = Data[](Data(1.0, 2.0), Data(3.0, 4.0));
|
||||||
|
data2 = Data[](Data(4.0, 2.0), Data(3.0, 5.0));
|
||||||
|
Data param = data[gl_LocalInvocationID.x];
|
||||||
|
Data param_1 = data2[gl_LocalInvocationID.x];
|
||||||
|
Data _73 = combine(param, param_1);
|
||||||
|
_53.outdata[gl_WorkGroupID.x].a = _73.a;
|
||||||
|
_53.outdata[gl_WorkGroupID.x].b = _73.b;
|
||||||
|
}
|
||||||
|
|
29
shaders-hlsl/comp/composite-array-initialization.comp
Normal file
29
shaders-hlsl/comp/composite-array-initialization.comp
Normal file
@ -0,0 +1,29 @@
|
|||||||
|
#version 450
|
||||||
|
layout(local_size_x = 2) in;
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
layout(std430, binding = 0) buffer SSBO
|
||||||
|
{
|
||||||
|
Data outdata[];
|
||||||
|
};
|
||||||
|
|
||||||
|
layout(constant_id = 0) const float X = 4.0;
|
||||||
|
|
||||||
|
Data data[2] = Data[](Data(1.0, 2.0), Data(3.0, 4.0));
|
||||||
|
Data data2[2] = Data[](Data(X, 2.0), Data(3.0, 5.0));
|
||||||
|
|
||||||
|
Data combine(Data a, Data b)
|
||||||
|
{
|
||||||
|
return Data(a.a + b.a, a.b + b.b);
|
||||||
|
}
|
||||||
|
|
||||||
|
void main()
|
||||||
|
{
|
||||||
|
if (gl_LocalInvocationIndex == 0u)
|
||||||
|
outdata[gl_WorkGroupID.x] = combine(data[gl_LocalInvocationID.x], data2[gl_LocalInvocationID.x]);
|
||||||
|
}
|
28
shaders-msl/comp/composite-array-initialization.comp
Normal file
28
shaders-msl/comp/composite-array-initialization.comp
Normal file
@ -0,0 +1,28 @@
|
|||||||
|
#version 450
|
||||||
|
layout(local_size_x = 2) in;
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
layout(std430, binding = 0) buffer SSBO
|
||||||
|
{
|
||||||
|
Data outdata[];
|
||||||
|
};
|
||||||
|
|
||||||
|
layout(constant_id = 0) const float X = 4.0;
|
||||||
|
|
||||||
|
Data data[2] = Data[](Data(1.0, 2.0), Data(3.0, 4.0));
|
||||||
|
Data data2[2] = Data[](Data(X, 2.0), Data(3.0, 5.0));
|
||||||
|
|
||||||
|
Data combine(Data a, Data b)
|
||||||
|
{
|
||||||
|
return Data(a.a + b.a, a.b + b.b);
|
||||||
|
}
|
||||||
|
|
||||||
|
void main()
|
||||||
|
{
|
||||||
|
outdata[gl_WorkGroupID.x] = combine(data[gl_LocalInvocationID.x], data2[gl_LocalInvocationID.x]);
|
||||||
|
}
|
29
shaders/comp/composite-array-initialization.comp
Normal file
29
shaders/comp/composite-array-initialization.comp
Normal file
@ -0,0 +1,29 @@
|
|||||||
|
#version 310 es
|
||||||
|
#extension GL_EXT_shader_non_constant_global_initializers : require
|
||||||
|
layout(local_size_x = 2) in;
|
||||||
|
|
||||||
|
struct Data
|
||||||
|
{
|
||||||
|
float a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
layout(std430, binding = 0) buffer SSBO
|
||||||
|
{
|
||||||
|
Data outdata[];
|
||||||
|
};
|
||||||
|
|
||||||
|
layout(constant_id = 0) const float X = 4.0;
|
||||||
|
|
||||||
|
Data data[2] = Data[](Data(1.0, 2.0), Data(3.0, 4.0));
|
||||||
|
Data data2[2] = Data[](Data(X, 2.0), Data(3.0, 5.0));
|
||||||
|
|
||||||
|
Data combine(Data a, Data b)
|
||||||
|
{
|
||||||
|
return Data(a.a + b.a, a.b + b.b);
|
||||||
|
}
|
||||||
|
|
||||||
|
void main()
|
||||||
|
{
|
||||||
|
outdata[gl_WorkGroupID.x] = combine(data[gl_LocalInvocationID.x], data2[gl_LocalInvocationID.x]);
|
||||||
|
}
|
@ -2709,12 +2709,23 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c)
|
|||||||
{
|
{
|
||||||
if (!c.subconstants.empty())
|
if (!c.subconstants.empty())
|
||||||
{
|
{
|
||||||
|
auto &type = get<SPIRType>(c.constant_type);
|
||||||
|
|
||||||
// Handles Arrays and structures.
|
// Handles Arrays and structures.
|
||||||
string res;
|
string res;
|
||||||
if (backend.use_initializer_list)
|
if (backend.use_initializer_list && backend.use_typed_initializer_list && type.basetype == SPIRType::Struct &&
|
||||||
|
type.array.empty())
|
||||||
|
{
|
||||||
|
res = type_to_glsl_constructor(type) + "{ ";
|
||||||
|
}
|
||||||
|
else if (backend.use_initializer_list)
|
||||||
|
{
|
||||||
res = "{ ";
|
res = "{ ";
|
||||||
|
}
|
||||||
else
|
else
|
||||||
res = type_to_glsl_constructor(get<SPIRType>(c.constant_type)) + "(";
|
{
|
||||||
|
res = type_to_glsl_constructor(type) + "(";
|
||||||
|
}
|
||||||
|
|
||||||
for (auto &elem : c.subconstants)
|
for (auto &elem : c.subconstants)
|
||||||
{
|
{
|
||||||
@ -6478,7 +6489,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||||||
{
|
{
|
||||||
// Only use this path if we are building composites.
|
// Only use this path if we are building composites.
|
||||||
// This path cannot be used for arithmetic.
|
// This path cannot be used for arithmetic.
|
||||||
if (backend.use_typed_initializer_list && out_type.basetype == SPIRType::Struct)
|
if (backend.use_typed_initializer_list && out_type.basetype == SPIRType::Struct && out_type.array.empty())
|
||||||
constructor_op += type_to_glsl_constructor(get<SPIRType>(result_type));
|
constructor_op += type_to_glsl_constructor(get<SPIRType>(result_type));
|
||||||
constructor_op += "{ ";
|
constructor_op += "{ ";
|
||||||
if (type_is_empty(out_type) && !backend.supports_empty_struct)
|
if (type_is_empty(out_type) && !backend.supports_empty_struct)
|
||||||
|
@ -227,7 +227,7 @@ protected:
|
|||||||
virtual void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index,
|
virtual void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index,
|
||||||
const std::string &qualifier = "", uint32_t base_offset = 0);
|
const std::string &qualifier = "", uint32_t base_offset = 0);
|
||||||
virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0);
|
virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0);
|
||||||
virtual std::string constant_expression(const SPIRConstant &c);
|
std::string constant_expression(const SPIRConstant &c);
|
||||||
std::string constant_op_expression(const SPIRConstantOp &cop);
|
std::string constant_op_expression(const SPIRConstantOp &cop);
|
||||||
virtual std::string constant_expression_vector(const SPIRConstant &c, uint32_t vector);
|
virtual std::string constant_expression_vector(const SPIRConstant &c, uint32_t vector);
|
||||||
virtual void emit_fixup();
|
virtual void emit_fixup();
|
||||||
|
@ -3527,6 +3527,8 @@ void CompilerHLSL::emit_access_chain(const Instruction &instruction)
|
|||||||
bool need_transpose;
|
bool need_transpose;
|
||||||
base = access_chain(ops[2], &ops[3], to_plain_buffer_length, get<SPIRType>(ops[0]), &need_transpose);
|
base = access_chain(ops[2], &ops[3], to_plain_buffer_length, get<SPIRType>(ops[0]), &need_transpose);
|
||||||
}
|
}
|
||||||
|
else if (chain)
|
||||||
|
base = chain->base;
|
||||||
else
|
else
|
||||||
base = to_expression(ops[2]);
|
base = to_expression(ops[2]);
|
||||||
|
|
||||||
|
@ -3303,39 +3303,6 @@ uint32_t CompilerMSL::get_ordered_member_location(uint32_t type_id, uint32_t ind
|
|||||||
return index;
|
return index;
|
||||||
}
|
}
|
||||||
|
|
||||||
string CompilerMSL::constant_expression(const SPIRConstant &c)
|
|
||||||
{
|
|
||||||
if (!c.subconstants.empty())
|
|
||||||
{
|
|
||||||
// Handles Arrays and structures.
|
|
||||||
string res = "{";
|
|
||||||
for (auto &elem : c.subconstants)
|
|
||||||
{
|
|
||||||
res += constant_expression(get<SPIRConstant>(elem));
|
|
||||||
if (&elem != &c.subconstants.back())
|
|
||||||
res += ", ";
|
|
||||||
}
|
|
||||||
res += "}";
|
|
||||||
return res;
|
|
||||||
}
|
|
||||||
else if (c.columns() == 1)
|
|
||||||
{
|
|
||||||
return constant_expression_vector(c, 0);
|
|
||||||
}
|
|
||||||
else
|
|
||||||
{
|
|
||||||
string res = type_to_glsl(get<SPIRType>(c.constant_type)) + "(";
|
|
||||||
for (uint32_t col = 0; col < c.columns(); col++)
|
|
||||||
{
|
|
||||||
res += constant_expression_vector(c, col);
|
|
||||||
if (col + 1 < c.columns())
|
|
||||||
res += ", ";
|
|
||||||
}
|
|
||||||
res += ")";
|
|
||||||
return res;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
// Returns the type declaration for a function, including the
|
// Returns the type declaration for a function, including the
|
||||||
// entry type if the current function is the entry point function
|
// entry type if the current function is the entry point function
|
||||||
string CompilerMSL::func_type_decl(SPIRType &type)
|
string CompilerMSL::func_type_decl(SPIRType &type)
|
||||||
|
@ -290,7 +290,6 @@ protected:
|
|||||||
std::string image_type_glsl(const SPIRType &type, uint32_t id = 0) override;
|
std::string image_type_glsl(const SPIRType &type, uint32_t id = 0) override;
|
||||||
std::string sampler_type(const SPIRType &type);
|
std::string sampler_type(const SPIRType &type);
|
||||||
std::string builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClass storage) override;
|
std::string builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClass storage) override;
|
||||||
std::string constant_expression(const SPIRConstant &c) override;
|
|
||||||
size_t get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const override;
|
size_t get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const override;
|
||||||
std::string to_func_call_arg(uint32_t id) override;
|
std::string to_func_call_arg(uint32_t id) override;
|
||||||
std::string to_name(uint32_t id, bool allow_alias = true) const override;
|
std::string to_name(uint32_t id, bool allow_alias = true) const override;
|
||||||
|
Loading…
Reference in New Issue
Block a user