From e50eecfeebca4a5bfb9f340be08b8cd4326826fc Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 7 Nov 2018 00:31:11 -0600 Subject: [PATCH] MSL: Also pack members at unaligned offsets. This is necessary to support `VK_KHR_relaxed_block_layout`. --- .../asm/comp/relaxed-block-layout.asm.comp | 17 ++++++ .../asm/comp/relaxed-block-layout.asm.comp | 17 ++++++ .../asm/comp/relaxed-block-layout.asm.comp | 56 +++++++++++++++++++ spirv_msl.cpp | 6 +- 4 files changed, 95 insertions(+), 1 deletion(-) create mode 100644 reference/opt/shaders-msl/asm/comp/relaxed-block-layout.asm.comp create mode 100644 reference/shaders-msl/asm/comp/relaxed-block-layout.asm.comp create mode 100644 shaders-msl/asm/comp/relaxed-block-layout.asm.comp diff --git a/reference/opt/shaders-msl/asm/comp/relaxed-block-layout.asm.comp b/reference/opt/shaders-msl/asm/comp/relaxed-block-layout.asm.comp new file mode 100644 index 00000000..12547276 --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/relaxed-block-layout.asm.comp @@ -0,0 +1,17 @@ +#include +#include + +using namespace metal; + +struct foo +{ + uint bar; + packed_float3 baz; +}; + +kernel void main0(device foo& _6 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _6.bar = gl_LocalInvocationID.x; + _6.baz = float3(gl_GlobalInvocationID); +} + diff --git a/reference/shaders-msl/asm/comp/relaxed-block-layout.asm.comp b/reference/shaders-msl/asm/comp/relaxed-block-layout.asm.comp new file mode 100644 index 00000000..12547276 --- /dev/null +++ b/reference/shaders-msl/asm/comp/relaxed-block-layout.asm.comp @@ -0,0 +1,17 @@ +#include +#include + +using namespace metal; + +struct foo +{ + uint bar; + packed_float3 baz; +}; + +kernel void main0(device foo& _6 [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _6.bar = gl_LocalInvocationID.x; + _6.baz = float3(gl_GlobalInvocationID); +} + diff --git a/shaders-msl/asm/comp/relaxed-block-layout.asm.comp b/shaders-msl/asm/comp/relaxed-block-layout.asm.comp new file mode 100644 index 00000000..2545cc7b --- /dev/null +++ b/shaders-msl/asm/comp/relaxed-block-layout.asm.comp @@ -0,0 +1,56 @@ +; SPIR-V +; Version: 1.3 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 29 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_LocalInvocationID %gl_GlobalInvocationID + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %main "main" + OpName %foo "foo" + OpMemberName %foo 0 "bar" + OpMemberName %foo 1 "baz" + OpName %_ "" + OpName %gl_LocalInvocationID "gl_LocalInvocationID" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpMemberDecorate %foo 0 Offset 0 + OpMemberDecorate %foo 1 Offset 4 + OpDecorate %foo BufferBlock + OpDecorate %_ DescriptorSet 0 + OpDecorate %_ Binding 0 + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + %void = OpTypeVoid + %3 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %float = OpTypeFloat 32 + %v3float = OpTypeVector %float 3 + %foo = OpTypeStruct %uint %v3float +%_ptr_Uniform_foo = OpTypePointer Uniform %foo + %_ = OpVariable %_ptr_Uniform_foo Uniform + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %uint_0 = OpConstant %uint 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %int_1 = OpConstant %int 1 +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float + %main = OpFunction %void None %3 + %5 = OpLabel + %19 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %20 = OpLoad %uint %19 + %22 = OpAccessChain %_ptr_Uniform_uint %_ %int_0 + OpStore %22 %20 + %25 = OpLoad %v3uint %gl_GlobalInvocationID + %26 = OpConvertUToF %v3float %25 + %28 = OpAccessChain %_ptr_Uniform_v3float %_ %int_1 + OpStore %28 %26 + OpReturn + OpFunctionEnd diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 589914c9..5217bc0f 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -1256,11 +1256,15 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index) } else { + uint32_t mbr_offset_curr = get_member_decoration(ib_type.self, index, DecorationOffset); + // For vectors, pack if the member's offset doesn't conform to the + // type's usual alignment. For example, a float3 at offset 4. + if (!is_matrix(mbr_type) && (mbr_offset_curr % unpacked_mbr_size)) + return true; // Pack if there is not enough space between this member and next. // If last member, only pack if it's a row-major matrix. if (index < ib_type.member_types.size() - 1) { - uint32_t mbr_offset_curr = get_member_decoration(ib_type.self, index, DecorationOffset); uint32_t mbr_offset_next = get_member_decoration(ib_type.self, index + 1, DecorationOffset); return unpacked_mbr_size > mbr_offset_next - mbr_offset_curr; }