AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

MSL: Handle constant construct of block-like array types.

Browse Source 
Need this to be context sensitive, since array of block-like struct is
template, but struct of block-like array is C-style.

Also, test a mix and match, so we have constant array of block-like
struct with array inside. :v
This commit is contained in:
Hans-Kristian Arntzen 2022-01-17 17:33:57 +01:00
parent 79b13813c6
commit 9b25581d49
7 changed files with 342 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

77
reference/shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp Normal file
View File

@ -0,0 +1,77 @@
#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 type_CommonConstants
{
uint g_count;
packed_uint3 g_padding4;
};
struct MyStruct
{
float4 m_coefficients[4];
};
struct type_RWStructuredBuffer_MyStruct
{
MyStruct _m0[1];
};
constant spvUnsafeArray<float4, 4> _27 = spvUnsafeArray<float4, 4>({ float4(0.0), float4(0.0), float4(0.0), float4(0.0) });
kernel void main0(constant type_CommonConstants& CommonConstants [[buffer(0)]], device type_RWStructuredBuffer_MyStruct& g_data [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
do
{
if (gl_GlobalInvocationID.x >= CommonConstants.g_count)
{
break;
}
g_data._m0[gl_GlobalInvocationID.x] = MyStruct{ { float4(0.0), float4(0.0), float4(0.0), float4(0.0) } };
break;
} while(false);
}

77
reference/shaders-msl-no-opt/asm/comp/block-like-array-type-construct.asm.comp Normal file
View File

@ -0,0 +1,77 @@
#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 _10
{
float _m0[4];
float _m1[4];
};
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
struct SSBO
{
uint a;
int b;
};
constant spvUnsafeArray<float, 4> _31 = spvUnsafeArray<float, 4>({ 1.0, 2.0, 3.0, 4.0 });
kernel void main0()
{
spvUnsafeArray<_10, 2> _34 = spvUnsafeArray<_10, 2>({ _10{ { 1.0, 2.0, 3.0, 4.0 }, { 1.0, 2.0, 3.0, 4.0 } }, _10{ { 1.0, 2.0, 3.0, 4.0 }, { 1.0, 2.0, 3.0, 4.0 } } });
spvUnsafeArray<float, 4> foo;
foo[0] = 1.0;
foo = _31;
foo[1] = 2.0;
foo[2] = 3.0;
foo[3] = 4.0;
spvUnsafeArray<float, 4> foo2;
foo2 = foo;
_10 _37 = _10{ { foo[0], foo[1], foo[2], foo[3] }, { foo2[0], foo2[1], foo2[2], foo2[3] } };
}

85
shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp Normal file
View File

@ -0,0 +1,85 @@
; SPIR-V
; Version: 1.3
; Generator: Google spiregg; 0
; Bound: 40
; Schema: 0
OpCapability Shader
OpExtension "SPV_GOOGLE_hlsl_functionality1"
OpExtension "SPV_GOOGLE_user_type"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %csMainClear "main" %gl_GlobalInvocationID
OpExecutionMode %csMainClear LocalSize 64 1 1
OpSource HLSL 600
OpName %type_CommonConstants "type.CommonConstants"
OpMemberName %type_CommonConstants 0 "g_count"
OpMemberName %type_CommonConstants 1 "g_padding4"
OpName %CommonConstants "CommonConstants"
OpName %type_RWStructuredBuffer_MyStruct "type.RWStructuredBuffer.MyStruct"
OpName %MyStruct "MyStruct"
OpMemberName %MyStruct 0 "m_coefficients"
OpName %g_data "g_data"
OpName %csMainClear "csMainClear"
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorateString %gl_GlobalInvocationID UserSemantic "SV_DispatchThreadID"
OpDecorate %CommonConstants DescriptorSet 0
OpDecorate %CommonConstants Binding 0
OpDecorate %g_data DescriptorSet 0
OpDecorate %g_data Binding 0
OpMemberDecorate %type_CommonConstants 0 Offset 0
OpMemberDecorate %type_CommonConstants 1 Offset 4
OpDecorate %type_CommonConstants Block
OpDecorateString %CommonConstants UserTypeGOOGLE "cbuffer"
OpDecorate %_arr_v4float_uint_4 ArrayStride 16
OpMemberDecorate %MyStruct 0 Offset 0
OpDecorate %_runtimearr_MyStruct ArrayStride 64
OpMemberDecorate %type_RWStructuredBuffer_MyStruct 0 Offset 0
OpDecorate %type_RWStructuredBuffer_MyStruct BufferBlock
OpDecorateString %g_data UserTypeGOOGLE "rwstructuredbuffer"
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%v3uint = OpTypeVector %uint 3
%type_CommonConstants = OpTypeStruct %uint %v3uint
%_ptr_Uniform_type_CommonConstants = OpTypePointer Uniform %type_CommonConstants
%_arr_v4float_uint_4 = OpTypeArray %v4float %uint_4
%MyStruct = OpTypeStruct %_arr_v4float_uint_4
%_runtimearr_MyStruct = OpTypeRuntimeArray %MyStruct
%type_RWStructuredBuffer_MyStruct = OpTypeStruct %_runtimearr_MyStruct
%_ptr_Uniform_type_RWStructuredBuffer_MyStruct = OpTypePointer Uniform %type_RWStructuredBuffer_MyStruct
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%21 = OpTypeFunction %void
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%bool = OpTypeBool
%_ptr_Uniform_MyStruct = OpTypePointer Uniform %MyStruct
%CommonConstants = OpVariable %_ptr_Uniform_type_CommonConstants Uniform
%g_data = OpVariable %_ptr_Uniform_type_RWStructuredBuffer_MyStruct Uniform
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
%uint_0 = OpConstant %uint 0
%26 = OpConstantNull %v4float
%27 = OpConstantComposite %_arr_v4float_uint_4 %26 %26 %26 %26
%28 = OpConstantComposite %MyStruct %27
%csMainClear = OpFunction %void None %21
%29 = OpLabel
%30 = OpLoad %v3uint %gl_GlobalInvocationID
OpSelectionMerge %31 None
OpSwitch %uint_0 %32
%32 = OpLabel
%33 = OpCompositeExtract %uint %30 0
%34 = OpAccessChain %_ptr_Uniform_uint %CommonConstants %int_0
%35 = OpLoad %uint %34
%36 = OpUGreaterThanEqual %bool %33 %35
OpSelectionMerge %37 DontFlatten
OpBranchConditional %36 %38 %37
%38 = OpLabel
OpBranch %31
%37 = OpLabel
%39 = OpAccessChain %_ptr_Uniform_MyStruct %g_data %int_0 %33
OpStore %39 %28
OpBranch %31
%31 = OpLabel
OpReturn
OpFunctionEnd

80
shaders-msl-no-opt/asm/comp/block-like-array-type-construct.asm.comp Normal file
View File

@ -0,0 +1,80 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 10
; Bound: 32
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpName %main "main"
OpName %foo "foo"
OpName %foo2 "foo2"
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "a"
OpMemberName %SSBO 1 "b"
OpName %_ ""
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
OpMemberDecorate %SSBO 0 Offset 0
OpMemberDecorate %SSBO 1 Offset 4
OpDecorate %SSBO BufferBlock
OpDecorate %_ DescriptorSet 0
OpDecorate %_ Binding 0
OpDecorate %_arr_float_uint_4 ArrayStride 4
OpDecorate %struct_arr ArrayStride 32
OpMemberDecorate %struct 0 Offset 0
OpMemberDecorate %struct 1 Offset 16
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2
%uint_4 = OpConstant %uint 4
%_arr_float_uint_4 = OpTypeArray %float %uint_4
%_ptr_Private__arr_float_uint_4 = OpTypePointer Private %_arr_float_uint_4
%foo = OpVariable %_ptr_Private__arr_float_uint_4 Private
%foo2 = OpVariable %_ptr_Private__arr_float_uint_4 Private
%int = OpTypeInt 32 1
%int_0 = OpConstant %int 0
%float_1 = OpConstant %float 1
%struct = OpTypeStruct %_arr_float_uint_4 %_arr_float_uint_4
%struct_arr = OpTypeArray %struct %uint_2
%ptr_struct = OpTypePointer Function %struct
%_ptr_Private_float = OpTypePointer Private %float
%int_1 = OpConstant %int 1
%float_2 = OpConstant %float 2
%int_2 = OpConstant %int 2
%float_3 = OpConstant %float 3
%int_3 = OpConstant %int 3
%float_4 = OpConstant %float 4
%v3uint = OpTypeVector %uint 3
%uint_1 = OpConstant %uint 1
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
%carr = OpConstantComposite %_arr_float_uint_4 %float_1 %float_2 %float_3 %float_4
%struct_constant_0 = OpConstantComposite %struct %carr %carr
%struct_constant_1 = OpConstantComposite %struct %carr %carr
%struct_arr_constant = OpConstantComposite %struct_arr %struct_constant_0 %struct_constant_1
%SSBO = OpTypeStruct %uint %int
%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO
%_ = OpVariable %_ptr_Uniform_SSBO Uniform
%main = OpFunction %void None %3
%5 = OpLabel
%struct_var = OpVariable %ptr_struct Function
%16 = OpAccessChain %_ptr_Private_float %foo %int_0
OpStore %16 %float_1
OpStore %foo %carr
%19 = OpAccessChain %_ptr_Private_float %foo %int_1
OpStore %19 %float_2
%22 = OpAccessChain %_ptr_Private_float %foo %int_2
OpStore %22 %float_3
%25 = OpAccessChain %_ptr_Private_float %foo %int_3
OpStore %25 %float_4
OpCopyMemory %foo2 %foo
%l0 = OpLoad %_arr_float_uint_4 %foo
%l1 = OpLoad %_arr_float_uint_4 %foo2
%struct0 = OpCompositeConstruct %struct %l0 %l1
OpStore %struct_var %struct0
OpReturn
OpFunctionEnd

24
spirv_glsl.cpp
View File

@ -4559,12 +4559,13 @@ string CompilerGLSL::to_rerolled_array_expression(const string &base_expr, const
return expr;
}
string CompilerGLSL::to_composite_constructor_expression(uint32_t id, bool uses_buffer_offset)
string CompilerGLSL::to_composite_constructor_expression(uint32_t id, bool block_like_type)
{
auto &type = expression_type(id);
bool reroll_array = !type.array.empty() && (!backend.array_is_value_type ||
(uses_buffer_offset && !backend.buffer_offset_array_is_value_type));
bool reroll_array = !type.array.empty() &&
(!backend.array_is_value_type ||
(block_like_type && !backend.array_is_value_type_in_buffer_blocks));
if (reroll_array)
{
@ -4966,7 +4967,7 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
}
}
string CompilerGLSL::constant_expression(const SPIRConstant &c)
string CompilerGLSL::constant_expression(const SPIRConstant &c, bool inside_block_like_struct_scope)
{
auto &type = get<SPIRType>(c.constant_type);
@ -4979,6 +4980,15 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c)
// Handles Arrays and structures.
string res;
// Only consider the decay if we are inside a struct scope.
// Outside a struct declaration, we can always bind to a constant array with templated type.
bool array_type_decays = inside_block_like_struct_scope &&
!type.array.empty() && !backend.array_is_value_type_in_buffer_blocks &&
has_decoration(c.constant_type, DecorationArrayStride);
if (type.array.empty() && type.basetype == SPIRType::Struct && type_is_block_like(type))
inside_block_like_struct_scope = true;
// Allow Metal to use the array<T> template to make arrays a value type
bool needs_trailing_tracket = false;
if (backend.use_initializer_list && backend.use_typed_initializer_list && type.basetype == SPIRType::Struct &&
@ -4987,7 +4997,7 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c)
res = type_to_glsl_constructor(type) + "{ ";
}
else if (backend.use_initializer_list && backend.use_typed_initializer_list && backend.array_is_value_type &&
!type.array.empty())
!type.array.empty() && !array_type_decays)
{
res = type_to_glsl_constructor(type) + "({ ";
needs_trailing_tracket = true;
@ -5007,7 +5017,7 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c)
if (subc.specialization)
res += to_name(elem);
else
res += constant_expression(subc);
res += constant_expression(subc, inside_block_like_struct_scope);
if (&elem != &c.subconstants.back())
res += ", ";
@ -6965,7 +6975,7 @@ bool CompilerGLSL::expression_is_non_value_type_array(uint32_t ptr)
return false;
auto &backed_type = get<SPIRType>(var->basetype);
return !backend.buffer_offset_array_is_value_type && backed_type.basetype == SPIRType::Struct &&
return !backend.array_is_value_type_in_buffer_blocks && backed_type.basetype == SPIRType::Struct &&
has_member_decoration(backed_type.self, 0, DecorationOffset);
}

6
spirv_glsl.hpp
View File

@ -385,7 +385,7 @@ protected:
const std::string &qualifier = "", uint32_t base_offset = 0);
virtual void emit_struct_padding_target(const SPIRType &type);
virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0);
std::string constant_expression(const SPIRConstant &c);
std::string constant_expression(const SPIRConstant &c, bool inside_block_like_struct_scope = false);
virtual std::string constant_op_expression(const SPIRConstantOp &cop);
virtual std::string constant_expression_vector(const SPIRConstant &c, uint32_t vector);
virtual void emit_fixup();
@ -577,7 +577,7 @@ protected:
bool supports_extensions = false;
bool supports_empty_struct = false;
bool array_is_value_type = true;
bool buffer_offset_array_is_value_type = true;
bool array_is_value_type_in_buffer_blocks = true;
bool comparison_image_samples_scalar = false;
bool native_pointers = false;
bool support_small_type_sampling_result = false;
@ -718,7 +718,7 @@ protected:
void append_global_func_args(const SPIRFunction &func, uint32_t index, SmallVector<std::string> &arglist);
std::string to_non_uniform_aware_expression(uint32_t id);
std::string to_expression(uint32_t id, bool register_expression_read = true);
std::string to_composite_constructor_expression(uint32_t id, bool uses_buffer_offset);
std::string to_composite_constructor_expression(uint32_t id, bool block_like_type);
std::string to_rerolled_array_expression(const std::string &expr, const SPIRType &type);
std::string to_enclosed_expression(uint32_t id, bool register_expression_read = true);
std::string to_unpacked_expression(uint32_t id, bool register_expression_read = true);

6
spirv_msl.cpp
View File

@ -1360,8 +1360,8 @@ string CompilerMSL::compile()
// Allow Metal to use the array<T> template unless we force it off.
backend.can_return_array = !msl_options.force_native_arrays;
backend.array_is_value_type = !msl_options.force_native_arrays;
// Arrays which are part of buffer objects are never considered to be native arrays.
backend.buffer_offset_array_is_value_type = false;
// Arrays which are part of buffer objects are never considered to be value types (just plain C-style).
backend.array_is_value_type_in_buffer_blocks = false;
backend.support_pointer_to_pointer = true;
capture_output_to_buffer = msl_options.capture_output_to_buffer;
@ -13578,7 +13578,7 @@ std::string CompilerMSL::variable_decl(const SPIRVariable &variable)
if (variable_decl_is_remapped_storage(variable, StorageClassWorkgroup))
is_using_builtin_array = true;
std::string expr = CompilerGLSL::variable_decl(variable);
auto expr = CompilerGLSL::variable_decl(variable);
is_using_builtin_array = old_is_using_builtin_array;
return expr;
}