From 7920249ed1e444eb33eba0c8d2af472ce76d2b36 Mon Sep 17 00:00:00 2001 From: Samuliak Date: Sat, 4 May 2024 14:49:33 +0200 Subject: [PATCH] move input generation to context --- .../backend/msl/emit_msl.cpp | 206 +--------------- .../backend/msl/msl_emit_context.cpp | 226 +++++++++++++++++- .../backend/msl/msl_emit_context.h | 3 + 3 files changed, 224 insertions(+), 211 deletions(-) diff --git a/src/shader_recompiler/backend/msl/emit_msl.cpp b/src/shader_recompiler/backend/msl/emit_msl.cpp index 45c2141528..ce6b902c53 100644 --- a/src/shader_recompiler/backend/msl/emit_msl.cpp +++ b/src/shader_recompiler/backend/msl/emit_msl.cpp @@ -205,213 +205,12 @@ void DefineVariables(const EmitContext& ctx, std::string& header) { header += fmt::format("int loop{}=0x2000;", i); } } - -// TODO -std::string_view DepthSamplerType(TextureType type) { - switch (type) { - case TextureType::Color1D: - return "sampler1DShadow"; - case TextureType::ColorArray1D: - return "sampler1DArrayShadow"; - case TextureType::Color2D: - return "sampler2DShadow"; - case TextureType::ColorArray2D: - return "sampler2DArrayShadow"; - case TextureType::ColorCube: - return "samplerCubeShadow"; - case TextureType::ColorArrayCube: - return "samplerCubeArrayShadow"; - default: - throw NotImplementedException("Texture type: {}", type); - } -} - -// TODO: emit sampler as well -// TODO: handle multisample -// TODO: handle texture buffer -std::string_view ColorSamplerType(TextureType type, bool is_multisample = false) { - if (is_multisample) { - ASSERT(type == TextureType::Color2D || type == TextureType::ColorArray2D); - } - switch (type) { - case TextureType::Color1D: - return "texture1d"; - case TextureType::ColorArray1D: - return "texture1d_array"; - case TextureType::Color2D: - case TextureType::Color2DRect: - return "texture2d"; - case TextureType::ColorArray2D: - return "texture2d_array"; - case TextureType::Color3D: - return "texture3d"; - case TextureType::ColorCube: - return "texturecube"; - case TextureType::ColorArrayCube: - return "texturecube_array"; - default: - throw NotImplementedException("Texture type: {}", type); - } -} - -// TODO: handle texture buffer -std::string_view ImageType(TextureType type) { - switch (type) { - case TextureType::Color1D: - return "texture1d"; - case TextureType::ColorArray1D: - return "texture1d_array"; - case TextureType::Color2D: - return "texture2d"; - case TextureType::ColorArray2D: - return "texture2d_array"; - case TextureType::Color3D: - return "texture3d"; - case TextureType::ColorCube: - return "texturecube"; - case TextureType::ColorArrayCube: - return "texturecube_array"; - default: - throw NotImplementedException("Image type: {}", type); - } -} - -// TODO: is this needed? -/* -std::string_view ImageFormatString(ImageFormat format) { - switch (format) { - case ImageFormat::Typeless: - return ""; - case ImageFormat::R8_UINT: - return ",r8ui"; - case ImageFormat::R8_SINT: - return ",r8i"; - case ImageFormat::R16_UINT: - return ",r16ui"; - case ImageFormat::R16_SINT: - return ",r16i"; - case ImageFormat::R32_UINT: - return ",r32ui"; - case ImageFormat::R32G32_UINT: - return ",rg32ui"; - case ImageFormat::R32G32B32A32_UINT: - return ",rgba32ui"; - default: - throw NotImplementedException("Image format: {}", format); - } -} -*/ - -std::string_view ImageAccessQualifier(bool is_written, bool is_read) { - if (is_written && is_read) { - return "access::read, access::write"; - } - if (is_written) { - return "access::write"; - } - if (is_read) { - return "access::read"; - } - return ""; -} - -void DefineInputs(EmitContext& ctx, std::string& header, Bindings& bindings) { - bool added{false}; - - // Constant buffers - for (const auto& desc : ctx.info.constant_buffer_descriptors) { - const u32 cbuf_used_size{ - Common::DivCeil(ctx.info.constant_buffer_used_sizes[desc.index], 16U)}; - const u32 cbuf_binding_size{ctx.info.uses_global_memory ? 0x1000U : cbuf_used_size}; - if (added) - header += ","; - header += fmt::format("constant float4& cbuf{}[{}] [[buffer({})]]", desc.index, - cbuf_binding_size, bindings.uniform_buffer); - bindings.uniform_buffer += desc.count; - added = true; - } - - // Constant buffer indirect - // TODO - - // Storage space buffers - u32 index{}; - for (const auto& desc : ctx.info.storage_buffers_descriptors) { - if (added) - header += ","; - header += - fmt::format("device uint& ssbo{}[] [[buffer({})]]", index, bindings.storage_buffer); - bindings.storage_buffer += desc.count; - index += desc.count; - added = true; - } - - // Images - // TODO - /* - image_buffers.reserve(info.image_buffer_descriptors.size()); - for (const auto& desc : info.image_buffer_descriptors) { - image_buffers.push_back({bindings.image, desc.count}); - const auto format{ImageFormatString(desc.format)}; - const auto qualifier{ImageAccessQualifier(desc.is_written, desc.is_read)}; - const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""}; - header += fmt::format("layout(binding={}{}) uniform {}uimageBuffer img{}{};", - bindings.image, format, qualifier, bindings.image, array_decorator); - bindings.image += desc.count; - } - */ - ctx.images.reserve(ctx.info.image_descriptors.size()); - for (const auto& desc : ctx.info.image_descriptors) { - ctx.images.push_back({bindings.image, desc.count}); - // TODO: do we need format? - // const auto format{ImageFormatString(desc.format)}; - const auto image_type{ImageType(desc.type)}; - const auto qualifier{ImageAccessQualifier(desc.is_written, desc.is_read)}; - const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""}; - if (added) - header += ","; - header += fmt::format("{}<{}> img{}{} [[texture({})]]", qualifier, image_type, - bindings.image, array_decorator, bindings.image); - bindings.image += desc.count; - added = true; - } - - // Textures - // TODO - /* - texture_buffers.reserve(info.texture_buffer_descriptors.size()); - for (const auto& desc : info.texture_buffer_descriptors) { - texture_buffers.push_back({bindings.texture, desc.count}); - const auto sampler_type{ColorSamplerType(TextureType::Buffer)}; - const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""}; - header += fmt::format("layout(binding={}) uniform {} tex{}{};", bindings.texture, - sampler_type, bindings.texture, array_decorator); - bindings.texture += desc.count; - } - */ - ctx.textures.reserve(ctx.info.texture_descriptors.size()); - for (const auto& desc : ctx.info.texture_descriptors) { - ctx.textures.push_back({bindings.texture, desc.count}); - const auto texture_type{desc.is_depth ? DepthSamplerType(desc.type) - : ColorSamplerType(desc.type, desc.is_multisample)}; - const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""}; - if (added) - header += ","; - header += fmt::format("{} tex{}{} [[texture({})]]", texture_type, bindings.texture, - array_decorator, bindings.texture); - header += fmt::format(",sampler samp{}{} [[sampler({})]]", bindings.texture, - array_decorator, bindings.texture); - bindings.texture += desc.count; - added = true; - } -} } // Anonymous namespace std::string EmitMSL(const Profile& profile, const RuntimeInfo& runtime_info, IR::Program& program, Bindings& bindings) { EmitContext ctx{program, bindings, profile, runtime_info}; std::string inputs; - DefineInputs(ctx, inputs, bindings); Precolor(program); EmitCode(ctx, program); ctx.header.insert(0, "#include \nusing namespace metal;\n"); @@ -426,8 +225,9 @@ std::string EmitMSL(const Profile& profile, const RuntimeInfo& runtime_info, IR: const auto smem_size{needs_clamp ? max_size : requested_size}; ctx.header += fmt::format("shared uint smem[{}];", Common::DivCeil(smem_size, 4U)); } - ctx.header += "void main_("; - ctx.header += inputs; + ctx.header += ctx.stage_name; + ctx.header += " __Output main_("; + ctx.header += ctx.input_str; ctx.header += "){\n"; if (program.local_memory_size > 0) { ctx.header += fmt::format("uint lmem[{}];", Common::DivCeil(program.local_memory_size, 4U)); diff --git a/src/shader_recompiler/backend/msl/msl_emit_context.cpp b/src/shader_recompiler/backend/msl/msl_emit_context.cpp index 9dc25df7a9..203b87fdfd 100644 --- a/src/shader_recompiler/backend/msl/msl_emit_context.cpp +++ b/src/shader_recompiler/backend/msl/msl_emit_context.cpp @@ -107,6 +107,115 @@ std::string_view OutputPrimitive(OutputTopology topology) { } throw InvalidArgument("Invalid output topology {}", topology); } + +// TODO +std::string_view DepthSamplerType(TextureType type) { + switch (type) { + case TextureType::Color1D: + return "sampler1DShadow"; + case TextureType::ColorArray1D: + return "sampler1DArrayShadow"; + case TextureType::Color2D: + return "sampler2DShadow"; + case TextureType::ColorArray2D: + return "sampler2DArrayShadow"; + case TextureType::ColorCube: + return "samplerCubeShadow"; + case TextureType::ColorArrayCube: + return "samplerCubeArrayShadow"; + default: + throw NotImplementedException("Texture type: {}", type); + } +} + +// TODO: emit sampler as well +// TODO: handle multisample +// TODO: handle texture buffer +std::string_view ColorSamplerType(TextureType type, bool is_multisample = false) { + if (is_multisample) { + ASSERT(type == TextureType::Color2D || type == TextureType::ColorArray2D); + } + switch (type) { + case TextureType::Color1D: + return "texture1d"; + case TextureType::ColorArray1D: + return "texture1d_array"; + case TextureType::Color2D: + case TextureType::Color2DRect: + return "texture2d"; + case TextureType::ColorArray2D: + return "texture2d_array"; + case TextureType::Color3D: + return "texture3d"; + case TextureType::ColorCube: + return "texturecube"; + case TextureType::ColorArrayCube: + return "texturecube_array"; + default: + throw NotImplementedException("Texture type: {}", type); + } +} + +// TODO: handle texture buffer +std::string_view ImageType(TextureType type) { + switch (type) { + case TextureType::Color1D: + return "texture1d"; + case TextureType::ColorArray1D: + return "texture1d_array"; + case TextureType::Color2D: + return "texture2d"; + case TextureType::ColorArray2D: + return "texture2d_array"; + case TextureType::Color3D: + return "texture3d"; + case TextureType::ColorCube: + return "texturecube"; + case TextureType::ColorArrayCube: + return "texturecube_array"; + default: + throw NotImplementedException("Image type: {}", type); + } +} + +// TODO: is this needed? +/* +std::string_view ImageFormatString(ImageFormat format) { + switch (format) { + case ImageFormat::Typeless: + return ""; + case ImageFormat::R8_UINT: + return ",r8ui"; + case ImageFormat::R8_SINT: + return ",r8i"; + case ImageFormat::R16_UINT: + return ",r16ui"; + case ImageFormat::R16_SINT: + return ",r16i"; + case ImageFormat::R32_UINT: + return ",r32ui"; + case ImageFormat::R32G32_UINT: + return ",rg32ui"; + case ImageFormat::R32G32B32A32_UINT: + return ",rgba32ui"; + default: + throw NotImplementedException("Image format: {}", format); + } +} +*/ + +std::string_view ImageAccessQualifier(bool is_written, bool is_read) { + if (is_written && is_read) { + return "access::read, access::write"; + } + if (is_written) { + return "access::write"; + } + if (is_read) { + return "access::read"; + } + return ""; +} } // Anonymous namespace EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, @@ -120,20 +229,20 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile switch (program.stage) { case Stage::VertexA: case Stage::VertexB: - stage_name = "vs"; + stage_name = "vertex"; break; case Stage::TessellationControl: - stage_name = "tcs"; + stage_name = "kernel"; header += fmt::format("layout(vertices={})out;", program.invocations); break; case Stage::TessellationEval: - stage_name = "tes"; + stage_name = "vertex"; header += fmt::format("layout({},{},{})in;", GetTessMode(runtime_info.tess_primitive), GetTessSpacing(runtime_info.tess_spacing), runtime_info.tess_clockwise ? "cw" : "ccw"); break; case Stage::Geometry: - stage_name = "gs"; + stage_name = "vertex"; header += fmt::format("layout({})in;", InputPrimitive(runtime_info.input_topology)); if (uses_geometry_passthrough) { header += "layout(passthrough)in gl_PerVertex{vec4 gl_Position;};"; @@ -147,13 +256,13 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile OutputPrimitive(program.output_topology), program.output_vertices); break; case Stage::Fragment: - stage_name = "fs"; + stage_name = "fragment"; if (runtime_info.force_early_z) { header += "layout(early_fragment_tests)in;"; } break; case Stage::Compute: - stage_name = "cs"; + stage_name = "kernel"; const u32 local_x{std::max(program.workgroup_size[0], 1u)}; const u32 local_y{std::max(program.workgroup_size[1], 1u)}; const u32 local_z{std::max(program.workgroup_size[2], 1u)}; @@ -190,21 +299,122 @@ EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile header += fmt::format("layout(location={})out vec4 frag_color{};", index, index); } } + header += "struct __Output {\n"; for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { if (info.stores.Generic(index)) { DefineGenericOutput(index, program.invocations); } } + header += "};\n"; + bool added = DefineInputs(bindings); if (info.uses_rescaling_uniform) { - header += "layout(location=0) uniform vec4 scaling;"; + if (added) + input_str += ","; + input_str += "constant float4& scaling"; + added = true; } if (info.uses_render_area) { - header += "layout(location=1) uniform vec4 render_area;"; + if (added) + input_str += ","; + input_str += "constant float4& render_area"; + added = true; } DefineHelperFunctions(); DefineConstants(); } +bool EmitContext::DefineInputs(Bindings& bindings) { + bool added{false}; + + // Constant buffers + for (const auto& desc : info.constant_buffer_descriptors) { + const u32 cbuf_used_size{Common::DivCeil(info.constant_buffer_used_sizes[desc.index], 16U)}; + const u32 cbuf_binding_size{info.uses_global_memory ? 0x1000U : cbuf_used_size}; + if (added) + input_str += ","; + input_str += fmt::format("constant float4& cbuf{}[{}] [[buffer({})]]", desc.index, + cbuf_binding_size, bindings.uniform_buffer); + bindings.uniform_buffer += desc.count; + added = true; + } + + // Constant buffer indirect + // TODO + + // Storage space buffers + u32 index{}; + for (const auto& desc : info.storage_buffers_descriptors) { + if (added) + input_str += ","; + input_str += + fmt::format("device uint& ssbo{}[] [[buffer({})]]", index, bindings.storage_buffer); + bindings.storage_buffer += desc.count; + index += desc.count; + added = true; + } + + // Images + // TODO + /* + image_buffers.reserve(info.image_buffer_descriptors.size()); + for (const auto& desc : info.image_buffer_descriptors) { + image_buffers.push_back({bindings.image, desc.count}); + const auto format{ImageFormatString(desc.format)}; + const auto qualifier{ImageAccessQualifier(desc.is_written, desc.is_read)}; + const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""}; + input_str += fmt::format("layout(binding={}{}) uniform {}uimageBuffer img{}{};", + bindings.image, format, qualifier, bindings.image, array_decorator); + bindings.image += desc.count; + } + */ + images.reserve(info.image_descriptors.size()); + for (const auto& desc : info.image_descriptors) { + images.push_back({bindings.image, desc.count}); + // TODO: do we need format? + // const auto format{ImageFormatString(desc.format)}; + const auto image_type{ImageType(desc.type)}; + const auto qualifier{ImageAccessQualifier(desc.is_written, desc.is_read)}; + const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""}; + if (added) + input_str += ","; + input_str += fmt::format("{}<{}> img{}{} [[texture({})]]", qualifier, image_type, + bindings.image, array_decorator, bindings.image); + bindings.image += desc.count; + added = true; + } + + // Textures + // TODO + /* + texture_buffers.reserve(info.texture_buffer_descriptors.size()); + for (const auto& desc : info.texture_buffer_descriptors) { + texture_buffers.push_back({bindings.texture, desc.count}); + const auto sampler_type{ColorSamplerType(TextureType::Buffer)}; + const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""}; + input_str += fmt::format("layout(binding={}) uniform {} tex{}{};", bindings.texture, + sampler_type, bindings.texture, array_decorator); + bindings.texture += desc.count; + } + */ + textures.reserve(info.texture_descriptors.size()); + for (const auto& desc : info.texture_descriptors) { + textures.push_back({bindings.texture, desc.count}); + const auto texture_type{desc.is_depth ? DepthSamplerType(desc.type) + : ColorSamplerType(desc.type, desc.is_multisample)}; + const auto array_decorator{desc.count > 1 ? fmt::format("[{}]", desc.count) : ""}; + if (added) + input_str += ","; + input_str += fmt::format("{} tex{}{} [[texture({})]]", texture_type, bindings.texture, + array_decorator, bindings.texture); + input_str += fmt::format(",sampler samp{}{} [[sampler({})]]", bindings.texture, + array_decorator, bindings.texture); + bindings.texture += desc.count; + added = true; + } + + return added; +} + // TODO void EmitContext::DefineGenericOutput(size_t index, u32 invocations) { static constexpr std::string_view swizzle{"xyzw"}; diff --git a/src/shader_recompiler/backend/msl/msl_emit_context.h b/src/shader_recompiler/backend/msl/msl_emit_context.h index 233b4bd606..830e97ce9a 100644 --- a/src/shader_recompiler/backend/msl/msl_emit_context.h +++ b/src/shader_recompiler/backend/msl/msl_emit_context.h @@ -136,6 +136,7 @@ public: } std::string header; + std::string input_str; std::string code; VarAlloc var_alloc; const Info& info; @@ -158,6 +159,8 @@ public: bool uses_geometry_passthrough{}; private: + // TODO: break down into smaller functions + bool DefineInputs(Bindings& bindings); void DefineGenericOutput(size_t index, u32 invocations); void DefineHelperFunctions(); void DefineConstants();