Disabled spvUnsafeArray<> type for packed vectors and added test cases for those arrays.
This commit is contained in:
parent
ffbd801853
commit
0853bcaee1
@ -50,7 +50,7 @@ struct spvUnsafeArray
|
|||||||
|
|
||||||
struct S0
|
struct S0
|
||||||
{
|
{
|
||||||
packed_spvUnsafeArray<float2, 1> a;
|
packed_float2 a[1];
|
||||||
float b;
|
float b;
|
||||||
};
|
};
|
||||||
|
|
||||||
@ -62,7 +62,7 @@ struct S1
|
|||||||
|
|
||||||
struct S2
|
struct S2
|
||||||
{
|
{
|
||||||
packed_spvUnsafeArray<float3, 1> a;
|
packed_float3 a[1];
|
||||||
float b;
|
float b;
|
||||||
};
|
};
|
||||||
|
|
||||||
@ -166,22 +166,22 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|||||||
|
|
||||||
kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
|
kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
|
||||||
{
|
{
|
||||||
spvUnsafeArray<float2, 1>(ssbo_scalar.content.m0s[0].a[0]) = ssbo_140.content.m0s[0].a[0].xy;
|
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.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].a = float3(ssbo_140.content.m1s[0].a);
|
||||||
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
||||||
spvUnsafeArray<float3, 1>(ssbo_scalar.content.m2s[0].a[0]) = ssbo_140.content.m2s[0].a[0];
|
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.m2s[0].b = ssbo_140.content.m2s[0].b;
|
||||||
spvUnsafeArray<float2, 1>(ssbo_scalar.content.m0.a[0]) = ssbo_140.content.m0.a[0].xy;
|
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.m0.b = ssbo_140.content.m0.b;
|
||||||
ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
|
ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
|
||||||
ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;
|
ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;
|
||||||
spvUnsafeArray<float3, 1>(ssbo_scalar.content.m2.a[0]) = ssbo_140.content.m2.a[0];
|
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.m2.b = ssbo_140.content.m2.b;
|
||||||
ssbo_scalar.content.m3.a = ssbo_140.content.m3.a;
|
ssbo_scalar.content.m3.a = ssbo_140.content.m3.a;
|
||||||
ssbo_scalar.content.m3.b = ssbo_140.content.m3.b;
|
ssbo_scalar.content.m3.b = ssbo_140.content.m3.b;
|
||||||
ssbo_scalar.content.m4 = ssbo_140.content.m4;
|
ssbo_scalar.content.m4 = ssbo_140.content.m4;
|
||||||
ssbo_scalar.content.m1.a = float2x3(float3(float3x2(float2(ssbo_scalar.m2[1][0]), float2(ssbo_scalar.m2[1][1]), float2(ssbo_scalar.m2[1][2]))[0]), float3(float3x2(float2(ssbo_scalar.m2[1][0]), float2(ssbo_scalar.m2[1][1]), float2(ssbo_scalar.m2[1][2]))[1])) * float2(spvUnsafeArray<float2, 1>(ssbo_scalar.content.m0.a[0]));
|
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_scalar.m0 = float2x2(float2(ssbo_scalar2.m1[0]), float2(ssbo_scalar2.m1[1]));
|
||||||
ssbo_scalar2.m1[0] = float2(ssbo_scalar.m4[0][0], ssbo_scalar.m4[1][0]);
|
ssbo_scalar2.m1[0] = float2(ssbo_scalar.m4[0][0], ssbo_scalar.m4[1][0]);
|
||||||
ssbo_scalar2.m1[1] = float2(ssbo_scalar.m4[0][1], ssbo_scalar.m4[1][1]);
|
ssbo_scalar2.m1[1] = float2(ssbo_scalar.m4[0][1], ssbo_scalar.m4[1][1]);
|
||||||
|
17
reference/shaders-msl-no-opt/packing/array-of-vec3.comp
Normal file
17
reference/shaders-msl-no-opt/packing/array-of-vec3.comp
Normal file
@ -0,0 +1,17 @@
|
|||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
struct SSBO
|
||||||
|
{
|
||||||
|
packed_float3 v[16];
|
||||||
|
};
|
||||||
|
|
||||||
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||||
|
|
||||||
|
kernel void main0(device SSBO& _13 [[buffer(0)]])
|
||||||
|
{
|
||||||
|
_13.v[1] = float3(_13.v[0]);
|
||||||
|
}
|
||||||
|
|
58
reference/shaders-msl-no-opt/packing/array-of-vec4.comp
Normal file
58
reference/shaders-msl-no-opt/packing/array-of-vec4.comp
Normal file
@ -0,0 +1,58 @@
|
|||||||
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||||
|
#pragma clang diagnostic ignored "-Wmissing-braces"
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
template<typename T, size_t Num>
|
||||||
|
struct spvUnsafeArray
|
||||||
|
{
|
||||||
|
T elements[Num ? Num : 1];
|
||||||
|
|
||||||
|
thread T& operator [] (size_t pos) thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const thread T& operator [] (size_t pos) const thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
device T& operator [] (size_t pos) device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const device T& operator [] (size_t pos) const device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
constexpr const constant T& operator [] (size_t pos) const constant
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
threadgroup T& operator [] (size_t pos) threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
struct SSBO
|
||||||
|
{
|
||||||
|
spvUnsafeArray<float4, 16> v;
|
||||||
|
};
|
||||||
|
|
||||||
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
||||||
|
|
||||||
|
kernel void main0(device SSBO& _13 [[buffer(0)]])
|
||||||
|
{
|
||||||
|
_13.v[1] = _13.v[0];
|
||||||
|
}
|
||||||
|
|
@ -50,7 +50,7 @@ struct spvUnsafeArray
|
|||||||
|
|
||||||
struct S0
|
struct S0
|
||||||
{
|
{
|
||||||
packed_spvUnsafeArray<float2, 1> a;
|
packed_float2 a[1];
|
||||||
float b;
|
float b;
|
||||||
};
|
};
|
||||||
|
|
||||||
@ -62,7 +62,7 @@ struct S1
|
|||||||
|
|
||||||
struct S2
|
struct S2
|
||||||
{
|
{
|
||||||
packed_spvUnsafeArray<float3, 1> a;
|
packed_float3 a[1];
|
||||||
float b;
|
float b;
|
||||||
};
|
};
|
||||||
|
|
||||||
@ -166,22 +166,22 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|||||||
|
|
||||||
kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
|
kernel void main0(device SSBO1& ssbo_scalar [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]], device SSBO2& ssbo_scalar2 [[buffer(2)]])
|
||||||
{
|
{
|
||||||
spvUnsafeArray<float2, 1>(ssbo_scalar.content.m0s[0].a[0]) = ssbo_140.content.m0s[0].a[0].xy;
|
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.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].a = float3(ssbo_140.content.m1s[0].a);
|
||||||
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
|
||||||
spvUnsafeArray<float3, 1>(ssbo_scalar.content.m2s[0].a[0]) = ssbo_140.content.m2s[0].a[0];
|
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.m2s[0].b = ssbo_140.content.m2s[0].b;
|
||||||
spvUnsafeArray<float2, 1>(ssbo_scalar.content.m0.a[0]) = ssbo_140.content.m0.a[0].xy;
|
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.m0.b = ssbo_140.content.m0.b;
|
||||||
ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
|
ssbo_scalar.content.m1.a = float3(ssbo_140.content.m1.a);
|
||||||
ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;
|
ssbo_scalar.content.m1.b = ssbo_140.content.m1.b;
|
||||||
spvUnsafeArray<float3, 1>(ssbo_scalar.content.m2.a[0]) = ssbo_140.content.m2.a[0];
|
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.m2.b = ssbo_140.content.m2.b;
|
||||||
ssbo_scalar.content.m3.a = ssbo_140.content.m3.a;
|
ssbo_scalar.content.m3.a = ssbo_140.content.m3.a;
|
||||||
ssbo_scalar.content.m3.b = ssbo_140.content.m3.b;
|
ssbo_scalar.content.m3.b = ssbo_140.content.m3.b;
|
||||||
ssbo_scalar.content.m4 = ssbo_140.content.m4;
|
ssbo_scalar.content.m4 = ssbo_140.content.m4;
|
||||||
ssbo_scalar.content.m1.a = float2x3(float3(float3x2(float2(ssbo_scalar.m2[1][0]), float2(ssbo_scalar.m2[1][1]), float2(ssbo_scalar.m2[1][2]))[0]), float3(float3x2(float2(ssbo_scalar.m2[1][0]), float2(ssbo_scalar.m2[1][1]), float2(ssbo_scalar.m2[1][2]))[1])) * float2(spvUnsafeArray<float2, 1>(ssbo_scalar.content.m0.a[0]));
|
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_scalar.m0 = float2x2(float2(ssbo_scalar2.m1[0]), float2(ssbo_scalar2.m1[1]));
|
||||||
ssbo_scalar2.m1[0] = float2(ssbo_scalar.m4[0][0], ssbo_scalar.m4[1][0]);
|
ssbo_scalar2.m1[0] = float2(ssbo_scalar.m4[0][0], ssbo_scalar.m4[1][0]);
|
||||||
ssbo_scalar2.m1[1] = float2(ssbo_scalar.m4[0][1], ssbo_scalar.m4[1][1]);
|
ssbo_scalar2.m1[1] = float2(ssbo_scalar.m4[0][1], ssbo_scalar.m4[1][1]);
|
||||||
|
13
shaders-msl-no-opt/packing/array-of-vec3.comp
Normal file
13
shaders-msl-no-opt/packing/array-of-vec3.comp
Normal file
@ -0,0 +1,13 @@
|
|||||||
|
#version 450
|
||||||
|
#extension GL_EXT_scalar_block_layout : require
|
||||||
|
layout(local_size_x = 1) in;
|
||||||
|
|
||||||
|
layout(set = 0, binding = 0, scalar) buffer SSBO
|
||||||
|
{
|
||||||
|
vec3 v[16];
|
||||||
|
};
|
||||||
|
|
||||||
|
void main()
|
||||||
|
{
|
||||||
|
v[1] = v[0];
|
||||||
|
}
|
13
shaders-msl-no-opt/packing/array-of-vec4.comp
Normal file
13
shaders-msl-no-opt/packing/array-of-vec4.comp
Normal file
@ -0,0 +1,13 @@
|
|||||||
|
#version 450
|
||||||
|
#extension GL_EXT_scalar_block_layout : require
|
||||||
|
layout(local_size_x = 1) in;
|
||||||
|
|
||||||
|
layout(set = 0, binding = 0, scalar) buffer SSBO
|
||||||
|
{
|
||||||
|
vec4 v[16];
|
||||||
|
};
|
||||||
|
|
||||||
|
void main()
|
||||||
|
{
|
||||||
|
v[1] = v[0];
|
||||||
|
}
|
@ -3249,7 +3249,14 @@ string CompilerMSL::unpack_expression_type(string expr_str, const SPIRType &type
|
|||||||
return unpack_expr;
|
return unpack_expr;
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
return join(type_to_glsl(type), "(", expr_str, ")");
|
{
|
||||||
|
// Don't expose "spvUnsafeArray" when unpacking expressions,
|
||||||
|
// the input "type" will be the unpacked type and might also appear in l-value expressions
|
||||||
|
use_builtin_array = true;
|
||||||
|
string unpack_expr = join(type_to_glsl(type), "(", expr_str, ")");
|
||||||
|
use_builtin_array = false;
|
||||||
|
return unpack_expr;
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
// Emits the file header info
|
// Emits the file header info
|
||||||
@ -8068,8 +8075,11 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
|
|||||||
add_typedef_line(td_line);
|
add_typedef_line(td_line);
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
|
{
|
||||||
|
use_builtin_array = true;
|
||||||
pack_pfx = "packed_";
|
pack_pfx = "packed_";
|
||||||
}
|
}
|
||||||
|
}
|
||||||
else if (row_major)
|
else if (row_major)
|
||||||
{
|
{
|
||||||
// Need to declare type with flipped vecsize/columns.
|
// Need to declare type with flipped vecsize/columns.
|
||||||
@ -8092,7 +8102,8 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
|
|||||||
{
|
{
|
||||||
// Force the use of C style array declaration.
|
// Force the use of C style array declaration.
|
||||||
BuiltIn builtin = BuiltInMax;
|
BuiltIn builtin = BuiltInMax;
|
||||||
use_builtin_array = is_member_builtin(type, index, &builtin);
|
if (is_member_builtin(type, index, &builtin))
|
||||||
|
use_builtin_array = true;
|
||||||
array_type = type_to_array_glsl(physical_type);
|
array_type = type_to_array_glsl(physical_type);
|
||||||
}
|
}
|
||||||
|
|
||||||
@ -12027,14 +12038,6 @@ std::string CompilerMSL::access_chain_internal(uint32_t base, const uint32_t *in
|
|||||||
// Sample mask input for Metal is not an array
|
// Sample mask input for Metal is not an array
|
||||||
else if (ir.meta[base].decoration.builtin_type != BuiltInSampleMask)
|
else if (ir.meta[base].decoration.builtin_type != BuiltInSampleMask)
|
||||||
{
|
{
|
||||||
if (is_packed)
|
|
||||||
{
|
|
||||||
if (!remove_duplicate_swizzle(expr))
|
|
||||||
remove_unity_swizzle(base, expr);
|
|
||||||
append_index(index);
|
|
||||||
expr = unpack_expression_type(expr, *type, physical_type, is_packed, true);
|
|
||||||
}
|
|
||||||
else
|
|
||||||
append_index(index);
|
append_index(index);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
Loading…
Reference in New Issue
Block a user