From 9012a39b6090ee2237ca56f25c99368a487613fc Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Mon, 6 Jan 2020 11:47:26 +0100 Subject: [PATCH 1/4] Basic implementation of OpCopyLogical. --- .../asm/comp/copy-logical.spv14.asm.comp | 47 +++++++++++++ .../asm/comp/copy-logical.spv14.asm.comp | 45 ++++++++++++ .../asm/comp/copy-logical.spv14.asm.comp | 69 +++++++++++++++++++ .../asm/comp/copy-logical.spv14.asm.comp | 69 +++++++++++++++++++ spirv_glsl.cpp | 52 ++++++++++++++ spirv_glsl.hpp | 3 + test_shaders.py | 12 ++-- 7 files changed, 293 insertions(+), 4 deletions(-) create mode 100644 reference/shaders-msl-no-opt/asm/comp/copy-logical.spv14.asm.comp create mode 100644 reference/shaders-no-opt/asm/comp/copy-logical.spv14.asm.comp create mode 100644 shaders-msl-no-opt/asm/comp/copy-logical.spv14.asm.comp create mode 100644 shaders-no-opt/asm/comp/copy-logical.spv14.asm.comp diff --git a/reference/shaders-msl-no-opt/asm/comp/copy-logical.spv14.asm.comp b/reference/shaders-msl-no-opt/asm/comp/copy-logical.spv14.asm.comp new file mode 100644 index 00000000..22259815 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/copy-logical.spv14.asm.comp @@ -0,0 +1,47 @@ +#include +#include + +using namespace metal; + +struct B2 +{ + float4 elem2; +}; + +struct C +{ + float4 c; + B2 b2; + B2 b2_array[4]; +}; + +struct B1 +{ + float4 elem1; +}; + +struct A +{ + float4 a; + B1 b1; + B1 b1_array[4]; +}; + +struct _8 +{ + A a_block; + C c_block; +}; + +kernel void main0(device _8& _3 [[buffer(0)]]) +{ + A _27; + _27.a = _3.c_block.c; + _27.b1.elem1 = _3.c_block.b2.elem2; + _27.b1_array[0].elem1 = _3.c_block.b2_array[0].elem2; + _27.b1_array[1].elem1 = _3.c_block.b2_array[1].elem2; + _27.b1_array[2].elem1 = _3.c_block.b2_array[2].elem2; + _27.b1_array[3].elem1 = _3.c_block.b2_array[3].elem2; + _3.a_block = _27; +} + diff --git a/reference/shaders-no-opt/asm/comp/copy-logical.spv14.asm.comp b/reference/shaders-no-opt/asm/comp/copy-logical.spv14.asm.comp new file mode 100644 index 00000000..28b2d1d0 --- /dev/null +++ b/reference/shaders-no-opt/asm/comp/copy-logical.spv14.asm.comp @@ -0,0 +1,45 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct B2 +{ + vec4 elem2; +}; + +struct C +{ + vec4 c; + B2 b2; + B2 b2_array[4]; +}; + +struct B1 +{ + vec4 elem1; +}; + +struct A +{ + vec4 a; + B1 b1; + B1 b1_array[4]; +}; + +layout(binding = 0, std430) buffer _8_3 +{ + A a_block; + C c_block; +} _3; + +void main() +{ + A _27; + _27.a = _3.c_block.c; + _27.b1.elem1 = _3.c_block.b2.elem2; + _27.b1_array[0].elem1 = _3.c_block.b2_array[0].elem2; + _27.b1_array[1].elem1 = _3.c_block.b2_array[1].elem2; + _27.b1_array[2].elem1 = _3.c_block.b2_array[2].elem2; + _27.b1_array[3].elem1 = _3.c_block.b2_array[3].elem2; + _3.a_block = _27; +} + diff --git a/shaders-msl-no-opt/asm/comp/copy-logical.spv14.asm.comp b/shaders-msl-no-opt/asm/comp/copy-logical.spv14.asm.comp new file mode 100644 index 00000000..20fa0b09 --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/copy-logical.spv14.asm.comp @@ -0,0 +1,69 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 8 +; Bound: 48 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %ssbo + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %B1 "B1" + OpName %A "A" + OpName %C "C" + OpName %B2 "B2" + OpMemberName %A 0 "a" + OpMemberName %A 1 "b1" + OpMemberName %A 2 "b1_array" + OpMemberName %C 0 "c" + OpMemberName %C 1 "b2" + OpMemberName %C 2 "b2_array" + OpMemberName %B1 0 "elem1" + OpMemberName %B2 0 "elem2" + OpMemberName %SSBO 0 "a_block" + OpMemberName %SSBO 1 "c_block" + OpDecorate %B1Array ArrayStride 16 + OpDecorate %B2Array ArrayStride 16 + OpMemberDecorate %B1 0 Offset 0 + OpMemberDecorate %A 0 Offset 0 + OpMemberDecorate %A 1 Offset 16 + OpMemberDecorate %A 2 Offset 32 + OpMemberDecorate %B2 0 Offset 0 + OpMemberDecorate %C 0 Offset 0 + OpMemberDecorate %C 1 Offset 16 + OpMemberDecorate %C 2 Offset 32 + OpMemberDecorate %SSBO 0 Offset 0 + OpMemberDecorate %SSBO 1 Offset 96 + OpDecorate %SSBO Block + OpDecorate %ssbo DescriptorSet 0 + OpDecorate %ssbo Binding 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %float = OpTypeFloat 32 + %uint = OpTypeInt 32 0 + %uint_4 = OpConstant %uint 4 + %v4float = OpTypeVector %float 4 + %B2 = OpTypeStruct %v4float + %B2Array = OpTypeArray %B2 %uint_4 + %C = OpTypeStruct %v4float %B2 %B2Array + %B1 = OpTypeStruct %v4float + %B1Array = OpTypeArray %B1 %uint_4 + %A = OpTypeStruct %v4float %B1 %B1Array + %SSBO = OpTypeStruct %A %C +%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO + %ssbo = OpVariable %_ptr_Uniform_SSBO StorageBuffer + %int = OpTypeInt 32 1 + %int_1 = OpConstant %int 1 +%_ptr_Uniform_C = OpTypePointer StorageBuffer %C + %int_0 = OpConstant %int 0 +%_ptr_Uniform_A = OpTypePointer StorageBuffer %A + %main = OpFunction %void None %3 + %5 = OpLabel + %22 = OpAccessChain %_ptr_Uniform_C %ssbo %int_1 + %39 = OpAccessChain %_ptr_Uniform_A %ssbo %int_0 + %23 = OpLoad %C %22 + %24 = OpCopyLogical %A %23 + OpStore %39 %24 + OpReturn + OpFunctionEnd diff --git a/shaders-no-opt/asm/comp/copy-logical.spv14.asm.comp b/shaders-no-opt/asm/comp/copy-logical.spv14.asm.comp new file mode 100644 index 00000000..20fa0b09 --- /dev/null +++ b/shaders-no-opt/asm/comp/copy-logical.spv14.asm.comp @@ -0,0 +1,69 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 8 +; Bound: 48 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %ssbo + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %B1 "B1" + OpName %A "A" + OpName %C "C" + OpName %B2 "B2" + OpMemberName %A 0 "a" + OpMemberName %A 1 "b1" + OpMemberName %A 2 "b1_array" + OpMemberName %C 0 "c" + OpMemberName %C 1 "b2" + OpMemberName %C 2 "b2_array" + OpMemberName %B1 0 "elem1" + OpMemberName %B2 0 "elem2" + OpMemberName %SSBO 0 "a_block" + OpMemberName %SSBO 1 "c_block" + OpDecorate %B1Array ArrayStride 16 + OpDecorate %B2Array ArrayStride 16 + OpMemberDecorate %B1 0 Offset 0 + OpMemberDecorate %A 0 Offset 0 + OpMemberDecorate %A 1 Offset 16 + OpMemberDecorate %A 2 Offset 32 + OpMemberDecorate %B2 0 Offset 0 + OpMemberDecorate %C 0 Offset 0 + OpMemberDecorate %C 1 Offset 16 + OpMemberDecorate %C 2 Offset 32 + OpMemberDecorate %SSBO 0 Offset 0 + OpMemberDecorate %SSBO 1 Offset 96 + OpDecorate %SSBO Block + OpDecorate %ssbo DescriptorSet 0 + OpDecorate %ssbo Binding 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %float = OpTypeFloat 32 + %uint = OpTypeInt 32 0 + %uint_4 = OpConstant %uint 4 + %v4float = OpTypeVector %float 4 + %B2 = OpTypeStruct %v4float + %B2Array = OpTypeArray %B2 %uint_4 + %C = OpTypeStruct %v4float %B2 %B2Array + %B1 = OpTypeStruct %v4float + %B1Array = OpTypeArray %B1 %uint_4 + %A = OpTypeStruct %v4float %B1 %B1Array + %SSBO = OpTypeStruct %A %C +%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO + %ssbo = OpVariable %_ptr_Uniform_SSBO StorageBuffer + %int = OpTypeInt 32 1 + %int_1 = OpConstant %int 1 +%_ptr_Uniform_C = OpTypePointer StorageBuffer %C + %int_0 = OpConstant %int 0 +%_ptr_Uniform_A = OpTypePointer StorageBuffer %A + %main = OpFunction %void None %3 + %5 = OpLabel + %22 = OpAccessChain %_ptr_Uniform_C %ssbo %int_1 + %39 = OpAccessChain %_ptr_Uniform_A %ssbo %int_0 + %23 = OpLoad %C %22 + %24 = OpCopyLogical %A %23 + OpStore %39 %24 + OpReturn + OpFunctionEnd diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index f6a1b8b6..7f9681e2 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -8555,6 +8555,22 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) break; } + case OpCopyLogical: + { + // This is used for copying object of different types, arrays and structs. + // We need to unroll the copy, element-by-element. + uint32_t result_type = ops[0]; + uint32_t id = ops[1]; + uint32_t rhs = ops[2]; + + emit_uninitialized_temporary_expression(result_type, id); + + auto &input_type = expression_type(rhs); + auto &output_type = get(result_type); + emit_copy_logical_type(to_expression(id), output_type, to_unpacked_expression(rhs), input_type); + break; + } + case OpCopyObject: { uint32_t result_type = ops[0]; @@ -13098,3 +13114,39 @@ void CompilerGLSL::propagate_nonuniform_qualifier(uint32_t id) propagate_nonuniform_qualifier(expr); } } + +void CompilerGLSL::emit_copy_logical_type(const std::string &lhs, const SPIRType &lhs_type, + const std::string &rhs, const SPIRType &rhs_type) +{ + if (!lhs_type.array.empty()) + { + // Could use a loop here to support specialization constants, but it gets rather complicated with nested array types, + // and this is a rather obscure opcode anyways, keep it simple unless we are forced to. + uint32_t array_size = to_array_size_literal(lhs_type); + auto &lhs_parent_type = get(lhs_type.parent_type); + auto &rhs_parent_type = get(rhs_type.parent_type); + + for (uint32_t i = 0; i < array_size; i++) + { + emit_copy_logical_type(join(lhs, "[", i, "]"), + lhs_parent_type, + join(rhs, "[", i, "]"), + rhs_parent_type); + } + } + else if (lhs_type.basetype == SPIRType::Struct) + { + uint32_t member_count = uint32_t(lhs_type.member_types.size()); + for (uint32_t i = 0; i < member_count; i++) + { + emit_copy_logical_type(join(lhs, ".", to_member_name(lhs_type, i)), + get(lhs_type.member_types[i]), + join(rhs, ".", to_member_name(rhs_type, i)), + get(rhs_type.member_types[i])); + } + } + else + { + statement(lhs, " = ", rhs, ";"); + } +} diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index 3326a244..9f982b16 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -275,6 +275,9 @@ protected: virtual bool builtin_translates_to_nonarray(spv::BuiltIn builtin) const; + void emit_copy_logical_type(const std::string &lhs, const SPIRType &lhs_type, + const std::string &rhs, const SPIRType &rhs_type); + StringStream<> buffer; template diff --git a/test_shaders.py b/test_shaders.py index 52a344f5..da94b7cc 100755 --- a/test_shaders.py +++ b/test_shaders.py @@ -178,7 +178,9 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths): spirv_path = create_temporary() msl_path = create_temporary(os.path.basename(shader)) - spirv_cmd = [paths.spirv_as, '--target-env', 'vulkan1.1', '-o', spirv_path, shader] + spirv_env = 'vulkan1.1spv1.4' if ('.spv14.' in shader) else 'vulkan1.1' + + spirv_cmd = [paths.spirv_as, '--target-env', spirv_env, '-o', spirv_path, shader] if '.preserve.' in shader: spirv_cmd.append('--preserve-numeric-ids') @@ -249,7 +251,7 @@ def cross_compile_msl(shader, spirv, opt, iterations, paths): subprocess.check_call(msl_args) if not shader_is_invalid_spirv(msl_path): - subprocess.check_call([paths.spirv_val, '--scalar-block-layout', '--target-env', 'vulkan1.1', spirv_path]) + subprocess.check_call([paths.spirv_val, '--scalar-block-layout', '--target-env', spirv_env, spirv_path]) return (spirv_path, msl_path) @@ -388,10 +390,12 @@ def cross_compile(shader, vulkan, spirv, invalid_spirv, eliminate, is_legacy, fl spirv_path = create_temporary() glsl_path = create_temporary(os.path.basename(shader)) + spirv_env = 'vulkan1.1spv1.4' if ('.spv14.' in shader) else 'vulkan1.1' + if vulkan or spirv: vulkan_glsl_path = create_temporary('vk' + os.path.basename(shader)) - spirv_cmd = [paths.spirv_as, '--target-env', 'vulkan1.1', '-o', spirv_path, shader] + spirv_cmd = [paths.spirv_as, '--target-env', spirv_env, '-o', spirv_path, shader] if '.preserve.' in shader: spirv_cmd.append('--preserve-numeric-ids') @@ -404,7 +408,7 @@ def cross_compile(shader, vulkan, spirv, invalid_spirv, eliminate, is_legacy, fl subprocess.check_call([paths.spirv_opt, '--skip-validation', '-O', '-o', spirv_path, spirv_path]) if not invalid_spirv: - subprocess.check_call([paths.spirv_val, '--scalar-block-layout', '--target-env', 'vulkan1.1', spirv_path]) + subprocess.check_call([paths.spirv_val, '--scalar-block-layout', '--target-env', spirv_env, spirv_path]) extra_args = ['--iterations', str(iterations)] if eliminate: From cf725b4c635b1a4a6efec7881307a9c260106fb3 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Mon, 6 Jan 2020 12:29:44 +0100 Subject: [PATCH 2/4] Go through access chain path for OpCopyLogical. We will need to deal with packing/unpacking data when copying from/to complex types in MSL. --- spirv_cross.cpp | 7 ++++- spirv_glsl.cpp | 72 +++++++++++++++++++++++++++++++++++++------------ spirv_glsl.hpp | 5 ++-- 3 files changed, 64 insertions(+), 20 deletions(-) diff --git a/spirv_cross.cpp b/spirv_cross.cpp index 286b4504..6e2cc95f 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -317,6 +317,8 @@ void Compiler::register_write(uint32_t chain) var = maybe_get(access_chain->loaded_from); } + auto &chain_type = expression_type(chain); + if (var) { bool check_argument_storage_qualifier = true; @@ -359,7 +361,7 @@ void Compiler::register_write(uint32_t chain) force_recompile(); } } - else + else if (chain_type.pointer) { // If we stored through a variable pointer, then we don't know which // variable we stored to. So *all* expressions after this point need to @@ -368,6 +370,9 @@ void Compiler::register_write(uint32_t chain) // only certain variables, we can invalidate only those. flush_all_active_variables(); } + + // If chain_type.pointer is false, we're not writing to memory backed variables, but temporaries instead. + // This can happen in copy_logical_type where we unroll complex reads and writes to temporaries. } void Compiler::flush_dependees(SPIRVariable &var) diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 7f9681e2..9fbb079e 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -8564,10 +8564,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) uint32_t rhs = ops[2]; emit_uninitialized_temporary_expression(result_type, id); - - auto &input_type = expression_type(rhs); - auto &output_type = get(result_type); - emit_copy_logical_type(to_expression(id), output_type, to_unpacked_expression(rhs), input_type); + emit_copy_logical_type(id, result_type, rhs, expression_type_id(rhs), {}); break; } @@ -13115,38 +13112,79 @@ void CompilerGLSL::propagate_nonuniform_qualifier(uint32_t id) } } -void CompilerGLSL::emit_copy_logical_type(const std::string &lhs, const SPIRType &lhs_type, - const std::string &rhs, const SPIRType &rhs_type) +void CompilerGLSL::emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id, + uint32_t rhs_id, uint32_t rhs_type_id, + SmallVector chain) { + // Fully unroll all member/array indices one by one. + + auto &lhs_type = get(lhs_type_id); + auto &rhs_type = get(rhs_type_id); + if (!lhs_type.array.empty()) { // Could use a loop here to support specialization constants, but it gets rather complicated with nested array types, // and this is a rather obscure opcode anyways, keep it simple unless we are forced to. uint32_t array_size = to_array_size_literal(lhs_type); - auto &lhs_parent_type = get(lhs_type.parent_type); - auto &rhs_parent_type = get(rhs_type.parent_type); + chain.push_back(0); for (uint32_t i = 0; i < array_size; i++) { - emit_copy_logical_type(join(lhs, "[", i, "]"), - lhs_parent_type, - join(rhs, "[", i, "]"), - rhs_parent_type); + chain.back() = i; + emit_copy_logical_type(lhs_id, lhs_type.parent_type, rhs_id, rhs_type.parent_type, chain); } } else if (lhs_type.basetype == SPIRType::Struct) { + chain.push_back(0); uint32_t member_count = uint32_t(lhs_type.member_types.size()); for (uint32_t i = 0; i < member_count; i++) { - emit_copy_logical_type(join(lhs, ".", to_member_name(lhs_type, i)), - get(lhs_type.member_types[i]), - join(rhs, ".", to_member_name(rhs_type, i)), - get(rhs_type.member_types[i])); + chain.back() = i; + emit_copy_logical_type(lhs_id, lhs_type.member_types[i], rhs_id, rhs_type.member_types[i], chain); } } else { - statement(lhs, " = ", rhs, ";"); + // Need to handle unpack/packing fixups since this can differ wildly between the logical types, + // particularly in MSL. + // To deal with this, we emit access chains and go through emit_store_statement + // to deal with all the special cases we can encounter. + + AccessChainMeta lhs_meta, rhs_meta; + auto lhs = access_chain_internal(lhs_id, chain.data(), chain.size(), ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &lhs_meta); + auto rhs = access_chain_internal(rhs_id, chain.data(), chain.size(), ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &rhs_meta); + + uint32_t id = ir.increase_bound_by(2); + lhs_id = id; + rhs_id = id + 1; + + { + auto &lhs_expr = set(lhs_id, move(lhs), lhs_type_id, true); + lhs_expr.need_transpose = lhs_meta.need_transpose; + + if (lhs_meta.storage_is_packed) + set_extended_decoration(lhs_id, SPIRVCrossDecorationPhysicalTypePacked); + if (lhs_meta.storage_physical_type != 0) + set_extended_decoration(lhs_id, SPIRVCrossDecorationPhysicalTypeID, lhs_meta.storage_physical_type); + + forwarded_temporaries.insert(lhs_id); + suppressed_usage_tracking.insert(lhs_id); + } + + { + auto &rhs_expr = set(rhs_id, move(rhs), rhs_type_id, true); + rhs_expr.need_transpose = rhs_meta.need_transpose; + + if (rhs_meta.storage_is_packed) + set_extended_decoration(rhs_id, SPIRVCrossDecorationPhysicalTypePacked); + if (rhs_meta.storage_physical_type != 0) + set_extended_decoration(rhs_id, SPIRVCrossDecorationPhysicalTypeID, rhs_meta.storage_physical_type); + + forwarded_temporaries.insert(rhs_id); + suppressed_usage_tracking.insert(rhs_id); + } + + emit_store_statement(lhs_id, rhs_id); } } diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index 9f982b16..0a1cb50f 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -275,8 +275,9 @@ protected: virtual bool builtin_translates_to_nonarray(spv::BuiltIn builtin) const; - void emit_copy_logical_type(const std::string &lhs, const SPIRType &lhs_type, - const std::string &rhs, const SPIRType &rhs_type); + void emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id, + uint32_t rhs_id, uint32_t rhs_type_id, + SmallVector chain); StringStream<> buffer; From 8bef6ff16787b11a35d72bc029f7b07df6c15c8c Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Mon, 6 Jan 2020 12:44:18 +0100 Subject: [PATCH 3/4] Add test shader for OpCopyLogical with packing/unpacking. --- .../asm/comp/copy-logical-2.spv14.asm.comp | 60 ++++++++++++++ .../asm/comp/copy-logical-2.spv14.asm.comp | 81 +++++++++++++++++++ 2 files changed, 141 insertions(+) create mode 100644 reference/shaders-msl-no-opt/asm/comp/copy-logical-2.spv14.asm.comp create mode 100644 shaders-msl-no-opt/asm/comp/copy-logical-2.spv14.asm.comp diff --git a/reference/shaders-msl-no-opt/asm/comp/copy-logical-2.spv14.asm.comp b/reference/shaders-msl-no-opt/asm/comp/copy-logical-2.spv14.asm.comp new file mode 100644 index 00000000..09a31d68 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/copy-logical-2.spv14.asm.comp @@ -0,0 +1,60 @@ +#include +#include + +using namespace metal; + +struct _11 +{ + float2x2 _m0; +}; + +struct _12 +{ + float2x4 _m0; +}; + +struct B2 +{ + float4 elem2; +}; + +struct C +{ + float4 c; + B2 b2; + B2 b2_array[4]; + _12 _m3; +}; + +struct B1 +{ + float4 elem1; +}; + +struct A +{ + float4 a; + B1 b1; + B1 b1_array[4]; + _11 _m3; +}; + +struct _8 +{ + A a_block; + C c_block; +}; + +kernel void main0(device _8& _3 [[buffer(0)]]) +{ + A _31; + _31.a = _3.c_block.c; + _31.b1.elem1 = _3.c_block.b2.elem2; + _31.b1_array[0].elem1 = _3.c_block.b2_array[0].elem2; + _31.b1_array[1].elem1 = _3.c_block.b2_array[1].elem2; + _31.b1_array[2].elem1 = _3.c_block.b2_array[2].elem2; + _31.b1_array[3].elem1 = _3.c_block.b2_array[3].elem2; + _31._m3._m0 = transpose(float2x2(_3.c_block._m3._m0[0].xy, _3.c_block._m3._m0[1].xy)); + _3.a_block = _31; +} + diff --git a/shaders-msl-no-opt/asm/comp/copy-logical-2.spv14.asm.comp b/shaders-msl-no-opt/asm/comp/copy-logical-2.spv14.asm.comp new file mode 100644 index 00000000..6a7065a6 --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/copy-logical-2.spv14.asm.comp @@ -0,0 +1,81 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 8 +; Bound: 48 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %ssbo + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %B1 "B1" + OpName %A "A" + OpName %C "C" + OpName %B2 "B2" + OpMemberName %A 0 "a" + OpMemberName %A 1 "b1" + OpMemberName %A 2 "b1_array" + OpMemberName %C 0 "c" + OpMemberName %C 1 "b2" + OpMemberName %C 2 "b2_array" + OpMemberName %B1 0 "elem1" + OpMemberName %B2 0 "elem2" + OpMemberName %SSBO 0 "a_block" + OpMemberName %SSBO 1 "c_block" + OpDecorate %B1Array ArrayStride 16 + OpDecorate %B2Array ArrayStride 16 + OpMemberDecorate %B1 0 Offset 0 + OpMemberDecorate %A 0 Offset 0 + OpMemberDecorate %A 1 Offset 16 + OpMemberDecorate %A 2 Offset 32 + OpMemberDecorate %A 3 Offset 96 + OpMemberDecorate %B2 0 Offset 0 + OpMemberDecorate %C 0 Offset 0 + OpMemberDecorate %C 1 Offset 16 + OpMemberDecorate %C 2 Offset 32 + OpMemberDecorate %C 3 Offset 96 + OpMemberDecorate %SSBO 0 Offset 0 + OpMemberDecorate %SSBO 1 Offset 112 + OpMemberDecorate %A0 0 Offset 0 + OpMemberDecorate %C0 0 Offset 0 + OpMemberDecorate %A0 0 RowMajor + OpMemberDecorate %A0 0 MatrixStride 8 + OpMemberDecorate %C0 0 ColMajor + OpMemberDecorate %C0 0 MatrixStride 16 + OpDecorate %SSBO Block + OpDecorate %ssbo DescriptorSet 0 + OpDecorate %ssbo Binding 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %float = OpTypeFloat 32 + %uint = OpTypeInt 32 0 + %uint_4 = OpConstant %uint 4 + %v4float = OpTypeVector %float 4 + %v2float = OpTypeVector %float 2 + %m2float = OpTypeMatrix %v2float 2 + %A0 = OpTypeStruct %m2float + %C0 = OpTypeStruct %m2float + %B2 = OpTypeStruct %v4float + %B2Array = OpTypeArray %B2 %uint_4 + %C = OpTypeStruct %v4float %B2 %B2Array %C0 + %B1 = OpTypeStruct %v4float + %B1Array = OpTypeArray %B1 %uint_4 + %A = OpTypeStruct %v4float %B1 %B1Array %A0 + %SSBO = OpTypeStruct %A %C +%_ptr_Uniform_SSBO = OpTypePointer StorageBuffer %SSBO + %ssbo = OpVariable %_ptr_Uniform_SSBO StorageBuffer + %int = OpTypeInt 32 1 + %int_1 = OpConstant %int 1 +%_ptr_Uniform_C = OpTypePointer StorageBuffer %C + %int_0 = OpConstant %int 0 +%_ptr_Uniform_A = OpTypePointer StorageBuffer %A + %main = OpFunction %void None %3 + %5 = OpLabel + %22 = OpAccessChain %_ptr_Uniform_C %ssbo %int_1 + %39 = OpAccessChain %_ptr_Uniform_A %ssbo %int_0 + %23 = OpLoad %C %22 + %24 = OpCopyLogical %A %23 + OpStore %39 %24 + OpReturn + OpFunctionEnd From 91e917b8ca272608da148737d4ed47dcae386b5d Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Mon, 6 Jan 2020 13:40:33 +0100 Subject: [PATCH 4/4] Fix Clang warnings. --- spirv_glsl.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 9fbb079e..26eec84a 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -13152,8 +13152,8 @@ void CompilerGLSL::emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id, // to deal with all the special cases we can encounter. AccessChainMeta lhs_meta, rhs_meta; - auto lhs = access_chain_internal(lhs_id, chain.data(), chain.size(), ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &lhs_meta); - auto rhs = access_chain_internal(rhs_id, chain.data(), chain.size(), ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &rhs_meta); + auto lhs = access_chain_internal(lhs_id, chain.data(), uint32_t(chain.size()), ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &lhs_meta); + auto rhs = access_chain_internal(rhs_id, chain.data(), uint32_t(chain.size()), ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &rhs_meta); uint32_t id = ir.increase_bound_by(2); lhs_id = id;