MSL: Always emit block variable for block types.

This commit is contained in:
Hans-Kristian Arntzen 2021-04-07 10:55:40 +02:00
parent ae7bb41ef4
commit 74b2acab9b
5 changed files with 5 additions and 1 deletions

View File

@ -21,6 +21,7 @@ struct main0_out
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
{
V _22 = {};
device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
if (any(gl_GlobalInvocationID >= spvStageInputSize))
return;

View File

@ -21,6 +21,7 @@ struct main0_out
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
{
V _22 = {};
device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
if (any(gl_GlobalInvocationID >= spvStageInputSize))
return;

View File

@ -21,6 +21,7 @@ struct main0_out
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
{
V _22 = {};
device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
if (any(gl_GlobalInvocationID >= spvStageInputSize))
return;

View File

@ -21,6 +21,7 @@ struct main0_out
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
{
V _22 = {};
device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
if (any(gl_GlobalInvocationID >= spvStageInputSize))
return;

View File

@ -2907,7 +2907,7 @@ void CompilerMSL::add_variable_to_interface_block(StorageClass storage, const st
if (var_type.basetype == SPIRType::Struct)
{
if (!is_builtin_type(var_type) && (!capture_output_to_buffer || storage == StorageClassInput) &&
if (!is_builtin_type(var_type) && (!capture_output_to_buffer || is_block || storage == StorageClassInput) &&
!meta.strip_array)
{
// For I/O blocks or structs, we will need to pass the block itself around