diff --git a/reference/opt/shaders-msl/tesc/reload-tess-level.multi-patch.tesc b/reference/opt/shaders-msl/tesc/reload-tess-level.multi-patch.tesc new file mode 100644 index 00000000..a55755e0 --- /dev/null +++ b/reference/opt/shaders-msl/tesc/reload-tess-level.multi-patch.tesc @@ -0,0 +1,35 @@ +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 gl_Position; +}; + +struct main0_in +{ + uint3 m_82; + ushort2 m_86; + float4 gl_Position; +}; + +kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]]) +{ + device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4]; + device main0_in* gl_in = &spvIn[min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1) * spvIndirectParams[0]]; + uint gl_InvocationID = gl_GlobalInvocationID.x % 4; + uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1]); + if (gl_InvocationID == 0) + { + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(2.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(3.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(4.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(5.0); + spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(mix(float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0]), float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3]), 0.5)); + spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(mix(float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2]), float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1]), 0.5)); + } + gl_out[gl_InvocationID].gl_Position = gl_in[gl_InvocationID].gl_Position; +} + diff --git a/reference/opt/shaders-msl/tesc/reload-tess-level.tesc b/reference/opt/shaders-msl/tesc/reload-tess-level.tesc new file mode 100644 index 00000000..eafc5060 --- /dev/null +++ b/reference/opt/shaders-msl/tesc/reload-tess-level.tesc @@ -0,0 +1,35 @@ +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 gl_Position; +}; + +struct main0_in +{ + float4 gl_Position [[attribute(0)]]; +}; + +kernel void main0(main0_in in [[stage_in]], 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 MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]]) +{ + device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4]; + if (gl_InvocationID < spvIndirectParams[0]) + gl_in[gl_InvocationID] = in; + threadgroup_barrier(mem_flags::mem_threadgroup); + if (gl_InvocationID >= 4) + return; + if (gl_InvocationID == 0) + { + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(2.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(3.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(4.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(5.0); + spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(mix(float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0]), float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3]), 0.5)); + spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(mix(float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2]), float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1]), 0.5)); + } + gl_out[gl_InvocationID].gl_Position = gl_in[gl_InvocationID].gl_Position; +} + diff --git a/reference/shaders-msl/tesc/reload-tess-level.multi-patch.tesc b/reference/shaders-msl/tesc/reload-tess-level.multi-patch.tesc new file mode 100644 index 00000000..a55755e0 --- /dev/null +++ b/reference/shaders-msl/tesc/reload-tess-level.multi-patch.tesc @@ -0,0 +1,35 @@ +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 gl_Position; +}; + +struct main0_in +{ + uint3 m_82; + ushort2 m_86; + float4 gl_Position; +}; + +kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], device main0_out* spvOut [[buffer(28)]], constant uint* spvIndirectParams [[buffer(29)]], device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], device main0_in* spvIn [[buffer(22)]]) +{ + device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 4]; + device main0_in* gl_in = &spvIn[min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1] - 1) * spvIndirectParams[0]]; + uint gl_InvocationID = gl_GlobalInvocationID.x % 4; + uint gl_PrimitiveID = min(gl_GlobalInvocationID.x / 4, spvIndirectParams[1]); + if (gl_InvocationID == 0) + { + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(2.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(3.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(4.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(5.0); + spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(mix(float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0]), float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3]), 0.5)); + spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(mix(float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2]), float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1]), 0.5)); + } + gl_out[gl_InvocationID].gl_Position = gl_in[gl_InvocationID].gl_Position; +} + diff --git a/reference/shaders-msl/tesc/reload-tess-level.tesc b/reference/shaders-msl/tesc/reload-tess-level.tesc new file mode 100644 index 00000000..eafc5060 --- /dev/null +++ b/reference/shaders-msl/tesc/reload-tess-level.tesc @@ -0,0 +1,35 @@ +#include +#include + +using namespace metal; + +struct main0_out +{ + float4 gl_Position; +}; + +struct main0_in +{ + float4 gl_Position [[attribute(0)]]; +}; + +kernel void main0(main0_in in [[stage_in]], 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 MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]], threadgroup main0_in* gl_in [[threadgroup(0)]]) +{ + device main0_out* gl_out = &spvOut[gl_PrimitiveID * 4]; + if (gl_InvocationID < spvIndirectParams[0]) + gl_in[gl_InvocationID] = in; + threadgroup_barrier(mem_flags::mem_threadgroup); + if (gl_InvocationID >= 4) + return; + if (gl_InvocationID == 0) + { + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0] = half(2.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1] = half(3.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2] = half(4.0); + spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3] = half(5.0); + spvTessLevel[gl_PrimitiveID].insideTessellationFactor[0] = half(mix(float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[0]), float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[3]), 0.5)); + spvTessLevel[gl_PrimitiveID].insideTessellationFactor[1] = half(mix(float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[2]), float(spvTessLevel[gl_PrimitiveID].edgeTessellationFactor[1]), 0.5)); + } + gl_out[gl_InvocationID].gl_Position = gl_in[gl_InvocationID].gl_Position; +} + diff --git a/shaders-msl/tesc/reload-tess-level.multi-patch.tesc b/shaders-msl/tesc/reload-tess-level.multi-patch.tesc new file mode 100644 index 00000000..c3f0195c --- /dev/null +++ b/shaders-msl/tesc/reload-tess-level.multi-patch.tesc @@ -0,0 +1,17 @@ +#version 450 +layout(vertices = 4) out; + +void main() +{ + if (gl_InvocationID == 0) + { + gl_TessLevelOuter[0] = 2.0; + gl_TessLevelOuter[1] = 3.0; + gl_TessLevelOuter[2] = 4.0; + gl_TessLevelOuter[3] = 5.0; + gl_TessLevelInner[0] = mix(gl_TessLevelOuter[0], gl_TessLevelOuter[3], 0.5); + gl_TessLevelInner[1] = mix(gl_TessLevelOuter[2], gl_TessLevelOuter[1], 0.5); + } + + gl_out[gl_InvocationID].gl_Position = gl_in[gl_InvocationID].gl_Position; +} diff --git a/shaders-msl/tesc/reload-tess-level.tesc b/shaders-msl/tesc/reload-tess-level.tesc new file mode 100644 index 00000000..c3f0195c --- /dev/null +++ b/shaders-msl/tesc/reload-tess-level.tesc @@ -0,0 +1,17 @@ +#version 450 +layout(vertices = 4) out; + +void main() +{ + if (gl_InvocationID == 0) + { + gl_TessLevelOuter[0] = 2.0; + gl_TessLevelOuter[1] = 3.0; + gl_TessLevelOuter[2] = 4.0; + gl_TessLevelOuter[3] = 5.0; + gl_TessLevelInner[0] = mix(gl_TessLevelOuter[0], gl_TessLevelOuter[3], 0.5); + gl_TessLevelInner[1] = mix(gl_TessLevelOuter[2], gl_TessLevelOuter[1], 0.5); + } + + gl_out[gl_InvocationID].gl_Position = gl_in[gl_InvocationID].gl_Position; +} diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 130c2884..5589d889 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -9313,8 +9313,8 @@ void CompilerGLSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_ex auto lhs = to_dereferenced_expression(lhs_expression); - // We might need to bitcast in order to store to a builtin. - bitcast_to_builtin_store(lhs_expression, rhs, expression_type(rhs_expression)); + // We might need to cast in order to store to a builtin. + cast_to_builtin_store(lhs_expression, rhs, expression_type(rhs_expression)); // Tries to optimize assignments like " = op expr". // While this is purely cosmetic, this is important for legacy ESSL where loop @@ -9477,8 +9477,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) if (expr_type.vecsize > type.vecsize) expr = enclose_expression(expr + vector_swizzle(type.vecsize, 0)); - // We might need to bitcast in order to load from a builtin. - bitcast_from_builtin_load(ptr, expr, type); + // We might need to cast in order to load from a builtin. + cast_from_builtin_load(ptr, expr, type); // We might be trying to load a gl_Position[N], where we should be // doing float4[](gl_in[i].gl_Position, ...) instead. @@ -14484,7 +14484,7 @@ void CompilerGLSL::unroll_array_from_complex_load(uint32_t target_id, uint32_t s } } -void CompilerGLSL::bitcast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) +void CompilerGLSL::cast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) { auto *var = maybe_get_backing_variable(source_id); if (var) @@ -14536,7 +14536,7 @@ void CompilerGLSL::bitcast_from_builtin_load(uint32_t source_id, std::string &ex expr = bitcast_expression(expr_type, expected_type, expr); } -void CompilerGLSL::bitcast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) +void CompilerGLSL::cast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) { // Only interested in standalone builtin variables. if (!has_decoration(target_id, DecorationBuiltIn)) diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index 33625f38..d8125456 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -854,9 +854,9 @@ protected: // Builtins in GLSL are always specific signedness, but the SPIR-V can declare them // as either unsigned or signed. - // Sometimes we will need to automatically perform bitcasts on load and store to make this work. - virtual void bitcast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type); - virtual void bitcast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type); + // Sometimes we will need to automatically perform casts on load and store to make this work. + virtual void cast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type); + virtual void cast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type); void unroll_array_from_complex_load(uint32_t target_id, uint32_t source_id, std::string &expr); void convert_non_uniform_expression(const SPIRType &type, std::string &expr); diff --git a/spirv_msl.cpp b/spirv_msl.cpp index ddef2ce3..39798466 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -13299,7 +13299,7 @@ void CompilerMSL::remap_constexpr_sampler_by_binding(uint32_t desc_set, uint32_t constexpr_samplers_by_binding[{ desc_set, binding }] = sampler; } -void CompilerMSL::bitcast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) +void CompilerMSL::cast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) { auto *var = maybe_get_backing_variable(source_id); if (var) @@ -13311,6 +13311,7 @@ void CompilerMSL::bitcast_from_builtin_load(uint32_t source_id, std::string &exp auto builtin = static_cast(get_decoration(source_id, DecorationBuiltIn)); auto expected_type = expr_type.basetype; + auto expected_width = expr_type.width; switch (builtin) { case BuiltInGlobalInvocationId: @@ -13331,12 +13332,16 @@ void CompilerMSL::bitcast_from_builtin_load(uint32_t source_id, std::string &exp case BuiltInBaseInstance: case BuiltInBaseVertex: expected_type = SPIRType::UInt; + expected_width = 32; break; case BuiltInTessLevelInner: case BuiltInTessLevelOuter: if (get_execution_model() == ExecutionModelTessellationControl) + { expected_type = SPIRType::Half; + expected_width = 16; + } break; default: @@ -13344,7 +13349,17 @@ void CompilerMSL::bitcast_from_builtin_load(uint32_t source_id, std::string &exp } if (expected_type != expr_type.basetype) - expr = bitcast_expression(expr_type, expected_type, expr); + { + if (expected_width != expr_type.width) + { + // These are of different widths, so we cannot do a straight bitcast. + expr = join(type_to_glsl(expr_type), "(", expr, ")"); + } + else + { + expr = bitcast_expression(expr_type, expected_type, expr); + } + } if (builtin == BuiltInTessCoord && get_entry_point().flags.get(ExecutionModeQuads) && expr_type.vecsize == 3) { @@ -13354,7 +13369,7 @@ void CompilerMSL::bitcast_from_builtin_load(uint32_t source_id, std::string &exp } } -void CompilerMSL::bitcast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) +void CompilerMSL::cast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) { auto *var = maybe_get_backing_variable(target_id); if (var) @@ -13366,6 +13381,7 @@ void CompilerMSL::bitcast_to_builtin_store(uint32_t target_id, std::string &expr auto builtin = static_cast(get_decoration(target_id, DecorationBuiltIn)); auto expected_type = expr_type.basetype; + auto expected_width = expr_type.width; switch (builtin) { case BuiltInLayer: @@ -13374,11 +13390,13 @@ void CompilerMSL::bitcast_to_builtin_store(uint32_t target_id, std::string &expr case BuiltInPrimitiveId: case BuiltInViewIndex: expected_type = SPIRType::UInt; + expected_width = 32; break; case BuiltInTessLevelInner: case BuiltInTessLevelOuter: expected_type = SPIRType::Half; + expected_width = 16; break; default: @@ -13387,10 +13405,13 @@ void CompilerMSL::bitcast_to_builtin_store(uint32_t target_id, std::string &expr if (expected_type != expr_type.basetype) { - if (expected_type == SPIRType::Half && expr_type.basetype == SPIRType::Float) + if (expected_width != expr_type.width) { // These are of different widths, so we cannot do a straight bitcast. - expr = join("half(", expr, ")"); + auto type = expr_type; + type.basetype = expected_type; + type.width = expected_width; + expr = join(type_to_glsl(type), "(", expr, ")"); } else { diff --git a/spirv_msl.hpp b/spirv_msl.hpp index c8c9d8f1..949e30ec 100644 --- a/spirv_msl.hpp +++ b/spirv_msl.hpp @@ -844,8 +844,8 @@ protected: bool does_shader_write_sample_mask = false; - void bitcast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) override; - void bitcast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) override; + void cast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) override; + void cast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) override; void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression) override; void analyze_sampled_image_usage();