From 5286a7bc4ccf8da0827b0352f40dbce651b57d09 Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Sun, 5 Dec 2021 16:33:44 -0500 Subject: shader_recompiler: Rename backend emit_context files --- src/shader_recompiler/CMakeLists.txt | 12 +- .../backend/glasm/emit_context.cpp | 156 -- src/shader_recompiler/backend/glasm/emit_context.h | 80 - .../backend/glasm/glasm_emit_context.cpp | 156 ++ .../backend/glasm/glasm_emit_context.h | 80 + .../backend/glsl/emit_context.cpp | 718 --------- src/shader_recompiler/backend/glsl/emit_context.h | 174 --- .../backend/glsl/glsl_emit_context.cpp | 718 +++++++++ .../backend/glsl/glsl_emit_context.h | 174 +++ .../backend/spirv/emit_context.cpp | 1585 -------------------- src/shader_recompiler/backend/spirv/emit_context.h | 335 ----- .../backend/spirv/spirv_emit_context.cpp | 1585 ++++++++++++++++++++ .../backend/spirv/spirv_emit_context.h | 335 +++++ 13 files changed, 3054 insertions(+), 3054 deletions(-) delete mode 100644 src/shader_recompiler/backend/glasm/emit_context.cpp delete mode 100644 src/shader_recompiler/backend/glasm/emit_context.h create mode 100644 src/shader_recompiler/backend/glasm/glasm_emit_context.cpp create mode 100644 src/shader_recompiler/backend/glasm/glasm_emit_context.h delete mode 100644 src/shader_recompiler/backend/glsl/emit_context.cpp delete mode 100644 src/shader_recompiler/backend/glsl/emit_context.h create mode 100644 src/shader_recompiler/backend/glsl/glsl_emit_context.cpp create mode 100644 src/shader_recompiler/backend/glsl/glsl_emit_context.h delete mode 100644 src/shader_recompiler/backend/spirv/emit_context.cpp delete mode 100644 src/shader_recompiler/backend/spirv/emit_context.h create mode 100644 src/shader_recompiler/backend/spirv/spirv_emit_context.cpp create mode 100644 src/shader_recompiler/backend/spirv/spirv_emit_context.h (limited to 'src') diff --git a/src/shader_recompiler/CMakeLists.txt b/src/shader_recompiler/CMakeLists.txt index bc3df80c8..4c76ce1ea 100644 --- a/src/shader_recompiler/CMakeLists.txt +++ b/src/shader_recompiler/CMakeLists.txt @@ -1,7 +1,5 @@ add_library(shader_recompiler STATIC backend/bindings.h - backend/glasm/emit_context.cpp - backend/glasm/emit_context.h backend/glasm/emit_glasm.cpp backend/glasm/emit_glasm.h backend/glasm/emit_glasm_barriers.cpp @@ -22,10 +20,10 @@ add_library(shader_recompiler STATIC backend/glasm/emit_glasm_special.cpp backend/glasm/emit_glasm_undefined.cpp backend/glasm/emit_glasm_warp.cpp + backend/glasm/glasm_emit_context.cpp + backend/glasm/glasm_emit_context.h backend/glasm/reg_alloc.cpp backend/glasm/reg_alloc.h - backend/glsl/emit_context.cpp - backend/glsl/emit_context.h backend/glsl/emit_glsl.cpp backend/glsl/emit_glsl.h backend/glsl/emit_glsl_atomic.cpp @@ -47,10 +45,10 @@ add_library(shader_recompiler STATIC backend/glsl/emit_glsl_special.cpp backend/glsl/emit_glsl_undefined.cpp backend/glsl/emit_glsl_warp.cpp + backend/glsl/glsl_emit_context.cpp + backend/glsl/glsl_emit_context.h backend/glsl/var_alloc.cpp backend/glsl/var_alloc.h - backend/spirv/emit_context.cpp - backend/spirv/emit_context.h backend/spirv/emit_spirv.cpp backend/spirv/emit_spirv.h backend/spirv/emit_spirv_atomic.cpp @@ -72,6 +70,8 @@ add_library(shader_recompiler STATIC backend/spirv/emit_spirv_special.cpp backend/spirv/emit_spirv_undefined.cpp backend/spirv/emit_spirv_warp.cpp + backend/spirv/spirv_emit_context.cpp + backend/spirv/spirv_emit_context.h environment.h exception.h frontend/ir/abstract_syntax_list.h diff --git a/src/shader_recompiler/backend/glasm/emit_context.cpp b/src/shader_recompiler/backend/glasm/emit_context.cpp deleted file mode 100644 index 8fd459dfe..000000000 --- a/src/shader_recompiler/backend/glasm/emit_context.cpp +++ /dev/null @@ -1,156 +0,0 @@ -// Copyright 2021 yuzu Emulator Project -// Licensed under GPLv2 or any later version -// Refer to the license.txt file included. - -#include - -#include "shader_recompiler/backend/bindings.h" -#include "shader_recompiler/backend/glasm/emit_context.h" -#include "shader_recompiler/backend/glasm/emit_glasm.h" -#include "shader_recompiler/frontend/ir/program.h" -#include "shader_recompiler/profile.h" -#include "shader_recompiler/runtime_info.h" - -namespace Shader::Backend::GLASM { -namespace { -std::string_view InterpDecorator(Interpolation interp) { - switch (interp) { - case Interpolation::Smooth: - return ""; - case Interpolation::Flat: - return "FLAT "; - case Interpolation::NoPerspective: - return "NOPERSPECTIVE "; - } - throw InvalidArgument("Invalid interpolation {}", interp); -} - -bool IsInputArray(Stage stage) { - return stage == Stage::Geometry || stage == Stage::TessellationControl || - stage == Stage::TessellationEval; -} -} // Anonymous namespace - -EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, - const RuntimeInfo& runtime_info_) - : info{program.info}, profile{profile_}, runtime_info{runtime_info_} { - // FIXME: Temporary partial implementation - u32 cbuf_index{}; - for (const auto& desc : info.constant_buffer_descriptors) { - if (desc.count != 1) { - throw NotImplementedException("Constant buffer descriptor array"); - } - Add("CBUFFER c{}[]={{program.buffer[{}]}};", desc.index, cbuf_index); - ++cbuf_index; - } - u32 ssbo_index{}; - for (const auto& desc : info.storage_buffers_descriptors) { - if (desc.count != 1) { - throw NotImplementedException("Storage buffer descriptor array"); - } - if (runtime_info.glasm_use_storage_buffers) { - Add("STORAGE ssbo{}[]={{program.storage[{}]}};", ssbo_index, bindings.storage_buffer); - ++bindings.storage_buffer; - ++ssbo_index; - } - } - if (!runtime_info.glasm_use_storage_buffers) { - if (const size_t num = info.storage_buffers_descriptors.size(); num > 0) { - const size_t index{num + PROGRAM_LOCAL_PARAMETER_STORAGE_BUFFER_BASE}; - Add("PARAM c[{}]={{program.local[0..{}]}};", index, index - 1); - } - } - stage = program.stage; - switch (program.stage) { - case Stage::VertexA: - case Stage::VertexB: - stage_name = "vertex"; - attrib_name = "vertex"; - break; - case Stage::TessellationControl: - case Stage::TessellationEval: - stage_name = "primitive"; - attrib_name = "primitive"; - break; - case Stage::Geometry: - stage_name = "primitive"; - attrib_name = "vertex"; - break; - case Stage::Fragment: - stage_name = "fragment"; - attrib_name = "fragment"; - break; - case Stage::Compute: - stage_name = "invocation"; - break; - } - const std::string_view attr_stage{stage == Stage::Fragment ? "fragment" : "vertex"}; - const VaryingState loads{info.loads.mask | info.passthrough.mask}; - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - if (loads.Generic(index)) { - Add("{}ATTRIB in_attr{}[]={{{}.attrib[{}..{}]}};", - InterpDecorator(info.interpolation[index]), index, attr_stage, index, index); - } - } - if (IsInputArray(stage) && loads.AnyComponent(IR::Attribute::PositionX)) { - Add("ATTRIB vertex_position=vertex.position;"); - } - if (info.uses_invocation_id) { - Add("ATTRIB primitive_invocation=primitive.invocation;"); - } - if (info.stores_tess_level_outer) { - Add("OUTPUT result_patch_tessouter[]={{result.patch.tessouter[0..3]}};"); - } - if (info.stores_tess_level_inner) { - Add("OUTPUT result_patch_tessinner[]={{result.patch.tessinner[0..1]}};"); - } - if (info.stores.ClipDistances()) { - Add("OUTPUT result_clip[]={{result.clip[0..7]}};"); - } - for (size_t index = 0; index < info.uses_patches.size(); ++index) { - if (!info.uses_patches[index]) { - continue; - } - if (stage == Stage::TessellationControl) { - Add("OUTPUT result_patch_attrib{}[]={{result.patch.attrib[{}..{}]}};" - "ATTRIB primitive_out_patch_attrib{}[]={{primitive.out.patch.attrib[{}..{}]}};", - index, index, index, index, index, index); - } else { - Add("ATTRIB primitive_patch_attrib{}[]={{primitive.patch.attrib[{}..{}]}};", index, - index, index); - } - } - if (stage == Stage::Fragment) { - Add("OUTPUT frag_color0=result.color;"); - for (size_t index = 1; index < info.stores_frag_color.size(); ++index) { - Add("OUTPUT frag_color{}=result.color[{}];", index, index); - } - } - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - if (info.stores.Generic(index)) { - Add("OUTPUT out_attr{}[]={{result.attrib[{}..{}]}};", index, index, index); - } - } - image_buffer_bindings.reserve(info.image_buffer_descriptors.size()); - for (const auto& desc : info.image_buffer_descriptors) { - image_buffer_bindings.push_back(bindings.image); - bindings.image += desc.count; - } - image_bindings.reserve(info.image_descriptors.size()); - for (const auto& desc : info.image_descriptors) { - image_bindings.push_back(bindings.image); - bindings.image += desc.count; - } - texture_buffer_bindings.reserve(info.texture_buffer_descriptors.size()); - for (const auto& desc : info.texture_buffer_descriptors) { - texture_buffer_bindings.push_back(bindings.texture); - bindings.texture += desc.count; - } - texture_bindings.reserve(info.texture_descriptors.size()); - for (const auto& desc : info.texture_descriptors) { - texture_bindings.push_back(bindings.texture); - bindings.texture += desc.count; - } -} - -} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/emit_context.h b/src/shader_recompiler/backend/glasm/emit_context.h deleted file mode 100644 index 8433e5c00..000000000 --- a/src/shader_recompiler/backend/glasm/emit_context.h +++ /dev/null @@ -1,80 +0,0 @@ -// Copyright 2021 yuzu Emulator Project -// Licensed under GPLv2 or any later version -// Refer to the license.txt file included. - -#pragma once - -#include -#include -#include - -#include - -#include "shader_recompiler/backend/glasm/reg_alloc.h" -#include "shader_recompiler/stage.h" - -namespace Shader { -struct Info; -struct Profile; -struct RuntimeInfo; -} // namespace Shader - -namespace Shader::Backend { -struct Bindings; -} - -namespace Shader::IR { -class Inst; -struct Program; -} // namespace Shader::IR - -namespace Shader::Backend::GLASM { - -class EmitContext { -public: - explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, - const RuntimeInfo& runtime_info_); - - template - void Add(const char* format_str, IR::Inst& inst, Args&&... args) { - code += fmt::format(fmt::runtime(format_str), reg_alloc.Define(inst), - std::forward(args)...); - // TODO: Remove this - code += '\n'; - } - - template - void LongAdd(const char* format_str, IR::Inst& inst, Args&&... args) { - code += fmt::format(fmt::runtime(format_str), reg_alloc.LongDefine(inst), - std::forward(args)...); - // TODO: Remove this - code += '\n'; - } - - template - void Add(const char* format_str, Args&&... args) { - code += fmt::format(fmt::runtime(format_str), std::forward(args)...); - // TODO: Remove this - code += '\n'; - } - - std::string code; - RegAlloc reg_alloc{}; - const Info& info; - const Profile& profile; - const RuntimeInfo& runtime_info; - - std::vector texture_buffer_bindings; - std::vector image_buffer_bindings; - std::vector texture_bindings; - std::vector image_bindings; - - Stage stage{}; - std::string_view stage_name = "invalid"; - std::string_view attrib_name = "invalid"; - - u32 num_safety_loop_vars{}; - bool uses_y_direction{}; -}; - -} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp b/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp new file mode 100644 index 000000000..8fd459dfe --- /dev/null +++ b/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp @@ -0,0 +1,156 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include + +#include "shader_recompiler/backend/bindings.h" +#include "shader_recompiler/backend/glasm/emit_context.h" +#include "shader_recompiler/backend/glasm/emit_glasm.h" +#include "shader_recompiler/frontend/ir/program.h" +#include "shader_recompiler/profile.h" +#include "shader_recompiler/runtime_info.h" + +namespace Shader::Backend::GLASM { +namespace { +std::string_view InterpDecorator(Interpolation interp) { + switch (interp) { + case Interpolation::Smooth: + return ""; + case Interpolation::Flat: + return "FLAT "; + case Interpolation::NoPerspective: + return "NOPERSPECTIVE "; + } + throw InvalidArgument("Invalid interpolation {}", interp); +} + +bool IsInputArray(Stage stage) { + return stage == Stage::Geometry || stage == Stage::TessellationControl || + stage == Stage::TessellationEval; +} +} // Anonymous namespace + +EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, + const RuntimeInfo& runtime_info_) + : info{program.info}, profile{profile_}, runtime_info{runtime_info_} { + // FIXME: Temporary partial implementation + u32 cbuf_index{}; + for (const auto& desc : info.constant_buffer_descriptors) { + if (desc.count != 1) { + throw NotImplementedException("Constant buffer descriptor array"); + } + Add("CBUFFER c{}[]={{program.buffer[{}]}};", desc.index, cbuf_index); + ++cbuf_index; + } + u32 ssbo_index{}; + for (const auto& desc : info.storage_buffers_descriptors) { + if (desc.count != 1) { + throw NotImplementedException("Storage buffer descriptor array"); + } + if (runtime_info.glasm_use_storage_buffers) { + Add("STORAGE ssbo{}[]={{program.storage[{}]}};", ssbo_index, bindings.storage_buffer); + ++bindings.storage_buffer; + ++ssbo_index; + } + } + if (!runtime_info.glasm_use_storage_buffers) { + if (const size_t num = info.storage_buffers_descriptors.size(); num > 0) { + const size_t index{num + PROGRAM_LOCAL_PARAMETER_STORAGE_BUFFER_BASE}; + Add("PARAM c[{}]={{program.local[0..{}]}};", index, index - 1); + } + } + stage = program.stage; + switch (program.stage) { + case Stage::VertexA: + case Stage::VertexB: + stage_name = "vertex"; + attrib_name = "vertex"; + break; + case Stage::TessellationControl: + case Stage::TessellationEval: + stage_name = "primitive"; + attrib_name = "primitive"; + break; + case Stage::Geometry: + stage_name = "primitive"; + attrib_name = "vertex"; + break; + case Stage::Fragment: + stage_name = "fragment"; + attrib_name = "fragment"; + break; + case Stage::Compute: + stage_name = "invocation"; + break; + } + const std::string_view attr_stage{stage == Stage::Fragment ? "fragment" : "vertex"}; + const VaryingState loads{info.loads.mask | info.passthrough.mask}; + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + if (loads.Generic(index)) { + Add("{}ATTRIB in_attr{}[]={{{}.attrib[{}..{}]}};", + InterpDecorator(info.interpolation[index]), index, attr_stage, index, index); + } + } + if (IsInputArray(stage) && loads.AnyComponent(IR::Attribute::PositionX)) { + Add("ATTRIB vertex_position=vertex.position;"); + } + if (info.uses_invocation_id) { + Add("ATTRIB primitive_invocation=primitive.invocation;"); + } + if (info.stores_tess_level_outer) { + Add("OUTPUT result_patch_tessouter[]={{result.patch.tessouter[0..3]}};"); + } + if (info.stores_tess_level_inner) { + Add("OUTPUT result_patch_tessinner[]={{result.patch.tessinner[0..1]}};"); + } + if (info.stores.ClipDistances()) { + Add("OUTPUT result_clip[]={{result.clip[0..7]}};"); + } + for (size_t index = 0; index < info.uses_patches.size(); ++index) { + if (!info.uses_patches[index]) { + continue; + } + if (stage == Stage::TessellationControl) { + Add("OUTPUT result_patch_attrib{}[]={{result.patch.attrib[{}..{}]}};" + "ATTRIB primitive_out_patch_attrib{}[]={{primitive.out.patch.attrib[{}..{}]}};", + index, index, index, index, index, index); + } else { + Add("ATTRIB primitive_patch_attrib{}[]={{primitive.patch.attrib[{}..{}]}};", index, + index, index); + } + } + if (stage == Stage::Fragment) { + Add("OUTPUT frag_color0=result.color;"); + for (size_t index = 1; index < info.stores_frag_color.size(); ++index) { + Add("OUTPUT frag_color{}=result.color[{}];", index, index); + } + } + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + if (info.stores.Generic(index)) { + Add("OUTPUT out_attr{}[]={{result.attrib[{}..{}]}};", index, index, index); + } + } + image_buffer_bindings.reserve(info.image_buffer_descriptors.size()); + for (const auto& desc : info.image_buffer_descriptors) { + image_buffer_bindings.push_back(bindings.image); + bindings.image += desc.count; + } + image_bindings.reserve(info.image_descriptors.size()); + for (const auto& desc : info.image_descriptors) { + image_bindings.push_back(bindings.image); + bindings.image += desc.count; + } + texture_buffer_bindings.reserve(info.texture_buffer_descriptors.size()); + for (const auto& desc : info.texture_buffer_descriptors) { + texture_buffer_bindings.push_back(bindings.texture); + bindings.texture += desc.count; + } + texture_bindings.reserve(info.texture_descriptors.size()); + for (const auto& desc : info.texture_descriptors) { + texture_bindings.push_back(bindings.texture); + bindings.texture += desc.count; + } +} + +} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/glasm_emit_context.h b/src/shader_recompiler/backend/glasm/glasm_emit_context.h new file mode 100644 index 000000000..8433e5c00 --- /dev/null +++ b/src/shader_recompiler/backend/glasm/glasm_emit_context.h @@ -0,0 +1,80 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#pragma once + +#include +#include +#include + +#include + +#include "shader_recompiler/backend/glasm/reg_alloc.h" +#include "shader_recompiler/stage.h" + +namespace Shader { +struct Info; +struct Profile; +struct RuntimeInfo; +} // namespace Shader + +namespace Shader::Backend { +struct Bindings; +} + +namespace Shader::IR { +class Inst; +struct Program; +} // namespace Shader::IR + +namespace Shader::Backend::GLASM { + +class EmitContext { +public: + explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, + const RuntimeInfo& runtime_info_); + + template + void Add(const char* format_str, IR::Inst& inst, Args&&... args) { + code += fmt::format(fmt::runtime(format_str), reg_alloc.Define(inst), + std::forward(args)...); + // TODO: Remove this + code += '\n'; + } + + template + void LongAdd(const char* format_str, IR::Inst& inst, Args&&... args) { + code += fmt::format(fmt::runtime(format_str), reg_alloc.LongDefine(inst), + std::forward(args)...); + // TODO: Remove this + code += '\n'; + } + + template + void Add(const char* format_str, Args&&... args) { + code += fmt::format(fmt::runtime(format_str), std::forward(args)...); + // TODO: Remove this + code += '\n'; + } + + std::string code; + RegAlloc reg_alloc{}; + const Info& info; + const Profile& profile; + const RuntimeInfo& runtime_info; + + std::vector texture_buffer_bindings; + std::vector image_buffer_bindings; + std::vector texture_bindings; + std::vector image_bindings; + + Stage stage{}; + std::string_view stage_name = "invalid"; + std::string_view attrib_name = "invalid"; + + u32 num_safety_loop_vars{}; + bool uses_y_direction{}; +}; + +} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glsl/emit_context.cpp b/src/shader_recompiler/backend/glsl/emit_context.cpp deleted file mode 100644 index 97bd59302..000000000 --- a/src/shader_recompiler/backend/glsl/emit_context.cpp +++ /dev/null @@ -1,718 +0,0 @@ -// Copyright 2021 yuzu Emulator Project -// Licensed under GPLv2 or any later version -// Refer to the license.txt file included. - -#include "shader_recompiler/backend/bindings.h" -#include "shader_recompiler/backend/glsl/emit_context.h" -#include "shader_recompiler/frontend/ir/program.h" -#include "shader_recompiler/profile.h" -#include "shader_recompiler/runtime_info.h" - -namespace Shader::Backend::GLSL { -namespace { -u32 CbufIndex(size_t offset) { - return (offset / 4) % 4; -} - -char Swizzle(size_t offset) { - return "xyzw"[CbufIndex(offset)]; -} - -std::string_view InterpDecorator(Interpolation interp) { - switch (interp) { - case Interpolation::Smooth: - return ""; - case Interpolation::Flat: - return "flat "; - case Interpolation::NoPerspective: - return "noperspective "; - } - throw InvalidArgument("Invalid interpolation {}", interp); -} - -std::string_view InputArrayDecorator(Stage stage) { - switch (stage) { - case Stage::Geometry: - case Stage::TessellationControl: - case Stage::TessellationEval: - return "[]"; - default: - return ""; - } -} - -bool StoresPerVertexAttributes(Stage stage) { - switch (stage) { - case Stage::VertexA: - case Stage::VertexB: - case Stage::Geometry: - case Stage::TessellationEval: - return true; - default: - return false; - } -} - -std::string OutputDecorator(Stage stage, u32 size) { - switch (stage) { - case Stage::TessellationControl: - return fmt::format("[{}]", size); - default: - return ""; - } -} - -std::string_view SamplerType(TextureType type, bool is_depth) { - if (is_depth) { - 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); - } - } - switch (type) { - case TextureType::Color1D: - return "sampler1D"; - case TextureType::ColorArray1D: - return "sampler1DArray"; - case TextureType::Color2D: - return "sampler2D"; - case TextureType::ColorArray2D: - return "sampler2DArray"; - case TextureType::Color3D: - return "sampler3D"; - case TextureType::ColorCube: - return "samplerCube"; - case TextureType::ColorArrayCube: - return "samplerCubeArray"; - case TextureType::Buffer: - return "samplerBuffer"; - default: - throw NotImplementedException("Texture type: {}", type); - } -} - -std::string_view ImageType(TextureType type) { - switch (type) { - case TextureType::Color1D: - return "uimage1D"; - case TextureType::ColorArray1D: - return "uimage1DArray"; - case TextureType::Color2D: - return "uimage2D"; - case TextureType::ColorArray2D: - return "uimage2DArray"; - case TextureType::Color3D: - return "uimage3D"; - case TextureType::ColorCube: - return "uimageCube"; - case TextureType::ColorArrayCube: - return "uimageCubeArray"; - case TextureType::Buffer: - return "uimageBuffer"; - default: - throw NotImplementedException("Image type: {}", type); - } -} - -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 "writeonly "; - } - if (is_read && !is_written) { - return "readonly "; - } - return ""; -} - -std::string_view GetTessMode(TessPrimitive primitive) { - switch (primitive) { - case TessPrimitive::Triangles: - return "triangles"; - case TessPrimitive::Quads: - return "quads"; - case TessPrimitive::Isolines: - return "isolines"; - } - throw InvalidArgument("Invalid tessellation primitive {}", primitive); -} - -std::string_view GetTessSpacing(TessSpacing spacing) { - switch (spacing) { - case TessSpacing::Equal: - return "equal_spacing"; - case TessSpacing::FractionalOdd: - return "fractional_odd_spacing"; - case TessSpacing::FractionalEven: - return "fractional_even_spacing"; - } - throw InvalidArgument("Invalid tessellation spacing {}", spacing); -} - -std::string_view InputPrimitive(InputTopology topology) { - switch (topology) { - case InputTopology::Points: - return "points"; - case InputTopology::Lines: - return "lines"; - case InputTopology::LinesAdjacency: - return "lines_adjacency"; - case InputTopology::Triangles: - return "triangles"; - case InputTopology::TrianglesAdjacency: - return "triangles_adjacency"; - } - throw InvalidArgument("Invalid input topology {}", topology); -} - -std::string_view OutputPrimitive(OutputTopology topology) { - switch (topology) { - case OutputTopology::PointList: - return "points"; - case OutputTopology::LineStrip: - return "line_strip"; - case OutputTopology::TriangleStrip: - return "triangle_strip"; - } - throw InvalidArgument("Invalid output topology {}", topology); -} - -void SetupLegacyOutPerVertex(EmitContext& ctx, std::string& header) { - if (!ctx.info.stores.Legacy()) { - return; - } - if (ctx.info.stores.FixedFunctionTexture()) { - header += "vec4 gl_TexCoord[8];"; - } - if (ctx.info.stores.AnyComponent(IR::Attribute::ColorFrontDiffuseR)) { - header += "vec4 gl_FrontColor;"; - } - if (ctx.info.stores.AnyComponent(IR::Attribute::ColorFrontSpecularR)) { - header += "vec4 gl_FrontSecondaryColor;"; - } - if (ctx.info.stores.AnyComponent(IR::Attribute::ColorBackDiffuseR)) { - header += "vec4 gl_BackColor;"; - } - if (ctx.info.stores.AnyComponent(IR::Attribute::ColorBackSpecularR)) { - header += "vec4 gl_BackSecondaryColor;"; - } -} - -void SetupOutPerVertex(EmitContext& ctx, std::string& header) { - if (!StoresPerVertexAttributes(ctx.stage)) { - return; - } - if (ctx.uses_geometry_passthrough) { - return; - } - header += "out gl_PerVertex{vec4 gl_Position;"; - if (ctx.info.stores[IR::Attribute::PointSize]) { - header += "float gl_PointSize;"; - } - if (ctx.info.stores.ClipDistances()) { - header += "float gl_ClipDistance[];"; - } - if (ctx.info.stores[IR::Attribute::ViewportIndex] && - ctx.profile.support_viewport_index_layer_non_geometry && ctx.stage != Stage::Geometry) { - header += "int gl_ViewportIndex;"; - } - SetupLegacyOutPerVertex(ctx, header); - header += "};"; - if (ctx.info.stores[IR::Attribute::ViewportIndex] && ctx.stage == Stage::Geometry) { - header += "out int gl_ViewportIndex;"; - } -} - -void SetupInPerVertex(EmitContext& ctx, std::string& header) { - // Currently only required for TessellationControl to adhere to - // ARB_separate_shader_objects requirements - if (ctx.stage != Stage::TessellationControl) { - return; - } - const bool loads_position{ctx.info.loads.AnyComponent(IR::Attribute::PositionX)}; - const bool loads_point_size{ctx.info.loads[IR::Attribute::PointSize]}; - const bool loads_clip_distance{ctx.info.loads.ClipDistances()}; - const bool loads_per_vertex{loads_position || loads_point_size || loads_clip_distance}; - if (!loads_per_vertex) { - return; - } - header += "in gl_PerVertex{"; - if (loads_position) { - header += "vec4 gl_Position;"; - } - if (loads_point_size) { - header += "float gl_PointSize;"; - } - if (loads_clip_distance) { - header += "float gl_ClipDistance[];"; - } - header += "}gl_in[gl_MaxPatchVertices];"; -} - -void SetupLegacyInPerFragment(EmitContext& ctx, std::string& header) { - if (!ctx.info.loads.Legacy()) { - return; - } - header += "in gl_PerFragment{"; - if (ctx.info.loads.FixedFunctionTexture()) { - header += "vec4 gl_TexCoord[8];"; - } - if (ctx.info.loads.AnyComponent(IR::Attribute::ColorFrontDiffuseR)) { - header += "vec4 gl_Color;"; - } - header += "};"; -} - -} // Anonymous namespace - -EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, - const RuntimeInfo& runtime_info_) - : info{program.info}, profile{profile_}, runtime_info{runtime_info_}, stage{program.stage}, - uses_geometry_passthrough{program.is_geometry_passthrough && - profile.support_geometry_shader_passthrough} { - if (profile.need_fastmath_off) { - header += "#pragma optionNV(fastmath off)\n"; - } - SetupExtensions(); - switch (program.stage) { - case Stage::VertexA: - case Stage::VertexB: - stage_name = "vs"; - break; - case Stage::TessellationControl: - stage_name = "tcs"; - header += fmt::format("layout(vertices={})out;", program.invocations); - break; - case Stage::TessellationEval: - stage_name = "tes"; - 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"; - header += fmt::format("layout({})in;", InputPrimitive(runtime_info.input_topology)); - if (uses_geometry_passthrough) { - header += "layout(passthrough)in gl_PerVertex{vec4 gl_Position;};"; - break; - } else if (program.is_geometry_passthrough && - !profile.support_geometry_shader_passthrough) { - LOG_WARNING(Shader_GLSL, "Passthrough geometry program used but not supported"); - } - header += fmt::format( - "layout({},max_vertices={})out;in gl_PerVertex{{vec4 gl_Position;}}gl_in[];", - OutputPrimitive(program.output_topology), program.output_vertices); - break; - case Stage::Fragment: - stage_name = "fs"; - position_name = "gl_FragCoord"; - if (runtime_info.force_early_z) { - header += "layout(early_fragment_tests)in;"; - } - if (info.uses_sample_id) { - header += "in int gl_SampleID;"; - } - if (info.stores_sample_mask) { - header += "out int gl_SampleMask[];"; - } - break; - case Stage::Compute: - stage_name = "cs"; - 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)}; - header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;", - local_x, local_y, local_z); - break; - } - SetupOutPerVertex(*this, header); - SetupInPerVertex(*this, header); - SetupLegacyInPerFragment(*this, header); - - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - if (!info.loads.Generic(index) || !runtime_info.previous_stage_stores.Generic(index)) { - continue; - } - const auto qualifier{uses_geometry_passthrough ? "passthrough" - : fmt::format("location={}", index)}; - header += fmt::format("layout({}){}in vec4 in_attr{}{};", qualifier, - InterpDecorator(info.interpolation[index]), index, - InputArrayDecorator(stage)); - } - for (size_t index = 0; index < info.uses_patches.size(); ++index) { - if (!info.uses_patches[index]) { - continue; - } - const auto qualifier{stage == Stage::TessellationControl ? "out" : "in"}; - header += fmt::format("layout(location={})patch {} vec4 patch{};", index, qualifier, index); - } - if (stage == Stage::Fragment) { - for (size_t index = 0; index < info.stores_frag_color.size(); ++index) { - if (!info.stores_frag_color[index] && !profile.need_declared_frag_colors) { - continue; - } - header += fmt::format("layout(location={})out vec4 frag_color{};", index, index); - } - } - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - if (info.stores.Generic(index)) { - DefineGenericOutput(index, program.invocations); - } - } - if (info.uses_rescaling_uniform) { - header += "layout(location=0) uniform vec4 scaling;"; - } - DefineConstantBuffers(bindings); - DefineStorageBuffers(bindings); - SetupImages(bindings); - SetupTextures(bindings); - DefineHelperFunctions(); - DefineConstants(); -} - -void EmitContext::SetupExtensions() { - header += "#extension GL_ARB_separate_shader_objects : enable\n"; - if (info.uses_shadow_lod && profile.support_gl_texture_shadow_lod) { - header += "#extension GL_EXT_texture_shadow_lod : enable\n"; - } - if (info.uses_int64 && profile.support_int64) { - header += "#extension GL_ARB_gpu_shader_int64 : enable\n"; - } - if (info.uses_int64_bit_atomics) { - header += "#extension GL_NV_shader_atomic_int64 : enable\n"; - } - if (info.uses_atomic_f32_add) { - header += "#extension GL_NV_shader_atomic_float : enable\n"; - } - if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) { - header += "#extension GL_NV_shader_atomic_fp16_vector : enable\n"; - } - if (info.uses_fp16) { - if (profile.support_gl_nv_gpu_shader_5) { - header += "#extension GL_NV_gpu_shader5 : enable\n"; - } - if (profile.support_gl_amd_gpu_shader_half_float) { - header += "#extension GL_AMD_gpu_shader_half_float : enable\n"; - } - } - if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote || - info.uses_subgroup_shuffles || info.uses_fswzadd) { - header += "#extension GL_ARB_shader_ballot : enable\n" - "#extension GL_ARB_shader_group_vote : enable\n"; - if (!info.uses_int64 && profile.support_int64) { - header += "#extension GL_ARB_gpu_shader_int64 : enable\n"; - } - if (profile.support_gl_warp_intrinsics) { - header += "#extension GL_NV_shader_thread_shuffle : enable\n"; - } - } - if ((info.stores[IR::Attribute::ViewportIndex] || info.stores[IR::Attribute::Layer]) && - profile.support_viewport_index_layer_non_geometry && stage != Stage::Geometry) { - header += "#extension GL_ARB_shader_viewport_layer_array : enable\n"; - } - if (info.uses_sparse_residency && profile.support_gl_sparse_textures) { - header += "#extension GL_ARB_sparse_texture2 : enable\n"; - } - if (info.stores[IR::Attribute::ViewportMask] && profile.support_viewport_mask) { - header += "#extension GL_NV_viewport_array2 : enable\n"; - } - if (info.uses_typeless_image_reads) { - header += "#extension GL_EXT_shader_image_load_formatted : enable\n"; - } - if (info.uses_derivatives && profile.support_gl_derivative_control) { - header += "#extension GL_ARB_derivative_control : enable\n"; - } - if (uses_geometry_passthrough) { - header += "#extension GL_NV_geometry_shader_passthrough : enable\n"; - } -} - -void EmitContext::DefineConstantBuffers(Bindings& bindings) { - if (info.constant_buffer_descriptors.empty()) { - return; - } - for (const auto& desc : info.constant_buffer_descriptors) { - header += fmt::format( - "layout(std140,binding={}) uniform {}_cbuf_{}{{vec4 {}_cbuf{}[{}];}};", - bindings.uniform_buffer, stage_name, desc.index, stage_name, desc.index, 4 * 1024); - bindings.uniform_buffer += desc.count; - } -} - -void EmitContext::DefineStorageBuffers(Bindings& bindings) { - if (info.storage_buffers_descriptors.empty()) { - return; - } - u32 index{}; - for (const auto& desc : info.storage_buffers_descriptors) { - header += fmt::format("layout(std430,binding={}) buffer {}_ssbo_{}{{uint {}_ssbo{}[];}};", - bindings.storage_buffer, stage_name, bindings.storage_buffer, - stage_name, index); - bindings.storage_buffer += desc.count; - index += desc.count; - } -} - -void EmitContext::DefineGenericOutput(size_t index, u32 invocations) { - static constexpr std::string_view swizzle{"xyzw"}; - const size_t base_index{static_cast(IR::Attribute::Generic0X) + index * 4}; - u32 element{0}; - while (element < 4) { - std::string definition{fmt::format("layout(location={}", index)}; - const u32 remainder{4 - element}; - const TransformFeedbackVarying* xfb_varying{}; - if (!runtime_info.xfb_varyings.empty()) { - xfb_varying = &runtime_info.xfb_varyings[base_index + element]; - xfb_varying = xfb_varying && xfb_varying->components > 0 ? xfb_varying : nullptr; - } - const u32 num_components{xfb_varying ? xfb_varying->components : remainder}; - if (element > 0) { - definition += fmt::format(",component={}", element); - } - if (xfb_varying) { - definition += - fmt::format(",xfb_buffer={},xfb_stride={},xfb_offset={}", xfb_varying->buffer, - xfb_varying->stride, xfb_varying->offset); - } - std::string name{fmt::format("out_attr{}", index)}; - if (num_components < 4 || element > 0) { - name += fmt::format("_{}", swizzle.substr(element, num_components)); - } - const auto type{num_components == 1 ? "float" : fmt::format("vec{}", num_components)}; - definition += fmt::format(")out {} {}{};", type, name, OutputDecorator(stage, invocations)); - header += definition; - - const GenericElementInfo element_info{ - .name = name, - .first_element = element, - .num_components = num_components, - }; - std::fill_n(output_generics[index].begin() + element, num_components, element_info); - element += num_components; - } -} - -void EmitContext::DefineHelperFunctions() { - header += "\n#define ftoi floatBitsToInt\n#define ftou floatBitsToUint\n" - "#define itof intBitsToFloat\n#define utof uintBitsToFloat\n"; - if (info.uses_global_increment || info.uses_shared_increment) { - header += "uint CasIncrement(uint op_a,uint op_b){return op_a>=op_b?0u:(op_a+1u);}"; - } - if (info.uses_global_decrement || info.uses_shared_decrement) { - header += "uint CasDecrement(uint op_a,uint op_b){" - "return op_a==0||op_a>op_b?op_b:(op_a-1u);}"; - } - if (info.uses_atomic_f32_add) { - header += "uint CasFloatAdd(uint op_a,float op_b){" - "return ftou(utof(op_a)+op_b);}"; - } - if (info.uses_atomic_f32x2_add) { - header += "uint CasFloatAdd32x2(uint op_a,vec2 op_b){" - "return packHalf2x16(unpackHalf2x16(op_a)+op_b);}"; - } - if (info.uses_atomic_f32x2_min) { - header += "uint CasFloatMin32x2(uint op_a,vec2 op_b){return " - "packHalf2x16(min(unpackHalf2x16(op_a),op_b));}"; - } - if (info.uses_atomic_f32x2_max) { - header += "uint CasFloatMax32x2(uint op_a,vec2 op_b){return " - "packHalf2x16(max(unpackHalf2x16(op_a),op_b));}"; - } - if (info.uses_atomic_f16x2_add) { - header += "uint CasFloatAdd16x2(uint op_a,f16vec2 op_b){return " - "packFloat2x16(unpackFloat2x16(op_a)+op_b);}"; - } - if (info.uses_atomic_f16x2_min) { - header += "uint CasFloatMin16x2(uint op_a,f16vec2 op_b){return " - "packFloat2x16(min(unpackFloat2x16(op_a),op_b));}"; - } - if (info.uses_atomic_f16x2_max) { - header += "uint CasFloatMax16x2(uint op_a,f16vec2 op_b){return " - "packFloat2x16(max(unpackFloat2x16(op_a),op_b));}"; - } - if (info.uses_atomic_s32_min) { - header += "uint CasMinS32(uint op_a,uint op_b){return uint(min(int(op_a),int(op_b)));}"; - } - if (info.uses_atomic_s32_max) { - header += "uint CasMaxS32(uint op_a,uint op_b){return uint(max(int(op_a),int(op_b)));}"; - } - if (info.uses_global_memory && profile.support_int64) { - header += DefineGlobalMemoryFunctions(); - } - if (info.loads_indexed_attributes) { - const bool is_array{stage == Stage::Geometry}; - const auto vertex_arg{is_array ? ",uint vertex" : ""}; - std::string func{ - fmt::format("float IndexedAttrLoad(int offset{}){{int base_index=offset>>2;uint " - "masked_index=uint(base_index)&3u;switch(base_index>>2){{", - vertex_arg)}; - if (info.loads.AnyComponent(IR::Attribute::PositionX)) { - const auto position_idx{is_array ? "gl_in[vertex]." : ""}; - func += fmt::format("case {}:return {}{}[masked_index];", - static_cast(IR::Attribute::PositionX) >> 2, position_idx, - position_name); - } - const u32 base_attribute_value = static_cast(IR::Attribute::Generic0X) >> 2; - for (u32 index = 0; index < IR::NUM_GENERICS; ++index) { - if (!info.loads.Generic(index)) { - continue; - } - const auto vertex_idx{is_array ? "[vertex]" : ""}; - func += fmt::format("case {}:return in_attr{}{}[masked_index];", - base_attribute_value + index, index, vertex_idx); - } - func += "default: return 0.0;}}"; - header += func; - } - if (info.stores_indexed_attributes) { - // TODO - } -} - -std::string EmitContext::DefineGlobalMemoryFunctions() { - const auto define_body{[&](std::string& func, size_t index, std::string_view return_statement) { - const auto& ssbo{info.storage_buffers_descriptors[index]}; - const u32 size_cbuf_offset{ssbo.cbuf_offset + 8}; - const auto ssbo_addr{fmt::format("ssbo_addr{}", index)}; - const auto cbuf{fmt::format("{}_cbuf{}", stage_name, ssbo.cbuf_index)}; - std::array addr_xy; - std::array size_xy; - for (size_t i = 0; i < addr_xy.size(); ++i) { - const auto addr_loc{ssbo.cbuf_offset + 4 * i}; - const auto size_loc{size_cbuf_offset + 4 * i}; - addr_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, addr_loc / 16, Swizzle(addr_loc)); - size_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, size_loc / 16, Swizzle(size_loc)); - } - const auto addr_pack{fmt::format("packUint2x32(uvec2({},{}))", addr_xy[0], addr_xy[1])}; - const auto addr_statment{fmt::format("uint64_t {}={};", ssbo_addr, addr_pack)}; - func += addr_statment; - - const auto size_vec{fmt::format("uvec2({},{})", size_xy[0], size_xy[1])}; - const auto comp_lhs{fmt::format("(addr>={})", ssbo_addr)}; - const auto comp_rhs{fmt::format("(addr<({}+uint64_t({})))", ssbo_addr, size_vec)}; - const auto comparison{fmt::format("if({}&&{}){{", comp_lhs, comp_rhs)}; - func += comparison; - - const auto ssbo_name{fmt::format("{}_ssbo{}", stage_name, index)}; - func += fmt::format(fmt::runtime(return_statement), ssbo_name, ssbo_addr); - }}; - std::string write_func{"void WriteGlobal32(uint64_t addr,uint data){"}; - std::string write_func_64{"void WriteGlobal64(uint64_t addr,uvec2 data){"}; - std::string write_func_128{"void WriteGlobal128(uint64_t addr,uvec4 data){"}; - std::string load_func{"uint LoadGlobal32(uint64_t addr){"}; - std::string load_func_64{"uvec2 LoadGlobal64(uint64_t addr){"}; - std::string load_func_128{"uvec4 LoadGlobal128(uint64_t addr){"}; - const size_t num_buffers{info.storage_buffers_descriptors.size()}; - for (size_t index = 0; index < num_buffers; ++index) { - if (!info.nvn_buffer_used[index]) { - continue; - } - define_body(write_func, index, "{0}[uint(addr-{1})>>2]=data;return;}}"); - define_body(write_func_64, index, - "{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;return;}}"); - define_body(write_func_128, index, - "{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;{0}[uint(" - "addr-{1}+8)>>2]=data.z;{0}[uint(addr-{1}+12)>>2]=data.w;return;}}"); - define_body(load_func, index, "return {0}[uint(addr-{1})>>2];}}"); - define_body(load_func_64, index, - "return uvec2({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2]);}}"); - define_body(load_func_128, index, - "return uvec4({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2],{0}[" - "uint(addr-{1}+8)>>2],{0}[uint(addr-{1}+12)>>2]);}}"); - } - write_func += '}'; - write_func_64 += '}'; - write_func_128 += '}'; - load_func += "return 0u;}"; - load_func_64 += "return uvec2(0);}"; - load_func_128 += "return uvec4(0);}"; - return write_func + write_func_64 + write_func_128 + load_func + load_func_64 + load_func_128; -} - -void EmitContext::SetupImages(Bindings& bindings) { - 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; - } - images.reserve(info.image_descriptors.size()); - for (const auto& desc : info.image_descriptors) { - images.push_back({bindings.image, desc.count}); - 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) : ""}; - header += fmt::format("layout(binding={}{})uniform {}{} img{}{};", bindings.image, format, - qualifier, image_type, bindings.image, array_decorator); - bindings.image += desc.count; - } -} - -void EmitContext::SetupTextures(Bindings& bindings) { - 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{SamplerType(TextureType::Buffer, false)}; - 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; - } - textures.reserve(info.texture_descriptors.size()); - for (const auto& desc : info.texture_descriptors) { - textures.push_back({bindings.texture, desc.count}); - const auto sampler_type{SamplerType(desc.type, desc.is_depth)}; - 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; - } -} - -void EmitContext::DefineConstants() { - if (info.uses_fswzadd) { - header += "const float FSWZ_A[]=float[4](-1.f,1.f,-1.f,0.f);" - "const float FSWZ_B[]=float[4](-1.f,-1.f,1.f,-1.f);"; - } -} - -} // namespace Shader::Backend::GLSL diff --git a/src/shader_recompiler/backend/glsl/emit_context.h b/src/shader_recompiler/backend/glsl/emit_context.h deleted file mode 100644 index d9b639d29..000000000 --- a/src/shader_recompiler/backend/glsl/emit_context.h +++ /dev/null @@ -1,174 +0,0 @@ -// Copyright 2021 yuzu Emulator Project -// Licensed under GPLv2 or any later version -// Refer to the license.txt file included. - -#pragma once - -#include -#include -#include - -#include - -#include "shader_recompiler/backend/glsl/var_alloc.h" -#include "shader_recompiler/stage.h" - -namespace Shader { -struct Info; -struct Profile; -struct RuntimeInfo; -} // namespace Shader - -namespace Shader::Backend { -struct Bindings; -} - -namespace Shader::IR { -class Inst; -struct Program; -} // namespace Shader::IR - -namespace Shader::Backend::GLSL { - -struct GenericElementInfo { - std::string name; - u32 first_element{}; - u32 num_components{}; -}; - -struct TextureImageDefinition { - u32 binding; - u32 count; -}; - -class EmitContext { -public: - explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, - const RuntimeInfo& runtime_info_); - - template - void Add(const char* format_str, IR::Inst& inst, Args&&... args) { - const auto var_def{var_alloc.AddDefine(inst, type)}; - if (var_def.empty()) { - // skip assigment. - code += fmt::format(fmt::runtime(format_str + 3), std::forward(args)...); - } else { - code += fmt::format(fmt::runtime(format_str), var_def, std::forward(args)...); - } - // TODO: Remove this - code += '\n'; - } - - template - void AddU1(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddF16x2(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddU32(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddF32(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddU64(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddF64(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddU32x2(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddF32x2(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddU32x3(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddF32x3(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddU32x4(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddF32x4(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddPrecF32(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void AddPrecF64(const char* format_str, IR::Inst& inst, Args&&... args) { - Add(format_str, inst, args...); - } - - template - void Add(const char* format_str, Args&&... args) { - code += fmt::format(fmt::runtime(format_str), std::forward(args)...); - // TODO: Remove this - code += '\n'; - } - - std::string header; - std::string code; - VarAlloc var_alloc; - const Info& info; - const Profile& profile; - const RuntimeInfo& runtime_info; - - Stage stage{}; - std::string_view stage_name = "invalid"; - std::string_view position_name = "gl_Position"; - - std::vector texture_buffers; - std::vector image_buffers; - std::vector textures; - std::vector images; - std::array, 32> output_generics{}; - - u32 num_safety_loop_vars{}; - - bool uses_y_direction{}; - bool uses_cc_carry{}; - bool uses_geometry_passthrough{}; - -private: - void SetupExtensions(); - void DefineConstantBuffers(Bindings& bindings); - void DefineStorageBuffers(Bindings& bindings); - void DefineGenericOutput(size_t index, u32 invocations); - void DefineHelperFunctions(); - void DefineConstants(); - std::string DefineGlobalMemoryFunctions(); - void SetupImages(Bindings& bindings); - void SetupTextures(Bindings& bindings); -}; - -} // namespace Shader::Backend::GLSL diff --git a/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp b/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp new file mode 100644 index 000000000..97bd59302 --- /dev/null +++ b/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp @@ -0,0 +1,718 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include "shader_recompiler/backend/bindings.h" +#include "shader_recompiler/backend/glsl/emit_context.h" +#include "shader_recompiler/frontend/ir/program.h" +#include "shader_recompiler/profile.h" +#include "shader_recompiler/runtime_info.h" + +namespace Shader::Backend::GLSL { +namespace { +u32 CbufIndex(size_t offset) { + return (offset / 4) % 4; +} + +char Swizzle(size_t offset) { + return "xyzw"[CbufIndex(offset)]; +} + +std::string_view InterpDecorator(Interpolation interp) { + switch (interp) { + case Interpolation::Smooth: + return ""; + case Interpolation::Flat: + return "flat "; + case Interpolation::NoPerspective: + return "noperspective "; + } + throw InvalidArgument("Invalid interpolation {}", interp); +} + +std::string_view InputArrayDecorator(Stage stage) { + switch (stage) { + case Stage::Geometry: + case Stage::TessellationControl: + case Stage::TessellationEval: + return "[]"; + default: + return ""; + } +} + +bool StoresPerVertexAttributes(Stage stage) { + switch (stage) { + case Stage::VertexA: + case Stage::VertexB: + case Stage::Geometry: + case Stage::TessellationEval: + return true; + default: + return false; + } +} + +std::string OutputDecorator(Stage stage, u32 size) { + switch (stage) { + case Stage::TessellationControl: + return fmt::format("[{}]", size); + default: + return ""; + } +} + +std::string_view SamplerType(TextureType type, bool is_depth) { + if (is_depth) { + 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); + } + } + switch (type) { + case TextureType::Color1D: + return "sampler1D"; + case TextureType::ColorArray1D: + return "sampler1DArray"; + case TextureType::Color2D: + return "sampler2D"; + case TextureType::ColorArray2D: + return "sampler2DArray"; + case TextureType::Color3D: + return "sampler3D"; + case TextureType::ColorCube: + return "samplerCube"; + case TextureType::ColorArrayCube: + return "samplerCubeArray"; + case TextureType::Buffer: + return "samplerBuffer"; + default: + throw NotImplementedException("Texture type: {}", type); + } +} + +std::string_view ImageType(TextureType type) { + switch (type) { + case TextureType::Color1D: + return "uimage1D"; + case TextureType::ColorArray1D: + return "uimage1DArray"; + case TextureType::Color2D: + return "uimage2D"; + case TextureType::ColorArray2D: + return "uimage2DArray"; + case TextureType::Color3D: + return "uimage3D"; + case TextureType::ColorCube: + return "uimageCube"; + case TextureType::ColorArrayCube: + return "uimageCubeArray"; + case TextureType::Buffer: + return "uimageBuffer"; + default: + throw NotImplementedException("Image type: {}", type); + } +} + +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 "writeonly "; + } + if (is_read && !is_written) { + return "readonly "; + } + return ""; +} + +std::string_view GetTessMode(TessPrimitive primitive) { + switch (primitive) { + case TessPrimitive::Triangles: + return "triangles"; + case TessPrimitive::Quads: + return "quads"; + case TessPrimitive::Isolines: + return "isolines"; + } + throw InvalidArgument("Invalid tessellation primitive {}", primitive); +} + +std::string_view GetTessSpacing(TessSpacing spacing) { + switch (spacing) { + case TessSpacing::Equal: + return "equal_spacing"; + case TessSpacing::FractionalOdd: + return "fractional_odd_spacing"; + case TessSpacing::FractionalEven: + return "fractional_even_spacing"; + } + throw InvalidArgument("Invalid tessellation spacing {}", spacing); +} + +std::string_view InputPrimitive(InputTopology topology) { + switch (topology) { + case InputTopology::Points: + return "points"; + case InputTopology::Lines: + return "lines"; + case InputTopology::LinesAdjacency: + return "lines_adjacency"; + case InputTopology::Triangles: + return "triangles"; + case InputTopology::TrianglesAdjacency: + return "triangles_adjacency"; + } + throw InvalidArgument("Invalid input topology {}", topology); +} + +std::string_view OutputPrimitive(OutputTopology topology) { + switch (topology) { + case OutputTopology::PointList: + return "points"; + case OutputTopology::LineStrip: + return "line_strip"; + case OutputTopology::TriangleStrip: + return "triangle_strip"; + } + throw InvalidArgument("Invalid output topology {}", topology); +} + +void SetupLegacyOutPerVertex(EmitContext& ctx, std::string& header) { + if (!ctx.info.stores.Legacy()) { + return; + } + if (ctx.info.stores.FixedFunctionTexture()) { + header += "vec4 gl_TexCoord[8];"; + } + if (ctx.info.stores.AnyComponent(IR::Attribute::ColorFrontDiffuseR)) { + header += "vec4 gl_FrontColor;"; + } + if (ctx.info.stores.AnyComponent(IR::Attribute::ColorFrontSpecularR)) { + header += "vec4 gl_FrontSecondaryColor;"; + } + if (ctx.info.stores.AnyComponent(IR::Attribute::ColorBackDiffuseR)) { + header += "vec4 gl_BackColor;"; + } + if (ctx.info.stores.AnyComponent(IR::Attribute::ColorBackSpecularR)) { + header += "vec4 gl_BackSecondaryColor;"; + } +} + +void SetupOutPerVertex(EmitContext& ctx, std::string& header) { + if (!StoresPerVertexAttributes(ctx.stage)) { + return; + } + if (ctx.uses_geometry_passthrough) { + return; + } + header += "out gl_PerVertex{vec4 gl_Position;"; + if (ctx.info.stores[IR::Attribute::PointSize]) { + header += "float gl_PointSize;"; + } + if (ctx.info.stores.ClipDistances()) { + header += "float gl_ClipDistance[];"; + } + if (ctx.info.stores[IR::Attribute::ViewportIndex] && + ctx.profile.support_viewport_index_layer_non_geometry && ctx.stage != Stage::Geometry) { + header += "int gl_ViewportIndex;"; + } + SetupLegacyOutPerVertex(ctx, header); + header += "};"; + if (ctx.info.stores[IR::Attribute::ViewportIndex] && ctx.stage == Stage::Geometry) { + header += "out int gl_ViewportIndex;"; + } +} + +void SetupInPerVertex(EmitContext& ctx, std::string& header) { + // Currently only required for TessellationControl to adhere to + // ARB_separate_shader_objects requirements + if (ctx.stage != Stage::TessellationControl) { + return; + } + const bool loads_position{ctx.info.loads.AnyComponent(IR::Attribute::PositionX)}; + const bool loads_point_size{ctx.info.loads[IR::Attribute::PointSize]}; + const bool loads_clip_distance{ctx.info.loads.ClipDistances()}; + const bool loads_per_vertex{loads_position || loads_point_size || loads_clip_distance}; + if (!loads_per_vertex) { + return; + } + header += "in gl_PerVertex{"; + if (loads_position) { + header += "vec4 gl_Position;"; + } + if (loads_point_size) { + header += "float gl_PointSize;"; + } + if (loads_clip_distance) { + header += "float gl_ClipDistance[];"; + } + header += "}gl_in[gl_MaxPatchVertices];"; +} + +void SetupLegacyInPerFragment(EmitContext& ctx, std::string& header) { + if (!ctx.info.loads.Legacy()) { + return; + } + header += "in gl_PerFragment{"; + if (ctx.info.loads.FixedFunctionTexture()) { + header += "vec4 gl_TexCoord[8];"; + } + if (ctx.info.loads.AnyComponent(IR::Attribute::ColorFrontDiffuseR)) { + header += "vec4 gl_Color;"; + } + header += "};"; +} + +} // Anonymous namespace + +EmitContext::EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, + const RuntimeInfo& runtime_info_) + : info{program.info}, profile{profile_}, runtime_info{runtime_info_}, stage{program.stage}, + uses_geometry_passthrough{program.is_geometry_passthrough && + profile.support_geometry_shader_passthrough} { + if (profile.need_fastmath_off) { + header += "#pragma optionNV(fastmath off)\n"; + } + SetupExtensions(); + switch (program.stage) { + case Stage::VertexA: + case Stage::VertexB: + stage_name = "vs"; + break; + case Stage::TessellationControl: + stage_name = "tcs"; + header += fmt::format("layout(vertices={})out;", program.invocations); + break; + case Stage::TessellationEval: + stage_name = "tes"; + 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"; + header += fmt::format("layout({})in;", InputPrimitive(runtime_info.input_topology)); + if (uses_geometry_passthrough) { + header += "layout(passthrough)in gl_PerVertex{vec4 gl_Position;};"; + break; + } else if (program.is_geometry_passthrough && + !profile.support_geometry_shader_passthrough) { + LOG_WARNING(Shader_GLSL, "Passthrough geometry program used but not supported"); + } + header += fmt::format( + "layout({},max_vertices={})out;in gl_PerVertex{{vec4 gl_Position;}}gl_in[];", + OutputPrimitive(program.output_topology), program.output_vertices); + break; + case Stage::Fragment: + stage_name = "fs"; + position_name = "gl_FragCoord"; + if (runtime_info.force_early_z) { + header += "layout(early_fragment_tests)in;"; + } + if (info.uses_sample_id) { + header += "in int gl_SampleID;"; + } + if (info.stores_sample_mask) { + header += "out int gl_SampleMask[];"; + } + break; + case Stage::Compute: + stage_name = "cs"; + 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)}; + header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;", + local_x, local_y, local_z); + break; + } + SetupOutPerVertex(*this, header); + SetupInPerVertex(*this, header); + SetupLegacyInPerFragment(*this, header); + + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + if (!info.loads.Generic(index) || !runtime_info.previous_stage_stores.Generic(index)) { + continue; + } + const auto qualifier{uses_geometry_passthrough ? "passthrough" + : fmt::format("location={}", index)}; + header += fmt::format("layout({}){}in vec4 in_attr{}{};", qualifier, + InterpDecorator(info.interpolation[index]), index, + InputArrayDecorator(stage)); + } + for (size_t index = 0; index < info.uses_patches.size(); ++index) { + if (!info.uses_patches[index]) { + continue; + } + const auto qualifier{stage == Stage::TessellationControl ? "out" : "in"}; + header += fmt::format("layout(location={})patch {} vec4 patch{};", index, qualifier, index); + } + if (stage == Stage::Fragment) { + for (size_t index = 0; index < info.stores_frag_color.size(); ++index) { + if (!info.stores_frag_color[index] && !profile.need_declared_frag_colors) { + continue; + } + header += fmt::format("layout(location={})out vec4 frag_color{};", index, index); + } + } + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + if (info.stores.Generic(index)) { + DefineGenericOutput(index, program.invocations); + } + } + if (info.uses_rescaling_uniform) { + header += "layout(location=0) uniform vec4 scaling;"; + } + DefineConstantBuffers(bindings); + DefineStorageBuffers(bindings); + SetupImages(bindings); + SetupTextures(bindings); + DefineHelperFunctions(); + DefineConstants(); +} + +void EmitContext::SetupExtensions() { + header += "#extension GL_ARB_separate_shader_objects : enable\n"; + if (info.uses_shadow_lod && profile.support_gl_texture_shadow_lod) { + header += "#extension GL_EXT_texture_shadow_lod : enable\n"; + } + if (info.uses_int64 && profile.support_int64) { + header += "#extension GL_ARB_gpu_shader_int64 : enable\n"; + } + if (info.uses_int64_bit_atomics) { + header += "#extension GL_NV_shader_atomic_int64 : enable\n"; + } + if (info.uses_atomic_f32_add) { + header += "#extension GL_NV_shader_atomic_float : enable\n"; + } + if (info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max) { + header += "#extension GL_NV_shader_atomic_fp16_vector : enable\n"; + } + if (info.uses_fp16) { + if (profile.support_gl_nv_gpu_shader_5) { + header += "#extension GL_NV_gpu_shader5 : enable\n"; + } + if (profile.support_gl_amd_gpu_shader_half_float) { + header += "#extension GL_AMD_gpu_shader_half_float : enable\n"; + } + } + if (info.uses_subgroup_invocation_id || info.uses_subgroup_mask || info.uses_subgroup_vote || + info.uses_subgroup_shuffles || info.uses_fswzadd) { + header += "#extension GL_ARB_shader_ballot : enable\n" + "#extension GL_ARB_shader_group_vote : enable\n"; + if (!info.uses_int64 && profile.support_int64) { + header += "#extension GL_ARB_gpu_shader_int64 : enable\n"; + } + if (profile.support_gl_warp_intrinsics) { + header += "#extension GL_NV_shader_thread_shuffle : enable\n"; + } + } + if ((info.stores[IR::Attribute::ViewportIndex] || info.stores[IR::Attribute::Layer]) && + profile.support_viewport_index_layer_non_geometry && stage != Stage::Geometry) { + header += "#extension GL_ARB_shader_viewport_layer_array : enable\n"; + } + if (info.uses_sparse_residency && profile.support_gl_sparse_textures) { + header += "#extension GL_ARB_sparse_texture2 : enable\n"; + } + if (info.stores[IR::Attribute::ViewportMask] && profile.support_viewport_mask) { + header += "#extension GL_NV_viewport_array2 : enable\n"; + } + if (info.uses_typeless_image_reads) { + header += "#extension GL_EXT_shader_image_load_formatted : enable\n"; + } + if (info.uses_derivatives && profile.support_gl_derivative_control) { + header += "#extension GL_ARB_derivative_control : enable\n"; + } + if (uses_geometry_passthrough) { + header += "#extension GL_NV_geometry_shader_passthrough : enable\n"; + } +} + +void EmitContext::DefineConstantBuffers(Bindings& bindings) { + if (info.constant_buffer_descriptors.empty()) { + return; + } + for (const auto& desc : info.constant_buffer_descriptors) { + header += fmt::format( + "layout(std140,binding={}) uniform {}_cbuf_{}{{vec4 {}_cbuf{}[{}];}};", + bindings.uniform_buffer, stage_name, desc.index, stage_name, desc.index, 4 * 1024); + bindings.uniform_buffer += desc.count; + } +} + +void EmitContext::DefineStorageBuffers(Bindings& bindings) { + if (info.storage_buffers_descriptors.empty()) { + return; + } + u32 index{}; + for (const auto& desc : info.storage_buffers_descriptors) { + header += fmt::format("layout(std430,binding={}) buffer {}_ssbo_{}{{uint {}_ssbo{}[];}};", + bindings.storage_buffer, stage_name, bindings.storage_buffer, + stage_name, index); + bindings.storage_buffer += desc.count; + index += desc.count; + } +} + +void EmitContext::DefineGenericOutput(size_t index, u32 invocations) { + static constexpr std::string_view swizzle{"xyzw"}; + const size_t base_index{static_cast(IR::Attribute::Generic0X) + index * 4}; + u32 element{0}; + while (element < 4) { + std::string definition{fmt::format("layout(location={}", index)}; + const u32 remainder{4 - element}; + const TransformFeedbackVarying* xfb_varying{}; + if (!runtime_info.xfb_varyings.empty()) { + xfb_varying = &runtime_info.xfb_varyings[base_index + element]; + xfb_varying = xfb_varying && xfb_varying->components > 0 ? xfb_varying : nullptr; + } + const u32 num_components{xfb_varying ? xfb_varying->components : remainder}; + if (element > 0) { + definition += fmt::format(",component={}", element); + } + if (xfb_varying) { + definition += + fmt::format(",xfb_buffer={},xfb_stride={},xfb_offset={}", xfb_varying->buffer, + xfb_varying->stride, xfb_varying->offset); + } + std::string name{fmt::format("out_attr{}", index)}; + if (num_components < 4 || element > 0) { + name += fmt::format("_{}", swizzle.substr(element, num_components)); + } + const auto type{num_components == 1 ? "float" : fmt::format("vec{}", num_components)}; + definition += fmt::format(")out {} {}{};", type, name, OutputDecorator(stage, invocations)); + header += definition; + + const GenericElementInfo element_info{ + .name = name, + .first_element = element, + .num_components = num_components, + }; + std::fill_n(output_generics[index].begin() + element, num_components, element_info); + element += num_components; + } +} + +void EmitContext::DefineHelperFunctions() { + header += "\n#define ftoi floatBitsToInt\n#define ftou floatBitsToUint\n" + "#define itof intBitsToFloat\n#define utof uintBitsToFloat\n"; + if (info.uses_global_increment || info.uses_shared_increment) { + header += "uint CasIncrement(uint op_a,uint op_b){return op_a>=op_b?0u:(op_a+1u);}"; + } + if (info.uses_global_decrement || info.uses_shared_decrement) { + header += "uint CasDecrement(uint op_a,uint op_b){" + "return op_a==0||op_a>op_b?op_b:(op_a-1u);}"; + } + if (info.uses_atomic_f32_add) { + header += "uint CasFloatAdd(uint op_a,float op_b){" + "return ftou(utof(op_a)+op_b);}"; + } + if (info.uses_atomic_f32x2_add) { + header += "uint CasFloatAdd32x2(uint op_a,vec2 op_b){" + "return packHalf2x16(unpackHalf2x16(op_a)+op_b);}"; + } + if (info.uses_atomic_f32x2_min) { + header += "uint CasFloatMin32x2(uint op_a,vec2 op_b){return " + "packHalf2x16(min(unpackHalf2x16(op_a),op_b));}"; + } + if (info.uses_atomic_f32x2_max) { + header += "uint CasFloatMax32x2(uint op_a,vec2 op_b){return " + "packHalf2x16(max(unpackHalf2x16(op_a),op_b));}"; + } + if (info.uses_atomic_f16x2_add) { + header += "uint CasFloatAdd16x2(uint op_a,f16vec2 op_b){return " + "packFloat2x16(unpackFloat2x16(op_a)+op_b);}"; + } + if (info.uses_atomic_f16x2_min) { + header += "uint CasFloatMin16x2(uint op_a,f16vec2 op_b){return " + "packFloat2x16(min(unpackFloat2x16(op_a),op_b));}"; + } + if (info.uses_atomic_f16x2_max) { + header += "uint CasFloatMax16x2(uint op_a,f16vec2 op_b){return " + "packFloat2x16(max(unpackFloat2x16(op_a),op_b));}"; + } + if (info.uses_atomic_s32_min) { + header += "uint CasMinS32(uint op_a,uint op_b){return uint(min(int(op_a),int(op_b)));}"; + } + if (info.uses_atomic_s32_max) { + header += "uint CasMaxS32(uint op_a,uint op_b){return uint(max(int(op_a),int(op_b)));}"; + } + if (info.uses_global_memory && profile.support_int64) { + header += DefineGlobalMemoryFunctions(); + } + if (info.loads_indexed_attributes) { + const bool is_array{stage == Stage::Geometry}; + const auto vertex_arg{is_array ? ",uint vertex" : ""}; + std::string func{ + fmt::format("float IndexedAttrLoad(int offset{}){{int base_index=offset>>2;uint " + "masked_index=uint(base_index)&3u;switch(base_index>>2){{", + vertex_arg)}; + if (info.loads.AnyComponent(IR::Attribute::PositionX)) { + const auto position_idx{is_array ? "gl_in[vertex]." : ""}; + func += fmt::format("case {}:return {}{}[masked_index];", + static_cast(IR::Attribute::PositionX) >> 2, position_idx, + position_name); + } + const u32 base_attribute_value = static_cast(IR::Attribute::Generic0X) >> 2; + for (u32 index = 0; index < IR::NUM_GENERICS; ++index) { + if (!info.loads.Generic(index)) { + continue; + } + const auto vertex_idx{is_array ? "[vertex]" : ""}; + func += fmt::format("case {}:return in_attr{}{}[masked_index];", + base_attribute_value + index, index, vertex_idx); + } + func += "default: return 0.0;}}"; + header += func; + } + if (info.stores_indexed_attributes) { + // TODO + } +} + +std::string EmitContext::DefineGlobalMemoryFunctions() { + const auto define_body{[&](std::string& func, size_t index, std::string_view return_statement) { + const auto& ssbo{info.storage_buffers_descriptors[index]}; + const u32 size_cbuf_offset{ssbo.cbuf_offset + 8}; + const auto ssbo_addr{fmt::format("ssbo_addr{}", index)}; + const auto cbuf{fmt::format("{}_cbuf{}", stage_name, ssbo.cbuf_index)}; + std::array addr_xy; + std::array size_xy; + for (size_t i = 0; i < addr_xy.size(); ++i) { + const auto addr_loc{ssbo.cbuf_offset + 4 * i}; + const auto size_loc{size_cbuf_offset + 4 * i}; + addr_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, addr_loc / 16, Swizzle(addr_loc)); + size_xy[i] = fmt::format("ftou({}[{}].{})", cbuf, size_loc / 16, Swizzle(size_loc)); + } + const auto addr_pack{fmt::format("packUint2x32(uvec2({},{}))", addr_xy[0], addr_xy[1])}; + const auto addr_statment{fmt::format("uint64_t {}={};", ssbo_addr, addr_pack)}; + func += addr_statment; + + const auto size_vec{fmt::format("uvec2({},{})", size_xy[0], size_xy[1])}; + const auto comp_lhs{fmt::format("(addr>={})", ssbo_addr)}; + const auto comp_rhs{fmt::format("(addr<({}+uint64_t({})))", ssbo_addr, size_vec)}; + const auto comparison{fmt::format("if({}&&{}){{", comp_lhs, comp_rhs)}; + func += comparison; + + const auto ssbo_name{fmt::format("{}_ssbo{}", stage_name, index)}; + func += fmt::format(fmt::runtime(return_statement), ssbo_name, ssbo_addr); + }}; + std::string write_func{"void WriteGlobal32(uint64_t addr,uint data){"}; + std::string write_func_64{"void WriteGlobal64(uint64_t addr,uvec2 data){"}; + std::string write_func_128{"void WriteGlobal128(uint64_t addr,uvec4 data){"}; + std::string load_func{"uint LoadGlobal32(uint64_t addr){"}; + std::string load_func_64{"uvec2 LoadGlobal64(uint64_t addr){"}; + std::string load_func_128{"uvec4 LoadGlobal128(uint64_t addr){"}; + const size_t num_buffers{info.storage_buffers_descriptors.size()}; + for (size_t index = 0; index < num_buffers; ++index) { + if (!info.nvn_buffer_used[index]) { + continue; + } + define_body(write_func, index, "{0}[uint(addr-{1})>>2]=data;return;}}"); + define_body(write_func_64, index, + "{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;return;}}"); + define_body(write_func_128, index, + "{0}[uint(addr-{1})>>2]=data.x;{0}[uint(addr-{1}+4)>>2]=data.y;{0}[uint(" + "addr-{1}+8)>>2]=data.z;{0}[uint(addr-{1}+12)>>2]=data.w;return;}}"); + define_body(load_func, index, "return {0}[uint(addr-{1})>>2];}}"); + define_body(load_func_64, index, + "return uvec2({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2]);}}"); + define_body(load_func_128, index, + "return uvec4({0}[uint(addr-{1})>>2],{0}[uint(addr-{1}+4)>>2],{0}[" + "uint(addr-{1}+8)>>2],{0}[uint(addr-{1}+12)>>2]);}}"); + } + write_func += '}'; + write_func_64 += '}'; + write_func_128 += '}'; + load_func += "return 0u;}"; + load_func_64 += "return uvec2(0);}"; + load_func_128 += "return uvec4(0);}"; + return write_func + write_func_64 + write_func_128 + load_func + load_func_64 + load_func_128; +} + +void EmitContext::SetupImages(Bindings& bindings) { + 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; + } + images.reserve(info.image_descriptors.size()); + for (const auto& desc : info.image_descriptors) { + images.push_back({bindings.image, desc.count}); + 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) : ""}; + header += fmt::format("layout(binding={}{})uniform {}{} img{}{};", bindings.image, format, + qualifier, image_type, bindings.image, array_decorator); + bindings.image += desc.count; + } +} + +void EmitContext::SetupTextures(Bindings& bindings) { + 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{SamplerType(TextureType::Buffer, false)}; + 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; + } + textures.reserve(info.texture_descriptors.size()); + for (const auto& desc : info.texture_descriptors) { + textures.push_back({bindings.texture, desc.count}); + const auto sampler_type{SamplerType(desc.type, desc.is_depth)}; + 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; + } +} + +void EmitContext::DefineConstants() { + if (info.uses_fswzadd) { + header += "const float FSWZ_A[]=float[4](-1.f,1.f,-1.f,0.f);" + "const float FSWZ_B[]=float[4](-1.f,-1.f,1.f,-1.f);"; + } +} + +} // namespace Shader::Backend::GLSL diff --git a/src/shader_recompiler/backend/glsl/glsl_emit_context.h b/src/shader_recompiler/backend/glsl/glsl_emit_context.h new file mode 100644 index 000000000..d9b639d29 --- /dev/null +++ b/src/shader_recompiler/backend/glsl/glsl_emit_context.h @@ -0,0 +1,174 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#pragma once + +#include +#include +#include + +#include + +#include "shader_recompiler/backend/glsl/var_alloc.h" +#include "shader_recompiler/stage.h" + +namespace Shader { +struct Info; +struct Profile; +struct RuntimeInfo; +} // namespace Shader + +namespace Shader::Backend { +struct Bindings; +} + +namespace Shader::IR { +class Inst; +struct Program; +} // namespace Shader::IR + +namespace Shader::Backend::GLSL { + +struct GenericElementInfo { + std::string name; + u32 first_element{}; + u32 num_components{}; +}; + +struct TextureImageDefinition { + u32 binding; + u32 count; +}; + +class EmitContext { +public: + explicit EmitContext(IR::Program& program, Bindings& bindings, const Profile& profile_, + const RuntimeInfo& runtime_info_); + + template + void Add(const char* format_str, IR::Inst& inst, Args&&... args) { + const auto var_def{var_alloc.AddDefine(inst, type)}; + if (var_def.empty()) { + // skip assigment. + code += fmt::format(fmt::runtime(format_str + 3), std::forward(args)...); + } else { + code += fmt::format(fmt::runtime(format_str), var_def, std::forward(args)...); + } + // TODO: Remove this + code += '\n'; + } + + template + void AddU1(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddF16x2(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddU32(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddF32(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddU64(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddF64(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddU32x2(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddF32x2(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddU32x3(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddF32x3(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddU32x4(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddF32x4(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddPrecF32(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void AddPrecF64(const char* format_str, IR::Inst& inst, Args&&... args) { + Add(format_str, inst, args...); + } + + template + void Add(const char* format_str, Args&&... args) { + code += fmt::format(fmt::runtime(format_str), std::forward(args)...); + // TODO: Remove this + code += '\n'; + } + + std::string header; + std::string code; + VarAlloc var_alloc; + const Info& info; + const Profile& profile; + const RuntimeInfo& runtime_info; + + Stage stage{}; + std::string_view stage_name = "invalid"; + std::string_view position_name = "gl_Position"; + + std::vector texture_buffers; + std::vector image_buffers; + std::vector textures; + std::vector images; + std::array, 32> output_generics{}; + + u32 num_safety_loop_vars{}; + + bool uses_y_direction{}; + bool uses_cc_carry{}; + bool uses_geometry_passthrough{}; + +private: + void SetupExtensions(); + void DefineConstantBuffers(Bindings& bindings); + void DefineStorageBuffers(Bindings& bindings); + void DefineGenericOutput(size_t index, u32 invocations); + void DefineHelperFunctions(); + void DefineConstants(); + std::string DefineGlobalMemoryFunctions(); + void SetupImages(Bindings& bindings); + void SetupTextures(Bindings& bindings); +}; + +} // namespace Shader::Backend::GLSL diff --git a/src/shader_recompiler/backend/spirv/emit_context.cpp b/src/shader_recompiler/backend/spirv/emit_context.cpp deleted file mode 100644 index 723455462..000000000 --- a/src/shader_recompiler/backend/spirv/emit_context.cpp +++ /dev/null @@ -1,1585 +0,0 @@ -// Copyright 2021 yuzu Emulator Project -// Licensed under GPLv2 or any later version -// Refer to the license.txt file included. - -#include -#include -#include -#include - -#include - -#include - -#include "common/common_types.h" -#include "common/div_ceil.h" -#include "shader_recompiler/backend/spirv/emit_context.h" -#include "shader_recompiler/backend/spirv/emit_spirv.h" - -namespace Shader::Backend::SPIRV { -namespace { -constexpr size_t NUM_FIXEDFNCTEXTURE = 10; - -enum class Operation { - Increment, - Decrement, - FPAdd, - FPMin, - FPMax, -}; - -struct AttrInfo { - Id pointer; - Id id; - bool needs_cast; -}; - -Id ImageType(EmitContext& ctx, const TextureDescriptor& desc) { - const spv::ImageFormat format{spv::ImageFormat::Unknown}; - const Id type{ctx.F32[1]}; - const bool depth{desc.is_depth}; - switch (desc.type) { - case TextureType::Color1D: - return ctx.TypeImage(type, spv::Dim::Dim1D, depth, false, false, 1, format); - case TextureType::ColorArray1D: - return ctx.TypeImage(type, spv::Dim::Dim1D, depth, true, false, 1, format); - case TextureType::Color2D: - return ctx.TypeImage(type, spv::Dim::Dim2D, depth, false, false, 1, format); - case TextureType::ColorArray2D: - return ctx.TypeImage(type, spv::Dim::Dim2D, depth, true, false, 1, format); - case TextureType::Color3D: - return ctx.TypeImage(type, spv::Dim::Dim3D, depth, false, false, 1, format); - case TextureType::ColorCube: - return ctx.TypeImage(type, spv::Dim::Cube, depth, false, false, 1, format); - case TextureType::ColorArrayCube: - return ctx.TypeImage(type, spv::Dim::Cube, depth, true, false, 1, format); - case TextureType::Buffer: - break; - } - throw InvalidArgument("Invalid texture type {}", desc.type); -} - -spv::ImageFormat GetImageFormat(ImageFormat format) { - switch (format) { - case ImageFormat::Typeless: - return spv::ImageFormat::Unknown; - case ImageFormat::R8_UINT: - return spv::ImageFormat::R8ui; - case ImageFormat::R8_SINT: - return spv::ImageFormat::R8i; - case ImageFormat::R16_UINT: - return spv::ImageFormat::R16ui; - case ImageFormat::R16_SINT: - return spv::ImageFormat::R16i; - case ImageFormat::R32_UINT: - return spv::ImageFormat::R32ui; - case ImageFormat::R32G32_UINT: - return spv::ImageFormat::Rg32ui; - case ImageFormat::R32G32B32A32_UINT: - return spv::ImageFormat::Rgba32ui; - } - throw InvalidArgument("Invalid image format {}", format); -} - -Id ImageType(EmitContext& ctx, const ImageDescriptor& desc) { - const spv::ImageFormat format{GetImageFormat(desc.format)}; - const Id type{ctx.U32[1]}; - switch (desc.type) { - case TextureType::Color1D: - return ctx.TypeImage(type, spv::Dim::Dim1D, false, false, false, 2, format); - case TextureType::ColorArray1D: - return ctx.TypeImage(type, spv::Dim::Dim1D, false, true, false, 2, format); - case TextureType::Color2D: - return ctx.TypeImage(type, spv::Dim::Dim2D, false, false, false, 2, format); - case TextureType::ColorArray2D: - return ctx.TypeImage(type, spv::Dim::Dim2D, false, true, false, 2, format); - case TextureType::Color3D: - return ctx.TypeImage(type, spv::Dim::Dim3D, false, false, false, 2, format); - case TextureType::Buffer: - throw NotImplementedException("Image buffer"); - default: - break; - } - throw InvalidArgument("Invalid texture type {}", desc.type); -} - -Id DefineVariable(EmitContext& ctx, Id type, std::optional builtin, - spv::StorageClass storage_class) { - const Id pointer_type{ctx.TypePointer(storage_class, type)}; - const Id id{ctx.AddGlobalVariable(pointer_type, storage_class)}; - if (builtin) { - ctx.Decorate(id, spv::Decoration::BuiltIn, *builtin); - } - ctx.interfaces.push_back(id); - return id; -} - -u32 NumVertices(InputTopology input_topology) { - switch (input_topology) { - case InputTopology::Points: - return 1; - case InputTopology::Lines: - return 2; - case InputTopology::LinesAdjacency: - return 4; - case InputTopology::Triangles: - return 3; - case InputTopology::TrianglesAdjacency: - return 6; - } - throw InvalidArgument("Invalid input topology {}", input_topology); -} - -Id DefineInput(EmitContext& ctx, Id type, bool per_invocation, - std::optional builtin = std::nullopt) { - switch (ctx.stage) { - case Stage::TessellationControl: - case Stage::TessellationEval: - if (per_invocation) { - type = ctx.TypeArray(type, ctx.Const(32u)); - } - break; - case Stage::Geometry: - if (per_invocation) { - const u32 num_vertices{NumVertices(ctx.runtime_info.input_topology)}; - type = ctx.TypeArray(type, ctx.Const(num_vertices)); - } - break; - default: - break; - } - return DefineVariable(ctx, type, builtin, spv::StorageClass::Input); -} - -Id DefineOutput(EmitContext& ctx, Id type, std::optional invocations, - std::optional builtin = std::nullopt) { - if (invocations && ctx.stage == Stage::TessellationControl) { - type = ctx.TypeArray(type, ctx.Const(*invocations)); - } - return DefineVariable(ctx, type, builtin, spv::StorageClass::Output); -} - -void DefineGenericOutput(EmitContext& ctx, size_t index, std::optional invocations) { - static constexpr std::string_view swizzle{"xyzw"}; - const size_t base_attr_index{static_cast(IR::Attribute::Generic0X) + index * 4}; - u32 element{0}; - while (element < 4) { - const u32 remainder{4 - element}; - const TransformFeedbackVarying* xfb_varying{}; - if (!ctx.runtime_info.xfb_varyings.empty()) { - xfb_varying = &ctx.runtime_info.xfb_varyings[base_attr_index + element]; - xfb_varying = xfb_varying && xfb_varying->components > 0 ? xfb_varying : nullptr; - } - const u32 num_components{xfb_varying ? xfb_varying->components : remainder}; - - const Id id{DefineOutput(ctx, ctx.F32[num_components], invocations)}; - ctx.Decorate(id, spv::Decoration::Location, static_cast(index)); - if (element > 0) { - ctx.Decorate(id, spv::Decoration::Component, element); - } - if (xfb_varying) { - ctx.Decorate(id, spv::Decoration::XfbBuffer, xfb_varying->buffer); - ctx.Decorate(id, spv::Decoration::XfbStride, xfb_varying->stride); - ctx.Decorate(id, spv::Decoration::Offset, xfb_varying->offset); - } - if (num_components < 4 || element > 0) { - const std::string_view subswizzle{swizzle.substr(element, num_components)}; - ctx.Name(id, fmt::format("out_attr{}_{}", index, subswizzle)); - } else { - ctx.Name(id, fmt::format("out_attr{}", index)); - } - const GenericElementInfo info{ - .id = id, - .first_element = element, - .num_components = num_components, - }; - std::fill_n(ctx.output_generics[index].begin() + element, num_components, info); - element += num_components; - } -} - -Id GetAttributeType(EmitContext& ctx, AttributeType type) { - switch (type) { - case AttributeType::Float: - return ctx.F32[4]; - case AttributeType::SignedInt: - return ctx.TypeVector(ctx.TypeInt(32, true), 4); - case AttributeType::UnsignedInt: - return ctx.U32[4]; - case AttributeType::Disabled: - break; - } - throw InvalidArgument("Invalid attribute type {}", type); -} - -std::optional AttrTypes(EmitContext& ctx, u32 index) { - const AttributeType type{ctx.runtime_info.generic_input_types.at(index)}; - switch (type) { - case AttributeType::Float: - return AttrInfo{ctx.input_f32, ctx.F32[1], false}; - case AttributeType::UnsignedInt: - return AttrInfo{ctx.input_u32, ctx.U32[1], true}; - case AttributeType::SignedInt: - return AttrInfo{ctx.input_s32, ctx.TypeInt(32, true), true}; - case AttributeType::Disabled: - return std::nullopt; - } - throw InvalidArgument("Invalid attribute type {}", type); -} - -std::string_view StageName(Stage stage) { - switch (stage) { - case Stage::VertexA: - return "vs_a"; - case Stage::VertexB: - return "vs"; - case Stage::TessellationControl: - return "tcs"; - case Stage::TessellationEval: - return "tes"; - case Stage::Geometry: - return "gs"; - case Stage::Fragment: - return "fs"; - case Stage::Compute: - return "cs"; - } - throw InvalidArgument("Invalid stage {}", stage); -} - -template -void Name(EmitContext& ctx, Id object, std::string_view format_str, Args&&... args) { - ctx.Name(object, fmt::format(fmt::runtime(format_str), StageName(ctx.stage), - std::forward(args)...) - .c_str()); -} - -void DefineConstBuffers(EmitContext& ctx, const Info& info, Id UniformDefinitions::*member_type, - u32 binding, Id type, char type_char, u32 element_size) { - const Id array_type{ctx.TypeArray(type, ctx.Const(65536U / element_size))}; - ctx.Decorate(array_type, spv::Decoration::ArrayStride, element_size); - - const Id struct_type{ctx.TypeStruct(array_type)}; - Name(ctx, struct_type, "{}_cbuf_block_{}{}", ctx.stage, type_char, element_size * CHAR_BIT); - ctx.Decorate(struct_type, spv::Decoration::Block); - ctx.MemberName(struct_type, 0, "data"); - ctx.MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U); - - const Id struct_pointer_type{ctx.TypePointer(spv::StorageClass::Uniform, struct_type)}; - const Id uniform_type{ctx.TypePointer(spv::StorageClass::Uniform, type)}; - ctx.uniform_types.*member_type = uniform_type; - - for (const ConstantBufferDescriptor& desc : info.constant_buffer_descriptors) { - const Id id{ctx.AddGlobalVariable(struct_pointer_type, spv::StorageClass::Uniform)}; - ctx.Decorate(id, spv::Decoration::Binding, binding); - ctx.Decorate(id, spv::Decoration::DescriptorSet, 0U); - ctx.Name(id, fmt::format("c{}", desc.index)); - for (size_t i = 0; i < desc.count; ++i) { - ctx.cbufs[desc.index + i].*member_type = id; - } - if (ctx.profile.supported_spirv >= 0x00010400) { - ctx.interfaces.push_back(id); - } - binding += desc.count; - } -} - -void DefineSsbos(EmitContext& ctx, StorageTypeDefinition& type_def, - Id StorageDefinitions::*member_type, const Info& info, u32 binding, Id type, - u32 stride) { - const Id array_type{ctx.TypeRuntimeArray(type)}; - ctx.Decorate(array_type, spv::Decoration::ArrayStride, stride); - - const Id struct_type{ctx.TypeStruct(array_type)}; - ctx.Decorate(struct_type, spv::Decoration::Block); - ctx.MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U); - - const Id struct_pointer{ctx.TypePointer(spv::StorageClass::StorageBuffer, struct_type)}; - type_def.array = struct_pointer; - type_def.element = ctx.TypePointer(spv::StorageClass::StorageBuffer, type); - - u32 index{}; - for (const StorageBufferDescriptor& desc : info.storage_buffers_descriptors) { - const Id id{ctx.AddGlobalVariable(struct_pointer, spv::StorageClass::StorageBuffer)}; - ctx.Decorate(id, spv::Decoration::Binding, binding); - ctx.Decorate(id, spv::Decoration::DescriptorSet, 0U); - ctx.Name(id, fmt::format("ssbo{}", index)); - if (ctx.profile.supported_spirv >= 0x00010400) { - ctx.interfaces.push_back(id); - } - for (size_t i = 0; i < desc.count; ++i) { - ctx.ssbos[index + i].*member_type = id; - } - index += desc.count; - binding += desc.count; - } -} - -Id CasFunction(EmitContext& ctx, Operation operation, Id value_type) { - const Id func_type{ctx.TypeFunction(value_type, value_type, value_type)}; - const Id func{ctx.OpFunction(value_type, spv::FunctionControlMask::MaskNone, func_type)}; - const Id op_a{ctx.OpFunctionParameter(value_type)}; - const Id op_b{ctx.OpFunctionParameter(value_type)}; - ctx.AddLabel(); - Id result{}; - switch (operation) { - case Operation::Increment: { - const Id pred{ctx.OpUGreaterThanEqual(ctx.U1, op_a, op_b)}; - const Id incr{ctx.OpIAdd(value_type, op_a, ctx.Constant(value_type, 1))}; - result = ctx.OpSelect(value_type, pred, ctx.u32_zero_value, incr); - break; - } - case Operation::Decrement: { - const Id lhs{ctx.OpIEqual(ctx.U1, op_a, ctx.Constant(value_type, 0u))}; - const Id rhs{ctx.OpUGreaterThan(ctx.U1, op_a, op_b)}; - const Id pred{ctx.OpLogicalOr(ctx.U1, lhs, rhs)}; - const Id decr{ctx.OpISub(value_type, op_a, ctx.Constant(value_type, 1))}; - result = ctx.OpSelect(value_type, pred, op_b, decr); - break; - } - case Operation::FPAdd: - result = ctx.OpFAdd(value_type, op_a, op_b); - break; - case Operation::FPMin: - result = ctx.OpFMin(value_type, op_a, op_b); - break; - case Operation::FPMax: - result = ctx.OpFMax(value_type, op_a, op_b); - break; - default: - break; - } - ctx.OpReturnValue(result); - ctx.OpFunctionEnd(); - return func; -} - -Id CasLoop(EmitContext& ctx, Operation operation, Id array_pointer, Id element_pointer, - Id value_type, Id memory_type, spv::Scope scope) { - const bool is_shared{scope == spv::Scope::Workgroup}; - const bool is_struct{!is_shared || ctx.profile.support_explicit_workgroup_layout}; - const Id cas_func{CasFunction(ctx, operation, value_type)}; - const Id zero{ctx.u32_zero_value}; - const Id scope_id{ctx.Const(static_cast(scope))}; - - const Id loop_header{ctx.OpLabel()}; - const Id continue_block{ctx.OpLabel()}; - const Id merge_block{ctx.OpLabel()}; - const Id func_type{is_shared - ? ctx.TypeFunction(value_type, ctx.U32[1], value_type) - : ctx.TypeFunction(value_type, ctx.U32[1], value_type, array_pointer)}; - - const Id func{ctx.OpFunction(value_type, spv::FunctionControlMask::MaskNone, func_type)}; - const Id index{ctx.OpFunctionParameter(ctx.U32[1])}; - const Id op_b{ctx.OpFunctionParameter(value_type)}; - const Id base{is_shared ? ctx.shared_memory_u32 : ctx.OpFunctionParameter(array_pointer)}; - ctx.AddLabel(); - ctx.OpBranch(loop_header); - ctx.AddLabel(loop_header); - - ctx.OpLoopMerge(merge_block, continue_block, spv::LoopControlMask::MaskNone); - ctx.OpBranch(continue_block); - - ctx.AddLabel(continue_block); - const Id word_pointer{is_struct ? ctx.OpAccessChain(element_pointer, base, zero, index) - : ctx.OpAccessChain(element_pointer, base, index)}; - if (value_type.value == ctx.F32[2].value) { - const Id u32_value{ctx.OpLoad(ctx.U32[1], word_pointer)}; - const Id value{ctx.OpUnpackHalf2x16(ctx.F32[2], u32_value)}; - const Id new_value{ctx.OpFunctionCall(value_type, cas_func, value, op_b)}; - const Id u32_new_value{ctx.OpPackHalf2x16(ctx.U32[1], new_value)}; - const Id atomic_res{ctx.OpAtomicCompareExchange(ctx.U32[1], word_pointer, scope_id, zero, - zero, u32_new_value, u32_value)}; - const Id success{ctx.OpIEqual(ctx.U1, atomic_res, u32_value)}; - ctx.OpBranchConditional(success, merge_block, loop_header); - - ctx.AddLabel(merge_block); - ctx.OpReturnValue(ctx.OpUnpackHalf2x16(ctx.F32[2], atomic_res)); - } else { - const Id value{ctx.OpLoad(memory_type, word_pointer)}; - const bool matching_type{value_type.value == memory_type.value}; - const Id bitcast_value{matching_type ? value : ctx.OpBitcast(value_type, value)}; - const Id cal_res{ctx.OpFunctionCall(value_type, cas_func, bitcast_value, op_b)}; - const Id new_value{matching_type ? cal_res : ctx.OpBitcast(memory_type, cal_res)}; - const Id atomic_res{ctx.OpAtomicCompareExchange(ctx.U32[1], word_pointer, scope_id, zero, - zero, new_value, value)}; - const Id success{ctx.OpIEqual(ctx.U1, atomic_res, value)}; - ctx.OpBranchConditional(success, merge_block, loop_header); - - ctx.AddLabel(merge_block); - ctx.OpReturnValue(ctx.OpBitcast(value_type, atomic_res)); - } - ctx.OpFunctionEnd(); - return func; -} - -template -std::string NameOf(Stage stage, const Desc& desc, std::string_view prefix) { - if (desc.count > 1) { - return fmt::format("{}_{}{}_{:02x}x{}", StageName(stage), prefix, desc.cbuf_index, - desc.cbuf_offset, desc.count); - } else { - return fmt::format("{}_{}{}_{:02x}", StageName(stage), prefix, desc.cbuf_index, - desc.cbuf_offset); - } -} - -Id DescType(EmitContext& ctx, Id sampled_type, Id pointer_type, u32 count) { - if (count > 1) { - const Id array_type{ctx.TypeArray(sampled_type, ctx.Const(count))}; - return ctx.TypePointer(spv::StorageClass::UniformConstant, array_type); - } else { - return pointer_type; - } -} - -size_t FindAndSetNextUnusedLocation(std::bitset& used_locations, - size_t& start_offset) { - for (size_t location = start_offset; location < used_locations.size(); ++location) { - if (!used_locations.test(location)) { - start_offset = location; - used_locations.set(location); - return location; - } - } - throw RuntimeError("Unable to get an unused location for legacy attribute"); -} - -Id DefineLegacyInput(EmitContext& ctx, std::bitset& used_locations, - size_t& start_offset) { - const Id id{DefineInput(ctx, ctx.F32[4], true)}; - const size_t location = FindAndSetNextUnusedLocation(used_locations, start_offset); - ctx.Decorate(id, spv::Decoration::Location, location); - return id; -} - -Id DefineLegacyOutput(EmitContext& ctx, std::bitset& used_locations, - size_t& start_offset, std::optional invocations) { - const Id id{DefineOutput(ctx, ctx.F32[4], invocations)}; - const size_t location = FindAndSetNextUnusedLocation(used_locations, start_offset); - ctx.Decorate(id, spv::Decoration::Location, location); - return id; -} -} // Anonymous namespace - -void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_view name) { - defs[0] = sirit_ctx.Name(base_type, name); - - std::array def_name; - for (int i = 1; i < 4; ++i) { - const std::string_view def_name_view( - def_name.data(), - fmt::format_to_n(def_name.data(), def_name.size(), "{}x{}", name, i + 1).size); - defs[static_cast(i)] = - sirit_ctx.Name(sirit_ctx.TypeVector(base_type, i + 1), def_name_view); - } -} - -EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_info_, - IR::Program& program, Bindings& bindings) - : Sirit::Module(profile_.supported_spirv), profile{profile_}, runtime_info{runtime_info_}, - stage{program.stage}, texture_rescaling_index{bindings.texture_scaling_index}, - image_rescaling_index{bindings.image_scaling_index} { - const bool is_unified{profile.unified_descriptor_binding}; - u32& uniform_binding{is_unified ? bindings.unified : bindings.uniform_buffer}; - u32& storage_binding{is_unified ? bindings.unified : bindings.storage_buffer}; - u32& texture_binding{is_unified ? bindings.unified : bindings.texture}; - u32& image_binding{is_unified ? bindings.unified : bindings.image}; - AddCapability(spv::Capability::Shader); - DefineCommonTypes(program.info); - DefineCommonConstants(); - DefineInterfaces(program); - DefineLocalMemory(program); - DefineSharedMemory(program); - DefineSharedMemoryFunctions(program); - DefineConstantBuffers(program.info, uniform_binding); - DefineStorageBuffers(program.info, storage_binding); - DefineTextureBuffers(program.info, texture_binding); - DefineImageBuffers(program.info, image_binding); - DefineTextures(program.info, texture_binding, bindings.texture_scaling_index); - DefineImages(program.info, image_binding, bindings.image_scaling_index); - DefineAttributeMemAccess(program.info); - DefineGlobalMemoryFunctions(program.info); - DefineRescalingInput(program.info); -} - -EmitContext::~EmitContext() = default; - -Id EmitContext::Def(const IR::Value& value) { - if (!value.IsImmediate()) { - return value.InstRecursive()->Definition(); - } - switch (value.Type()) { - case IR::Type::Void: - // Void instructions are used for optional arguments (e.g. texture offsets) - // They are not meant to be used in the SPIR-V module - return Id{}; - case IR::Type::U1: - return value.U1() ? true_value : false_value; - case IR::Type::U32: - return Const(value.U32()); - case IR::Type::U64: - return Constant(U64, value.U64()); - case IR::Type::F32: - return Const(value.F32()); - case IR::Type::F64: - return Constant(F64[1], value.F64()); - default: - throw NotImplementedException("Immediate type {}", value.Type()); - } -} - -Id EmitContext::BitOffset8(const IR::Value& offset) { - if (offset.IsImmediate()) { - return Const((offset.U32() % 4) * 8); - } - return OpBitwiseAnd(U32[1], OpShiftLeftLogical(U32[1], Def(offset), Const(3u)), Const(24u)); -} - -Id EmitContext::BitOffset16(const IR::Value& offset) { - if (offset.IsImmediate()) { - return Const(((offset.U32() / 2) % 2) * 16); - } - return OpBitwiseAnd(U32[1], OpShiftLeftLogical(U32[1], Def(offset), Const(3u)), Const(16u)); -} - -Id EmitContext::InputLegacyAttribute(IR::Attribute attribute) { - if (attribute >= IR::Attribute::ColorFrontDiffuseR && - attribute <= IR::Attribute::ColorFrontDiffuseA) { - return input_front_color; - } - if (attribute >= IR::Attribute::ColorFrontSpecularR && - attribute <= IR::Attribute::ColorFrontSpecularA) { - return input_front_secondary_color; - } - if (attribute >= IR::Attribute::ColorBackDiffuseR && - attribute <= IR::Attribute::ColorBackDiffuseA) { - return input_back_color; - } - if (attribute >= IR::Attribute::ColorBackSpecularR && - attribute <= IR::Attribute::ColorBackSpecularA) { - return input_back_secondary_color; - } - if (attribute == IR::Attribute::FogCoordinate) { - return input_fog_frag_coord; - } - if (attribute >= IR::Attribute::FixedFncTexture0S && - attribute <= IR::Attribute::FixedFncTexture9Q) { - u32 index = - (static_cast(attribute) - static_cast(IR::Attribute::FixedFncTexture0S)) / 4; - return input_fixed_fnc_textures[index]; - } - throw InvalidArgument("Attribute is not legacy attribute {}", attribute); -} - -Id EmitContext::OutputLegacyAttribute(IR::Attribute attribute) { - if (attribute >= IR::Attribute::ColorFrontDiffuseR && - attribute <= IR::Attribute::ColorFrontDiffuseA) { - return output_front_color; - } - if (attribute >= IR::Attribute::ColorFrontSpecularR && - attribute <= IR::Attribute::ColorFrontSpecularA) { - return output_front_secondary_color; - } - if (attribute >= IR::Attribute::ColorBackDiffuseR && - attribute <= IR::Attribute::ColorBackDiffuseA) { - return output_back_color; - } - if (attribute >= IR::Attribute::ColorBackSpecularR && - attribute <= IR::Attribute::ColorBackSpecularA) { - return output_back_secondary_color; - } - if (attribute == IR::Attribute::FogCoordinate) { - return output_fog_frag_coord; - } - if (attribute >= IR::Attribute::FixedFncTexture0S && - attribute <= IR::Attribute::FixedFncTexture9Q) { - u32 index = - (static_cast(attribute) - static_cast(IR::Attribute::FixedFncTexture0S)) / 4; - return output_fixed_fnc_textures[index]; - } - throw InvalidArgument("Attribute is not legacy attribute {}", attribute); -} - -void EmitContext::DefineCommonTypes(const Info& info) { - void_id = TypeVoid(); - - U1 = Name(TypeBool(), "u1"); - - F32.Define(*this, TypeFloat(32), "f32"); - U32.Define(*this, TypeInt(32, false), "u32"); - S32.Define(*this, TypeInt(32, true), "s32"); - - private_u32 = Name(TypePointer(spv::StorageClass::Private, U32[1]), "private_u32"); - - input_f32 = Name(TypePointer(spv::StorageClass::Input, F32[1]), "input_f32"); - input_u32 = Name(TypePointer(spv::StorageClass::Input, U32[1]), "input_u32"); - input_s32 = Name(TypePointer(spv::StorageClass::Input, TypeInt(32, true)), "input_s32"); - - output_f32 = Name(TypePointer(spv::StorageClass::Output, F32[1]), "output_f32"); - output_u32 = Name(TypePointer(spv::StorageClass::Output, U32[1]), "output_u32"); - - if (info.uses_int8 && profile.support_int8) { - AddCapability(spv::Capability::Int8); - U8 = Name(TypeInt(8, false), "u8"); - S8 = Name(TypeInt(8, true), "s8"); - } - if (info.uses_int16 && profile.support_int16) { - AddCapability(spv::Capability::Int16); - U16 = Name(TypeInt(16, false), "u16"); - S16 = Name(TypeInt(16, true), "s16"); - } - if (info.uses_int64) { - AddCapability(spv::Capability::Int64); - U64 = Name(TypeInt(64, false), "u64"); - } - if (info.uses_fp16) { - AddCapability(spv::Capability::Float16); - F16.Define(*this, TypeFloat(16), "f16"); - } - if (info.uses_fp64) { - AddCapability(spv::Capability::Float64); - F64.Define(*this, TypeFloat(64), "f64"); - } -} - -void EmitContext::DefineCommonConstants() { - true_value = ConstantTrue(U1); - false_value = ConstantFalse(U1); - u32_zero_value = Const(0U); - f32_zero_value = Const(0.0f); -} - -void EmitContext::DefineInterfaces(const IR::Program& program) { - DefineInputs(program); - DefineOutputs(program); -} - -void EmitContext::DefineLocalMemory(const IR::Program& program) { - if (program.local_memory_size == 0) { - return; - } - const u32 num_elements{Common::DivCeil(program.local_memory_size, 4U)}; - const Id type{TypeArray(U32[1], Const(num_elements))}; - const Id pointer{TypePointer(spv::StorageClass::Private, type)}; - local_memory = AddGlobalVariable(pointer, spv::StorageClass::Private); - if (profile.supported_spirv >= 0x00010400) { - interfaces.push_back(local_memory); - } -} - -void EmitContext::DefineSharedMemory(const IR::Program& program) { - if (program.shared_memory_size == 0) { - return; - } - const auto make{[&](Id element_type, u32 element_size) { - const u32 num_elements{Common::DivCeil(program.shared_memory_size, element_size)}; - const Id array_type{TypeArray(element_type, Const(num_elements))}; - Decorate(array_type, spv::Decoration::ArrayStride, element_size); - - const Id struct_type{TypeStruct(array_type)}; - MemberDecorate(struct_type, 0U, spv::Decoration::Offset, 0U); - Decorate(struct_type, spv::Decoration::Block); - - const Id pointer{TypePointer(spv::StorageClass::Workgroup, struct_type)}; - const Id element_pointer{TypePointer(spv::StorageClass::Workgroup, element_type)}; - const Id variable{AddGlobalVariable(pointer, spv::StorageClass::Workgroup)}; - Decorate(variable, spv::Decoration::Aliased); - interfaces.push_back(variable); - - return std::make_tuple(variable, element_pointer, pointer); - }}; - if (profile.support_explicit_workgroup_layout) { - AddExtension("SPV_KHR_workgroup_memory_explicit_layout"); - AddCapability(spv::Capability::WorkgroupMemoryExplicitLayoutKHR); - if (program.info.uses_int8) { - AddCapability(spv::Capability::WorkgroupMemoryExplicitLayout8BitAccessKHR); - std::tie(shared_memory_u8, shared_u8, std::ignore) = make(U8, 1); - } - if (program.info.uses_int16) { - AddCapability(spv::Capability::WorkgroupMemoryExplicitLayout16BitAccessKHR); - std::tie(shared_memory_u16, shared_u16, std::ignore) = make(U16, 2); - } - if (program.info.uses_int64) { - std::tie(shared_memory_u64, shared_u64, std::ignore) = make(U64, 8); - } - std::tie(shared_memory_u32, shared_u32, shared_memory_u32_type) = make(U32[1], 4); - std::tie(shared_memory_u32x2, shared_u32x2, std::ignore) = make(U32[2], 8); - std::tie(shared_memory_u32x4, shared_u32x4, std::ignore) = make(U32[4], 16); - return; - } - const u32 num_elements{Common::DivCeil(program.shared_memory_size, 4U)}; - const Id type{TypeArray(U32[1], Const(num_elements))}; - shared_memory_u32_type = TypePointer(spv::StorageClass::Workgroup, type); - - shared_u32 = TypePointer(spv::StorageClass::Workgroup, U32[1]); - shared_memory_u32 = AddGlobalVariable(shared_memory_u32_type, spv::StorageClass::Workgroup); - interfaces.push_back(shared_memory_u32); - - const Id func_type{TypeFunction(void_id, U32[1], U32[1])}; - const auto make_function{[&](u32 mask, u32 size) { - const Id loop_header{OpLabel()}; - const Id continue_block{OpLabel()}; - const Id merge_block{OpLabel()}; - - const Id func{OpFunction(void_id, spv::FunctionControlMask::MaskNone, func_type)}; - const Id offset{OpFunctionParameter(U32[1])}; - const Id insert_value{OpFunctionParameter(U32[1])}; - AddLabel(); - OpBranch(loop_header); - - AddLabel(loop_header); - const Id word_offset{OpShiftRightArithmetic(U32[1], offset, Const(2U))}; - const Id shift_offset{OpShiftLeftLogical(U32[1], offset, Const(3U))}; - const Id bit_offset{OpBitwiseAnd(U32[1], shift_offset, Const(mask))}; - const Id count{Const(size)}; - OpLoopMerge(merge_block, continue_block, spv::LoopControlMask::MaskNone); - OpBranch(continue_block); - - AddLabel(continue_block); - const Id word_pointer{OpAccessChain(shared_u32, shared_memory_u32, word_offset)}; - const Id old_value{OpLoad(U32[1], word_pointer)}; - const Id new_value{OpBitFieldInsert(U32[1], old_value, insert_value, bit_offset, count)}; - const Id atomic_res{OpAtomicCompareExchange(U32[1], word_pointer, Const(1U), u32_zero_value, - u32_zero_value, new_value, old_value)}; - const Id success{OpIEqual(U1, atomic_res, old_value)}; - OpBranchConditional(success, merge_block, loop_header); - - AddLabel(merge_block); - OpReturn(); - OpFunctionEnd(); - return func; - }}; - if (program.info.uses_int8) { - shared_store_u8_func = make_function(24, 8); - } - if (program.info.uses_int16) { - shared_store_u16_func = make_function(16, 16); - } -} - -void EmitContext::DefineSharedMemoryFunctions(const IR::Program& program) { - if (program.info.uses_shared_increment) { - increment_cas_shared = CasLoop(*this, Operation::Increment, shared_memory_u32_type, - shared_u32, U32[1], U32[1], spv::Scope::Workgroup); - } - if (program.info.uses_shared_decrement) { - decrement_cas_shared = CasLoop(*this, Operation::Decrement, shared_memory_u32_type, - shared_u32, U32[1], U32[1], spv::Scope::Workgroup); - } -} - -void EmitContext::DefineAttributeMemAccess(const Info& info) { - const auto make_load{[&] { - const bool is_array{stage == Stage::Geometry}; - const Id end_block{OpLabel()}; - const Id default_label{OpLabel()}; - - const Id func_type_load{is_array ? TypeFunction(F32[1], U32[1], U32[1]) - : TypeFunction(F32[1], U32[1])}; - const Id func{OpFunction(F32[1], spv::FunctionControlMask::MaskNone, func_type_load)}; - const Id offset{OpFunctionParameter(U32[1])}; - const Id vertex{is_array ? OpFunctionParameter(U32[1]) : Id{}}; - - AddLabel(); - const Id base_index{OpShiftRightArithmetic(U32[1], offset, Const(2U))}; - const Id masked_index{OpBitwiseAnd(U32[1], base_index, Const(3U))}; - const Id compare_index{OpShiftRightArithmetic(U32[1], base_index, Const(2U))}; - std::vector literals; - std::vector labels; - if (info.loads.AnyComponent(IR::Attribute::PositionX)) { - literals.push_back(static_cast(IR::Attribute::PositionX) >> 2); - labels.push_back(OpLabel()); - } - const u32 base_attribute_value = static_cast(IR::Attribute::Generic0X) >> 2; - for (u32 index = 0; index < static_cast(IR::NUM_GENERICS); ++index) { - if (!info.loads.Generic(index)) { - continue; - } - literals.push_back(base_attribute_value + index); - labels.push_back(OpLabel()); - } - OpSelectionMerge(end_block, spv::SelectionControlMask::MaskNone); - OpSwitch(compare_index, default_label, literals, labels); - AddLabel(default_label); - OpReturnValue(Const(0.0f)); - size_t label_index{0}; - if (info.loads.AnyComponent(IR::Attribute::PositionX)) { - AddLabel(labels[label_index]); - const Id pointer{is_array - ? OpAccessChain(input_f32, input_position, vertex, masked_index) - : OpAccessChain(input_f32, input_position, masked_index)}; - const Id result{OpLoad(F32[1], pointer)}; - OpReturnValue(result); - ++label_index; - } - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - if (!info.loads.Generic(index)) { - continue; - } - AddLabel(labels[label_index]); - const auto type{AttrTypes(*this, static_cast(index))}; - if (!type) { - OpReturnValue(Const(0.0f)); - ++label_index; - continue; - } - const Id generic_id{input_generics.at(index)}; - const Id pointer{is_array - ? OpAccessChain(type->pointer, generic_id, vertex, masked_index) - : OpAccessChain(type->pointer, generic_id, masked_index)}; - const Id value{OpLoad(type->id, pointer)}; - const Id result{type->needs_cast ? OpBitcast(F32[1], value) : value}; - OpReturnValue(result); - ++label_index; - } - AddLabel(end_block); - OpUnreachable(); - OpFunctionEnd(); - return func; - }}; - const auto make_store{[&] { - const Id end_block{OpLabel()}; - const Id default_label{OpLabel()}; - - const Id func_type_store{TypeFunction(void_id, U32[1], F32[1])}; - const Id func{OpFunction(void_id, spv::FunctionControlMask::MaskNone, func_type_store)}; - const Id offset{OpFunctionParameter(U32[1])}; - const Id store_value{OpFunctionParameter(F32[1])}; - AddLabel(); - const Id base_index{OpShiftRightArithmetic(U32[1], offset, Const(2U))}; - const Id masked_index{OpBitwiseAnd(U32[1], base_index, Const(3U))}; - const Id compare_index{OpShiftRightArithmetic(U32[1], base_index, Const(2U))}; - std::vector literals; - std::vector labels; - if (info.stores.AnyComponent(IR::Attribute::PositionX)) { - literals.push_back(static_cast(IR::Attribute::PositionX) >> 2); - labels.push_back(OpLabel()); - } - const u32 base_attribute_value = static_cast(IR::Attribute::Generic0X) >> 2; - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - if (!info.stores.Generic(index)) { - continue; - } - literals.push_back(base_attribute_value + static_cast(index)); - labels.push_back(OpLabel()); - } - if (info.stores.ClipDistances()) { - literals.push_back(static_cast(IR::Attribute::ClipDistance0) >> 2); - labels.push_back(OpLabel()); - literals.push_back(static_cast(IR::Attribute::ClipDistance4) >> 2); - labels.push_back(OpLabel()); - } - OpSelectionMerge(end_block, spv::SelectionControlMask::MaskNone); - OpSwitch(compare_index, default_label, literals, labels); - AddLabel(default_label); - OpReturn(); - size_t label_index{0}; - if (info.stores.AnyComponent(IR::Attribute::PositionX)) { - AddLabel(labels[label_index]); - const Id pointer{OpAccessChain(output_f32, output_position, masked_index)}; - OpStore(pointer, store_value); - OpReturn(); - ++label_index; - } - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - if (!info.stores.Generic(index)) { - continue; - } - if (output_generics[index][0].num_components != 4) { - throw NotImplementedException("Physical stores and transform feedbacks"); - } - AddLabel(labels[label_index]); - const Id generic_id{output_generics[index][0].id}; - const Id pointer{OpAccessChain(output_f32, generic_id, masked_index)}; - OpStore(pointer, store_value); - OpReturn(); - ++label_index; - } - if (info.stores.ClipDistances()) { - AddLabel(labels[label_index]); - const Id pointer{OpAccessChain(output_f32, clip_distances, masked_index)}; - OpStore(pointer, store_value); - OpReturn(); - ++label_index; - AddLabel(labels[label_index]); - const Id fixed_index{OpIAdd(U32[1], masked_index, Const(4U))}; - const Id pointer2{OpAccessChain(output_f32, clip_distances, fixed_index)}; - OpStore(pointer2, store_value); - OpReturn(); - ++label_index; - } - AddLabel(end_block); - OpUnreachable(); - OpFunctionEnd(); - return func; - }}; - if (info.loads_indexed_attributes) { - indexed_load_func = make_load(); - } - if (info.stores_indexed_attributes) { - indexed_store_func = make_store(); - } -} - -void EmitContext::DefineGlobalMemoryFunctions(const Info& info) { - if (!info.uses_global_memory || !profile.support_int64) { - return; - } - using DefPtr = Id StorageDefinitions::*; - const Id zero{u32_zero_value}; - const auto define_body{[&](DefPtr ssbo_member, Id addr, Id element_pointer, u32 shift, - auto&& callback) { - AddLabel(); - const size_t num_buffers{info.storage_buffers_descriptors.size()}; - for (size_t index = 0; index < num_buffers; ++index) { - if (!info.nvn_buffer_used[index]) { - continue; - } - const auto& ssbo{info.storage_buffers_descriptors[index]}; - const Id ssbo_addr_cbuf_offset{Const(ssbo.cbuf_offset / 8)}; - const Id ssbo_size_cbuf_offset{Const(ssbo.cbuf_offset / 4 + 2)}; - const Id ssbo_addr_pointer{OpAccessChain( - uniform_types.U32x2, cbufs[ssbo.cbuf_index].U32x2, zero, ssbo_addr_cbuf_offset)}; - const Id ssbo_size_pointer{OpAccessChain(uniform_types.U32, cbufs[ssbo.cbuf_index].U32, - zero, ssbo_size_cbuf_offset)}; - - const Id ssbo_addr{OpBitcast(U64, OpLoad(U32[2], ssbo_addr_pointer))}; - const Id ssbo_size{OpUConvert(U64, OpLoad(U32[1], ssbo_size_pointer))}; - const Id ssbo_end{OpIAdd(U64, ssbo_addr, ssbo_size)}; - const Id cond{OpLogicalAnd(U1, OpUGreaterThanEqual(U1, addr, ssbo_addr), - OpULessThan(U1, addr, ssbo_end))}; - const Id then_label{OpLabel()}; - const Id else_label{OpLabel()}; - OpSelectionMerge(else_label, spv::SelectionControlMask::MaskNone); - OpBranchConditional(cond, then_label, else_label); - AddLabel(then_label); - const Id ssbo_id{ssbos[index].*ssbo_member}; - const Id ssbo_offset{OpUConvert(U32[1], OpISub(U64, addr, ssbo_addr))}; - const Id ssbo_index{OpShiftRightLogical(U32[1], ssbo_offset, Const(shift))}; - const Id ssbo_pointer{OpAccessChain(element_pointer, ssbo_id, zero, ssbo_index)}; - callback(ssbo_pointer); - AddLabel(else_label); - } - }}; - const auto define_load{[&](DefPtr ssbo_member, Id element_pointer, Id type, u32 shift) { - const Id function_type{TypeFunction(type, U64)}; - const Id func_id{OpFunction(type, spv::FunctionControlMask::MaskNone, function_type)}; - const Id addr{OpFunctionParameter(U64)}; - define_body(ssbo_member, addr, element_pointer, shift, - [&](Id ssbo_pointer) { OpReturnValue(OpLoad(type, ssbo_pointer)); }); - OpReturnValue(ConstantNull(type)); - OpFunctionEnd(); - return func_id; - }}; - const auto define_write{[&](DefPtr ssbo_member, Id element_pointer, Id type, u32 shift) { - const Id function_type{TypeFunction(void_id, U64, type)}; - const Id func_id{OpFunction(void_id, spv::FunctionControlMask::MaskNone, function_type)}; - const Id addr{OpFunctionParameter(U64)}; - const Id data{OpFunctionParameter(type)}; - define_body(ssbo_member, addr, element_pointer, shift, [&](Id ssbo_pointer) { - OpStore(ssbo_pointer, data); - OpReturn(); - }); - OpReturn(); - OpFunctionEnd(); - return func_id; - }}; - const auto define{ - [&](DefPtr ssbo_member, const StorageTypeDefinition& type_def, Id type, size_t size) { - const Id element_type{type_def.element}; - const u32 shift{static_cast(std::countr_zero(size))}; - const Id load_func{define_load(ssbo_member, element_type, type, shift)}; - const Id write_func{define_write(ssbo_member, element_type, type, shift)}; - return std::make_pair(load_func, write_func); - }}; - std::tie(load_global_func_u32, write_global_func_u32) = - define(&StorageDefinitions::U32, storage_types.U32, U32[1], sizeof(u32)); - std::tie(load_global_func_u32x2, write_global_func_u32x2) = - define(&StorageDefinitions::U32x2, storage_types.U32x2, U32[2], sizeof(u32[2])); - std::tie(load_global_func_u32x4, write_global_func_u32x4) = - define(&StorageDefinitions::U32x4, storage_types.U32x4, U32[4], sizeof(u32[4])); -} - -void EmitContext::DefineRescalingInput(const Info& info) { - if (!info.uses_rescaling_uniform) { - return; - } - if (profile.unified_descriptor_binding) { - DefineRescalingInputPushConstant(); - } else { - DefineRescalingInputUniformConstant(); - } -} - -void EmitContext::DefineRescalingInputPushConstant() { - boost::container::static_vector members{}; - u32 member_index{0}; - - rescaling_textures_type = TypeArray(U32[1], Const(4u)); - Decorate(rescaling_textures_type, spv::Decoration::ArrayStride, 4u); - members.push_back(rescaling_textures_type); - rescaling_textures_member_index = member_index++; - - rescaling_images_type = TypeArray(U32[1], Const(NUM_IMAGE_SCALING_WORDS)); - Decorate(rescaling_images_type, spv::Decoration::ArrayStride, 4u); - members.push_back(rescaling_images_type); - rescaling_images_member_index = member_index++; - - if (stage != Stage::Compute) { - members.push_back(F32[1]); - rescaling_downfactor_member_index = member_index++; - } - const Id push_constant_struct{TypeStruct(std::span(members.data(), members.size()))}; - Decorate(push_constant_struct, spv::Decoration::Block); - Name(push_constant_struct, "ResolutionInfo"); - - MemberDecorate(push_constant_struct, rescaling_textures_member_index, spv::Decoration::Offset, - static_cast(offsetof(RescalingLayout, rescaling_textures))); - MemberName(push_constant_struct, rescaling_textures_member_index, "rescaling_textures"); - - MemberDecorate(push_constant_struct, rescaling_images_member_index, spv::Decoration::Offset, - static_cast(offsetof(RescalingLayout, rescaling_images))); - MemberName(push_constant_struct, rescaling_images_member_index, "rescaling_images"); - - if (stage != Stage::Compute) { - MemberDecorate(push_constant_struct, rescaling_downfactor_member_index, - spv::Decoration::Offset, - static_cast(offsetof(RescalingLayout, down_factor))); - MemberName(push_constant_struct, rescaling_downfactor_member_index, "down_factor"); - } - const Id pointer_type{TypePointer(spv::StorageClass::PushConstant, push_constant_struct)}; - rescaling_push_constants = AddGlobalVariable(pointer_type, spv::StorageClass::PushConstant); - Name(rescaling_push_constants, "rescaling_push_constants"); - - if (profile.supported_spirv >= 0x00010400) { - interfaces.push_back(rescaling_push_constants); - } -} - -void EmitContext::DefineRescalingInputUniformConstant() { - const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, F32[4])}; - rescaling_uniform_constant = - AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant); - Decorate(rescaling_uniform_constant, spv::Decoration::Location, 0u); - - if (profile.supported_spirv >= 0x00010400) { - interfaces.push_back(rescaling_uniform_constant); - } -} - -void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) { - if (info.constant_buffer_descriptors.empty()) { - return; - } - if (!profile.support_descriptor_aliasing) { - DefineConstBuffers(*this, info, &UniformDefinitions::U32x4, binding, U32[4], 'u', - sizeof(u32[4])); - for (const ConstantBufferDescriptor& desc : info.constant_buffer_descriptors) { - binding += desc.count; - } - return; - } - IR::Type types{info.used_constant_buffer_types}; - if (True(types & IR::Type::U8)) { - if (profile.support_int8) { - DefineConstBuffers(*this, info, &UniformDefinitions::U8, binding, U8, 'u', sizeof(u8)); - DefineConstBuffers(*this, info, &UniformDefinitions::S8, binding, S8, 's', sizeof(s8)); - } else { - types |= IR::Type::U32; - } - } - if (True(types & IR::Type::U16)) { - if (profile.support_int16) { - DefineConstBuffers(*this, info, &UniformDefinitions::U16, binding, U16, 'u', - sizeof(u16)); - DefineConstBuffers(*this, info, &UniformDefinitions::S16, binding, S16, 's', - sizeof(s16)); - } else { - types |= IR::Type::U32; - } - } - if (True(types & IR::Type::U32)) { - DefineConstBuffers(*this, info, &UniformDefinitions::U32, binding, U32[1], 'u', - sizeof(u32)); - } - if (True(types & IR::Type::F32)) { - DefineConstBuffers(*this, info, &UniformDefinitions::F32, binding, F32[1], 'f', - sizeof(f32)); - } - if (True(types & IR::Type::U32x2)) { - DefineConstBuffers(*this, info, &UniformDefinitions::U32x2, binding, U32[2], 'u', - sizeof(u32[2])); - } - binding += static_cast(info.constant_buffer_descriptors.size()); -} - -void EmitContext::DefineStorageBuffers(const Info& info, u32& binding) { - if (info.storage_buffers_descriptors.empty()) { - return; - } - AddExtension("SPV_KHR_storage_buffer_storage_class"); - - const IR::Type used_types{profile.support_descriptor_aliasing ? info.used_storage_buffer_types - : IR::Type::U32}; - if (profile.support_int8 && True(used_types & IR::Type::U8)) { - DefineSsbos(*this, storage_types.U8, &StorageDefinitions::U8, info, binding, U8, - sizeof(u8)); - DefineSsbos(*this, storage_types.S8, &StorageDefinitions::S8, info, binding, S8, - sizeof(u8)); - } - if (profile.support_int16 && True(used_types & IR::Type::U16)) { - DefineSsbos(*this, storage_types.U16, &StorageDefinitions::U16, info, binding, U16, - sizeof(u16)); - DefineSsbos(*this, storage_types.S16, &StorageDefinitions::S16, info, binding, S16, - sizeof(u16)); - } - if (True(used_types & IR::Type::U32)) { - DefineSsbos(*this, storage_types.U32, &StorageDefinitions::U32, info, binding, U32[1], - sizeof(u32)); - } - if (True(used_types & IR::Type::F32)) { - DefineSsbos(*this, storage_types.F32, &StorageDefinitions::F32, info, binding, F32[1], - sizeof(f32)); - } - if (True(used_types & IR::Type::U64)) { - DefineSsbos(*this, storage_types.U64, &StorageDefinitions::U64, info, binding, U64, - sizeof(u64)); - } - if (True(used_types & IR::Type::U32x2)) { - DefineSsbos(*this, storage_types.U32x2, &StorageDefinitions::U32x2, info, binding, U32[2], - sizeof(u32[2])); - } - if (True(used_types & IR::Type::U32x4)) { - DefineSsbos(*this, storage_types.U32x4, &StorageDefinitions::U32x4, info, binding, U32[4], - sizeof(u32[4])); - } - for (const StorageBufferDescriptor& desc : info.storage_buffers_descriptors) { - binding += desc.count; - } - const bool needs_function{ - info.uses_global_increment || info.uses_global_decrement || info.uses_atomic_f32_add || - info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max || - info.uses_atomic_f32x2_add || info.uses_atomic_f32x2_min || info.uses_atomic_f32x2_max}; - if (needs_function) { - AddCapability(spv::Capability::VariablePointersStorageBuffer); - } - if (info.uses_global_increment) { - increment_cas_ssbo = CasLoop(*this, Operation::Increment, storage_types.U32.array, - storage_types.U32.element, U32[1], U32[1], spv::Scope::Device); - } - if (info.uses_global_decrement) { - decrement_cas_ssbo = CasLoop(*this, Operation::Decrement, storage_types.U32.array, - storage_types.U32.element, U32[1], U32[1], spv::Scope::Device); - } - if (info.uses_atomic_f32_add) { - f32_add_cas = CasLoop(*this, Operation::FPAdd, storage_types.U32.array, - storage_types.U32.element, F32[1], U32[1], spv::Scope::Device); - } - if (info.uses_atomic_f16x2_add) { - f16x2_add_cas = CasLoop(*this, Operation::FPAdd, storage_types.U32.array, - storage_types.U32.element, F16[2], F16[2], spv::Scope::Device); - } - if (info.uses_atomic_f16x2_min) { - f16x2_min_cas = CasLoop(*this, Operation::FPMin, storage_types.U32.array, - storage_types.U32.element, F16[2], F16[2], spv::Scope::Device); - } - if (info.uses_atomic_f16x2_max) { - f16x2_max_cas = CasLoop(*this, Operation::FPMax, storage_types.U32.array, - storage_types.U32.element, F16[2], F16[2], spv::Scope::Device); - } - if (info.uses_atomic_f32x2_add) { - f32x2_add_cas = CasLoop(*this, Operation::FPAdd, storage_types.U32.array, - storage_types.U32.element, F32[2], F32[2], spv::Scope::Device); - } - if (info.uses_atomic_f32x2_min) { - f32x2_min_cas = CasLoop(*this, Operation::FPMin, storage_types.U32.array, - storage_types.U32.element, F32[2], F32[2], spv::Scope::Device); - } - if (info.uses_atomic_f32x2_max) { - f32x2_max_cas = CasLoop(*this, Operation::FPMax, storage_types.U32.array, - storage_types.U32.element, F32[2], F32[2], spv::Scope::Device); - } -} - -void EmitContext::DefineTextureBuffers(const Info& info, u32& binding) { - if (info.texture_buffer_descriptors.empty()) { - return; - } - const spv::ImageFormat format{spv::ImageFormat::Unknown}; - image_buffer_type = TypeImage(F32[1], spv::Dim::Buffer, 0U, false, false, 1, format); - sampled_texture_buffer_type = TypeSampledImage(image_buffer_type); - - const Id type{TypePointer(spv::StorageClass::UniformConstant, sampled_texture_buffer_type)}; - texture_buffers.reserve(info.texture_buffer_descriptors.size()); - for (const TextureBufferDescriptor& desc : info.texture_buffer_descriptors) { - if (desc.count != 1) { - throw NotImplementedException("Array of texture buffers"); - } - const Id id{AddGlobalVariable(type, spv::StorageClass::UniformConstant)}; - Decorate(id, spv::Decoration::Binding, binding); - Decorate(id, spv::Decoration::DescriptorSet, 0U); - Name(id, NameOf(stage, desc, "texbuf")); - texture_buffers.push_back({ - .id = id, - .count = desc.count, - }); - if (profile.supported_spirv >= 0x00010400) { - interfaces.push_back(id); - } - ++binding; - } -} - -void EmitContext::DefineImageBuffers(const Info& info, u32& binding) { - image_buffers.reserve(info.image_buffer_descriptors.size()); - for (const ImageBufferDescriptor& desc : info.image_buffer_descriptors) { - if (desc.count != 1) { - throw NotImplementedException("Array of image buffers"); - } - const spv::ImageFormat format{GetImageFormat(desc.format)}; - const Id image_type{TypeImage(U32[1], spv::Dim::Buffer, false, false, false, 2, format)}; - const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, image_type)}; - const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant)}; - Decorate(id, spv::Decoration::Binding, binding); - Decorate(id, spv::Decoration::DescriptorSet, 0U); - Name(id, NameOf(stage, desc, "imgbuf")); - image_buffers.push_back({ - .id = id, - .image_type = image_type, - .count = desc.count, - }); - if (profile.supported_spirv >= 0x00010400) { - interfaces.push_back(id); - } - ++binding; - } -} - -void EmitContext::DefineTextures(const Info& info, u32& binding, u32& scaling_index) { - textures.reserve(info.texture_descriptors.size()); - for (const TextureDescriptor& desc : info.texture_descriptors) { - const Id image_type{ImageType(*this, desc)}; - const Id sampled_type{TypeSampledImage(image_type)}; - const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, sampled_type)}; - const Id desc_type{DescType(*this, sampled_type, pointer_type, desc.count)}; - const Id id{AddGlobalVariable(desc_type, spv::StorageClass::UniformConstant)}; - Decorate(id, spv::Decoration::Binding, binding); - Decorate(id, spv::Decoration::DescriptorSet, 0U); - Name(id, NameOf(stage, desc, "tex")); - textures.push_back({ - .id = id, - .sampled_type = sampled_type, - .pointer_type = pointer_type, - .image_type = image_type, - .count = desc.count, - }); - if (profile.supported_spirv >= 0x00010400) { - interfaces.push_back(id); - } - ++binding; - ++scaling_index; - } - if (info.uses_atomic_image_u32) { - image_u32 = TypePointer(spv::StorageClass::Image, U32[1]); - } -} - -void EmitContext::DefineImages(const Info& info, u32& binding, u32& scaling_index) { - images.reserve(info.image_descriptors.size()); - for (const ImageDescriptor& desc : info.image_descriptors) { - if (desc.count != 1) { - throw NotImplementedException("Array of images"); - } - const Id image_type{ImageType(*this, desc)}; - const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, image_type)}; - const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant)}; - Decorate(id, spv::Decoration::Binding, binding); - Decorate(id, spv::Decoration::DescriptorSet, 0U); - Name(id, NameOf(stage, desc, "img")); - images.push_back({ - .id = id, - .image_type = image_type, - .count = desc.count, - }); - if (profile.supported_spirv >= 0x00010400) { - interfaces.push_back(id); - } - ++binding; - ++scaling_index; - } -} - -void EmitContext::DefineInputs(const IR::Program& program) { - const Info& info{program.info}; - const VaryingState loads{info.loads.mask | info.passthrough.mask}; - - if (info.uses_workgroup_id) { - workgroup_id = DefineInput(*this, U32[3], false, spv::BuiltIn::WorkgroupId); - } - if (info.uses_local_invocation_id) { - local_invocation_id = DefineInput(*this, U32[3], false, spv::BuiltIn::LocalInvocationId); - } - if (info.uses_invocation_id) { - invocation_id = DefineInput(*this, U32[1], false, spv::BuiltIn::InvocationId); - } - if (info.uses_sample_id) { - sample_id = DefineInput(*this, U32[1], false, spv::BuiltIn::SampleId); - } - if (info.uses_is_helper_invocation) { - is_helper_invocation = DefineInput(*this, U1, false, spv::BuiltIn::HelperInvocation); - } - if (info.uses_subgroup_mask) { - subgroup_mask_eq = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupEqMaskKHR); - subgroup_mask_lt = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupLtMaskKHR); - subgroup_mask_le = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupLeMaskKHR); - subgroup_mask_gt = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupGtMaskKHR); - subgroup_mask_ge = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupGeMaskKHR); - } - if (info.uses_subgroup_invocation_id || info.uses_subgroup_shuffles || - (profile.warp_size_potentially_larger_than_guest && - (info.uses_subgroup_vote || info.uses_subgroup_mask))) { - subgroup_local_invocation_id = - DefineInput(*this, U32[1], false, spv::BuiltIn::SubgroupLocalInvocationId); - } - if (info.uses_fswzadd) { - const Id f32_one{Const(1.0f)}; - const Id f32_minus_one{Const(-1.0f)}; - const Id f32_zero{Const(0.0f)}; - fswzadd_lut_a = ConstantComposite(F32[4], f32_minus_one, f32_one, f32_minus_one, f32_zero); - fswzadd_lut_b = - ConstantComposite(F32[4], f32_minus_one, f32_minus_one, f32_one, f32_minus_one); - } - if (loads[IR::Attribute::PrimitiveId]) { - primitive_id = DefineInput(*this, U32[1], false, spv::BuiltIn::PrimitiveId); - } - if (loads.AnyComponent(IR::Attribute::PositionX)) { - const bool is_fragment{stage != Stage::Fragment}; - const spv::BuiltIn built_in{is_fragment ? spv::BuiltIn::Position : spv::BuiltIn::FragCoord}; - input_position = DefineInput(*this, F32[4], true, built_in); - if (profile.support_geometry_shader_passthrough) { - if (info.passthrough.AnyComponent(IR::Attribute::PositionX)) { - Decorate(input_position, spv::Decoration::PassthroughNV); - } - } - } - if (loads[IR::Attribute::InstanceId]) { - if (profile.support_vertex_instance_id) { - instance_id = DefineInput(*this, U32[1], true, spv::BuiltIn::InstanceId); - } else { - instance_index = DefineInput(*this, U32[1], true, spv::BuiltIn::InstanceIndex); - base_instance = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseInstance); - } - } - if (loads[IR::Attribute::VertexId]) { - if (profile.support_vertex_instance_id) { - vertex_id = DefineInput(*this, U32[1], true, spv::BuiltIn::VertexId); - } else { - vertex_index = DefineInput(*this, U32[1], true, spv::BuiltIn::VertexIndex); - base_vertex = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseVertex); - } - } - if (loads[IR::Attribute::FrontFace]) { - front_face = DefineInput(*this, U1, true, spv::BuiltIn::FrontFacing); - } - if (loads[IR::Attribute::PointSpriteS] || loads[IR::Attribute::PointSpriteT]) { - point_coord = DefineInput(*this, F32[2], true, spv::BuiltIn::PointCoord); - } - if (loads[IR::Attribute::TessellationEvaluationPointU] || - loads[IR::Attribute::TessellationEvaluationPointV]) { - tess_coord = DefineInput(*this, F32[3], false, spv::BuiltIn::TessCoord); - } - std::bitset used_locations{}; - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - const AttributeType input_type{runtime_info.generic_input_types[index]}; - if (!runtime_info.previous_stage_stores.Generic(index)) { - continue; - } - if (!loads.Generic(index)) { - continue; - } - if (input_type == AttributeType::Disabled) { - continue; - } - used_locations.set(index); - const Id type{GetAttributeType(*this, input_type)}; - const Id id{DefineInput(*this, type, true)}; - Decorate(id, spv::Decoration::Location, static_cast(index)); - Name(id, fmt::format("in_attr{}", index)); - input_generics[index] = id; - - if (info.passthrough.Generic(index) && profile.support_geometry_shader_passthrough) { - Decorate(id, spv::Decoration::PassthroughNV); - } - if (stage != Stage::Fragment) { - continue; - } - switch (info.interpolation[index]) { - case Interpolation::Smooth: - // Default - // Decorate(id, spv::Decoration::Smooth); - break; - case Interpolation::NoPerspective: - Decorate(id, spv::Decoration::NoPerspective); - break; - case Interpolation::Flat: - Decorate(id, spv::Decoration::Flat); - break; - } - } - size_t previous_unused_location = 0; - if (loads.AnyComponent(IR::Attribute::ColorFrontDiffuseR)) { - input_front_color = DefineLegacyInput(*this, used_locations, previous_unused_location); - } - if (loads.AnyComponent(IR::Attribute::ColorFrontSpecularR)) { - input_front_secondary_color = - DefineLegacyInput(*this, used_locations, previous_unused_location); - } - if (loads.AnyComponent(IR::Attribute::ColorBackDiffuseR)) { - input_back_color = DefineLegacyInput(*this, used_locations, previous_unused_location); - } - if (loads.AnyComponent(IR::Attribute::ColorBackSpecularR)) { - input_back_secondary_color = - DefineLegacyInput(*this, used_locations, previous_unused_location); - } - if (loads.AnyComponent(IR::Attribute::FogCoordinate)) { - input_fog_frag_coord = DefineLegacyInput(*this, used_locations, previous_unused_location); - } - for (size_t index = 0; index < NUM_FIXEDFNCTEXTURE; ++index) { - if (loads.AnyComponent(IR::Attribute::FixedFncTexture0S + index * 4)) { - input_fixed_fnc_textures[index] = - DefineLegacyInput(*this, used_locations, previous_unused_location); - } - } - if (stage == Stage::TessellationEval) { - for (size_t index = 0; index < info.uses_patches.size(); ++index) { - if (!info.uses_patches[index]) { - continue; - } - const Id id{DefineInput(*this, F32[4], false)}; - Decorate(id, spv::Decoration::Patch); - Decorate(id, spv::Decoration::Location, static_cast(index)); - patches[index] = id; - } - } -} - -void EmitContext::DefineOutputs(const IR::Program& program) { - const Info& info{program.info}; - const std::optional invocations{program.invocations}; - if (info.stores.AnyComponent(IR::Attribute::PositionX) || stage == Stage::VertexB) { - output_position = DefineOutput(*this, F32[4], invocations, spv::BuiltIn::Position); - } - if (info.stores[IR::Attribute::PointSize] || runtime_info.fixed_state_point_size) { - if (stage == Stage::Fragment) { - throw NotImplementedException("Storing PointSize in fragment stage"); - } - output_point_size = DefineOutput(*this, F32[1], invocations, spv::BuiltIn::PointSize); - } - if (info.stores.ClipDistances()) { - if (stage == Stage::Fragment) { - throw NotImplementedException("Storing ClipDistance in fragment stage"); - } - const Id type{TypeArray(F32[1], Const(8U))}; - clip_distances = DefineOutput(*this, type, invocations, spv::BuiltIn::ClipDistance); - } - if (info.stores[IR::Attribute::Layer] && - (profile.support_viewport_index_layer_non_geometry || stage == Stage::Geometry)) { - if (stage == Stage::Fragment) { - throw NotImplementedException("Storing Layer in fragment stage"); - } - layer = DefineOutput(*this, U32[1], invocations, spv::BuiltIn::Layer); - } - if (info.stores[IR::Attribute::ViewportIndex] && - (profile.support_viewport_index_layer_non_geometry || stage == Stage::Geometry)) { - if (stage == Stage::Fragment) { - throw NotImplementedException("Storing ViewportIndex in fragment stage"); - } - viewport_index = DefineOutput(*this, U32[1], invocations, spv::BuiltIn::ViewportIndex); - } - if (info.stores[IR::Attribute::ViewportMask] && profile.support_viewport_mask) { - viewport_mask = DefineOutput(*this, TypeArray(U32[1], Const(1u)), std::nullopt, - spv::BuiltIn::ViewportMaskNV); - } - std::bitset used_locations{}; - for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { - if (info.stores.Generic(index)) { - DefineGenericOutput(*this, index, invocations); - used_locations.set(index); - } - } - size_t previous_unused_location = 0; - if (info.stores.AnyComponent(IR::Attribute::ColorFrontDiffuseR)) { - output_front_color = - DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); - } - if (info.stores.AnyComponent(IR::Attribute::ColorFrontSpecularR)) { - output_front_secondary_color = - DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); - } - if (info.stores.AnyComponent(IR::Attribute::ColorBackDiffuseR)) { - output_back_color = - DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); - } - if (info.stores.AnyComponent(IR::Attribute::ColorBackSpecularR)) { - output_back_secondary_color = - DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); - } - if (info.stores.AnyComponent(IR::Attribute::FogCoordinate)) { - output_fog_frag_coord = - DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); - } - for (size_t index = 0; index < NUM_FIXEDFNCTEXTURE; ++index) { - if (info.stores.AnyComponent(IR::Attribute::FixedFncTexture0S + index * 4)) { - output_fixed_fnc_textures[index] = - DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); - } - } - switch (stage) { - case Stage::TessellationControl: - if (info.stores_tess_level_outer) { - const Id type{TypeArray(F32[1], Const(4U))}; - output_tess_level_outer = - DefineOutput(*this, type, std::nullopt, spv::BuiltIn::TessLevelOuter); - Decorate(output_tess_level_outer, spv::Decoration::Patch); - } - if (info.stores_tess_level_inner) { - const Id type{TypeArray(F32[1], Const(2U))}; - output_tess_level_inner = - DefineOutput(*this, type, std::nullopt, spv::BuiltIn::TessLevelInner); - Decorate(output_tess_level_inner, spv::Decoration::Patch); - } - for (size_t index = 0; index < info.uses_patches.size(); ++index) { - if (!info.uses_patches[index]) { - continue; - } - const Id id{DefineOutput(*this, F32[4], std::nullopt)}; - Decorate(id, spv::Decoration::Patch); - Decorate(id, spv::Decoration::Location, static_cast(index)); - patches[index] = id; - } - break; - case Stage::Fragment: - for (u32 index = 0; index < 8; ++index) { - if (!info.stores_frag_color[index] && !profile.need_declared_frag_colors) { - continue; - } - frag_color[index] = DefineOutput(*this, F32[4], std::nullopt); - Decorate(frag_color[index], spv::Decoration::Location, index); - Name(frag_color[index], fmt::format("frag_color{}", index)); - } - if (info.stores_frag_depth) { - frag_depth = DefineOutput(*this, F32[1], std::nullopt); - Decorate(frag_depth, spv::Decoration::BuiltIn, spv::BuiltIn::FragDepth); - } - if (info.stores_sample_mask) { - sample_mask = DefineOutput(*this, U32[1], std::nullopt); - Decorate(sample_mask, spv::Decoration::BuiltIn, spv::BuiltIn::SampleMask); - } - break; - default: - break; - } -} - -} // namespace Shader::Backend::SPIRV diff --git a/src/shader_recompiler/backend/spirv/emit_context.h b/src/shader_recompiler/backend/spirv/emit_context.h deleted file mode 100644 index 63f8185d9..000000000 --- a/src/shader_recompiler/backend/spirv/emit_context.h +++ /dev/null @@ -1,335 +0,0 @@ -// Copyright 2021 yuzu Emulator Project -// Licensed under GPLv2 or any later version -// Refer to the license.txt file included. - -#pragma once - -#include -#include - -#include - -#include "shader_recompiler/backend/bindings.h" -#include "shader_recompiler/frontend/ir/program.h" -#include "shader_recompiler/profile.h" -#include "shader_recompiler/runtime_info.h" -#include "shader_recompiler/shader_info.h" - -namespace Shader::Backend::SPIRV { - -using Sirit::Id; - -class VectorTypes { -public: - void Define(Sirit::Module& sirit_ctx, Id base_type, std::string_view name); - - [[nodiscard]] Id operator[](size_t size) const noexcept { - return defs[size - 1]; - } - -private: - std::array defs{}; -}; - -struct TextureDefinition { - Id id; - Id sampled_type; - Id pointer_type; - Id image_type; - u32 count; -}; - -struct TextureBufferDefinition { - Id id; - u32 count; -}; - -struct ImageBufferDefinition { - Id id; - Id image_type; - u32 count; -}; - -struct ImageDefinition { - Id id; - Id image_type; - u32 count; -}; - -struct UniformDefinitions { - Id U8{}; - Id S8{}; - Id U16{}; - Id S16{}; - Id U32{}; - Id F32{}; - Id U32x2{}; - Id U32x4{}; -}; - -struct StorageTypeDefinition { - Id array{}; - Id element{}; -}; - -struct StorageTypeDefinitions { - StorageTypeDefinition U8{}; - StorageTypeDefinition S8{}; - StorageTypeDefinition U16{}; - StorageTypeDefinition S16{}; - StorageTypeDefinition U32{}; - StorageTypeDefinition U64{}; - StorageTypeDefinition F32{}; - StorageTypeDefinition U32x2{}; - StorageTypeDefinition U32x4{}; -}; - -struct StorageDefinitions { - Id U8{}; - Id S8{}; - Id U16{}; - Id S16{}; - Id U32{}; - Id F32{}; - Id U64{}; - Id U32x2{}; - Id U32x4{}; -}; - -struct GenericElementInfo { - Id id{}; - u32 first_element{}; - u32 num_components{}; -}; - -class EmitContext final : public Sirit::Module { -public: - explicit EmitContext(const Profile& profile, const RuntimeInfo& runtime_info, - IR::Program& program, Bindings& binding); - ~EmitContext(); - - [[nodiscard]] Id Def(const IR::Value& value); - - [[nodiscard]] Id BitOffset8(const IR::Value& offset); - [[nodiscard]] Id BitOffset16(const IR::Value& offset); - - Id InputLegacyAttribute(IR::Attribute attribute); - Id OutputLegacyAttribute(IR::Attribute attribute); - - Id Const(u32 value) { - return Constant(U32[1], value); - } - - Id Const(u32 element_1, u32 element_2) { - return ConstantComposite(U32[2], Const(element_1), Const(element_2)); - } - - Id Const(u32 element_1, u32 element_2, u32 element_3) { - return ConstantComposite(U32[3], Const(element_1), Const(element_2), Const(element_3)); - } - - Id Const(u32 element_1, u32 element_2, u32 element_3, u32 element_4) { - return ConstantComposite(U32[4], Const(element_1), Const(element_2), Const(element_3), - Const(element_4)); - } - - Id SConst(s32 value) { - return Constant(S32[1], value); - } - - Id SConst(s32 element_1, s32 element_2) { - return ConstantComposite(S32[2], SConst(element_1), SConst(element_2)); - } - - Id SConst(s32 element_1, s32 element_2, s32 element_3) { - return ConstantComposite(S32[3], SConst(element_1), SConst(element_2), SConst(element_3)); - } - - Id SConst(s32 element_1, s32 element_2, s32 element_3, s32 element_4) { - return ConstantComposite(S32[4], SConst(element_1), SConst(element_2), SConst(element_3), - SConst(element_4)); - } - - Id Const(f32 value) { - return Constant(F32[1], value); - } - - const Profile& profile; - const RuntimeInfo& runtime_info; - Stage stage{}; - - Id void_id{}; - Id U1{}; - Id U8{}; - Id S8{}; - Id U16{}; - Id S16{}; - Id U64{}; - VectorTypes F32; - VectorTypes U32; - VectorTypes S32; - VectorTypes F16; - VectorTypes F64; - - Id true_value{}; - Id false_value{}; - Id u32_zero_value{}; - Id f32_zero_value{}; - - UniformDefinitions uniform_types; - StorageTypeDefinitions storage_types; - - Id private_u32{}; - - Id shared_u8{}; - Id shared_u16{}; - Id shared_u32{}; - Id shared_u64{}; - Id shared_u32x2{}; - Id shared_u32x4{}; - - Id input_f32{}; - Id input_u32{}; - Id input_s32{}; - - Id output_f32{}; - Id output_u32{}; - - Id image_buffer_type{}; - Id sampled_texture_buffer_type{}; - Id image_u32{}; - - std::array cbufs{}; - std::array ssbos{}; - std::vector texture_buffers; - std::vector image_buffers; - std::vector textures; - std::vector images; - - Id workgroup_id{}; - Id local_invocation_id{}; - Id invocation_id{}; - Id sample_id{}; - Id is_helper_invocation{}; - Id subgroup_local_invocation_id{}; - Id subgroup_mask_eq{}; - Id subgroup_mask_lt{}; - Id subgroup_mask_le{}; - Id subgroup_mask_gt{}; - Id subgroup_mask_ge{}; - Id instance_id{}; - Id instance_index{}; - Id base_instance{}; - Id vertex_id{}; - Id vertex_index{}; - Id base_vertex{}; - Id front_face{}; - Id point_coord{}; - Id tess_coord{}; - Id clip_distances{}; - Id layer{}; - Id viewport_index{}; - Id viewport_mask{}; - Id primitive_id{}; - - Id fswzadd_lut_a{}; - Id fswzadd_lut_b{}; - - Id indexed_load_func{}; - Id indexed_store_func{}; - - Id rescaling_uniform_constant{}; - Id rescaling_push_constants{}; - Id rescaling_textures_type{}; - Id rescaling_images_type{}; - u32 rescaling_textures_member_index{}; - u32 rescaling_images_member_index{}; - u32 rescaling_downfactor_member_index{}; - u32 texture_rescaling_index{}; - u32 image_rescaling_index{}; - - Id local_memory{}; - - Id shared_memory_u8{}; - Id shared_memory_u16{}; - Id shared_memory_u32{}; - Id shared_memory_u64{}; - Id shared_memory_u32x2{}; - Id shared_memory_u32x4{}; - - Id shared_memory_u32_type{}; - - Id shared_store_u8_func{}; - Id shared_store_u16_func{}; - Id increment_cas_shared{}; - Id increment_cas_ssbo{}; - Id decrement_cas_shared{}; - Id decrement_cas_ssbo{}; - Id f32_add_cas{}; - Id f16x2_add_cas{}; - Id f16x2_min_cas{}; - Id f16x2_max_cas{}; - Id f32x2_add_cas{}; - Id f32x2_min_cas{}; - Id f32x2_max_cas{}; - - Id load_global_func_u32{}; - Id load_global_func_u32x2{}; - Id load_global_func_u32x4{}; - Id write_global_func_u32{}; - Id write_global_func_u32x2{}; - Id write_global_func_u32x4{}; - - Id input_position{}; - Id input_front_color{}; - Id input_front_secondary_color{}; - Id input_back_color{}; - Id input_back_secondary_color{}; - Id input_fog_frag_coord{}; - std::array input_fixed_fnc_textures{}; - std::array input_generics{}; - - Id output_point_size{}; - Id output_position{}; - Id output_front_color{}; - Id output_front_secondary_color{}; - Id output_back_color{}; - Id output_back_secondary_color{}; - Id output_fog_frag_coord{}; - std::array output_fixed_fnc_textures{}; - std::array, 32> output_generics{}; - - Id output_tess_level_outer{}; - Id output_tess_level_inner{}; - std::array patches{}; - - std::array frag_color{}; - Id sample_mask{}; - Id frag_depth{}; - - std::vector interfaces; - -private: - void DefineCommonTypes(const Info& info); - void DefineCommonConstants(); - void DefineInterfaces(const IR::Program& program); - void DefineLocalMemory(const IR::Program& program); - void DefineSharedMemory(const IR::Program& program); - void DefineSharedMemoryFunctions(const IR::Program& program); - void DefineConstantBuffers(const Info& info, u32& binding); - void DefineStorageBuffers(const Info& info, u32& binding); - void DefineTextureBuffers(const Info& info, u32& binding); - void DefineImageBuffers(const Info& info, u32& binding); - void DefineTextures(const Info& info, u32& binding, u32& scaling_index); - void DefineImages(const Info& info, u32& binding, u32& scaling_index); - void DefineAttributeMemAccess(const Info& info); - void DefineGlobalMemoryFunctions(const Info& info); - void DefineRescalingInput(const Info& info); - void DefineRescalingInputPushConstant(); - void DefineRescalingInputUniformConstant(); - - void DefineInputs(const IR::Program& program); - void DefineOutputs(const IR::Program& program); -}; - -} // namespace Shader::Backend::SPIRV diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp new file mode 100644 index 000000000..723455462 --- /dev/null +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp @@ -0,0 +1,1585 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include +#include +#include +#include + +#include + +#include + +#include "common/common_types.h" +#include "common/div_ceil.h" +#include "shader_recompiler/backend/spirv/emit_context.h" +#include "shader_recompiler/backend/spirv/emit_spirv.h" + +namespace Shader::Backend::SPIRV { +namespace { +constexpr size_t NUM_FIXEDFNCTEXTURE = 10; + +enum class Operation { + Increment, + Decrement, + FPAdd, + FPMin, + FPMax, +}; + +struct AttrInfo { + Id pointer; + Id id; + bool needs_cast; +}; + +Id ImageType(EmitContext& ctx, const TextureDescriptor& desc) { + const spv::ImageFormat format{spv::ImageFormat::Unknown}; + const Id type{ctx.F32[1]}; + const bool depth{desc.is_depth}; + switch (desc.type) { + case TextureType::Color1D: + return ctx.TypeImage(type, spv::Dim::Dim1D, depth, false, false, 1, format); + case TextureType::ColorArray1D: + return ctx.TypeImage(type, spv::Dim::Dim1D, depth, true, false, 1, format); + case TextureType::Color2D: + return ctx.TypeImage(type, spv::Dim::Dim2D, depth, false, false, 1, format); + case TextureType::ColorArray2D: + return ctx.TypeImage(type, spv::Dim::Dim2D, depth, true, false, 1, format); + case TextureType::Color3D: + return ctx.TypeImage(type, spv::Dim::Dim3D, depth, false, false, 1, format); + case TextureType::ColorCube: + return ctx.TypeImage(type, spv::Dim::Cube, depth, false, false, 1, format); + case TextureType::ColorArrayCube: + return ctx.TypeImage(type, spv::Dim::Cube, depth, true, false, 1, format); + case TextureType::Buffer: + break; + } + throw InvalidArgument("Invalid texture type {}", desc.type); +} + +spv::ImageFormat GetImageFormat(ImageFormat format) { + switch (format) { + case ImageFormat::Typeless: + return spv::ImageFormat::Unknown; + case ImageFormat::R8_UINT: + return spv::ImageFormat::R8ui; + case ImageFormat::R8_SINT: + return spv::ImageFormat::R8i; + case ImageFormat::R16_UINT: + return spv::ImageFormat::R16ui; + case ImageFormat::R16_SINT: + return spv::ImageFormat::R16i; + case ImageFormat::R32_UINT: + return spv::ImageFormat::R32ui; + case ImageFormat::R32G32_UINT: + return spv::ImageFormat::Rg32ui; + case ImageFormat::R32G32B32A32_UINT: + return spv::ImageFormat::Rgba32ui; + } + throw InvalidArgument("Invalid image format {}", format); +} + +Id ImageType(EmitContext& ctx, const ImageDescriptor& desc) { + const spv::ImageFormat format{GetImageFormat(desc.format)}; + const Id type{ctx.U32[1]}; + switch (desc.type) { + case TextureType::Color1D: + return ctx.TypeImage(type, spv::Dim::Dim1D, false, false, false, 2, format); + case TextureType::ColorArray1D: + return ctx.TypeImage(type, spv::Dim::Dim1D, false, true, false, 2, format); + case TextureType::Color2D: + return ctx.TypeImage(type, spv::Dim::Dim2D, false, false, false, 2, format); + case TextureType::ColorArray2D: + return ctx.TypeImage(type, spv::Dim::Dim2D, false, true, false, 2, format); + case TextureType::Color3D: + return ctx.TypeImage(type, spv::Dim::Dim3D, false, false, false, 2, format); + case TextureType::Buffer: + throw NotImplementedException("Image buffer"); + default: + break; + } + throw InvalidArgument("Invalid texture type {}", desc.type); +} + +Id DefineVariable(EmitContext& ctx, Id type, std::optional builtin, + spv::StorageClass storage_class) { + const Id pointer_type{ctx.TypePointer(storage_class, type)}; + const Id id{ctx.AddGlobalVariable(pointer_type, storage_class)}; + if (builtin) { + ctx.Decorate(id, spv::Decoration::BuiltIn, *builtin); + } + ctx.interfaces.push_back(id); + return id; +} + +u32 NumVertices(InputTopology input_topology) { + switch (input_topology) { + case InputTopology::Points: + return 1; + case InputTopology::Lines: + return 2; + case InputTopology::LinesAdjacency: + return 4; + case InputTopology::Triangles: + return 3; + case InputTopology::TrianglesAdjacency: + return 6; + } + throw InvalidArgument("Invalid input topology {}", input_topology); +} + +Id DefineInput(EmitContext& ctx, Id type, bool per_invocation, + std::optional builtin = std::nullopt) { + switch (ctx.stage) { + case Stage::TessellationControl: + case Stage::TessellationEval: + if (per_invocation) { + type = ctx.TypeArray(type, ctx.Const(32u)); + } + break; + case Stage::Geometry: + if (per_invocation) { + const u32 num_vertices{NumVertices(ctx.runtime_info.input_topology)}; + type = ctx.TypeArray(type, ctx.Const(num_vertices)); + } + break; + default: + break; + } + return DefineVariable(ctx, type, builtin, spv::StorageClass::Input); +} + +Id DefineOutput(EmitContext& ctx, Id type, std::optional invocations, + std::optional builtin = std::nullopt) { + if (invocations && ctx.stage == Stage::TessellationControl) { + type = ctx.TypeArray(type, ctx.Const(*invocations)); + } + return DefineVariable(ctx, type, builtin, spv::StorageClass::Output); +} + +void DefineGenericOutput(EmitContext& ctx, size_t index, std::optional invocations) { + static constexpr std::string_view swizzle{"xyzw"}; + const size_t base_attr_index{static_cast(IR::Attribute::Generic0X) + index * 4}; + u32 element{0}; + while (element < 4) { + const u32 remainder{4 - element}; + const TransformFeedbackVarying* xfb_varying{}; + if (!ctx.runtime_info.xfb_varyings.empty()) { + xfb_varying = &ctx.runtime_info.xfb_varyings[base_attr_index + element]; + xfb_varying = xfb_varying && xfb_varying->components > 0 ? xfb_varying : nullptr; + } + const u32 num_components{xfb_varying ? xfb_varying->components : remainder}; + + const Id id{DefineOutput(ctx, ctx.F32[num_components], invocations)}; + ctx.Decorate(id, spv::Decoration::Location, static_cast(index)); + if (element > 0) { + ctx.Decorate(id, spv::Decoration::Component, element); + } + if (xfb_varying) { + ctx.Decorate(id, spv::Decoration::XfbBuffer, xfb_varying->buffer); + ctx.Decorate(id, spv::Decoration::XfbStride, xfb_varying->stride); + ctx.Decorate(id, spv::Decoration::Offset, xfb_varying->offset); + } + if (num_components < 4 || element > 0) { + const std::string_view subswizzle{swizzle.substr(element, num_components)}; + ctx.Name(id, fmt::format("out_attr{}_{}", index, subswizzle)); + } else { + ctx.Name(id, fmt::format("out_attr{}", index)); + } + const GenericElementInfo info{ + .id = id, + .first_element = element, + .num_components = num_components, + }; + std::fill_n(ctx.output_generics[index].begin() + element, num_components, info); + element += num_components; + } +} + +Id GetAttributeType(EmitContext& ctx, AttributeType type) { + switch (type) { + case AttributeType::Float: + return ctx.F32[4]; + case AttributeType::SignedInt: + return ctx.TypeVector(ctx.TypeInt(32, true), 4); + case AttributeType::UnsignedInt: + return ctx.U32[4]; + case AttributeType::Disabled: + break; + } + throw InvalidArgument("Invalid attribute type {}", type); +} + +std::optional AttrTypes(EmitContext& ctx, u32 index) { + const AttributeType type{ctx.runtime_info.generic_input_types.at(index)}; + switch (type) { + case AttributeType::Float: + return AttrInfo{ctx.input_f32, ctx.F32[1], false}; + case AttributeType::UnsignedInt: + return AttrInfo{ctx.input_u32, ctx.U32[1], true}; + case AttributeType::SignedInt: + return AttrInfo{ctx.input_s32, ctx.TypeInt(32, true), true}; + case AttributeType::Disabled: + return std::nullopt; + } + throw InvalidArgument("Invalid attribute type {}", type); +} + +std::string_view StageName(Stage stage) { + switch (stage) { + case Stage::VertexA: + return "vs_a"; + case Stage::VertexB: + return "vs"; + case Stage::TessellationControl: + return "tcs"; + case Stage::TessellationEval: + return "tes"; + case Stage::Geometry: + return "gs"; + case Stage::Fragment: + return "fs"; + case Stage::Compute: + return "cs"; + } + throw InvalidArgument("Invalid stage {}", stage); +} + +template +void Name(EmitContext& ctx, Id object, std::string_view format_str, Args&&... args) { + ctx.Name(object, fmt::format(fmt::runtime(format_str), StageName(ctx.stage), + std::forward(args)...) + .c_str()); +} + +void DefineConstBuffers(EmitContext& ctx, const Info& info, Id UniformDefinitions::*member_type, + u32 binding, Id type, char type_char, u32 element_size) { + const Id array_type{ctx.TypeArray(type, ctx.Const(65536U / element_size))}; + ctx.Decorate(array_type, spv::Decoration::ArrayStride, element_size); + + const Id struct_type{ctx.TypeStruct(array_type)}; + Name(ctx, struct_type, "{}_cbuf_block_{}{}", ctx.stage, type_char, element_size * CHAR_BIT); + ctx.Decorate(struct_type, spv::Decoration::Block); + ctx.MemberName(struct_type, 0, "data"); + ctx.MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U); + + const Id struct_pointer_type{ctx.TypePointer(spv::StorageClass::Uniform, struct_type)}; + const Id uniform_type{ctx.TypePointer(spv::StorageClass::Uniform, type)}; + ctx.uniform_types.*member_type = uniform_type; + + for (const ConstantBufferDescriptor& desc : info.constant_buffer_descriptors) { + const Id id{ctx.AddGlobalVariable(struct_pointer_type, spv::StorageClass::Uniform)}; + ctx.Decorate(id, spv::Decoration::Binding, binding); + ctx.Decorate(id, spv::Decoration::DescriptorSet, 0U); + ctx.Name(id, fmt::format("c{}", desc.index)); + for (size_t i = 0; i < desc.count; ++i) { + ctx.cbufs[desc.index + i].*member_type = id; + } + if (ctx.profile.supported_spirv >= 0x00010400) { + ctx.interfaces.push_back(id); + } + binding += desc.count; + } +} + +void DefineSsbos(EmitContext& ctx, StorageTypeDefinition& type_def, + Id StorageDefinitions::*member_type, const Info& info, u32 binding, Id type, + u32 stride) { + const Id array_type{ctx.TypeRuntimeArray(type)}; + ctx.Decorate(array_type, spv::Decoration::ArrayStride, stride); + + const Id struct_type{ctx.TypeStruct(array_type)}; + ctx.Decorate(struct_type, spv::Decoration::Block); + ctx.MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U); + + const Id struct_pointer{ctx.TypePointer(spv::StorageClass::StorageBuffer, struct_type)}; + type_def.array = struct_pointer; + type_def.element = ctx.TypePointer(spv::StorageClass::StorageBuffer, type); + + u32 index{}; + for (const StorageBufferDescriptor& desc : info.storage_buffers_descriptors) { + const Id id{ctx.AddGlobalVariable(struct_pointer, spv::StorageClass::StorageBuffer)}; + ctx.Decorate(id, spv::Decoration::Binding, binding); + ctx.Decorate(id, spv::Decoration::DescriptorSet, 0U); + ctx.Name(id, fmt::format("ssbo{}", index)); + if (ctx.profile.supported_spirv >= 0x00010400) { + ctx.interfaces.push_back(id); + } + for (size_t i = 0; i < desc.count; ++i) { + ctx.ssbos[index + i].*member_type = id; + } + index += desc.count; + binding += desc.count; + } +} + +Id CasFunction(EmitContext& ctx, Operation operation, Id value_type) { + const Id func_type{ctx.TypeFunction(value_type, value_type, value_type)}; + const Id func{ctx.OpFunction(value_type, spv::FunctionControlMask::MaskNone, func_type)}; + const Id op_a{ctx.OpFunctionParameter(value_type)}; + const Id op_b{ctx.OpFunctionParameter(value_type)}; + ctx.AddLabel(); + Id result{}; + switch (operation) { + case Operation::Increment: { + const Id pred{ctx.OpUGreaterThanEqual(ctx.U1, op_a, op_b)}; + const Id incr{ctx.OpIAdd(value_type, op_a, ctx.Constant(value_type, 1))}; + result = ctx.OpSelect(value_type, pred, ctx.u32_zero_value, incr); + break; + } + case Operation::Decrement: { + const Id lhs{ctx.OpIEqual(ctx.U1, op_a, ctx.Constant(value_type, 0u))}; + const Id rhs{ctx.OpUGreaterThan(ctx.U1, op_a, op_b)}; + const Id pred{ctx.OpLogicalOr(ctx.U1, lhs, rhs)}; + const Id decr{ctx.OpISub(value_type, op_a, ctx.Constant(value_type, 1))}; + result = ctx.OpSelect(value_type, pred, op_b, decr); + break; + } + case Operation::FPAdd: + result = ctx.OpFAdd(value_type, op_a, op_b); + break; + case Operation::FPMin: + result = ctx.OpFMin(value_type, op_a, op_b); + break; + case Operation::FPMax: + result = ctx.OpFMax(value_type, op_a, op_b); + break; + default: + break; + } + ctx.OpReturnValue(result); + ctx.OpFunctionEnd(); + return func; +} + +Id CasLoop(EmitContext& ctx, Operation operation, Id array_pointer, Id element_pointer, + Id value_type, Id memory_type, spv::Scope scope) { + const bool is_shared{scope == spv::Scope::Workgroup}; + const bool is_struct{!is_shared || ctx.profile.support_explicit_workgroup_layout}; + const Id cas_func{CasFunction(ctx, operation, value_type)}; + const Id zero{ctx.u32_zero_value}; + const Id scope_id{ctx.Const(static_cast(scope))}; + + const Id loop_header{ctx.OpLabel()}; + const Id continue_block{ctx.OpLabel()}; + const Id merge_block{ctx.OpLabel()}; + const Id func_type{is_shared + ? ctx.TypeFunction(value_type, ctx.U32[1], value_type) + : ctx.TypeFunction(value_type, ctx.U32[1], value_type, array_pointer)}; + + const Id func{ctx.OpFunction(value_type, spv::FunctionControlMask::MaskNone, func_type)}; + const Id index{ctx.OpFunctionParameter(ctx.U32[1])}; + const Id op_b{ctx.OpFunctionParameter(value_type)}; + const Id base{is_shared ? ctx.shared_memory_u32 : ctx.OpFunctionParameter(array_pointer)}; + ctx.AddLabel(); + ctx.OpBranch(loop_header); + ctx.AddLabel(loop_header); + + ctx.OpLoopMerge(merge_block, continue_block, spv::LoopControlMask::MaskNone); + ctx.OpBranch(continue_block); + + ctx.AddLabel(continue_block); + const Id word_pointer{is_struct ? ctx.OpAccessChain(element_pointer, base, zero, index) + : ctx.OpAccessChain(element_pointer, base, index)}; + if (value_type.value == ctx.F32[2].value) { + const Id u32_value{ctx.OpLoad(ctx.U32[1], word_pointer)}; + const Id value{ctx.OpUnpackHalf2x16(ctx.F32[2], u32_value)}; + const Id new_value{ctx.OpFunctionCall(value_type, cas_func, value, op_b)}; + const Id u32_new_value{ctx.OpPackHalf2x16(ctx.U32[1], new_value)}; + const Id atomic_res{ctx.OpAtomicCompareExchange(ctx.U32[1], word_pointer, scope_id, zero, + zero, u32_new_value, u32_value)}; + const Id success{ctx.OpIEqual(ctx.U1, atomic_res, u32_value)}; + ctx.OpBranchConditional(success, merge_block, loop_header); + + ctx.AddLabel(merge_block); + ctx.OpReturnValue(ctx.OpUnpackHalf2x16(ctx.F32[2], atomic_res)); + } else { + const Id value{ctx.OpLoad(memory_type, word_pointer)}; + const bool matching_type{value_type.value == memory_type.value}; + const Id bitcast_value{matching_type ? value : ctx.OpBitcast(value_type, value)}; + const Id cal_res{ctx.OpFunctionCall(value_type, cas_func, bitcast_value, op_b)}; + const Id new_value{matching_type ? cal_res : ctx.OpBitcast(memory_type, cal_res)}; + const Id atomic_res{ctx.OpAtomicCompareExchange(ctx.U32[1], word_pointer, scope_id, zero, + zero, new_value, value)}; + const Id success{ctx.OpIEqual(ctx.U1, atomic_res, value)}; + ctx.OpBranchConditional(success, merge_block, loop_header); + + ctx.AddLabel(merge_block); + ctx.OpReturnValue(ctx.OpBitcast(value_type, atomic_res)); + } + ctx.OpFunctionEnd(); + return func; +} + +template +std::string NameOf(Stage stage, const Desc& desc, std::string_view prefix) { + if (desc.count > 1) { + return fmt::format("{}_{}{}_{:02x}x{}", StageName(stage), prefix, desc.cbuf_index, + desc.cbuf_offset, desc.count); + } else { + return fmt::format("{}_{}{}_{:02x}", StageName(stage), prefix, desc.cbuf_index, + desc.cbuf_offset); + } +} + +Id DescType(EmitContext& ctx, Id sampled_type, Id pointer_type, u32 count) { + if (count > 1) { + const Id array_type{ctx.TypeArray(sampled_type, ctx.Const(count))}; + return ctx.TypePointer(spv::StorageClass::UniformConstant, array_type); + } else { + return pointer_type; + } +} + +size_t FindAndSetNextUnusedLocation(std::bitset& used_locations, + size_t& start_offset) { + for (size_t location = start_offset; location < used_locations.size(); ++location) { + if (!used_locations.test(location)) { + start_offset = location; + used_locations.set(location); + return location; + } + } + throw RuntimeError("Unable to get an unused location for legacy attribute"); +} + +Id DefineLegacyInput(EmitContext& ctx, std::bitset& used_locations, + size_t& start_offset) { + const Id id{DefineInput(ctx, ctx.F32[4], true)}; + const size_t location = FindAndSetNextUnusedLocation(used_locations, start_offset); + ctx.Decorate(id, spv::Decoration::Location, location); + return id; +} + +Id DefineLegacyOutput(EmitContext& ctx, std::bitset& used_locations, + size_t& start_offset, std::optional invocations) { + const Id id{DefineOutput(ctx, ctx.F32[4], invocations)}; + const size_t location = FindAndSetNextUnusedLocation(used_locations, start_offset); + ctx.Decorate(id, spv::Decoration::Location, location); + return id; +} +} // Anonymous namespace + +void VectorTypes::Define(Sirit::Module& sirit_ctx, Id base_type, std::string_view name) { + defs[0] = sirit_ctx.Name(base_type, name); + + std::array def_name; + for (int i = 1; i < 4; ++i) { + const std::string_view def_name_view( + def_name.data(), + fmt::format_to_n(def_name.data(), def_name.size(), "{}x{}", name, i + 1).size); + defs[static_cast(i)] = + sirit_ctx.Name(sirit_ctx.TypeVector(base_type, i + 1), def_name_view); + } +} + +EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_info_, + IR::Program& program, Bindings& bindings) + : Sirit::Module(profile_.supported_spirv), profile{profile_}, runtime_info{runtime_info_}, + stage{program.stage}, texture_rescaling_index{bindings.texture_scaling_index}, + image_rescaling_index{bindings.image_scaling_index} { + const bool is_unified{profile.unified_descriptor_binding}; + u32& uniform_binding{is_unified ? bindings.unified : bindings.uniform_buffer}; + u32& storage_binding{is_unified ? bindings.unified : bindings.storage_buffer}; + u32& texture_binding{is_unified ? bindings.unified : bindings.texture}; + u32& image_binding{is_unified ? bindings.unified : bindings.image}; + AddCapability(spv::Capability::Shader); + DefineCommonTypes(program.info); + DefineCommonConstants(); + DefineInterfaces(program); + DefineLocalMemory(program); + DefineSharedMemory(program); + DefineSharedMemoryFunctions(program); + DefineConstantBuffers(program.info, uniform_binding); + DefineStorageBuffers(program.info, storage_binding); + DefineTextureBuffers(program.info, texture_binding); + DefineImageBuffers(program.info, image_binding); + DefineTextures(program.info, texture_binding, bindings.texture_scaling_index); + DefineImages(program.info, image_binding, bindings.image_scaling_index); + DefineAttributeMemAccess(program.info); + DefineGlobalMemoryFunctions(program.info); + DefineRescalingInput(program.info); +} + +EmitContext::~EmitContext() = default; + +Id EmitContext::Def(const IR::Value& value) { + if (!value.IsImmediate()) { + return value.InstRecursive()->Definition(); + } + switch (value.Type()) { + case IR::Type::Void: + // Void instructions are used for optional arguments (e.g. texture offsets) + // They are not meant to be used in the SPIR-V module + return Id{}; + case IR::Type::U1: + return value.U1() ? true_value : false_value; + case IR::Type::U32: + return Const(value.U32()); + case IR::Type::U64: + return Constant(U64, value.U64()); + case IR::Type::F32: + return Const(value.F32()); + case IR::Type::F64: + return Constant(F64[1], value.F64()); + default: + throw NotImplementedException("Immediate type {}", value.Type()); + } +} + +Id EmitContext::BitOffset8(const IR::Value& offset) { + if (offset.IsImmediate()) { + return Const((offset.U32() % 4) * 8); + } + return OpBitwiseAnd(U32[1], OpShiftLeftLogical(U32[1], Def(offset), Const(3u)), Const(24u)); +} + +Id EmitContext::BitOffset16(const IR::Value& offset) { + if (offset.IsImmediate()) { + return Const(((offset.U32() / 2) % 2) * 16); + } + return OpBitwiseAnd(U32[1], OpShiftLeftLogical(U32[1], Def(offset), Const(3u)), Const(16u)); +} + +Id EmitContext::InputLegacyAttribute(IR::Attribute attribute) { + if (attribute >= IR::Attribute::ColorFrontDiffuseR && + attribute <= IR::Attribute::ColorFrontDiffuseA) { + return input_front_color; + } + if (attribute >= IR::Attribute::ColorFrontSpecularR && + attribute <= IR::Attribute::ColorFrontSpecularA) { + return input_front_secondary_color; + } + if (attribute >= IR::Attribute::ColorBackDiffuseR && + attribute <= IR::Attribute::ColorBackDiffuseA) { + return input_back_color; + } + if (attribute >= IR::Attribute::ColorBackSpecularR && + attribute <= IR::Attribute::ColorBackSpecularA) { + return input_back_secondary_color; + } + if (attribute == IR::Attribute::FogCoordinate) { + return input_fog_frag_coord; + } + if (attribute >= IR::Attribute::FixedFncTexture0S && + attribute <= IR::Attribute::FixedFncTexture9Q) { + u32 index = + (static_cast(attribute) - static_cast(IR::Attribute::FixedFncTexture0S)) / 4; + return input_fixed_fnc_textures[index]; + } + throw InvalidArgument("Attribute is not legacy attribute {}", attribute); +} + +Id EmitContext::OutputLegacyAttribute(IR::Attribute attribute) { + if (attribute >= IR::Attribute::ColorFrontDiffuseR && + attribute <= IR::Attribute::ColorFrontDiffuseA) { + return output_front_color; + } + if (attribute >= IR::Attribute::ColorFrontSpecularR && + attribute <= IR::Attribute::ColorFrontSpecularA) { + return output_front_secondary_color; + } + if (attribute >= IR::Attribute::ColorBackDiffuseR && + attribute <= IR::Attribute::ColorBackDiffuseA) { + return output_back_color; + } + if (attribute >= IR::Attribute::ColorBackSpecularR && + attribute <= IR::Attribute::ColorBackSpecularA) { + return output_back_secondary_color; + } + if (attribute == IR::Attribute::FogCoordinate) { + return output_fog_frag_coord; + } + if (attribute >= IR::Attribute::FixedFncTexture0S && + attribute <= IR::Attribute::FixedFncTexture9Q) { + u32 index = + (static_cast(attribute) - static_cast(IR::Attribute::FixedFncTexture0S)) / 4; + return output_fixed_fnc_textures[index]; + } + throw InvalidArgument("Attribute is not legacy attribute {}", attribute); +} + +void EmitContext::DefineCommonTypes(const Info& info) { + void_id = TypeVoid(); + + U1 = Name(TypeBool(), "u1"); + + F32.Define(*this, TypeFloat(32), "f32"); + U32.Define(*this, TypeInt(32, false), "u32"); + S32.Define(*this, TypeInt(32, true), "s32"); + + private_u32 = Name(TypePointer(spv::StorageClass::Private, U32[1]), "private_u32"); + + input_f32 = Name(TypePointer(spv::StorageClass::Input, F32[1]), "input_f32"); + input_u32 = Name(TypePointer(spv::StorageClass::Input, U32[1]), "input_u32"); + input_s32 = Name(TypePointer(spv::StorageClass::Input, TypeInt(32, true)), "input_s32"); + + output_f32 = Name(TypePointer(spv::StorageClass::Output, F32[1]), "output_f32"); + output_u32 = Name(TypePointer(spv::StorageClass::Output, U32[1]), "output_u32"); + + if (info.uses_int8 && profile.support_int8) { + AddCapability(spv::Capability::Int8); + U8 = Name(TypeInt(8, false), "u8"); + S8 = Name(TypeInt(8, true), "s8"); + } + if (info.uses_int16 && profile.support_int16) { + AddCapability(spv::Capability::Int16); + U16 = Name(TypeInt(16, false), "u16"); + S16 = Name(TypeInt(16, true), "s16"); + } + if (info.uses_int64) { + AddCapability(spv::Capability::Int64); + U64 = Name(TypeInt(64, false), "u64"); + } + if (info.uses_fp16) { + AddCapability(spv::Capability::Float16); + F16.Define(*this, TypeFloat(16), "f16"); + } + if (info.uses_fp64) { + AddCapability(spv::Capability::Float64); + F64.Define(*this, TypeFloat(64), "f64"); + } +} + +void EmitContext::DefineCommonConstants() { + true_value = ConstantTrue(U1); + false_value = ConstantFalse(U1); + u32_zero_value = Const(0U); + f32_zero_value = Const(0.0f); +} + +void EmitContext::DefineInterfaces(const IR::Program& program) { + DefineInputs(program); + DefineOutputs(program); +} + +void EmitContext::DefineLocalMemory(const IR::Program& program) { + if (program.local_memory_size == 0) { + return; + } + const u32 num_elements{Common::DivCeil(program.local_memory_size, 4U)}; + const Id type{TypeArray(U32[1], Const(num_elements))}; + const Id pointer{TypePointer(spv::StorageClass::Private, type)}; + local_memory = AddGlobalVariable(pointer, spv::StorageClass::Private); + if (profile.supported_spirv >= 0x00010400) { + interfaces.push_back(local_memory); + } +} + +void EmitContext::DefineSharedMemory(const IR::Program& program) { + if (program.shared_memory_size == 0) { + return; + } + const auto make{[&](Id element_type, u32 element_size) { + const u32 num_elements{Common::DivCeil(program.shared_memory_size, element_size)}; + const Id array_type{TypeArray(element_type, Const(num_elements))}; + Decorate(array_type, spv::Decoration::ArrayStride, element_size); + + const Id struct_type{TypeStruct(array_type)}; + MemberDecorate(struct_type, 0U, spv::Decoration::Offset, 0U); + Decorate(struct_type, spv::Decoration::Block); + + const Id pointer{TypePointer(spv::StorageClass::Workgroup, struct_type)}; + const Id element_pointer{TypePointer(spv::StorageClass::Workgroup, element_type)}; + const Id variable{AddGlobalVariable(pointer, spv::StorageClass::Workgroup)}; + Decorate(variable, spv::Decoration::Aliased); + interfaces.push_back(variable); + + return std::make_tuple(variable, element_pointer, pointer); + }}; + if (profile.support_explicit_workgroup_layout) { + AddExtension("SPV_KHR_workgroup_memory_explicit_layout"); + AddCapability(spv::Capability::WorkgroupMemoryExplicitLayoutKHR); + if (program.info.uses_int8) { + AddCapability(spv::Capability::WorkgroupMemoryExplicitLayout8BitAccessKHR); + std::tie(shared_memory_u8, shared_u8, std::ignore) = make(U8, 1); + } + if (program.info.uses_int16) { + AddCapability(spv::Capability::WorkgroupMemoryExplicitLayout16BitAccessKHR); + std::tie(shared_memory_u16, shared_u16, std::ignore) = make(U16, 2); + } + if (program.info.uses_int64) { + std::tie(shared_memory_u64, shared_u64, std::ignore) = make(U64, 8); + } + std::tie(shared_memory_u32, shared_u32, shared_memory_u32_type) = make(U32[1], 4); + std::tie(shared_memory_u32x2, shared_u32x2, std::ignore) = make(U32[2], 8); + std::tie(shared_memory_u32x4, shared_u32x4, std::ignore) = make(U32[4], 16); + return; + } + const u32 num_elements{Common::DivCeil(program.shared_memory_size, 4U)}; + const Id type{TypeArray(U32[1], Const(num_elements))}; + shared_memory_u32_type = TypePointer(spv::StorageClass::Workgroup, type); + + shared_u32 = TypePointer(spv::StorageClass::Workgroup, U32[1]); + shared_memory_u32 = AddGlobalVariable(shared_memory_u32_type, spv::StorageClass::Workgroup); + interfaces.push_back(shared_memory_u32); + + const Id func_type{TypeFunction(void_id, U32[1], U32[1])}; + const auto make_function{[&](u32 mask, u32 size) { + const Id loop_header{OpLabel()}; + const Id continue_block{OpLabel()}; + const Id merge_block{OpLabel()}; + + const Id func{OpFunction(void_id, spv::FunctionControlMask::MaskNone, func_type)}; + const Id offset{OpFunctionParameter(U32[1])}; + const Id insert_value{OpFunctionParameter(U32[1])}; + AddLabel(); + OpBranch(loop_header); + + AddLabel(loop_header); + const Id word_offset{OpShiftRightArithmetic(U32[1], offset, Const(2U))}; + const Id shift_offset{OpShiftLeftLogical(U32[1], offset, Const(3U))}; + const Id bit_offset{OpBitwiseAnd(U32[1], shift_offset, Const(mask))}; + const Id count{Const(size)}; + OpLoopMerge(merge_block, continue_block, spv::LoopControlMask::MaskNone); + OpBranch(continue_block); + + AddLabel(continue_block); + const Id word_pointer{OpAccessChain(shared_u32, shared_memory_u32, word_offset)}; + const Id old_value{OpLoad(U32[1], word_pointer)}; + const Id new_value{OpBitFieldInsert(U32[1], old_value, insert_value, bit_offset, count)}; + const Id atomic_res{OpAtomicCompareExchange(U32[1], word_pointer, Const(1U), u32_zero_value, + u32_zero_value, new_value, old_value)}; + const Id success{OpIEqual(U1, atomic_res, old_value)}; + OpBranchConditional(success, merge_block, loop_header); + + AddLabel(merge_block); + OpReturn(); + OpFunctionEnd(); + return func; + }}; + if (program.info.uses_int8) { + shared_store_u8_func = make_function(24, 8); + } + if (program.info.uses_int16) { + shared_store_u16_func = make_function(16, 16); + } +} + +void EmitContext::DefineSharedMemoryFunctions(const IR::Program& program) { + if (program.info.uses_shared_increment) { + increment_cas_shared = CasLoop(*this, Operation::Increment, shared_memory_u32_type, + shared_u32, U32[1], U32[1], spv::Scope::Workgroup); + } + if (program.info.uses_shared_decrement) { + decrement_cas_shared = CasLoop(*this, Operation::Decrement, shared_memory_u32_type, + shared_u32, U32[1], U32[1], spv::Scope::Workgroup); + } +} + +void EmitContext::DefineAttributeMemAccess(const Info& info) { + const auto make_load{[&] { + const bool is_array{stage == Stage::Geometry}; + const Id end_block{OpLabel()}; + const Id default_label{OpLabel()}; + + const Id func_type_load{is_array ? TypeFunction(F32[1], U32[1], U32[1]) + : TypeFunction(F32[1], U32[1])}; + const Id func{OpFunction(F32[1], spv::FunctionControlMask::MaskNone, func_type_load)}; + const Id offset{OpFunctionParameter(U32[1])}; + const Id vertex{is_array ? OpFunctionParameter(U32[1]) : Id{}}; + + AddLabel(); + const Id base_index{OpShiftRightArithmetic(U32[1], offset, Const(2U))}; + const Id masked_index{OpBitwiseAnd(U32[1], base_index, Const(3U))}; + const Id compare_index{OpShiftRightArithmetic(U32[1], base_index, Const(2U))}; + std::vector literals; + std::vector labels; + if (info.loads.AnyComponent(IR::Attribute::PositionX)) { + literals.push_back(static_cast(IR::Attribute::PositionX) >> 2); + labels.push_back(OpLabel()); + } + const u32 base_attribute_value = static_cast(IR::Attribute::Generic0X) >> 2; + for (u32 index = 0; index < static_cast(IR::NUM_GENERICS); ++index) { + if (!info.loads.Generic(index)) { + continue; + } + literals.push_back(base_attribute_value + index); + labels.push_back(OpLabel()); + } + OpSelectionMerge(end_block, spv::SelectionControlMask::MaskNone); + OpSwitch(compare_index, default_label, literals, labels); + AddLabel(default_label); + OpReturnValue(Const(0.0f)); + size_t label_index{0}; + if (info.loads.AnyComponent(IR::Attribute::PositionX)) { + AddLabel(labels[label_index]); + const Id pointer{is_array + ? OpAccessChain(input_f32, input_position, vertex, masked_index) + : OpAccessChain(input_f32, input_position, masked_index)}; + const Id result{OpLoad(F32[1], pointer)}; + OpReturnValue(result); + ++label_index; + } + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + if (!info.loads.Generic(index)) { + continue; + } + AddLabel(labels[label_index]); + const auto type{AttrTypes(*this, static_cast(index))}; + if (!type) { + OpReturnValue(Const(0.0f)); + ++label_index; + continue; + } + const Id generic_id{input_generics.at(index)}; + const Id pointer{is_array + ? OpAccessChain(type->pointer, generic_id, vertex, masked_index) + : OpAccessChain(type->pointer, generic_id, masked_index)}; + const Id value{OpLoad(type->id, pointer)}; + const Id result{type->needs_cast ? OpBitcast(F32[1], value) : value}; + OpReturnValue(result); + ++label_index; + } + AddLabel(end_block); + OpUnreachable(); + OpFunctionEnd(); + return func; + }}; + const auto make_store{[&] { + const Id end_block{OpLabel()}; + const Id default_label{OpLabel()}; + + const Id func_type_store{TypeFunction(void_id, U32[1], F32[1])}; + const Id func{OpFunction(void_id, spv::FunctionControlMask::MaskNone, func_type_store)}; + const Id offset{OpFunctionParameter(U32[1])}; + const Id store_value{OpFunctionParameter(F32[1])}; + AddLabel(); + const Id base_index{OpShiftRightArithmetic(U32[1], offset, Const(2U))}; + const Id masked_index{OpBitwiseAnd(U32[1], base_index, Const(3U))}; + const Id compare_index{OpShiftRightArithmetic(U32[1], base_index, Const(2U))}; + std::vector literals; + std::vector labels; + if (info.stores.AnyComponent(IR::Attribute::PositionX)) { + literals.push_back(static_cast(IR::Attribute::PositionX) >> 2); + labels.push_back(OpLabel()); + } + const u32 base_attribute_value = static_cast(IR::Attribute::Generic0X) >> 2; + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + if (!info.stores.Generic(index)) { + continue; + } + literals.push_back(base_attribute_value + static_cast(index)); + labels.push_back(OpLabel()); + } + if (info.stores.ClipDistances()) { + literals.push_back(static_cast(IR::Attribute::ClipDistance0) >> 2); + labels.push_back(OpLabel()); + literals.push_back(static_cast(IR::Attribute::ClipDistance4) >> 2); + labels.push_back(OpLabel()); + } + OpSelectionMerge(end_block, spv::SelectionControlMask::MaskNone); + OpSwitch(compare_index, default_label, literals, labels); + AddLabel(default_label); + OpReturn(); + size_t label_index{0}; + if (info.stores.AnyComponent(IR::Attribute::PositionX)) { + AddLabel(labels[label_index]); + const Id pointer{OpAccessChain(output_f32, output_position, masked_index)}; + OpStore(pointer, store_value); + OpReturn(); + ++label_index; + } + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + if (!info.stores.Generic(index)) { + continue; + } + if (output_generics[index][0].num_components != 4) { + throw NotImplementedException("Physical stores and transform feedbacks"); + } + AddLabel(labels[label_index]); + const Id generic_id{output_generics[index][0].id}; + const Id pointer{OpAccessChain(output_f32, generic_id, masked_index)}; + OpStore(pointer, store_value); + OpReturn(); + ++label_index; + } + if (info.stores.ClipDistances()) { + AddLabel(labels[label_index]); + const Id pointer{OpAccessChain(output_f32, clip_distances, masked_index)}; + OpStore(pointer, store_value); + OpReturn(); + ++label_index; + AddLabel(labels[label_index]); + const Id fixed_index{OpIAdd(U32[1], masked_index, Const(4U))}; + const Id pointer2{OpAccessChain(output_f32, clip_distances, fixed_index)}; + OpStore(pointer2, store_value); + OpReturn(); + ++label_index; + } + AddLabel(end_block); + OpUnreachable(); + OpFunctionEnd(); + return func; + }}; + if (info.loads_indexed_attributes) { + indexed_load_func = make_load(); + } + if (info.stores_indexed_attributes) { + indexed_store_func = make_store(); + } +} + +void EmitContext::DefineGlobalMemoryFunctions(const Info& info) { + if (!info.uses_global_memory || !profile.support_int64) { + return; + } + using DefPtr = Id StorageDefinitions::*; + const Id zero{u32_zero_value}; + const auto define_body{[&](DefPtr ssbo_member, Id addr, Id element_pointer, u32 shift, + auto&& callback) { + AddLabel(); + const size_t num_buffers{info.storage_buffers_descriptors.size()}; + for (size_t index = 0; index < num_buffers; ++index) { + if (!info.nvn_buffer_used[index]) { + continue; + } + const auto& ssbo{info.storage_buffers_descriptors[index]}; + const Id ssbo_addr_cbuf_offset{Const(ssbo.cbuf_offset / 8)}; + const Id ssbo_size_cbuf_offset{Const(ssbo.cbuf_offset / 4 + 2)}; + const Id ssbo_addr_pointer{OpAccessChain( + uniform_types.U32x2, cbufs[ssbo.cbuf_index].U32x2, zero, ssbo_addr_cbuf_offset)}; + const Id ssbo_size_pointer{OpAccessChain(uniform_types.U32, cbufs[ssbo.cbuf_index].U32, + zero, ssbo_size_cbuf_offset)}; + + const Id ssbo_addr{OpBitcast(U64, OpLoad(U32[2], ssbo_addr_pointer))}; + const Id ssbo_size{OpUConvert(U64, OpLoad(U32[1], ssbo_size_pointer))}; + const Id ssbo_end{OpIAdd(U64, ssbo_addr, ssbo_size)}; + const Id cond{OpLogicalAnd(U1, OpUGreaterThanEqual(U1, addr, ssbo_addr), + OpULessThan(U1, addr, ssbo_end))}; + const Id then_label{OpLabel()}; + const Id else_label{OpLabel()}; + OpSelectionMerge(else_label, spv::SelectionControlMask::MaskNone); + OpBranchConditional(cond, then_label, else_label); + AddLabel(then_label); + const Id ssbo_id{ssbos[index].*ssbo_member}; + const Id ssbo_offset{OpUConvert(U32[1], OpISub(U64, addr, ssbo_addr))}; + const Id ssbo_index{OpShiftRightLogical(U32[1], ssbo_offset, Const(shift))}; + const Id ssbo_pointer{OpAccessChain(element_pointer, ssbo_id, zero, ssbo_index)}; + callback(ssbo_pointer); + AddLabel(else_label); + } + }}; + const auto define_load{[&](DefPtr ssbo_member, Id element_pointer, Id type, u32 shift) { + const Id function_type{TypeFunction(type, U64)}; + const Id func_id{OpFunction(type, spv::FunctionControlMask::MaskNone, function_type)}; + const Id addr{OpFunctionParameter(U64)}; + define_body(ssbo_member, addr, element_pointer, shift, + [&](Id ssbo_pointer) { OpReturnValue(OpLoad(type, ssbo_pointer)); }); + OpReturnValue(ConstantNull(type)); + OpFunctionEnd(); + return func_id; + }}; + const auto define_write{[&](DefPtr ssbo_member, Id element_pointer, Id type, u32 shift) { + const Id function_type{TypeFunction(void_id, U64, type)}; + const Id func_id{OpFunction(void_id, spv::FunctionControlMask::MaskNone, function_type)}; + const Id addr{OpFunctionParameter(U64)}; + const Id data{OpFunctionParameter(type)}; + define_body(ssbo_member, addr, element_pointer, shift, [&](Id ssbo_pointer) { + OpStore(ssbo_pointer, data); + OpReturn(); + }); + OpReturn(); + OpFunctionEnd(); + return func_id; + }}; + const auto define{ + [&](DefPtr ssbo_member, const StorageTypeDefinition& type_def, Id type, size_t size) { + const Id element_type{type_def.element}; + const u32 shift{static_cast(std::countr_zero(size))}; + const Id load_func{define_load(ssbo_member, element_type, type, shift)}; + const Id write_func{define_write(ssbo_member, element_type, type, shift)}; + return std::make_pair(load_func, write_func); + }}; + std::tie(load_global_func_u32, write_global_func_u32) = + define(&StorageDefinitions::U32, storage_types.U32, U32[1], sizeof(u32)); + std::tie(load_global_func_u32x2, write_global_func_u32x2) = + define(&StorageDefinitions::U32x2, storage_types.U32x2, U32[2], sizeof(u32[2])); + std::tie(load_global_func_u32x4, write_global_func_u32x4) = + define(&StorageDefinitions::U32x4, storage_types.U32x4, U32[4], sizeof(u32[4])); +} + +void EmitContext::DefineRescalingInput(const Info& info) { + if (!info.uses_rescaling_uniform) { + return; + } + if (profile.unified_descriptor_binding) { + DefineRescalingInputPushConstant(); + } else { + DefineRescalingInputUniformConstant(); + } +} + +void EmitContext::DefineRescalingInputPushConstant() { + boost::container::static_vector members{}; + u32 member_index{0}; + + rescaling_textures_type = TypeArray(U32[1], Const(4u)); + Decorate(rescaling_textures_type, spv::Decoration::ArrayStride, 4u); + members.push_back(rescaling_textures_type); + rescaling_textures_member_index = member_index++; + + rescaling_images_type = TypeArray(U32[1], Const(NUM_IMAGE_SCALING_WORDS)); + Decorate(rescaling_images_type, spv::Decoration::ArrayStride, 4u); + members.push_back(rescaling_images_type); + rescaling_images_member_index = member_index++; + + if (stage != Stage::Compute) { + members.push_back(F32[1]); + rescaling_downfactor_member_index = member_index++; + } + const Id push_constant_struct{TypeStruct(std::span(members.data(), members.size()))}; + Decorate(push_constant_struct, spv::Decoration::Block); + Name(push_constant_struct, "ResolutionInfo"); + + MemberDecorate(push_constant_struct, rescaling_textures_member_index, spv::Decoration::Offset, + static_cast(offsetof(RescalingLayout, rescaling_textures))); + MemberName(push_constant_struct, rescaling_textures_member_index, "rescaling_textures"); + + MemberDecorate(push_constant_struct, rescaling_images_member_index, spv::Decoration::Offset, + static_cast(offsetof(RescalingLayout, rescaling_images))); + MemberName(push_constant_struct, rescaling_images_member_index, "rescaling_images"); + + if (stage != Stage::Compute) { + MemberDecorate(push_constant_struct, rescaling_downfactor_member_index, + spv::Decoration::Offset, + static_cast(offsetof(RescalingLayout, down_factor))); + MemberName(push_constant_struct, rescaling_downfactor_member_index, "down_factor"); + } + const Id pointer_type{TypePointer(spv::StorageClass::PushConstant, push_constant_struct)}; + rescaling_push_constants = AddGlobalVariable(pointer_type, spv::StorageClass::PushConstant); + Name(rescaling_push_constants, "rescaling_push_constants"); + + if (profile.supported_spirv >= 0x00010400) { + interfaces.push_back(rescaling_push_constants); + } +} + +void EmitContext::DefineRescalingInputUniformConstant() { + const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, F32[4])}; + rescaling_uniform_constant = + AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant); + Decorate(rescaling_uniform_constant, spv::Decoration::Location, 0u); + + if (profile.supported_spirv >= 0x00010400) { + interfaces.push_back(rescaling_uniform_constant); + } +} + +void EmitContext::DefineConstantBuffers(const Info& info, u32& binding) { + if (info.constant_buffer_descriptors.empty()) { + return; + } + if (!profile.support_descriptor_aliasing) { + DefineConstBuffers(*this, info, &UniformDefinitions::U32x4, binding, U32[4], 'u', + sizeof(u32[4])); + for (const ConstantBufferDescriptor& desc : info.constant_buffer_descriptors) { + binding += desc.count; + } + return; + } + IR::Type types{info.used_constant_buffer_types}; + if (True(types & IR::Type::U8)) { + if (profile.support_int8) { + DefineConstBuffers(*this, info, &UniformDefinitions::U8, binding, U8, 'u', sizeof(u8)); + DefineConstBuffers(*this, info, &UniformDefinitions::S8, binding, S8, 's', sizeof(s8)); + } else { + types |= IR::Type::U32; + } + } + if (True(types & IR::Type::U16)) { + if (profile.support_int16) { + DefineConstBuffers(*this, info, &UniformDefinitions::U16, binding, U16, 'u', + sizeof(u16)); + DefineConstBuffers(*this, info, &UniformDefinitions::S16, binding, S16, 's', + sizeof(s16)); + } else { + types |= IR::Type::U32; + } + } + if (True(types & IR::Type::U32)) { + DefineConstBuffers(*this, info, &UniformDefinitions::U32, binding, U32[1], 'u', + sizeof(u32)); + } + if (True(types & IR::Type::F32)) { + DefineConstBuffers(*this, info, &UniformDefinitions::F32, binding, F32[1], 'f', + sizeof(f32)); + } + if (True(types & IR::Type::U32x2)) { + DefineConstBuffers(*this, info, &UniformDefinitions::U32x2, binding, U32[2], 'u', + sizeof(u32[2])); + } + binding += static_cast(info.constant_buffer_descriptors.size()); +} + +void EmitContext::DefineStorageBuffers(const Info& info, u32& binding) { + if (info.storage_buffers_descriptors.empty()) { + return; + } + AddExtension("SPV_KHR_storage_buffer_storage_class"); + + const IR::Type used_types{profile.support_descriptor_aliasing ? info.used_storage_buffer_types + : IR::Type::U32}; + if (profile.support_int8 && True(used_types & IR::Type::U8)) { + DefineSsbos(*this, storage_types.U8, &StorageDefinitions::U8, info, binding, U8, + sizeof(u8)); + DefineSsbos(*this, storage_types.S8, &StorageDefinitions::S8, info, binding, S8, + sizeof(u8)); + } + if (profile.support_int16 && True(used_types & IR::Type::U16)) { + DefineSsbos(*this, storage_types.U16, &StorageDefinitions::U16, info, binding, U16, + sizeof(u16)); + DefineSsbos(*this, storage_types.S16, &StorageDefinitions::S16, info, binding, S16, + sizeof(u16)); + } + if (True(used_types & IR::Type::U32)) { + DefineSsbos(*this, storage_types.U32, &StorageDefinitions::U32, info, binding, U32[1], + sizeof(u32)); + } + if (True(used_types & IR::Type::F32)) { + DefineSsbos(*this, storage_types.F32, &StorageDefinitions::F32, info, binding, F32[1], + sizeof(f32)); + } + if (True(used_types & IR::Type::U64)) { + DefineSsbos(*this, storage_types.U64, &StorageDefinitions::U64, info, binding, U64, + sizeof(u64)); + } + if (True(used_types & IR::Type::U32x2)) { + DefineSsbos(*this, storage_types.U32x2, &StorageDefinitions::U32x2, info, binding, U32[2], + sizeof(u32[2])); + } + if (True(used_types & IR::Type::U32x4)) { + DefineSsbos(*this, storage_types.U32x4, &StorageDefinitions::U32x4, info, binding, U32[4], + sizeof(u32[4])); + } + for (const StorageBufferDescriptor& desc : info.storage_buffers_descriptors) { + binding += desc.count; + } + const bool needs_function{ + info.uses_global_increment || info.uses_global_decrement || info.uses_atomic_f32_add || + info.uses_atomic_f16x2_add || info.uses_atomic_f16x2_min || info.uses_atomic_f16x2_max || + info.uses_atomic_f32x2_add || info.uses_atomic_f32x2_min || info.uses_atomic_f32x2_max}; + if (needs_function) { + AddCapability(spv::Capability::VariablePointersStorageBuffer); + } + if (info.uses_global_increment) { + increment_cas_ssbo = CasLoop(*this, Operation::Increment, storage_types.U32.array, + storage_types.U32.element, U32[1], U32[1], spv::Scope::Device); + } + if (info.uses_global_decrement) { + decrement_cas_ssbo = CasLoop(*this, Operation::Decrement, storage_types.U32.array, + storage_types.U32.element, U32[1], U32[1], spv::Scope::Device); + } + if (info.uses_atomic_f32_add) { + f32_add_cas = CasLoop(*this, Operation::FPAdd, storage_types.U32.array, + storage_types.U32.element, F32[1], U32[1], spv::Scope::Device); + } + if (info.uses_atomic_f16x2_add) { + f16x2_add_cas = CasLoop(*this, Operation::FPAdd, storage_types.U32.array, + storage_types.U32.element, F16[2], F16[2], spv::Scope::Device); + } + if (info.uses_atomic_f16x2_min) { + f16x2_min_cas = CasLoop(*this, Operation::FPMin, storage_types.U32.array, + storage_types.U32.element, F16[2], F16[2], spv::Scope::Device); + } + if (info.uses_atomic_f16x2_max) { + f16x2_max_cas = CasLoop(*this, Operation::FPMax, storage_types.U32.array, + storage_types.U32.element, F16[2], F16[2], spv::Scope::Device); + } + if (info.uses_atomic_f32x2_add) { + f32x2_add_cas = CasLoop(*this, Operation::FPAdd, storage_types.U32.array, + storage_types.U32.element, F32[2], F32[2], spv::Scope::Device); + } + if (info.uses_atomic_f32x2_min) { + f32x2_min_cas = CasLoop(*this, Operation::FPMin, storage_types.U32.array, + storage_types.U32.element, F32[2], F32[2], spv::Scope::Device); + } + if (info.uses_atomic_f32x2_max) { + f32x2_max_cas = CasLoop(*this, Operation::FPMax, storage_types.U32.array, + storage_types.U32.element, F32[2], F32[2], spv::Scope::Device); + } +} + +void EmitContext::DefineTextureBuffers(const Info& info, u32& binding) { + if (info.texture_buffer_descriptors.empty()) { + return; + } + const spv::ImageFormat format{spv::ImageFormat::Unknown}; + image_buffer_type = TypeImage(F32[1], spv::Dim::Buffer, 0U, false, false, 1, format); + sampled_texture_buffer_type = TypeSampledImage(image_buffer_type); + + const Id type{TypePointer(spv::StorageClass::UniformConstant, sampled_texture_buffer_type)}; + texture_buffers.reserve(info.texture_buffer_descriptors.size()); + for (const TextureBufferDescriptor& desc : info.texture_buffer_descriptors) { + if (desc.count != 1) { + throw NotImplementedException("Array of texture buffers"); + } + const Id id{AddGlobalVariable(type, spv::StorageClass::UniformConstant)}; + Decorate(id, spv::Decoration::Binding, binding); + Decorate(id, spv::Decoration::DescriptorSet, 0U); + Name(id, NameOf(stage, desc, "texbuf")); + texture_buffers.push_back({ + .id = id, + .count = desc.count, + }); + if (profile.supported_spirv >= 0x00010400) { + interfaces.push_back(id); + } + ++binding; + } +} + +void EmitContext::DefineImageBuffers(const Info& info, u32& binding) { + image_buffers.reserve(info.image_buffer_descriptors.size()); + for (const ImageBufferDescriptor& desc : info.image_buffer_descriptors) { + if (desc.count != 1) { + throw NotImplementedException("Array of image buffers"); + } + const spv::ImageFormat format{GetImageFormat(desc.format)}; + const Id image_type{TypeImage(U32[1], spv::Dim::Buffer, false, false, false, 2, format)}; + const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, image_type)}; + const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant)}; + Decorate(id, spv::Decoration::Binding, binding); + Decorate(id, spv::Decoration::DescriptorSet, 0U); + Name(id, NameOf(stage, desc, "imgbuf")); + image_buffers.push_back({ + .id = id, + .image_type = image_type, + .count = desc.count, + }); + if (profile.supported_spirv >= 0x00010400) { + interfaces.push_back(id); + } + ++binding; + } +} + +void EmitContext::DefineTextures(const Info& info, u32& binding, u32& scaling_index) { + textures.reserve(info.texture_descriptors.size()); + for (const TextureDescriptor& desc : info.texture_descriptors) { + const Id image_type{ImageType(*this, desc)}; + const Id sampled_type{TypeSampledImage(image_type)}; + const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, sampled_type)}; + const Id desc_type{DescType(*this, sampled_type, pointer_type, desc.count)}; + const Id id{AddGlobalVariable(desc_type, spv::StorageClass::UniformConstant)}; + Decorate(id, spv::Decoration::Binding, binding); + Decorate(id, spv::Decoration::DescriptorSet, 0U); + Name(id, NameOf(stage, desc, "tex")); + textures.push_back({ + .id = id, + .sampled_type = sampled_type, + .pointer_type = pointer_type, + .image_type = image_type, + .count = desc.count, + }); + if (profile.supported_spirv >= 0x00010400) { + interfaces.push_back(id); + } + ++binding; + ++scaling_index; + } + if (info.uses_atomic_image_u32) { + image_u32 = TypePointer(spv::StorageClass::Image, U32[1]); + } +} + +void EmitContext::DefineImages(const Info& info, u32& binding, u32& scaling_index) { + images.reserve(info.image_descriptors.size()); + for (const ImageDescriptor& desc : info.image_descriptors) { + if (desc.count != 1) { + throw NotImplementedException("Array of images"); + } + const Id image_type{ImageType(*this, desc)}; + const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, image_type)}; + const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant)}; + Decorate(id, spv::Decoration::Binding, binding); + Decorate(id, spv::Decoration::DescriptorSet, 0U); + Name(id, NameOf(stage, desc, "img")); + images.push_back({ + .id = id, + .image_type = image_type, + .count = desc.count, + }); + if (profile.supported_spirv >= 0x00010400) { + interfaces.push_back(id); + } + ++binding; + ++scaling_index; + } +} + +void EmitContext::DefineInputs(const IR::Program& program) { + const Info& info{program.info}; + const VaryingState loads{info.loads.mask | info.passthrough.mask}; + + if (info.uses_workgroup_id) { + workgroup_id = DefineInput(*this, U32[3], false, spv::BuiltIn::WorkgroupId); + } + if (info.uses_local_invocation_id) { + local_invocation_id = DefineInput(*this, U32[3], false, spv::BuiltIn::LocalInvocationId); + } + if (info.uses_invocation_id) { + invocation_id = DefineInput(*this, U32[1], false, spv::BuiltIn::InvocationId); + } + if (info.uses_sample_id) { + sample_id = DefineInput(*this, U32[1], false, spv::BuiltIn::SampleId); + } + if (info.uses_is_helper_invocation) { + is_helper_invocation = DefineInput(*this, U1, false, spv::BuiltIn::HelperInvocation); + } + if (info.uses_subgroup_mask) { + subgroup_mask_eq = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupEqMaskKHR); + subgroup_mask_lt = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupLtMaskKHR); + subgroup_mask_le = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupLeMaskKHR); + subgroup_mask_gt = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupGtMaskKHR); + subgroup_mask_ge = DefineInput(*this, U32[4], false, spv::BuiltIn::SubgroupGeMaskKHR); + } + if (info.uses_subgroup_invocation_id || info.uses_subgroup_shuffles || + (profile.warp_size_potentially_larger_than_guest && + (info.uses_subgroup_vote || info.uses_subgroup_mask))) { + subgroup_local_invocation_id = + DefineInput(*this, U32[1], false, spv::BuiltIn::SubgroupLocalInvocationId); + } + if (info.uses_fswzadd) { + const Id f32_one{Const(1.0f)}; + const Id f32_minus_one{Const(-1.0f)}; + const Id f32_zero{Const(0.0f)}; + fswzadd_lut_a = ConstantComposite(F32[4], f32_minus_one, f32_one, f32_minus_one, f32_zero); + fswzadd_lut_b = + ConstantComposite(F32[4], f32_minus_one, f32_minus_one, f32_one, f32_minus_one); + } + if (loads[IR::Attribute::PrimitiveId]) { + primitive_id = DefineInput(*this, U32[1], false, spv::BuiltIn::PrimitiveId); + } + if (loads.AnyComponent(IR::Attribute::PositionX)) { + const bool is_fragment{stage != Stage::Fragment}; + const spv::BuiltIn built_in{is_fragment ? spv::BuiltIn::Position : spv::BuiltIn::FragCoord}; + input_position = DefineInput(*this, F32[4], true, built_in); + if (profile.support_geometry_shader_passthrough) { + if (info.passthrough.AnyComponent(IR::Attribute::PositionX)) { + Decorate(input_position, spv::Decoration::PassthroughNV); + } + } + } + if (loads[IR::Attribute::InstanceId]) { + if (profile.support_vertex_instance_id) { + instance_id = DefineInput(*this, U32[1], true, spv::BuiltIn::InstanceId); + } else { + instance_index = DefineInput(*this, U32[1], true, spv::BuiltIn::InstanceIndex); + base_instance = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseInstance); + } + } + if (loads[IR::Attribute::VertexId]) { + if (profile.support_vertex_instance_id) { + vertex_id = DefineInput(*this, U32[1], true, spv::BuiltIn::VertexId); + } else { + vertex_index = DefineInput(*this, U32[1], true, spv::BuiltIn::VertexIndex); + base_vertex = DefineInput(*this, U32[1], true, spv::BuiltIn::BaseVertex); + } + } + if (loads[IR::Attribute::FrontFace]) { + front_face = DefineInput(*this, U1, true, spv::BuiltIn::FrontFacing); + } + if (loads[IR::Attribute::PointSpriteS] || loads[IR::Attribute::PointSpriteT]) { + point_coord = DefineInput(*this, F32[2], true, spv::BuiltIn::PointCoord); + } + if (loads[IR::Attribute::TessellationEvaluationPointU] || + loads[IR::Attribute::TessellationEvaluationPointV]) { + tess_coord = DefineInput(*this, F32[3], false, spv::BuiltIn::TessCoord); + } + std::bitset used_locations{}; + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + const AttributeType input_type{runtime_info.generic_input_types[index]}; + if (!runtime_info.previous_stage_stores.Generic(index)) { + continue; + } + if (!loads.Generic(index)) { + continue; + } + if (input_type == AttributeType::Disabled) { + continue; + } + used_locations.set(index); + const Id type{GetAttributeType(*this, input_type)}; + const Id id{DefineInput(*this, type, true)}; + Decorate(id, spv::Decoration::Location, static_cast(index)); + Name(id, fmt::format("in_attr{}", index)); + input_generics[index] = id; + + if (info.passthrough.Generic(index) && profile.support_geometry_shader_passthrough) { + Decorate(id, spv::Decoration::PassthroughNV); + } + if (stage != Stage::Fragment) { + continue; + } + switch (info.interpolation[index]) { + case Interpolation::Smooth: + // Default + // Decorate(id, spv::Decoration::Smooth); + break; + case Interpolation::NoPerspective: + Decorate(id, spv::Decoration::NoPerspective); + break; + case Interpolation::Flat: + Decorate(id, spv::Decoration::Flat); + break; + } + } + size_t previous_unused_location = 0; + if (loads.AnyComponent(IR::Attribute::ColorFrontDiffuseR)) { + input_front_color = DefineLegacyInput(*this, used_locations, previous_unused_location); + } + if (loads.AnyComponent(IR::Attribute::ColorFrontSpecularR)) { + input_front_secondary_color = + DefineLegacyInput(*this, used_locations, previous_unused_location); + } + if (loads.AnyComponent(IR::Attribute::ColorBackDiffuseR)) { + input_back_color = DefineLegacyInput(*this, used_locations, previous_unused_location); + } + if (loads.AnyComponent(IR::Attribute::ColorBackSpecularR)) { + input_back_secondary_color = + DefineLegacyInput(*this, used_locations, previous_unused_location); + } + if (loads.AnyComponent(IR::Attribute::FogCoordinate)) { + input_fog_frag_coord = DefineLegacyInput(*this, used_locations, previous_unused_location); + } + for (size_t index = 0; index < NUM_FIXEDFNCTEXTURE; ++index) { + if (loads.AnyComponent(IR::Attribute::FixedFncTexture0S + index * 4)) { + input_fixed_fnc_textures[index] = + DefineLegacyInput(*this, used_locations, previous_unused_location); + } + } + if (stage == Stage::TessellationEval) { + for (size_t index = 0; index < info.uses_patches.size(); ++index) { + if (!info.uses_patches[index]) { + continue; + } + const Id id{DefineInput(*this, F32[4], false)}; + Decorate(id, spv::Decoration::Patch); + Decorate(id, spv::Decoration::Location, static_cast(index)); + patches[index] = id; + } + } +} + +void EmitContext::DefineOutputs(const IR::Program& program) { + const Info& info{program.info}; + const std::optional invocations{program.invocations}; + if (info.stores.AnyComponent(IR::Attribute::PositionX) || stage == Stage::VertexB) { + output_position = DefineOutput(*this, F32[4], invocations, spv::BuiltIn::Position); + } + if (info.stores[IR::Attribute::PointSize] || runtime_info.fixed_state_point_size) { + if (stage == Stage::Fragment) { + throw NotImplementedException("Storing PointSize in fragment stage"); + } + output_point_size = DefineOutput(*this, F32[1], invocations, spv::BuiltIn::PointSize); + } + if (info.stores.ClipDistances()) { + if (stage == Stage::Fragment) { + throw NotImplementedException("Storing ClipDistance in fragment stage"); + } + const Id type{TypeArray(F32[1], Const(8U))}; + clip_distances = DefineOutput(*this, type, invocations, spv::BuiltIn::ClipDistance); + } + if (info.stores[IR::Attribute::Layer] && + (profile.support_viewport_index_layer_non_geometry || stage == Stage::Geometry)) { + if (stage == Stage::Fragment) { + throw NotImplementedException("Storing Layer in fragment stage"); + } + layer = DefineOutput(*this, U32[1], invocations, spv::BuiltIn::Layer); + } + if (info.stores[IR::Attribute::ViewportIndex] && + (profile.support_viewport_index_layer_non_geometry || stage == Stage::Geometry)) { + if (stage == Stage::Fragment) { + throw NotImplementedException("Storing ViewportIndex in fragment stage"); + } + viewport_index = DefineOutput(*this, U32[1], invocations, spv::BuiltIn::ViewportIndex); + } + if (info.stores[IR::Attribute::ViewportMask] && profile.support_viewport_mask) { + viewport_mask = DefineOutput(*this, TypeArray(U32[1], Const(1u)), std::nullopt, + spv::BuiltIn::ViewportMaskNV); + } + std::bitset used_locations{}; + for (size_t index = 0; index < IR::NUM_GENERICS; ++index) { + if (info.stores.Generic(index)) { + DefineGenericOutput(*this, index, invocations); + used_locations.set(index); + } + } + size_t previous_unused_location = 0; + if (info.stores.AnyComponent(IR::Attribute::ColorFrontDiffuseR)) { + output_front_color = + DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); + } + if (info.stores.AnyComponent(IR::Attribute::ColorFrontSpecularR)) { + output_front_secondary_color = + DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); + } + if (info.stores.AnyComponent(IR::Attribute::ColorBackDiffuseR)) { + output_back_color = + DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); + } + if (info.stores.AnyComponent(IR::Attribute::ColorBackSpecularR)) { + output_back_secondary_color = + DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); + } + if (info.stores.AnyComponent(IR::Attribute::FogCoordinate)) { + output_fog_frag_coord = + DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); + } + for (size_t index = 0; index < NUM_FIXEDFNCTEXTURE; ++index) { + if (info.stores.AnyComponent(IR::Attribute::FixedFncTexture0S + index * 4)) { + output_fixed_fnc_textures[index] = + DefineLegacyOutput(*this, used_locations, previous_unused_location, invocations); + } + } + switch (stage) { + case Stage::TessellationControl: + if (info.stores_tess_level_outer) { + const Id type{TypeArray(F32[1], Const(4U))}; + output_tess_level_outer = + DefineOutput(*this, type, std::nullopt, spv::BuiltIn::TessLevelOuter); + Decorate(output_tess_level_outer, spv::Decoration::Patch); + } + if (info.stores_tess_level_inner) { + const Id type{TypeArray(F32[1], Const(2U))}; + output_tess_level_inner = + DefineOutput(*this, type, std::nullopt, spv::BuiltIn::TessLevelInner); + Decorate(output_tess_level_inner, spv::Decoration::Patch); + } + for (size_t index = 0; index < info.uses_patches.size(); ++index) { + if (!info.uses_patches[index]) { + continue; + } + const Id id{DefineOutput(*this, F32[4], std::nullopt)}; + Decorate(id, spv::Decoration::Patch); + Decorate(id, spv::Decoration::Location, static_cast(index)); + patches[index] = id; + } + break; + case Stage::Fragment: + for (u32 index = 0; index < 8; ++index) { + if (!info.stores_frag_color[index] && !profile.need_declared_frag_colors) { + continue; + } + frag_color[index] = DefineOutput(*this, F32[4], std::nullopt); + Decorate(frag_color[index], spv::Decoration::Location, index); + Name(frag_color[index], fmt::format("frag_color{}", index)); + } + if (info.stores_frag_depth) { + frag_depth = DefineOutput(*this, F32[1], std::nullopt); + Decorate(frag_depth, spv::Decoration::BuiltIn, spv::BuiltIn::FragDepth); + } + if (info.stores_sample_mask) { + sample_mask = DefineOutput(*this, U32[1], std::nullopt); + Decorate(sample_mask, spv::Decoration::BuiltIn, spv::BuiltIn::SampleMask); + } + break; + default: + break; + } +} + +} // namespace Shader::Backend::SPIRV diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.h b/src/shader_recompiler/backend/spirv/spirv_emit_context.h new file mode 100644 index 000000000..63f8185d9 --- /dev/null +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h @@ -0,0 +1,335 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#pragma once + +#include +#include + +#include + +#include "shader_recompiler/backend/bindings.h" +#include "shader_recompiler/frontend/ir/program.h" +#include "shader_recompiler/profile.h" +#include "shader_recompiler/runtime_info.h" +#include "shader_recompiler/shader_info.h" + +namespace Shader::Backend::SPIRV { + +using Sirit::Id; + +class VectorTypes { +public: + void Define(Sirit::Module& sirit_ctx, Id base_type, std::string_view name); + + [[nodiscard]] Id operator[](size_t size) const noexcept { + return defs[size - 1]; + } + +private: + std::array defs{}; +}; + +struct TextureDefinition { + Id id; + Id sampled_type; + Id pointer_type; + Id image_type; + u32 count; +}; + +struct TextureBufferDefinition { + Id id; + u32 count; +}; + +struct ImageBufferDefinition { + Id id; + Id image_type; + u32 count; +}; + +struct ImageDefinition { + Id id; + Id image_type; + u32 count; +}; + +struct UniformDefinitions { + Id U8{}; + Id S8{}; + Id U16{}; + Id S16{}; + Id U32{}; + Id F32{}; + Id U32x2{}; + Id U32x4{}; +}; + +struct StorageTypeDefinition { + Id array{}; + Id element{}; +}; + +struct StorageTypeDefinitions { + StorageTypeDefinition U8{}; + StorageTypeDefinition S8{}; + StorageTypeDefinition U16{}; + StorageTypeDefinition S16{}; + StorageTypeDefinition U32{}; + StorageTypeDefinition U64{}; + StorageTypeDefinition F32{}; + StorageTypeDefinition U32x2{}; + StorageTypeDefinition U32x4{}; +}; + +struct StorageDefinitions { + Id U8{}; + Id S8{}; + Id U16{}; + Id S16{}; + Id U32{}; + Id F32{}; + Id U64{}; + Id U32x2{}; + Id U32x4{}; +}; + +struct GenericElementInfo { + Id id{}; + u32 first_element{}; + u32 num_components{}; +}; + +class EmitContext final : public Sirit::Module { +public: + explicit EmitContext(const Profile& profile, const RuntimeInfo& runtime_info, + IR::Program& program, Bindings& binding); + ~EmitContext(); + + [[nodiscard]] Id Def(const IR::Value& value); + + [[nodiscard]] Id BitOffset8(const IR::Value& offset); + [[nodiscard]] Id BitOffset16(const IR::Value& offset); + + Id InputLegacyAttribute(IR::Attribute attribute); + Id OutputLegacyAttribute(IR::Attribute attribute); + + Id Const(u32 value) { + return Constant(U32[1], value); + } + + Id Const(u32 element_1, u32 element_2) { + return ConstantComposite(U32[2], Const(element_1), Const(element_2)); + } + + Id Const(u32 element_1, u32 element_2, u32 element_3) { + return ConstantComposite(U32[3], Const(element_1), Const(element_2), Const(element_3)); + } + + Id Const(u32 element_1, u32 element_2, u32 element_3, u32 element_4) { + return ConstantComposite(U32[4], Const(element_1), Const(element_2), Const(element_3), + Const(element_4)); + } + + Id SConst(s32 value) { + return Constant(S32[1], value); + } + + Id SConst(s32 element_1, s32 element_2) { + return ConstantComposite(S32[2], SConst(element_1), SConst(element_2)); + } + + Id SConst(s32 element_1, s32 element_2, s32 element_3) { + return ConstantComposite(S32[3], SConst(element_1), SConst(element_2), SConst(element_3)); + } + + Id SConst(s32 element_1, s32 element_2, s32 element_3, s32 element_4) { + return ConstantComposite(S32[4], SConst(element_1), SConst(element_2), SConst(element_3), + SConst(element_4)); + } + + Id Const(f32 value) { + return Constant(F32[1], value); + } + + const Profile& profile; + const RuntimeInfo& runtime_info; + Stage stage{}; + + Id void_id{}; + Id U1{}; + Id U8{}; + Id S8{}; + Id U16{}; + Id S16{}; + Id U64{}; + VectorTypes F32; + VectorTypes U32; + VectorTypes S32; + VectorTypes F16; + VectorTypes F64; + + Id true_value{}; + Id false_value{}; + Id u32_zero_value{}; + Id f32_zero_value{}; + + UniformDefinitions uniform_types; + StorageTypeDefinitions storage_types; + + Id private_u32{}; + + Id shared_u8{}; + Id shared_u16{}; + Id shared_u32{}; + Id shared_u64{}; + Id shared_u32x2{}; + Id shared_u32x4{}; + + Id input_f32{}; + Id input_u32{}; + Id input_s32{}; + + Id output_f32{}; + Id output_u32{}; + + Id image_buffer_type{}; + Id sampled_texture_buffer_type{}; + Id image_u32{}; + + std::array cbufs{}; + std::array ssbos{}; + std::vector texture_buffers; + std::vector image_buffers; + std::vector textures; + std::vector images; + + Id workgroup_id{}; + Id local_invocation_id{}; + Id invocation_id{}; + Id sample_id{}; + Id is_helper_invocation{}; + Id subgroup_local_invocation_id{}; + Id subgroup_mask_eq{}; + Id subgroup_mask_lt{}; + Id subgroup_mask_le{}; + Id subgroup_mask_gt{}; + Id subgroup_mask_ge{}; + Id instance_id{}; + Id instance_index{}; + Id base_instance{}; + Id vertex_id{}; + Id vertex_index{}; + Id base_vertex{}; + Id front_face{}; + Id point_coord{}; + Id tess_coord{}; + Id clip_distances{}; + Id layer{}; + Id viewport_index{}; + Id viewport_mask{}; + Id primitive_id{}; + + Id fswzadd_lut_a{}; + Id fswzadd_lut_b{}; + + Id indexed_load_func{}; + Id indexed_store_func{}; + + Id rescaling_uniform_constant{}; + Id rescaling_push_constants{}; + Id rescaling_textures_type{}; + Id rescaling_images_type{}; + u32 rescaling_textures_member_index{}; + u32 rescaling_images_member_index{}; + u32 rescaling_downfactor_member_index{}; + u32 texture_rescaling_index{}; + u32 image_rescaling_index{}; + + Id local_memory{}; + + Id shared_memory_u8{}; + Id shared_memory_u16{}; + Id shared_memory_u32{}; + Id shared_memory_u64{}; + Id shared_memory_u32x2{}; + Id shared_memory_u32x4{}; + + Id shared_memory_u32_type{}; + + Id shared_store_u8_func{}; + Id shared_store_u16_func{}; + Id increment_cas_shared{}; + Id increment_cas_ssbo{}; + Id decrement_cas_shared{}; + Id decrement_cas_ssbo{}; + Id f32_add_cas{}; + Id f16x2_add_cas{}; + Id f16x2_min_cas{}; + Id f16x2_max_cas{}; + Id f32x2_add_cas{}; + Id f32x2_min_cas{}; + Id f32x2_max_cas{}; + + Id load_global_func_u32{}; + Id load_global_func_u32x2{}; + Id load_global_func_u32x4{}; + Id write_global_func_u32{}; + Id write_global_func_u32x2{}; + Id write_global_func_u32x4{}; + + Id input_position{}; + Id input_front_color{}; + Id input_front_secondary_color{}; + Id input_back_color{}; + Id input_back_secondary_color{}; + Id input_fog_frag_coord{}; + std::array input_fixed_fnc_textures{}; + std::array input_generics{}; + + Id output_point_size{}; + Id output_position{}; + Id output_front_color{}; + Id output_front_secondary_color{}; + Id output_back_color{}; + Id output_back_secondary_color{}; + Id output_fog_frag_coord{}; + std::array output_fixed_fnc_textures{}; + std::array, 32> output_generics{}; + + Id output_tess_level_outer{}; + Id output_tess_level_inner{}; + std::array patches{}; + + std::array frag_color{}; + Id sample_mask{}; + Id frag_depth{}; + + std::vector interfaces; + +private: + void DefineCommonTypes(const Info& info); + void DefineCommonConstants(); + void DefineInterfaces(const IR::Program& program); + void DefineLocalMemory(const IR::Program& program); + void DefineSharedMemory(const IR::Program& program); + void DefineSharedMemoryFunctions(const IR::Program& program); + void DefineConstantBuffers(const Info& info, u32& binding); + void DefineStorageBuffers(const Info& info, u32& binding); + void DefineTextureBuffers(const Info& info, u32& binding); + void DefineImageBuffers(const Info& info, u32& binding); + void DefineTextures(const Info& info, u32& binding, u32& scaling_index); + void DefineImages(const Info& info, u32& binding, u32& scaling_index); + void DefineAttributeMemAccess(const Info& info); + void DefineGlobalMemoryFunctions(const Info& info); + void DefineRescalingInput(const Info& info); + void DefineRescalingInputPushConstant(); + void DefineRescalingInputUniformConstant(); + + void DefineInputs(const IR::Program& program); + void DefineOutputs(const IR::Program& program); +}; + +} // namespace Shader::Backend::SPIRV -- cgit v1.2.3 From f32b2bcd200097659d7c6e0dfdef1e96e3f2f69e Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Sun, 5 Dec 2021 16:42:03 -0500 Subject: shader_recompiler: Adjust emit_context includes --- src/shader_recompiler/backend/glasm/emit_glasm.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_bitwise_conversion.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_composite.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_convert.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_image.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_select.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_shared_memory.cpp | 2 +- src/shader_recompiler/backend/glasm/emit_glasm_warp.cpp | 2 +- src/shader_recompiler/backend/glasm/glasm_emit_context.cpp | 2 +- src/shader_recompiler/backend/glasm/reg_alloc.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_barriers.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_bitwise_conversion.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_composite.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_control_flow.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_convert.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_floating_point.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_image.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_integer.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_logical.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_select.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_shared_memory.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_special.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_undefined.cpp | 2 +- src/shader_recompiler/backend/glsl/emit_glsl_warp.cpp | 2 +- src/shader_recompiler/backend/glsl/glsl_emit_context.cpp | 2 +- src/shader_recompiler/backend/spirv/emit_spirv.h | 2 +- src/shader_recompiler/backend/spirv/spirv_emit_context.cpp | 2 +- 37 files changed, 37 insertions(+), 37 deletions(-) (limited to 'src') diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp index 004658546..42eff443f 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp @@ -9,9 +9,9 @@ #include "common/div_ceil.h" #include "common/settings.h" #include "shader_recompiler/backend/bindings.h" -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/ir_emitter.h" #include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/profile.h" diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_bitwise_conversion.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_bitwise_conversion.cpp index 9201ccd39..3bfcbbe65 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_bitwise_conversion.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_bitwise_conversion.cpp @@ -2,8 +2,8 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLASM { diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_composite.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_composite.cpp index bff0b7c1c..babbe6654 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_composite.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_composite.cpp @@ -2,8 +2,8 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLASM { diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp index 02c9dc6d7..1b98e5b6c 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" #include "shader_recompiler/profile.h" #include "shader_recompiler/shader_info.h" diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_convert.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_convert.cpp index ccdf1cbc8..4cff70fe4 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_convert.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_convert.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" #include "shader_recompiler/frontend/ir/value.h" diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp index 4ed58619d..356640471 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" #include "shader_recompiler/frontend/ir/value.h" diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_image.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_image.cpp index d325d31c7..237a5af3f 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_image.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_image.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" #include "shader_recompiler/frontend/ir/value.h" diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp index 8aa494a4d..f698b8b9b 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp @@ -2,8 +2,8 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLASM { diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp index af9fac7c1..f135b67f5 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/frontend/ir/value.h" #include "shader_recompiler/runtime_info.h" diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp index 681aeda8d..b6a2161aa 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/frontend/ir/value.h" diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_select.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_select.cpp index 68fff613c..dc441c56d 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_select.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_select.cpp @@ -3,8 +3,8 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLASM { diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_shared_memory.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_shared_memory.cpp index c1498f449..39e1c6c3a 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_shared_memory.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_shared_memory.cpp @@ -3,8 +3,8 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLASM { diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_warp.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_warp.cpp index 544d475b4..32e0dd923 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_warp.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_warp.cpp @@ -2,8 +2,8 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" #include "shader_recompiler/profile.h" diff --git a/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp b/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp index 8fd459dfe..0401953f7 100644 --- a/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp +++ b/src/shader_recompiler/backend/glasm/glasm_emit_context.cpp @@ -5,8 +5,8 @@ #include #include "shader_recompiler/backend/bindings.h" -#include "shader_recompiler/backend/glasm/emit_context.h" #include "shader_recompiler/backend/glasm/emit_glasm.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/profile.h" #include "shader_recompiler/runtime_info.h" diff --git a/src/shader_recompiler/backend/glasm/reg_alloc.cpp b/src/shader_recompiler/backend/glasm/reg_alloc.cpp index 4c046db6e..201e428c1 100644 --- a/src/shader_recompiler/backend/glasm/reg_alloc.cpp +++ b/src/shader_recompiler/backend/glasm/reg_alloc.cpp @@ -6,7 +6,7 @@ #include -#include "shader_recompiler/backend/glasm/emit_context.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" #include "shader_recompiler/backend/glasm/reg_alloc.h" #include "shader_recompiler/exception.h" #include "shader_recompiler/frontend/ir/value.h" diff --git a/src/shader_recompiler/backend/glsl/emit_glsl.cpp b/src/shader_recompiler/backend/glsl/emit_glsl.cpp index 8a430d573..78b2eeaa2 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl.cpp @@ -9,9 +9,9 @@ #include "common/div_ceil.h" #include "common/settings.h" -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/ir_emitter.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp index 772acc5a4..dc377b053 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_atomic.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_barriers.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_barriers.cpp index e1d1b558e..8a9faa394 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_barriers.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_barriers.cpp @@ -2,8 +2,8 @@ // Licensed under GPLv2 or any later version // Refer to the license.txt file included. -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_bitwise_conversion.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_bitwise_conversion.cpp index 3c1714e89..0f2668d9e 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_bitwise_conversion.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_bitwise_conversion.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_composite.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_composite.cpp index 49a66e3ec..98cc57e58 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_composite.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_composite.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp index 4c26f3829..1920047f4 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_context_get_set.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" #include "shader_recompiler/profile.h" #include "shader_recompiler/runtime_info.h" diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_control_flow.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_control_flow.cpp index 53f8896be..c86465e8b 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_control_flow.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_control_flow.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/exception.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_convert.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_convert.cpp index eeae6562c..ce6ea1bb7 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_convert.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_convert.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_floating_point.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_floating_point.cpp index d423bfb1b..b765a251b 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_floating_point.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_floating_point.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" #include "shader_recompiler/frontend/ir/value.h" diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_image.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_image.cpp index 2f78d0267..fae2e397a 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_image.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_image.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" #include "shader_recompiler/frontend/ir/value.h" #include "shader_recompiler/profile.h" diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_integer.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_integer.cpp index 88c1d4c5e..44060df33 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_integer.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_integer.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_logical.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_logical.cpp index 338ff4bd6..742fec9cf 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_logical.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_logical.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp index e3957491f..9fd41b4fd 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_memory.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" #include "shader_recompiler/profile.h" diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp index f420fe388..4ebdfb3bc 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_not_implemented.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" #ifdef _MSC_VER diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_select.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_select.cpp index 49fba9073..b1e486e5f 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_select.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_select.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_shared_memory.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_shared_memory.cpp index 518b78f06..74ae345e5 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_shared_memory.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_shared_memory.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_special.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_special.cpp index 67f9dad68..b8ddafe48 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_special.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_special.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/frontend/ir/value.h" #include "shader_recompiler/profile.h" diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_undefined.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_undefined.cpp index 15bf02dd6..cace1db85 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_undefined.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_undefined.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" namespace Shader::Backend::GLSL { diff --git a/src/shader_recompiler/backend/glsl/emit_glsl_warp.cpp b/src/shader_recompiler/backend/glsl/emit_glsl_warp.cpp index cd285e2c8..6e01979b4 100644 --- a/src/shader_recompiler/backend/glsl/emit_glsl_warp.cpp +++ b/src/shader_recompiler/backend/glsl/emit_glsl_warp.cpp @@ -4,8 +4,8 @@ #include -#include "shader_recompiler/backend/glsl/emit_context.h" #include "shader_recompiler/backend/glsl/emit_glsl_instructions.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/value.h" #include "shader_recompiler/profile.h" diff --git a/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp b/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp index 97bd59302..1de017e76 100644 --- a/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp +++ b/src/shader_recompiler/backend/glsl/glsl_emit_context.cpp @@ -3,7 +3,7 @@ // Refer to the license.txt file included. #include "shader_recompiler/backend/bindings.h" -#include "shader_recompiler/backend/glsl/emit_context.h" +#include "shader_recompiler/backend/glsl/glsl_emit_context.h" #include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/profile.h" #include "shader_recompiler/runtime_info.h" diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.h b/src/shader_recompiler/backend/spirv/emit_spirv.h index 4b25534ce..63dea090d 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.h +++ b/src/shader_recompiler/backend/spirv/emit_spirv.h @@ -10,7 +10,7 @@ #include "common/common_types.h" #include "shader_recompiler/backend/bindings.h" -#include "shader_recompiler/backend/spirv/emit_context.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/profile.h" diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp index 723455462..4b6f792bf 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp @@ -13,8 +13,8 @@ #include "common/common_types.h" #include "common/div_ceil.h" -#include "shader_recompiler/backend/spirv/emit_context.h" #include "shader_recompiler/backend/spirv/emit_spirv.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { -- cgit v1.2.3 From 1e1f7b32341f6538fce07d0df415a4d494a2b209 Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Sun, 5 Dec 2021 16:58:18 -0500 Subject: glasm: Move implemented instructions from not_implemented.cpp --- .../backend/glasm/emit_glasm_barriers.cpp | 22 +++ .../backend/glasm/emit_glasm_context_get_set.cpp | 29 ++++ .../backend/glasm/emit_glasm_control_flow.cpp | 18 +++ .../backend/glasm/emit_glasm_logical.cpp | 26 ++++ .../backend/glasm/emit_glasm_not_implemented.cpp | 169 --------------------- .../backend/glasm/emit_glasm_special.cpp | 95 ++++++++++++ .../backend/glasm/emit_glasm_undefined.cpp | 30 ++++ 7 files changed, 220 insertions(+), 169 deletions(-) (limited to 'src') diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_barriers.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_barriers.cpp index e69de29bb..c0b97683e 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_barriers.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_barriers.cpp @@ -0,0 +1,22 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" + +namespace Shader::Backend::GLASM { + +void EmitBarrier(EmitContext& ctx) { + ctx.Add("BAR;"); +} + +void EmitWorkgroupMemoryBarrier(EmitContext& ctx) { + ctx.Add("MEMBAR.CTA;"); +} + +void EmitDeviceMemoryBarrier(EmitContext& ctx) { + ctx.Add("MEMBAR;"); +} + +} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp index 1b98e5b6c..081b2c8e0 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_context_get_set.cpp @@ -335,6 +335,35 @@ void EmitSetFragDepth(EmitContext& ctx, ScalarF32 value) { ctx.Add("MOV.F result.depth.z,{};", value); } +void EmitWorkgroupId(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {},invocation.groupid;", inst); +} + +void EmitLocalInvocationId(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {},invocation.localid;", inst); +} + +void EmitInvocationId(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {}.x,primitive_invocation.x;", inst); +} + +void EmitSampleId(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {}.x,fragment.sampleid.x;", inst); +} + +void EmitIsHelperInvocation(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {}.x,fragment.helperthread.x;", inst); +} + +void EmitYDirection(EmitContext& ctx, IR::Inst& inst) { + ctx.uses_y_direction = true; + ctx.Add("MOV.F {}.x,y_direction[0].w;", inst); +} + +void EmitResolutionDownFactor(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.F {}.x,scaling[0].z;", inst); +} + void EmitLoadLocal(EmitContext& ctx, IR::Inst& inst, ScalarU32 word_offset) { ctx.Add("MOV.U {},lmem[{}].x;", inst, word_offset); } diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_control_flow.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_control_flow.cpp index e69de29bb..8a14fc8d9 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_control_flow.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_control_flow.cpp @@ -0,0 +1,18 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" + +namespace Shader::Backend::GLASM { + +void EmitJoin(EmitContext&) { + throw NotImplementedException("Join shouldn't be emitted"); +} + +void EmitDemoteToHelperInvocation(EmitContext& ctx) { + ctx.Add("KIL TR.x;"); +} + +} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_logical.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_logical.cpp index e69de29bb..eed7bfec2 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_logical.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_logical.cpp @@ -0,0 +1,26 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" + +namespace Shader::Backend::GLASM { + +void EmitLogicalOr(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) { + ctx.Add("OR.S {},{},{};", inst, a, b); +} + +void EmitLogicalAnd(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) { + ctx.Add("AND.S {},{},{};", inst, a, b); +} + +void EmitLogicalXor(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) { + ctx.Add("XOR.S {},{},{};", inst, a, b); +} + +void EmitLogicalNot(EmitContext& ctx, IR::Inst& inst, ScalarS32 value) { + ctx.Add("SEQ.S {},{},0;", inst, value); +} + +} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp index b6a2161aa..86287ee3f 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp @@ -17,110 +17,6 @@ namespace Shader::Backend::GLASM { #define NotImplemented() throw NotImplementedException("GLASM instruction {}", __LINE__) -static void DefinePhi(EmitContext& ctx, IR::Inst& phi) { - switch (phi.Type()) { - case IR::Type::U1: - case IR::Type::U32: - case IR::Type::F32: - ctx.reg_alloc.Define(phi); - break; - case IR::Type::U64: - case IR::Type::F64: - ctx.reg_alloc.LongDefine(phi); - break; - default: - throw NotImplementedException("Phi node type {}", phi.Type()); - } -} - -void EmitPhi(EmitContext& ctx, IR::Inst& phi) { - const size_t num_args{phi.NumArgs()}; - for (size_t i = 0; i < num_args; ++i) { - ctx.reg_alloc.Consume(phi.Arg(i)); - } - if (!phi.Definition().is_valid) { - // The phi node wasn't forward defined - DefinePhi(ctx, phi); - } -} - -void EmitVoid(EmitContext&) {} - -void EmitReference(EmitContext& ctx, const IR::Value& value) { - ctx.reg_alloc.Consume(value); -} - -void EmitPhiMove(EmitContext& ctx, const IR::Value& phi_value, const IR::Value& value) { - IR::Inst& phi{RegAlloc::AliasInst(*phi_value.Inst())}; - if (!phi.Definition().is_valid) { - // The phi node wasn't forward defined - DefinePhi(ctx, phi); - } - const Register phi_reg{ctx.reg_alloc.Consume(IR::Value{&phi})}; - const Value eval_value{ctx.reg_alloc.Consume(value)}; - - if (phi_reg == eval_value) { - return; - } - switch (phi.Flags()) { - case IR::Type::U1: - case IR::Type::U32: - case IR::Type::F32: - ctx.Add("MOV.S {}.x,{};", phi_reg, ScalarS32{eval_value}); - break; - case IR::Type::U64: - case IR::Type::F64: - ctx.Add("MOV.U64 {}.x,{};", phi_reg, ScalarRegister{eval_value}); - break; - default: - throw NotImplementedException("Phi node type {}", phi.Type()); - } -} - -void EmitJoin(EmitContext& ctx) { - NotImplemented(); -} - -void EmitDemoteToHelperInvocation(EmitContext& ctx) { - ctx.Add("KIL TR.x;"); -} - -void EmitBarrier(EmitContext& ctx) { - ctx.Add("BAR;"); -} - -void EmitWorkgroupMemoryBarrier(EmitContext& ctx) { - ctx.Add("MEMBAR.CTA;"); -} - -void EmitDeviceMemoryBarrier(EmitContext& ctx) { - ctx.Add("MEMBAR;"); -} - -void EmitPrologue(EmitContext& ctx) { - // TODO -} - -void EmitEpilogue(EmitContext& ctx) { - // TODO -} - -void EmitEmitVertex(EmitContext& ctx, ScalarS32 stream) { - if (stream.type == Type::U32 && stream.imm_u32 == 0) { - ctx.Add("EMIT;"); - } else { - ctx.Add("EMITS {};", stream); - } -} - -void EmitEndPrimitive(EmitContext& ctx, const IR::Value& stream) { - if (!stream.IsImmediate()) { - LOG_WARNING(Shader_GLASM, "Stream is not immediate"); - } - ctx.reg_alloc.Consume(stream); - ctx.Add("ENDPRIM;"); -} - void EmitGetRegister(EmitContext& ctx) { NotImplemented(); } @@ -185,55 +81,6 @@ void EmitSetOFlag(EmitContext& ctx) { NotImplemented(); } -void EmitWorkgroupId(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {},invocation.groupid;", inst); -} - -void EmitLocalInvocationId(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {},invocation.localid;", inst); -} - -void EmitInvocationId(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {}.x,primitive_invocation.x;", inst); -} - -void EmitSampleId(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {}.x,fragment.sampleid.x;", inst); -} - -void EmitIsHelperInvocation(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {}.x,fragment.helperthread.x;", inst); -} - -void EmitYDirection(EmitContext& ctx, IR::Inst& inst) { - ctx.uses_y_direction = true; - ctx.Add("MOV.F {}.x,y_direction[0].w;", inst); -} - -void EmitResolutionDownFactor(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.F {}.x,scaling[0].z;", inst); -} - -void EmitUndefU1(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {}.x,0;", inst); -} - -void EmitUndefU8(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {}.x,0;", inst); -} - -void EmitUndefU16(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {}.x,0;", inst); -} - -void EmitUndefU32(EmitContext& ctx, IR::Inst& inst) { - ctx.Add("MOV.S {}.x,0;", inst); -} - -void EmitUndefU64(EmitContext& ctx, IR::Inst& inst) { - ctx.LongAdd("MOV.S64 {}.x,0;", inst); -} - void EmitGetZeroFromOp(EmitContext& ctx) { NotImplemented(); } @@ -258,20 +105,4 @@ void EmitGetInBoundsFromOp(EmitContext& ctx) { NotImplemented(); } -void EmitLogicalOr(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) { - ctx.Add("OR.S {},{},{};", inst, a, b); -} - -void EmitLogicalAnd(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) { - ctx.Add("AND.S {},{},{};", inst, a, b); -} - -void EmitLogicalXor(EmitContext& ctx, IR::Inst& inst, ScalarS32 a, ScalarS32 b) { - ctx.Add("XOR.S {},{},{};", inst, a, b); -} - -void EmitLogicalNot(EmitContext& ctx, IR::Inst& inst, ScalarS32 value) { - ctx.Add("SEQ.S {},{},0;", inst, value); -} - } // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_special.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_special.cpp index e69de29bb..e7a5fb13a 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_special.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_special.cpp @@ -0,0 +1,95 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" +#include "shader_recompiler/frontend/ir/value.h" + +namespace Shader::Backend::GLASM { + +static void DefinePhi(EmitContext& ctx, IR::Inst& phi) { + switch (phi.Type()) { + case IR::Type::U1: + case IR::Type::U32: + case IR::Type::F32: + ctx.reg_alloc.Define(phi); + break; + case IR::Type::U64: + case IR::Type::F64: + ctx.reg_alloc.LongDefine(phi); + break; + default: + throw NotImplementedException("Phi node type {}", phi.Type()); + } +} + +void EmitPhi(EmitContext& ctx, IR::Inst& phi) { + const size_t num_args{phi.NumArgs()}; + for (size_t i = 0; i < num_args; ++i) { + ctx.reg_alloc.Consume(phi.Arg(i)); + } + if (!phi.Definition().is_valid) { + // The phi node wasn't forward defined + DefinePhi(ctx, phi); + } +} + +void EmitVoid(EmitContext&) {} + +void EmitReference(EmitContext& ctx, const IR::Value& value) { + ctx.reg_alloc.Consume(value); +} + +void EmitPhiMove(EmitContext& ctx, const IR::Value& phi_value, const IR::Value& value) { + IR::Inst& phi{RegAlloc::AliasInst(*phi_value.Inst())}; + if (!phi.Definition().is_valid) { + // The phi node wasn't forward defined + DefinePhi(ctx, phi); + } + const Register phi_reg{ctx.reg_alloc.Consume(IR::Value{&phi})}; + const Value eval_value{ctx.reg_alloc.Consume(value)}; + + if (phi_reg == eval_value) { + return; + } + switch (phi.Flags()) { + case IR::Type::U1: + case IR::Type::U32: + case IR::Type::F32: + ctx.Add("MOV.S {}.x,{};", phi_reg, ScalarS32{eval_value}); + break; + case IR::Type::U64: + case IR::Type::F64: + ctx.Add("MOV.U64 {}.x,{};", phi_reg, ScalarRegister{eval_value}); + break; + default: + throw NotImplementedException("Phi node type {}", phi.Type()); + } +} + +void EmitPrologue(EmitContext&) { + // TODO +} + +void EmitEpilogue(EmitContext&) { + // TODO +} + +void EmitEmitVertex(EmitContext& ctx, ScalarS32 stream) { + if (stream.type == Type::U32 && stream.imm_u32 == 0) { + ctx.Add("EMIT;"); + } else { + ctx.Add("EMITS {};", stream); + } +} + +void EmitEndPrimitive(EmitContext& ctx, const IR::Value& stream) { + if (!stream.IsImmediate()) { + LOG_WARNING(Shader_GLASM, "Stream is not immediate"); + } + ctx.reg_alloc.Consume(stream); + ctx.Add("ENDPRIM;"); +} + +} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_undefined.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_undefined.cpp index e69de29bb..875e9d991 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_undefined.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_undefined.cpp @@ -0,0 +1,30 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/backend/glasm/glasm_emit_context.h" + +namespace Shader::Backend::GLASM { + +void EmitUndefU1(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {}.x,0;", inst); +} + +void EmitUndefU8(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {}.x,0;", inst); +} + +void EmitUndefU16(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {}.x,0;", inst); +} + +void EmitUndefU32(EmitContext& ctx, IR::Inst& inst) { + ctx.Add("MOV.S {}.x,0;", inst); +} + +void EmitUndefU64(EmitContext& ctx, IR::Inst& inst) { + ctx.LongAdd("MOV.S64 {}.x,0;", inst); +} + +} // namespace Shader::Backend::GLASM -- cgit v1.2.3 From 7105204a7ed0ac45701e57ff157c9e0c679c5ce9 Mon Sep 17 00:00:00 2001 From: ameerj <52414509+ameerj@users.noreply.github.com> Date: Sun, 5 Dec 2021 17:24:54 -0500 Subject: emit_spirv: Reduce emit_spirv.h include overhead emit_spirv.h is included in video_core, which was propagating further includes that video_core did not depend on. --- src/shader_recompiler/backend/spirv/emit_spirv.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv.h | 4 +--- src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_barriers.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_bitwise_conversion.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_composite.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_convert.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_image.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_image_atomic.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_integer.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_logical.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_select.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_special.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_undefined.cpp | 1 + src/shader_recompiler/backend/spirv/emit_spirv_warp.cpp | 1 + 20 files changed, 20 insertions(+), 3 deletions(-) (limited to 'src') diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index d7a86e270..6ce7ed12a 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -11,6 +11,7 @@ #include "common/settings.h" #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/basic_block.h" #include "shader_recompiler/frontend/ir/program.h" diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.h b/src/shader_recompiler/backend/spirv/emit_spirv.h index 63dea090d..b412957c7 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.h +++ b/src/shader_recompiler/backend/spirv/emit_spirv.h @@ -6,13 +6,11 @@ #include -#include - #include "common/common_types.h" #include "shader_recompiler/backend/bindings.h" -#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/program.h" #include "shader_recompiler/profile.h" +#include "shader_recompiler/runtime_info.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp index 9af8bb9e1..0d37b405c 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_barriers.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_barriers.cpp index e0b52a001..9ce95a41b 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_barriers.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_barriers.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_bitwise_conversion.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_bitwise_conversion.cpp index bb11f4f4e..02d1e63f7 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_bitwise_conversion.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_bitwise_conversion.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_composite.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_composite.cpp index 10ff4ecab..5c3e1ee2b 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_composite.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_composite.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp index bac683ae1..ad84966b5 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp @@ -7,6 +7,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp index d33486f28..1eca3aa85 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_control_flow.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_convert.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_convert.cpp index fd42b7a16..832de2452 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_convert.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_convert.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp index 61cf25f9c..0cdc46495 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_image.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_image.cpp index 4d168a96d..d18d5f1d5 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_image.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_image.cpp @@ -6,6 +6,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_image_atomic.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_image_atomic.cpp index d7f1a365a..a96190bc6 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_image_atomic.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_image_atomic.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/ir/modifiers.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_integer.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_integer.cpp index 50277eec3..44521f539 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_integer.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_integer.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_logical.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_logical.cpp index b9a9500fc..47745f7ee 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_logical.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_logical.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp index 679ee2684..175f4be19 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_memory.cpp @@ -6,6 +6,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_select.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_select.cpp index c5b4f4720..48caf1ffc 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_select.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_select.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp index 9a79fc7a2..330c9052c 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp index 9e7eb3cb1..d96a17583 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_undefined.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_undefined.cpp index c9f469e90..b5766fc52 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_undefined.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_undefined.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_warp.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_warp.cpp index cef52c56e..7034228bf 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_warp.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_warp.cpp @@ -4,6 +4,7 @@ #include "shader_recompiler/backend/spirv/emit_spirv.h" #include "shader_recompiler/backend/spirv/emit_spirv_instructions.h" +#include "shader_recompiler/backend/spirv/spirv_emit_context.h" namespace Shader::Backend::SPIRV { namespace { -- cgit v1.2.3