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

MSL: Test that we can mask location writes in TESC.

Browse Source
This commit is contained in:
Hans-Kristian Arntzen 2021-04-06 15:50:02 +02:00
parent 3255d6cef0
commit ae7bb41ef4
10 changed files with 240 additions and 15 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

34
reference/opt/shaders-msl/masking/write-outputs.mask-location-0.msl2.tesc Normal file
View File

@ -0,0 +1,34 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct main0_out
{
float4 gl_Position;
float gl_PointSize;
};
struct main0_patchOut
{
float4 v1;
};
kernel void main0(uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
{
threadgroup float4 v0[4];
device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
v0[gl_InvocationID] = float4(1.0);
((threadgroup float*)&v0[gl_InvocationID])[0u] = 2.0;
if (gl_InvocationID == 0)
{
patchOut.v1 = float4(2.0);
((device float*)&patchOut.v1)[3u] = 4.0;
}
gl_out[gl_InvocationID].gl_Position = float4(3.0);
gl_out[gl_InvocationID].gl_PointSize = 4.0;
gl_out[gl_InvocationID].gl_Position.z = 5.0;
gl_out[gl_InvocationID].gl_PointSize = 4.0;
}

33
reference/opt/shaders-msl/masking/write-outputs.mask-location-1.msl2.tesc Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct main0_out
{
float4 v0;
float4 gl_Position;
float gl_PointSize;
};
struct main0_patchOut
{
};
kernel void main0(uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
{
threadgroup float4 v1;
device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
gl_out[gl_InvocationID].v0 = float4(1.0);
gl_out[gl_InvocationID].v0.x = 2.0;
if (gl_InvocationID == 0)
{
v1 = float4(2.0);
((threadgroup float*)&v1)[3u] = 4.0;
}
gl_out[gl_InvocationID].gl_Position = float4(3.0);
gl_out[gl_InvocationID].gl_PointSize = 4.0;
gl_out[gl_InvocationID].gl_Position.z = 5.0;
gl_out[gl_InvocationID].gl_PointSize = 4.0;
}

42
reference/shaders-msl/masking/write-outputs.mask-location-0.msl2.tesc Normal file
View File

@ -0,0 +1,42 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct main0_out
{
float4 gl_Position;
float gl_PointSize;
};
struct main0_patchOut
{
float4 v1;
};
static inline __attribute__((always_inline))
void write_in_func(threadgroup float4 (&v0)[4], thread uint& gl_InvocationID, device float4& v1, device main0_out* thread & gl_out)
{
v0[gl_InvocationID] = float4(1.0);
((threadgroup float*)&v0[gl_InvocationID])[0u] = 2.0;
if (gl_InvocationID == 0)
{
v1 = float4(2.0);
((device float*)&v1)[3u] = 4.0;
}
gl_out[gl_InvocationID].gl_Position = float4(3.0);
gl_out[gl_InvocationID].gl_PointSize = 4.0;
gl_out[gl_InvocationID].gl_Position.z = 5.0;
gl_out[gl_InvocationID].gl_PointSize = 4.0;
}
kernel void main0(uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
{
threadgroup float4 v0[4];
device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
write_in_func(v0, gl_InvocationID, patchOut.v1, gl_out);
}

41
reference/shaders-msl/masking/write-outputs.mask-location-1.msl2.tesc Normal file
View File

@ -0,0 +1,41 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct main0_out
{
float4 v0;
float4 gl_Position;
float gl_PointSize;
};
struct main0_patchOut
{
};
static inline __attribute__((always_inline))
void write_in_func(device main0_out* thread & gl_out, thread uint& gl_InvocationID, threadgroup float4& v1)
{
gl_out[gl_InvocationID].v0 = float4(1.0);
gl_out[gl_InvocationID].v0.x = 2.0;
if (gl_InvocationID == 0)
{
v1 = float4(2.0);
((threadgroup float*)&v1)[3u] = 4.0;
}
gl_out[gl_InvocationID].gl_Position = float4(3.0);
gl_out[gl_InvocationID].gl_PointSize = 4.0;
gl_out[gl_InvocationID].gl_Position.z = 5.0;
gl_out[gl_InvocationID].gl_PointSize = 4.0;
}
kernel void main0(uint gl_InvocationID [[thread_index_in_threadgroup]], uint gl_PrimitiveID [[threadgroup_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device main0_patchOut* spvPatchOut [[buffer(27)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]])
{
threadgroup float4 v1;
device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4];
device main0_patchOut& patchOut = spvPatchOut[gl_PrimitiveID];
write_in_func(gl_out, gl_InvocationID, v1);
}

26
shaders-msl/masking/write-outputs.mask-location-0.msl2.tesc Normal file
View File

@ -0,0 +1,26 @@
#version 450
layout(vertices = 4) out;
layout(location = 0) out vec4 v0[];
layout(location = 1) patch out vec4 v1;
void write_in_func()
{
v0[gl_InvocationID] = vec4(1.0);
v0[gl_InvocationID][0] = 2.0;
if (gl_InvocationID == 0)
{
v1 = vec4(2.0);
v1[3] = 4.0;
}
gl_out[gl_InvocationID].gl_Position = vec4(3.0);
gl_out[gl_InvocationID].gl_PointSize = 4.0;
gl_out[gl_InvocationID].gl_Position[2] = 5.0;
gl_out[gl_InvocationID].gl_PointSize = 4.0;
}
void main()
{
write_in_func();
}

26
shaders-msl/masking/write-outputs.mask-location-1.msl2.tesc Normal file
View File

@ -0,0 +1,26 @@
#version 450
layout(vertices = 4) out;
layout(location = 0) out vec4 v0[];
layout(location = 1) patch out vec4 v1;
void write_in_func()
{
v0[gl_InvocationID] = vec4(1.0);
v0[gl_InvocationID][0] = 2.0;
if (gl_InvocationID == 0)
{
v1 = vec4(2.0);
v1[3] = 4.0;
}
gl_out[gl_InvocationID].gl_Position = vec4(3.0);
gl_out[gl_InvocationID].gl_PointSize = 4.0;
gl_out[gl_InvocationID].gl_Position[2] = 5.0;
gl_out[gl_InvocationID].gl_PointSize = 4.0;
}
void main()
{
write_in_func();
}

10
spirv_glsl.cpp
View File

@ -8656,8 +8656,10 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
if ((flags & ACCESS_CHAIN_FORCE_COMPOSITE_BIT) == 0)
{
auto *var = maybe_get_backing_variable(base);
if (var && variable_decl_is_threadgroup_like(*var))
if (var && variable_decl_is_remapped_storage(*var, StorageClassWorkgroup))
effective_storage = StorageClassWorkgroup;
else if (var && variable_decl_is_remapped_storage(*var, StorageClassStorageBuffer))
effective_storage = StorageClassStorageBuffer;
else if (expression_type(base).pointer)
effective_storage = get_expression_effective_storage_class(base);
}
@ -12650,9 +12652,9 @@ string CompilerGLSL::variable_decl(const SPIRType &type, const string &name, uin
return join(type_name, " ", name, type_to_array_glsl(type));
}
bool CompilerGLSL::variable_decl_is_threadgroup_like(const SPIRVariable &var) const
bool CompilerGLSL::variable_decl_is_remapped_storage(const SPIRVariable &var, StorageClass storage) const
{
return var.storage == StorageClassWorkgroup;
return var.storage == storage;
}
// Emit a structure member. Subclasses may override to modify output,
@ -13570,7 +13572,7 @@ void CompilerGLSL::emit_function(SPIRFunction &func, const Bitset &return_flags)
auto &var = get<SPIRVariable>(v);
var.deferred_declaration = false;
if (variable_decl_is_threadgroup_like(var))
if (variable_decl_is_remapped_storage(var, StorageClassWorkgroup))
{
// Special variable type which cannot have initializer,
// need to be declared as standalone variables.

2
spirv_glsl.hpp
View File

@ -386,7 +386,7 @@ protected:
virtual std::string constant_expression_vector(const SPIRConstant &c, uint32_t vector);
virtual void emit_fixup();
virtual std::string variable_decl(const SPIRType &type, const std::string &name, uint32_t id = 0);
virtual bool variable_decl_is_threadgroup_like(const SPIRVariable &var) const;
virtual bool variable_decl_is_remapped_storage(const SPIRVariable &var, spv::StorageClass storage) const;
virtual std::string to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id);
struct TextureFunctionBaseArguments

39
spirv_msl.cpp
View File

@ -2900,7 +2900,7 @@ void CompilerMSL::add_variable_to_interface_block(StorageClass storage, const st
// If we ignore an output, we must still emit it, since it might be used by app.
// Instead, just emit it as early declaration.
entry_func.add_local_variable(var.self);
if (!variable_decl_is_threadgroup_like(var))
if (!variable_decl_is_remapped_storage(var, StorageClassWorkgroup))
vars_needing_early_declaration.push_back(var.self);
return;
}
@ -12750,7 +12750,7 @@ string CompilerMSL::to_qualifiers_glsl(uint32_t id)
auto *var = maybe_get<SPIRVariable>(id);
auto &type = expression_type(id);
if (type.storage == StorageClassWorkgroup || (var && variable_decl_is_threadgroup_like(*var)))
if (type.storage == StorageClassWorkgroup || (var && variable_decl_is_remapped_storage(*var, StorageClassWorkgroup)))
quals += "threadgroup ";
return quals;
@ -12931,13 +12931,33 @@ string CompilerMSL::type_to_array_glsl(const SPIRType &type)
}
}
bool CompilerMSL::variable_decl_is_threadgroup_like(const SPIRVariable &variable) const
bool CompilerMSL::variable_decl_is_remapped_storage(const SPIRVariable &variable, spv::StorageClass storage) const
{
auto model = get_execution_model();
return variable.storage == StorageClassWorkgroup ||
(variable.storage == StorageClassOutput &&
model == ExecutionModelTessellationControl &&
is_stage_output_variable_masked(variable));
if (variable.storage == storage)
return true;
if (storage == StorageClassWorkgroup)
{
auto model = get_execution_model();
return variable.storage == StorageClassOutput &&
model == ExecutionModelTessellationControl &&
is_stage_output_variable_masked(variable);
}
else if (storage == StorageClassStorageBuffer)
{
// We won't be able to catch writes to control point outputs here since variable
// refers to a function local pointer.
// This is fine, as there cannot be concurrent writers to that memory anyways,
// so we just ignore that case.
return capture_output_to_buffer &&
variable.storage == StorageClassOutput &&
!is_stage_output_variable_masked(variable);
}
else
{
return false;
}
}
std::string CompilerMSL::variable_decl(const SPIRVariable &variable)
@ -12948,7 +12968,8 @@ std::string CompilerMSL::variable_decl(const SPIRVariable &variable)
// More special cases. ClipDistance and CullDistance are emitted as plain arrays in stage out,
// so preserve that property when emitting them as masked variables. Avoids lots of extra special casing
// in argument_decl(). Similar argument for TessLevels.
if (variable_decl_is_threadgroup_like(variable) || has_decoration(variable.self, DecorationBuiltIn))
if (variable_decl_is_remapped_storage(variable, StorageClassWorkgroup) ||
has_decoration(variable.self, DecorationBuiltIn))
is_using_builtin_array = true;
std::string expr = CompilerGLSL::variable_decl(variable);

2
spirv_msl.hpp
View File

@ -736,7 +736,7 @@ protected:
// Threadgroup arrays can't have a wrapper type
std::string variable_decl(const SPIRVariable &variable) override;
bool variable_decl_is_threadgroup_like(const SPIRVariable &variable) const override;
bool variable_decl_is_remapped_storage(const SPIRVariable &variable, spv::StorageClass storage) const override;
// GCC workaround of lambdas calling protected functions (for older GCC versions)
std::string variable_decl(const SPIRType &type, const std::string &name, uint32_t id = 0) override;