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

MSL: Also pack members at unaligned offsets.

Browse Source 
This is necessary to support `VK_KHR_relaxed_block_layout`.
This commit is contained in:
Chip Davis 2018-11-07 00:31:11 -06:00
parent daaffc4717
commit e50eecfeeb
4 changed files with 95 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
reference/opt/shaders-msl/asm/comp/relaxed-block-layout.asm.comp Normal file
View File

@ -0,0 +1,17 @@
#include <metal_stdlib>
#include <simd/simd.h>
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);
}

17
reference/shaders-msl/asm/comp/relaxed-block-layout.asm.comp Normal file
View File

@ -0,0 +1,17 @@
#include <metal_stdlib>
#include <simd/simd.h>
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);
}

56
shaders-msl/asm/comp/relaxed-block-layout.asm.comp Normal file
View File

@ -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

6
spirv_msl.cpp
View File

@ -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;
}