diff options
Diffstat (limited to 'src/shader_recompiler/backend')
9 files changed, 1167 insertions, 840 deletions
diff --git a/src/shader_recompiler/backend/glasm/emit_context.cpp b/src/shader_recompiler/backend/glasm/emit_context.cpp index b4db4ff8f..9f839f3bf 100644 --- a/src/shader_recompiler/backend/glasm/emit_context.cpp +++ b/src/shader_recompiler/backend/glasm/emit_context.cpp @@ -3,9 +3,28 @@ // Refer to the license.txt file included. #include "shader_recompiler/backend/glasm/emit_context.h" +#include "shader_recompiler/frontend/ir/program.h" namespace Shader::Backend::GLASM { -EmitContext::EmitContext() = default; +EmitContext::EmitContext(IR::Program& program) { + // FIXME: Temporary partial implementation + u32 cbuf_index{}; + for (const auto& desc : program.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; + } + for (const auto& desc : program.info.storage_buffers_descriptors) { + if (desc.count != 1) { + throw NotImplementedException("Storage buffer descriptor array"); + } + } + if (const size_t num = program.info.storage_buffers_descriptors.size(); num > 0) { + Add("PARAM c[{}]={{program.local[0..{}]}};", num, num - 1); + } +} } // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/emit_context.h b/src/shader_recompiler/backend/glasm/emit_context.h index cf66619de..4f98a9816 100644 --- a/src/shader_recompiler/backend/glasm/emit_context.h +++ b/src/shader_recompiler/backend/glasm/emit_context.h @@ -13,13 +13,14 @@ namespace Shader::IR { class Inst; -} +struct Program; +} // namespace Shader::IR namespace Shader::Backend::GLASM { class EmitContext { public: - explicit EmitContext(); + explicit EmitContext(IR::Program& program); template <typename... Args> void Add(const char* fmt, IR::Inst& inst, Args&&... args) { diff --git a/src/shader_recompiler/backend/glasm/emit_glasm.cpp b/src/shader_recompiler/backend/glasm/emit_glasm.cpp index 59d7c0f96..65600f58c 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm.cpp @@ -50,7 +50,7 @@ template <auto func, bool is_first_arg_inst, size_t... I> void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) { using Traits = FuncTraits<decltype(func)>; if constexpr (is_first_arg_inst) { - func(ctx, inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...); + func(ctx, *inst, Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...); } else { func(ctx, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...); } @@ -64,7 +64,7 @@ void Invoke(EmitContext& ctx, IR::Inst* inst) { Invoke<func, false>(ctx, inst, std::make_index_sequence<0>{}); } else { using FirstArgType = typename Traits::template ArgType<1>; - static constexpr bool is_first_arg_inst = std::is_same_v<FirstArgType, IR::Inst*>; + static constexpr bool is_first_arg_inst = std::is_same_v<FirstArgType, IR::Inst&>; using Indices = std::make_index_sequence<Traits::NUM_ARGS - (is_first_arg_inst ? 2 : 1)>; Invoke<func, is_first_arg_inst>(ctx, inst, Indices{}); } @@ -80,16 +80,76 @@ void EmitInst(EmitContext& ctx, IR::Inst* inst) { } throw LogicError("Invalid opcode {}", inst->GetOpcode()); } + +void Identity(IR::Inst& inst, const IR::Value& value) { + if (value.IsImmediate()) { + return; + } + IR::Inst* const value_inst{value.InstRecursive()}; + if (inst.GetOpcode() == IR::Opcode::Identity) { + value_inst->DestructiveAddUsage(inst.UseCount()); + value_inst->DestructiveRemoveUsage(); + } + inst.SetDefinition(value_inst->Definition<Id>()); +} } // Anonymous namespace std::string EmitGLASM(const Profile&, IR::Program& program, Bindings&) { - EmitContext ctx; + EmitContext ctx{program}; for (IR::Block* const block : program.blocks) { for (IR::Inst& inst : block->Instructions()) { EmitInst(ctx, &inst); } } + std::string header = "!!NVcp5.0\n" + "OPTION NV_internal;"; + switch (program.stage) { + case Stage::Compute: + header += fmt::format("GROUP_SIZE {} {} {};", program.workgroup_size[0], + program.workgroup_size[1], program.workgroup_size[2]); + break; + default: + break; + } + header += "TEMP "; + for (size_t index = 0; index < ctx.reg_alloc.NumUsedRegisters(); ++index) { + header += fmt::format("R{},", index); + } + header += "RC;"; + if (!program.info.storage_buffers_descriptors.empty()) { + header += "LONG TEMP LC;"; + } + ctx.code.insert(0, header); + ctx.code += "END"; return ctx.code; } +void EmitIdentity(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) { + Identity(inst, value); +} + +void EmitBitCastU16F16(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) { + Identity(inst, value); +} + +void EmitBitCastU32F32(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) { + Identity(inst, value); +} + +void EmitBitCastU64F64(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) { + Identity(inst, value); +} + +void EmitBitCastF16U16(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) { + Identity(inst, value); +} + +void EmitBitCastF32U32(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) { + Identity(inst, value); +} + +void EmitBitCastF64U64(EmitContext& ctx, IR::Inst& inst, const IR::Value& value) { + Identity(inst, value); +} + } // 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 e69de29bb..72733d1cf 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 @@ -0,0 +1,125 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include <string_view> + +#include "shader_recompiler/backend/glasm/emit_context.h" +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/frontend/ir/value.h" + +namespace Shader::Backend::GLASM { +namespace { +void GetCbuf(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, const IR::Value& offset, + std::string_view size) { + if (!binding.IsImmediate()) { + throw NotImplementedException("Indirect constant buffer loading"); + } + const std::string ret{ctx.reg_alloc.Define(inst)}; + ctx.Add("LDC.{} {},c{}[{}];", size, ret, binding.U32(), ctx.reg_alloc.Consume(offset)); +} +} // Anonymous namespace + +void EmitGetCbufU8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset) { + GetCbuf(ctx, inst, binding, offset, "U8"); +} + +void EmitGetCbufS8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset) { + GetCbuf(ctx, inst, binding, offset, "S8"); +} + +void EmitGetCbufU16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset) { + GetCbuf(ctx, inst, binding, offset, "U16"); +} + +void EmitGetCbufS16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset) { + GetCbuf(ctx, inst, binding, offset, "S16"); +} + +void EmitGetCbufU32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset) { + GetCbuf(ctx, inst, binding, offset, "U32"); +} + +void EmitGetCbufF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset) { + GetCbuf(ctx, inst, binding, offset, "F32"); +} + +void EmitGetCbufU32x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset) { + GetCbuf(ctx, inst, binding, offset, "U32X2"); +} + +void EmitGetAttribute(EmitContext& ctx, IR::Inst& inst, IR::Attribute attr, + [[maybe_unused]] std::string_view vertex) { + if (IR::IsGeneric(attr)) { + const u32 index{IR::GenericAttributeIndex(attr)}; + const u32 element{IR::GenericAttributeElement(attr)}; + ctx.Add("MOV.F {},in_attr{}.{};", inst, index, "xyzw"[element]); + return; + } + throw NotImplementedException("Get attribute {}", attr); +} + +void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view value, + [[maybe_unused]] std::string_view vertex) { + const u32 element{static_cast<u32>(attr) % 4}; + const char swizzle{"xyzw"[element]}; + if (IR::IsGeneric(attr)) { + const u32 index{IR::GenericAttributeIndex(attr)}; + ctx.Add("MOV.F out_attr{}.{},{};", index, swizzle, value); + return; + } + switch (attr) { + case IR::Attribute::PositionX: + case IR::Attribute::PositionY: + case IR::Attribute::PositionZ: + case IR::Attribute::PositionW: + ctx.Add("MOV.F result.position.{},{};", swizzle, value); + break; + default: + throw NotImplementedException("Set attribute {}", attr); + } +} + +void EmitGetAttributeIndexed([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view offset, + [[maybe_unused]] std::string_view vertex) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSetAttributeIndexed([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view offset, + [[maybe_unused]] std::string_view value, + [[maybe_unused]] std::string_view vertex) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitGetPatch([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Patch patch) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSetPatch([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Patch patch, + [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSetFragColor([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] u32 index, + [[maybe_unused]] u32 component, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSetSampleMask([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSetFragDepth([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +} // namespace Shader::Backend::GLASM 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 e69de29bb..db9dda261 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_floating_point.cpp @@ -0,0 +1,421 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include <string_view> + +#include "shader_recompiler/backend/glasm/emit_context.h" +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/frontend/ir/value.h" + +namespace Shader::Backend::GLASM { + +void EmitFPAbs16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPAbs32(EmitContext& ctx, IR::Inst& inst, std::string_view value) { + ctx.Add("MOV.F {},|{}|;", inst, value); +} + +void EmitFPAbs64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPAdd16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPAdd32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b) { + ctx.Add("ADD.F {},{},{};", inst, a, b); +} + +void EmitFPAdd64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPFma16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b, + [[maybe_unused]] std::string_view c) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPFma32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b, + std::string_view c) { + ctx.Add("MAD.F {},{},{},{};", inst, a, b, c); +} + +void EmitFPFma64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b, + [[maybe_unused]] std::string_view c) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPMax32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPMax64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPMin32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPMin64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPMul16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPMul32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b) { + ctx.Add("MUL.F {},{},{};", inst, a, b); +} + +void EmitFPMul64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPNeg16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPNeg32(EmitContext& ctx, IR::Inst& inst, std::string_view value) { + if (value[0] == '-') { + // Guard against negating a negative immediate + ctx.Add("MOV.F {},{};", inst, value.substr(1)); + } else { + ctx.Add("MOV.F {},-{};", inst, value); + } +} + +void EmitFPNeg64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPSin([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPCos([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPExp2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPLog2([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPRecip32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPRecip64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPRecipSqrt32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPRecipSqrt64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPSqrt([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPSaturate16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPSaturate32(EmitContext& ctx, IR::Inst& inst, std::string_view value) { + ctx.Add("MOV.F.SAT {},{};", inst, value); +} + +void EmitFPSaturate64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPClamp16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value, + [[maybe_unused]] std::string_view min_value, + [[maybe_unused]] std::string_view max_value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPClamp32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value, + [[maybe_unused]] std::string_view min_value, + [[maybe_unused]] std::string_view max_value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPClamp64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value, + [[maybe_unused]] std::string_view min_value, + [[maybe_unused]] std::string_view max_value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPRoundEven16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPRoundEven32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPRoundEven64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPFloor16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPFloor32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPFloor64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPCeil16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPCeil32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPCeil64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPTrunc16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPTrunc32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPTrunc64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdEqual16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdEqual32(EmitContext& ctx, IR::Inst& inst, std::string_view lhs, + std::string_view rhs) { + const std::string ret{ctx.reg_alloc.Define(inst)}; + ctx.Add("SEQ.F {},{},{};SNE.S {},{},0;", ret, lhs, rhs, ret, ret); +} + +void EmitFPOrdEqual64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordEqual16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordEqual32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordEqual64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdNotEqual16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdNotEqual32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdNotEqual64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordNotEqual16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordNotEqual32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordNotEqual64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdLessThan16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdLessThan32(EmitContext& ctx, IR::Inst& inst, std::string_view lhs, + std::string_view rhs) { + const std::string ret{ctx.reg_alloc.Define(inst)}; + ctx.Add("SLT.F {},{},{};SNE.S {},{},0;", ret, lhs, rhs, ret, ret); +} + +void EmitFPOrdLessThan64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordLessThan16([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordLessThan32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordLessThan64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdGreaterThan16([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdGreaterThan32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdGreaterThan64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordGreaterThan16([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordGreaterThan32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordGreaterThan64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdLessThanEqual16([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdLessThanEqual32(EmitContext& ctx, IR::Inst& inst, std::string_view lhs, + std::string_view rhs) { + const std::string ret{ctx.reg_alloc.Define(inst)}; + ctx.Add("SLE.F {},{},{};SNE.S {},{},0;", ret, lhs, rhs, ret, ret); +} + +void EmitFPOrdLessThanEqual64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordLessThanEqual16([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordLessThanEqual32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordLessThanEqual64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdGreaterThanEqual16([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdGreaterThanEqual32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPOrdGreaterThanEqual64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordGreaterThanEqual16([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordGreaterThanEqual32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFPUnordGreaterThanEqual64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +} // namespace Shader::Backend::GLASM diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_instructions.h b/src/shader_recompiler/backend/glasm/emit_glasm_instructions.h index 21d6af914..30cc6c2eb 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_instructions.h +++ b/src/shader_recompiler/backend/glasm/emit_glasm_instructions.h @@ -20,9 +20,9 @@ namespace Shader::Backend::GLASM { class EmitContext; // Microinstruction emitters -void EmitPhi(EmitContext& ctx, IR::Inst* inst); +void EmitPhi(EmitContext& ctx, IR::Inst& inst); void EmitVoid(EmitContext& ctx); -void EmitIdentity(EmitContext& ctx, const IR::Value& value); +void EmitIdentity(EmitContext& ctx, IR::Inst& inst, const IR::Value& value); void EmitBranch(EmitContext& ctx, std::string_view label); void EmitBranchConditional(EmitContext& ctx, std::string_view condition, std::string_view true_label, std::string_view false_label); @@ -47,14 +47,22 @@ void EmitSetGotoVariable(EmitContext& ctx); void EmitGetGotoVariable(EmitContext& ctx); void EmitSetIndirectBranchVariable(EmitContext& ctx); void EmitGetIndirectBranchVariable(EmitContext& ctx); -void EmitGetCbufU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitGetCbufS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitGetCbufU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitGetCbufS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitGetCbufU32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitGetCbufF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitGetCbufU32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitGetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view vertex); +void EmitGetCbufU8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset); +void EmitGetCbufS8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset); +void EmitGetCbufU16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset); +void EmitGetCbufS16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset); +void EmitGetCbufU32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset); +void EmitGetCbufF32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset); +void EmitGetCbufU32x2(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + const IR::Value& offset); +void EmitGetAttribute(EmitContext& ctx, IR::Inst& inst, IR::Attribute attr, + std::string_view vertex); void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view value, std::string_view vertex); void EmitGetAttributeIndexed(EmitContext& ctx, std::string_view offset, std::string_view vertex); @@ -100,26 +108,33 @@ void EmitWriteGlobalS16(EmitContext& ctx); void EmitWriteGlobal32(EmitContext& ctx, std::string_view address, std::string_view value); void EmitWriteGlobal64(EmitContext& ctx, std::string_view address, std::string_view value); void EmitWriteGlobal128(EmitContext& ctx, std::string_view address, std::string_view value); -void EmitLoadStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitLoadStorageS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitLoadStorageU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitLoadStorageS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitLoadStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitLoadStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset); -void EmitWriteStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, +void EmitLoadStorageU8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset); +void EmitLoadStorageS8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset); +void EmitLoadStorageU16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset); +void EmitLoadStorageS16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset); +void EmitLoadStorage32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset); +void EmitLoadStorage64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset); +void EmitLoadStorage128(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset); +void EmitWriteStorageU8(EmitContext& ctx, const IR::Value& binding, std::string_view offset, std::string_view value); -void EmitWriteStorageS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, +void EmitWriteStorageS8(EmitContext& ctx, const IR::Value& binding, std::string_view offset, std::string_view value); -void EmitWriteStorageU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, +void EmitWriteStorageU16(EmitContext& ctx, const IR::Value& binding, std::string_view offset, std::string_view value); -void EmitWriteStorageS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, +void EmitWriteStorageS16(EmitContext& ctx, const IR::Value& binding, std::string_view offset, std::string_view value); -void EmitWriteStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, +void EmitWriteStorage32(EmitContext& ctx, const IR::Value& binding, std::string_view offset, std::string_view value); -void EmitWriteStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, +void EmitWriteStorage64(EmitContext& ctx, const IR::Value& binding, std::string_view offset, std::string_view value); -void EmitWriteStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, +void EmitWriteStorage128(EmitContext& ctx, const IR::Value& binding, std::string_view offset, std::string_view value); void EmitLoadSharedU8(EmitContext& ctx, std::string_view offset); void EmitLoadSharedS8(EmitContext& ctx, std::string_view offset); @@ -203,12 +218,12 @@ void EmitSelectF32(EmitContext& ctx, std::string_view cond, std::string_view tru std::string_view false_value); void EmitSelectF64(EmitContext& ctx, std::string_view cond, std::string_view true_value, std::string_view false_value); -void EmitBitCastU16F16(EmitContext& ctx); -void EmitBitCastU32F32(EmitContext& ctx, std::string_view value); -void EmitBitCastU64F64(EmitContext& ctx); -void EmitBitCastF16U16(EmitContext& ctx); -void EmitBitCastF32U32(EmitContext& ctx, std::string_view value); -void EmitBitCastF64U64(EmitContext& ctx); +void EmitBitCastU16F16(EmitContext& ctx, IR::Inst& inst, const IR::Value& value); +void EmitBitCastU32F32(EmitContext& ctx, IR::Inst& inst, const IR::Value& value); +void EmitBitCastU64F64(EmitContext& ctx, IR::Inst& inst, const IR::Value& value); +void EmitBitCastF16U16(EmitContext& ctx, IR::Inst& inst, const IR::Value& value); +void EmitBitCastF32U32(EmitContext& ctx, IR::Inst& inst, const IR::Value& value); +void EmitBitCastF64U64(EmitContext& ctx, IR::Inst& inst, const IR::Value& value); void EmitPackUint2x32(EmitContext& ctx, std::string_view value); void EmitUnpackUint2x32(EmitContext& ctx, std::string_view value); void EmitPackFloat2x16(EmitContext& ctx, std::string_view value); @@ -224,26 +239,26 @@ void EmitGetOverflowFromOp(EmitContext& ctx); void EmitGetSparseFromOp(EmitContext& ctx); void EmitGetInBoundsFromOp(EmitContext& ctx); void EmitFPAbs16(EmitContext& ctx, std::string_view value); -void EmitFPAbs32(EmitContext& ctx, std::string_view value); +void EmitFPAbs32(EmitContext& ctx, IR::Inst& inst, std::string_view value); void EmitFPAbs64(EmitContext& ctx, std::string_view value); -void EmitFPAdd16(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); -void EmitFPAdd32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); -void EmitFPAdd64(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); -void EmitFPFma16(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b, +void EmitFPAdd16(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); +void EmitFPAdd32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); +void EmitFPAdd64(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); +void EmitFPFma16(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b, std::string_view c); -void EmitFPFma32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b, +void EmitFPFma32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b, std::string_view c); -void EmitFPFma64(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b, +void EmitFPFma64(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b, std::string_view c); void EmitFPMax32(EmitContext& ctx, std::string_view a, std::string_view b); void EmitFPMax64(EmitContext& ctx, std::string_view a, std::string_view b); void EmitFPMin32(EmitContext& ctx, std::string_view a, std::string_view b); void EmitFPMin64(EmitContext& ctx, std::string_view a, std::string_view b); -void EmitFPMul16(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); -void EmitFPMul32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); -void EmitFPMul64(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); +void EmitFPMul16(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); +void EmitFPMul32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); +void EmitFPMul64(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); void EmitFPNeg16(EmitContext& ctx, std::string_view value); -void EmitFPNeg32(EmitContext& ctx, std::string_view value); +void EmitFPNeg32(EmitContext& ctx, IR::Inst& inst, std::string_view value); void EmitFPNeg64(EmitContext& ctx, std::string_view value); void EmitFPSin(EmitContext& ctx, std::string_view value); void EmitFPCos(EmitContext& ctx, std::string_view value); @@ -255,7 +270,7 @@ void EmitFPRecipSqrt32(EmitContext& ctx, std::string_view value); void EmitFPRecipSqrt64(EmitContext& ctx, std::string_view value); void EmitFPSqrt(EmitContext& ctx, std::string_view value); void EmitFPSaturate16(EmitContext& ctx, std::string_view value); -void EmitFPSaturate32(EmitContext& ctx, std::string_view value); +void EmitFPSaturate32(EmitContext& ctx, IR::Inst& inst, std::string_view value); void EmitFPSaturate64(EmitContext& ctx, std::string_view value); void EmitFPClamp16(EmitContext& ctx, std::string_view value, std::string_view min_value, std::string_view max_value); @@ -276,7 +291,7 @@ void EmitFPTrunc16(EmitContext& ctx, std::string_view value); void EmitFPTrunc32(EmitContext& ctx, std::string_view value); void EmitFPTrunc64(EmitContext& ctx, std::string_view value); void EmitFPOrdEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs); -void EmitFPOrdEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs); +void EmitFPOrdEqual32(EmitContext& ctx, IR::Inst& inst, std::string_view lhs, std::string_view rhs); void EmitFPOrdEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPUnordEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPUnordEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs); @@ -288,7 +303,8 @@ void EmitFPUnordNotEqual16(EmitContext& ctx, std::string_view lhs, std::string_v void EmitFPUnordNotEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPUnordNotEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPOrdLessThan16(EmitContext& ctx, std::string_view lhs, std::string_view rhs); -void EmitFPOrdLessThan32(EmitContext& ctx, std::string_view lhs, std::string_view rhs); +void EmitFPOrdLessThan32(EmitContext& ctx, IR::Inst& inst, std::string_view lhs, + std::string_view rhs); void EmitFPOrdLessThan64(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPUnordLessThan16(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPUnordLessThan32(EmitContext& ctx, std::string_view lhs, std::string_view rhs); @@ -300,7 +316,8 @@ void EmitFPUnordGreaterThan16(EmitContext& ctx, std::string_view lhs, std::strin void EmitFPUnordGreaterThan32(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPUnordGreaterThan64(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPOrdLessThanEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs); -void EmitFPOrdLessThanEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs); +void EmitFPOrdLessThanEqual32(EmitContext& ctx, IR::Inst& inst, std::string_view lhs, + std::string_view rhs); void EmitFPOrdLessThanEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPUnordLessThanEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitFPUnordLessThanEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs); @@ -314,7 +331,7 @@ void EmitFPUnordGreaterThanEqual64(EmitContext& ctx, std::string_view lhs, std:: void EmitFPIsNan16(EmitContext& ctx, std::string_view value); void EmitFPIsNan32(EmitContext& ctx, std::string_view value); void EmitFPIsNan64(EmitContext& ctx, std::string_view value); -void EmitIAdd32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); +void EmitIAdd32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); void EmitIAdd64(EmitContext& ctx, std::string_view a, std::string_view b); void EmitISub32(EmitContext& ctx, std::string_view a, std::string_view b); void EmitISub64(EmitContext& ctx, std::string_view a, std::string_view b); @@ -329,14 +346,14 @@ void EmitShiftRightLogical32(EmitContext& ctx, std::string_view base, std::strin void EmitShiftRightLogical64(EmitContext& ctx, std::string_view base, std::string_view shift); void EmitShiftRightArithmetic32(EmitContext& ctx, std::string_view base, std::string_view shift); void EmitShiftRightArithmetic64(EmitContext& ctx, std::string_view base, std::string_view shift); -void EmitBitwiseAnd32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); -void EmitBitwiseOr32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); -void EmitBitwiseXor32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b); +void EmitBitwiseAnd32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); +void EmitBitwiseOr32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); +void EmitBitwiseXor32(EmitContext& ctx, IR::Inst& inst, std::string_view a, std::string_view b); void EmitBitFieldInsert(EmitContext& ctx, std::string_view base, std::string_view insert, std::string_view offset, std::string_view count); -void EmitBitFieldSExtract(EmitContext& ctx, IR::Inst* inst, std::string_view base, +void EmitBitFieldSExtract(EmitContext& ctx, IR::Inst& inst, std::string_view base, std::string_view offset, std::string_view count); -void EmitBitFieldUExtract(EmitContext& ctx, IR::Inst* inst, std::string_view base, +void EmitBitFieldUExtract(EmitContext& ctx, IR::Inst& inst, std::string_view base, std::string_view offset, std::string_view count); void EmitBitReverse32(EmitContext& ctx, std::string_view value); void EmitBitCount32(EmitContext& ctx, std::string_view value); @@ -347,9 +364,9 @@ void EmitSMin32(EmitContext& ctx, std::string_view a, std::string_view b); void EmitUMin32(EmitContext& ctx, std::string_view a, std::string_view b); void EmitSMax32(EmitContext& ctx, std::string_view a, std::string_view b); void EmitUMax32(EmitContext& ctx, std::string_view a, std::string_view b); -void EmitSClamp32(EmitContext& ctx, IR::Inst* inst, std::string_view value, std::string_view min, +void EmitSClamp32(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view min, std::string_view max); -void EmitUClamp32(EmitContext& ctx, IR::Inst* inst, std::string_view value, std::string_view min, +void EmitUClamp32(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view min, std::string_view max); void EmitSLessThan(EmitContext& ctx, std::string_view lhs, std::string_view rhs); void EmitULessThan(EmitContext& ctx, std::string_view lhs, std::string_view rhs); @@ -544,36 +561,36 @@ void EmitBoundImageQueryLod(EmitContext&); void EmitBoundImageGradient(EmitContext&); void EmitBoundImageRead(EmitContext&); void EmitBoundImageWrite(EmitContext&); -void EmitImageSampleImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageSampleImplicitLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view bias_lc, const IR::Value& offset); -void EmitImageSampleExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageSampleExplicitLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view lod_lc, const IR::Value& offset); -void EmitImageSampleDrefImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageSampleDrefImplicitLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view dref, std::string_view bias_lc, const IR::Value& offset); -void EmitImageSampleDrefExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageSampleDrefExplicitLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view dref, std::string_view lod_lc, const IR::Value& offset); -void EmitImageGather(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageGather(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, const IR::Value& offset, const IR::Value& offset2); -void EmitImageGatherDref(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageGatherDref(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, const IR::Value& offset, const IR::Value& offset2, std::string_view dref); -void EmitImageFetch(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageFetch(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view offset, std::string_view lod, std::string_view ms); -void EmitImageQueryDimensions(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageQueryDimensions(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view lod); -void EmitImageQueryLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageQueryLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords); -void EmitImageGradient(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageGradient(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view derivates, std::string_view offset, std::string_view lod_clamp); -void EmitImageRead(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageRead(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords); -void EmitImageWrite(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageWrite(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view color); void EmitBindlessImageAtomicIAdd32(EmitContext&); void EmitBindlessImageAtomicSMin32(EmitContext&); @@ -597,27 +614,27 @@ void EmitBoundImageAtomicAnd32(EmitContext&); void EmitBoundImageAtomicOr32(EmitContext&); void EmitBoundImageAtomicXor32(EmitContext&); void EmitBoundImageAtomicExchange32(EmitContext&); -void EmitImageAtomicIAdd32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicIAdd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicSMin32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicSMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicUMin32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicUMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicSMax32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicSMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicUMax32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicUMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicInc32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicInc32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicDec32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicDec32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicAnd32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicAnd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicOr32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicOr32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicXor32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicXor32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); -void EmitImageAtomicExchange32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicExchange32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value); void EmitLaneId(EmitContext& ctx); void EmitVoteAll(EmitContext& ctx, std::string_view pred); @@ -629,15 +646,15 @@ void EmitSubgroupLtMask(EmitContext& ctx); void EmitSubgroupLeMask(EmitContext& ctx); void EmitSubgroupGtMask(EmitContext& ctx); void EmitSubgroupGeMask(EmitContext& ctx); -void EmitShuffleIndex(EmitContext& ctx, IR::Inst* inst, std::string_view value, +void EmitShuffleIndex(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index, std::string_view clamp, std::string_view segmentation_mask); -void EmitShuffleUp(EmitContext& ctx, IR::Inst* inst, std::string_view value, std::string_view index, +void EmitShuffleUp(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index, std::string_view clamp, std::string_view segmentation_mask); -void EmitShuffleDown(EmitContext& ctx, IR::Inst* inst, std::string_view value, +void EmitShuffleDown(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index, std::string_view clamp, std::string_view segmentation_mask); -void EmitShuffleButterfly(EmitContext& ctx, IR::Inst* inst, std::string_view value, +void EmitShuffleButterfly(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index, std::string_view clamp, std::string_view segmentation_mask); void EmitFSwizzleAdd(EmitContext& ctx, std::string_view op_a, std::string_view op_b, diff --git a/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp b/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp index e69de29bb..e228fa072 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_integer.cpp @@ -0,0 +1,228 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include <string_view> + +#include "shader_recompiler/backend/glasm/emit_context.h" +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/frontend/ir/value.h" + +namespace Shader::Backend::GLASM { + +void EmitIAdd32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitIAdd64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitISub32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitISub64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitIMul32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitINeg32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitINeg64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitIAbs32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitIAbs64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitShiftLeftLogical32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view shift) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitShiftLeftLogical64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view shift) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitShiftRightLogical32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view shift) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitShiftRightLogical64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view shift) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitShiftRightArithmetic32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view shift) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitShiftRightArithmetic64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view shift) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitwiseAnd32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitwiseOr32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitwiseXor32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view a, [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitFieldInsert([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view insert, + [[maybe_unused]] std::string_view offset, + [[maybe_unused]] std::string_view count) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitFieldSExtract([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view offset, + [[maybe_unused]] std::string_view count) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitFieldUExtract([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view base, + [[maybe_unused]] std::string_view offset, + [[maybe_unused]] std::string_view count) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitReverse32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitCount32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitBitwiseNot32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFindSMsb32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitFindUMsb32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSMin32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitUMin32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSMax32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitUMax32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view a, + [[maybe_unused]] std::string_view b) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSClamp32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view value, [[maybe_unused]] std::string_view min, + [[maybe_unused]] std::string_view max) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitUClamp32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] IR::Inst& inst, + [[maybe_unused]] std::string_view value, [[maybe_unused]] std::string_view min, + [[maybe_unused]] std::string_view max) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSLessThan([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitULessThan([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitIEqual([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSLessThanEqual([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitULessThanEqual([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSGreaterThan([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitUGreaterThan([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitINotEqual([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitSGreaterThanEqual([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitUGreaterThanEqual([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view lhs, + [[maybe_unused]] std::string_view rhs) { + throw NotImplementedException("GLASM instruction"); +} + +} // 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 e69de29bb..9e38a1bdf 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_memory.cpp @@ -0,0 +1,178 @@ +// Copyright 2021 yuzu Emulator Project +// Licensed under GPLv2 or any later version +// Refer to the license.txt file included. + +#include <string_view> + +#include "shader_recompiler/backend/glasm/emit_context.h" +#include "shader_recompiler/backend/glasm/emit_glasm_instructions.h" +#include "shader_recompiler/frontend/ir/program.h" +#include "shader_recompiler/frontend/ir/value.h" + +namespace Shader::Backend::GLASM { +namespace { +void StorageOp(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view then_expr, std::string_view else_expr = {}) { + // Operate on bindless SSBO, call the expression with bounds checking + // address = c[binding].xy + // length = c[binding].z + const u32 sb_binding{binding.U32()}; + ctx.Add("PK64.U LC,c[{}];" // pointer = address + "CVT.U64.U32 LC.z,{};" // offset = uint64_t(offset) + "ADD.U64 LC.x,LC.x,LC.z;" // pointer += offset + "SLT.U.CC RC.x,{},c[{}].z;", // cc = offset < length + sb_binding, offset, offset, sb_binding); + if (else_expr.empty()) { + ctx.Add("{}", then_expr); + } else { + ctx.Add("IF NE.x;{}ELSE;{}ENDIF;", then_expr, else_expr); + } +} + +void Store(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view value, std::string_view size) { + StorageOp(ctx, binding, offset, fmt::format("STORE.{} {},LC.x;", size, value)); +} + +void Load(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, std::string_view offset, + std::string_view size) { + const std::string ret{ctx.reg_alloc.Define(inst)}; + StorageOp(ctx, binding, offset, fmt::format("STORE.{} {},LC.x;", size, ret), + fmt::format("MOV.U {},{{0,0,0,0}};", ret)); +} +} // Anonymous namespace + +void EmitLoadGlobalU8([[maybe_unused]] EmitContext& ctx) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitLoadGlobalS8([[maybe_unused]] EmitContext& ctx) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitLoadGlobalU16([[maybe_unused]] EmitContext& ctx) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitLoadGlobalS16([[maybe_unused]] EmitContext& ctx) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitLoadGlobal32([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view address) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitLoadGlobal64([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view address) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitLoadGlobal128([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view address) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitWriteGlobalU8([[maybe_unused]] EmitContext& ctx) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitWriteGlobalS8([[maybe_unused]] EmitContext& ctx) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitWriteGlobalU16([[maybe_unused]] EmitContext& ctx) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitWriteGlobalS16([[maybe_unused]] EmitContext& ctx) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitWriteGlobal32([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view address, + [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitWriteGlobal64([[maybe_unused]] EmitContext& ctx, [[maybe_unused]] std::string_view address, + [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitWriteGlobal128([[maybe_unused]] EmitContext& ctx, + [[maybe_unused]] std::string_view address, + [[maybe_unused]] std::string_view value) { + throw NotImplementedException("GLASM instruction"); +} + +void EmitLoadStorageU8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset) { + Load(ctx, inst, binding, offset, "U8"); +} + +void EmitLoadStorageS8(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset) { + Load(ctx, inst, binding, offset, "S8"); +} + +void EmitLoadStorageU16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset) { + Load(ctx, inst, binding, offset, "U16"); +} + +void EmitLoadStorageS16(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset) { + Load(ctx, inst, binding, offset, "S16"); +} + +void EmitLoadStorage32(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset) { + Load(ctx, inst, binding, offset, "U32"); +} + +void EmitLoadStorage64(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset) { + Load(ctx, inst, binding, offset, "U32X2"); +} + +void EmitLoadStorage128(EmitContext& ctx, IR::Inst& inst, const IR::Value& binding, + std::string_view offset) { + Load(ctx, inst, binding, offset, "U32X4"); +} + +void EmitWriteStorageU8(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view value) { + Store(ctx, binding, offset, value, "U8"); +} + +void EmitWriteStorageS8(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view value) { + Store(ctx, binding, offset, value, "S8"); +} + +void EmitWriteStorageU16(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view value) { + Store(ctx, binding, offset, value, "U16"); +} + +void EmitWriteStorageS16(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view value) { + Store(ctx, binding, offset, value, "S16"); +} + +void EmitWriteStorage32(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view value) { + Store(ctx, binding, offset, value, "U32"); +} + +void EmitWriteStorage64(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view value) { + Store(ctx, binding, offset, value, "U32X2"); +} + +void EmitWriteStorage128(EmitContext& ctx, const IR::Value& binding, std::string_view offset, + std::string_view value) { + Store(ctx, binding, offset, value, "U32X4"); +} + +} // 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 e90224e15..1337f4ae8 100644 --- a/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp +++ b/src/shader_recompiler/backend/glasm/emit_glasm_not_implemented.cpp @@ -15,11 +15,9 @@ namespace Shader::Backend::GLASM { -static void NotImplemented() { - throw NotImplementedException("GLASM instruction"); -} +#define NotImplemented() throw NotImplementedException("GLASM instruction {}", __LINE__) -void EmitPhi(EmitContext& ctx, IR::Inst* inst) { +void EmitPhi(EmitContext& ctx, IR::Inst& inst) { NotImplemented(); } @@ -27,10 +25,6 @@ void EmitVoid(EmitContext& ctx) { NotImplemented(); } -void EmitIdentity(EmitContext& ctx, const IR::Value& value) { - NotImplemented(); -} - void EmitBranch(EmitContext& ctx, std::string_view label) { NotImplemented(); } @@ -50,7 +44,7 @@ void EmitSelectionMerge(EmitContext& ctx, std::string_view merge_label) { } void EmitReturn(EmitContext& ctx) { - NotImplemented(); + ctx.Add("RET;"); } void EmitJoin(EmitContext& ctx) { @@ -78,11 +72,11 @@ void EmitDeviceMemoryBarrier(EmitContext& ctx) { } void EmitPrologue(EmitContext& ctx) { - NotImplemented(); + // TODO } void EmitEpilogue(EmitContext& ctx) { - NotImplemented(); + // TODO } void EmitEmitVertex(EmitContext& ctx, const IR::Value& stream) { @@ -125,72 +119,6 @@ void EmitGetIndirectBranchVariable(EmitContext& ctx) { NotImplemented(); } -void EmitGetCbufU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitGetCbufS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitGetCbufU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitGetCbufS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitGetCbufU32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitGetCbufF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitGetCbufU32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitGetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view vertex) { - NotImplemented(); -} - -void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, std::string_view value, - std::string_view vertex) { - NotImplemented(); -} - -void EmitGetAttributeIndexed(EmitContext& ctx, std::string_view offset, std::string_view vertex) { - NotImplemented(); -} - -void EmitSetAttributeIndexed(EmitContext& ctx, std::string_view offset, std::string_view value, - std::string_view vertex) { - NotImplemented(); -} - -void EmitGetPatch(EmitContext& ctx, IR::Patch patch) { - NotImplemented(); -} - -void EmitSetPatch(EmitContext& ctx, IR::Patch patch, std::string_view value) { - NotImplemented(); -} - -void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, std::string_view value) { - NotImplemented(); -} - -void EmitSetSampleMask(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitSetFragDepth(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - void EmitGetZFlag(EmitContext& ctx) { NotImplemented(); } @@ -275,125 +203,6 @@ void EmitUndefU64(EmitContext& ctx) { NotImplemented(); } -void EmitLoadGlobalU8(EmitContext& ctx) { - NotImplemented(); -} - -void EmitLoadGlobalS8(EmitContext& ctx) { - NotImplemented(); -} - -void EmitLoadGlobalU16(EmitContext& ctx) { - NotImplemented(); -} - -void EmitLoadGlobalS16(EmitContext& ctx) { - NotImplemented(); -} - -void EmitLoadGlobal32(EmitContext& ctx, std::string_view address) { - NotImplemented(); -} - -void EmitLoadGlobal64(EmitContext& ctx, std::string_view address) { - NotImplemented(); -} - -void EmitLoadGlobal128(EmitContext& ctx, std::string_view address) { - NotImplemented(); -} - -void EmitWriteGlobalU8(EmitContext& ctx) { - NotImplemented(); -} - -void EmitWriteGlobalS8(EmitContext& ctx) { - NotImplemented(); -} - -void EmitWriteGlobalU16(EmitContext& ctx) { - NotImplemented(); -} - -void EmitWriteGlobalS16(EmitContext& ctx) { - NotImplemented(); -} - -void EmitWriteGlobal32(EmitContext& ctx, std::string_view address, std::string_view value) { - NotImplemented(); -} - -void EmitWriteGlobal64(EmitContext& ctx, std::string_view address, std::string_view value) { - NotImplemented(); -} - -void EmitWriteGlobal128(EmitContext& ctx, std::string_view address, std::string_view value) { - NotImplemented(); -} - -void EmitLoadStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitLoadStorageS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitLoadStorageU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitLoadStorageS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitLoadStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitLoadStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitLoadStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset) { - NotImplemented(); -} - -void EmitWriteStorageU8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, - std::string_view value) { - NotImplemented(); -} - -void EmitWriteStorageS8(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, - std::string_view value) { - NotImplemented(); -} - -void EmitWriteStorageU16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, - std::string_view value) { - NotImplemented(); -} - -void EmitWriteStorageS16(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, - std::string_view value) { - NotImplemented(); -} - -void EmitWriteStorage32(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, - std::string_view value) { - NotImplemented(); -} - -void EmitWriteStorage64(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, - std::string_view value) { - NotImplemented(); -} - -void EmitWriteStorage128(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, - std::string_view value) { - NotImplemented(); -} - void EmitLoadSharedU8(EmitContext& ctx, std::string_view offset) { NotImplemented(); } @@ -644,30 +453,6 @@ void EmitSelectF64(EmitContext& ctx, std::string_view cond, std::string_view tru NotImplemented(); } -void EmitBitCastU16F16(EmitContext& ctx) { - NotImplemented(); -} - -void EmitBitCastU32F32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitBitCastU64F64(EmitContext& ctx) { - NotImplemented(); -} - -void EmitBitCastF16U16(EmitContext& ctx) { - NotImplemented(); -} - -void EmitBitCastF32U32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitBitCastF64U64(EmitContext& ctx) { - NotImplemented(); -} - void EmitPackUint2x32(EmitContext& ctx, std::string_view value) { NotImplemented(); } @@ -724,340 +509,6 @@ void EmitGetInBoundsFromOp(EmitContext& ctx) { NotImplemented(); } -void EmitFPAbs16(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPAbs32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPAbs64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPAdd16(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPAdd32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPAdd64(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPFma16(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b, - std::string_view c) { - NotImplemented(); -} - -void EmitFPFma32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b, - std::string_view c) { - NotImplemented(); -} - -void EmitFPFma64(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b, - std::string_view c) { - NotImplemented(); -} - -void EmitFPMax32(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPMax64(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPMin32(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPMin64(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPMul16(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPMul32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPMul64(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitFPNeg16(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPNeg32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPNeg64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPSin(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPCos(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPExp2(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPLog2(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPRecip32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPRecip64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPRecipSqrt32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPRecipSqrt64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPSqrt(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPSaturate16(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPSaturate32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPSaturate64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPClamp16(EmitContext& ctx, std::string_view value, std::string_view min_value, - std::string_view max_value) { - NotImplemented(); -} - -void EmitFPClamp32(EmitContext& ctx, std::string_view value, std::string_view min_value, - std::string_view max_value) { - NotImplemented(); -} - -void EmitFPClamp64(EmitContext& ctx, std::string_view value, std::string_view min_value, - std::string_view max_value) { - NotImplemented(); -} - -void EmitFPRoundEven16(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPRoundEven32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPRoundEven64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPFloor16(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPFloor32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPFloor64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPCeil16(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPCeil32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPCeil64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPTrunc16(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPTrunc32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPTrunc64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFPOrdEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdNotEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdNotEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdNotEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordNotEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordNotEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordNotEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdLessThan16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdLessThan32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdLessThan64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordLessThan16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordLessThan32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordLessThan64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdGreaterThan16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdGreaterThan32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdGreaterThan64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordGreaterThan16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordGreaterThan32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordGreaterThan64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdLessThanEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdLessThanEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdLessThanEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordLessThanEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordLessThanEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordLessThanEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdGreaterThanEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdGreaterThanEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPOrdGreaterThanEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordGreaterThanEqual16(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordGreaterThanEqual32(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitFPUnordGreaterThanEqual64(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - void EmitFPIsNan16(EmitContext& ctx, std::string_view value) { NotImplemented(); } @@ -1070,179 +521,6 @@ void EmitFPIsNan64(EmitContext& ctx, std::string_view value) { NotImplemented(); } -void EmitIAdd32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitIAdd64(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitISub32(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitISub64(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitIMul32(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitINeg32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitINeg64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitIAbs32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitIAbs64(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitShiftLeftLogical32(EmitContext& ctx, std::string_view base, std::string_view shift) { - NotImplemented(); -} - -void EmitShiftLeftLogical64(EmitContext& ctx, std::string_view base, std::string_view shift) { - NotImplemented(); -} - -void EmitShiftRightLogical32(EmitContext& ctx, std::string_view base, std::string_view shift) { - NotImplemented(); -} - -void EmitShiftRightLogical64(EmitContext& ctx, std::string_view base, std::string_view shift) { - NotImplemented(); -} - -void EmitShiftRightArithmetic32(EmitContext& ctx, std::string_view base, std::string_view shift) { - NotImplemented(); -} - -void EmitShiftRightArithmetic64(EmitContext& ctx, std::string_view base, std::string_view shift) { - NotImplemented(); -} - -void EmitBitwiseAnd32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitBitwiseOr32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitBitwiseXor32(EmitContext& ctx, IR::Inst* inst, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitBitFieldInsert(EmitContext& ctx, std::string_view base, std::string_view insert, - std::string_view offset, std::string_view count) { - NotImplemented(); -} - -void EmitBitFieldSExtract(EmitContext& ctx, IR::Inst* inst, std::string_view base, - std::string_view offset, std::string_view count) { - NotImplemented(); -} - -void EmitBitFieldUExtract(EmitContext& ctx, IR::Inst* inst, std::string_view base, - std::string_view offset, std::string_view count) { - NotImplemented(); -} - -void EmitBitReverse32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitBitCount32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitBitwiseNot32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFindSMsb32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitFindUMsb32(EmitContext& ctx, std::string_view value) { - NotImplemented(); -} - -void EmitSMin32(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitUMin32(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitSMax32(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitUMax32(EmitContext& ctx, std::string_view a, std::string_view b) { - NotImplemented(); -} - -void EmitSClamp32(EmitContext& ctx, IR::Inst* inst, std::string_view value, std::string_view min, - std::string_view max) { - NotImplemented(); -} - -void EmitUClamp32(EmitContext& ctx, IR::Inst* inst, std::string_view value, std::string_view min, - std::string_view max) { - NotImplemented(); -} - -void EmitSLessThan(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitULessThan(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitIEqual(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitSLessThanEqual(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitULessThanEqual(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitSGreaterThan(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitUGreaterThan(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitINotEqual(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitSGreaterThanEqual(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - -void EmitUGreaterThanEqual(EmitContext& ctx, std::string_view lhs, std::string_view rhs) { - NotImplemented(); -} - void EmitSharedAtomicIAdd32(EmitContext& ctx, std::string_view pointer_offset, std::string_view value) { NotImplemented(); @@ -1858,69 +1136,69 @@ void EmitBoundImageWrite(EmitContext&) { NotImplemented(); } -void EmitImageSampleImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageSampleImplicitLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view bias_lc, const IR::Value& offset) { NotImplemented(); } -void EmitImageSampleExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageSampleExplicitLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view lod_lc, const IR::Value& offset) { NotImplemented(); } -void EmitImageSampleDrefImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageSampleDrefImplicitLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view dref, std::string_view bias_lc, const IR::Value& offset) { NotImplemented(); } -void EmitImageSampleDrefExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageSampleDrefExplicitLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view dref, std::string_view lod_lc, const IR::Value& offset) { NotImplemented(); } -void EmitImageGather(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageGather(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, const IR::Value& offset, const IR::Value& offset2) { NotImplemented(); } -void EmitImageGatherDref(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageGatherDref(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, const IR::Value& offset, const IR::Value& offset2, std::string_view dref) { NotImplemented(); } -void EmitImageFetch(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageFetch(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view offset, std::string_view lod, std::string_view ms) { NotImplemented(); } -void EmitImageQueryDimensions(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageQueryDimensions(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view lod) { NotImplemented(); } -void EmitImageQueryLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageQueryLod(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords) { NotImplemented(); } -void EmitImageGradient(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageGradient(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view derivates, std::string_view offset, std::string_view lod_clamp) { NotImplemented(); } -void EmitImageRead(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageRead(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords) { NotImplemented(); } -void EmitImageWrite(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageWrite(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view color) { NotImplemented(); } @@ -2013,57 +1291,57 @@ void EmitBoundImageAtomicExchange32(EmitContext&) { NotImplemented(); } -void EmitImageAtomicIAdd32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicIAdd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicSMin32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicSMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicUMin32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicUMin32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicSMax32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicSMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicUMax32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicUMax32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicInc32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicInc32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicDec32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicDec32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicAnd32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicAnd32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicOr32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicOr32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicXor32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicXor32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } -void EmitImageAtomicExchange32(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, +void EmitImageAtomicExchange32(EmitContext& ctx, IR::Inst& inst, const IR::Value& index, std::string_view coords, std::string_view value) { NotImplemented(); } @@ -2108,24 +1386,24 @@ void EmitSubgroupGeMask(EmitContext& ctx) { NotImplemented(); } -void EmitShuffleIndex(EmitContext& ctx, IR::Inst* inst, std::string_view value, +void EmitShuffleIndex(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index, std::string_view clamp, std::string_view segmentation_mask) { NotImplemented(); } -void EmitShuffleUp(EmitContext& ctx, IR::Inst* inst, std::string_view value, std::string_view index, +void EmitShuffleUp(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index, std::string_view clamp, std::string_view segmentation_mask) { NotImplemented(); } -void EmitShuffleDown(EmitContext& ctx, IR::Inst* inst, std::string_view value, +void EmitShuffleDown(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index, std::string_view clamp, std::string_view segmentation_mask) { NotImplemented(); } -void EmitShuffleButterfly(EmitContext& ctx, IR::Inst* inst, std::string_view value, +void EmitShuffleButterfly(EmitContext& ctx, IR::Inst& inst, std::string_view value, std::string_view index, std::string_view clamp, std::string_view segmentation_mask) { NotImplemented(); |