diff options
author | ReinUsesLisp <reinuseslisp@airmail.cc> | 2021-05-08 21:28:52 +0200 |
---|---|---|
committer | ameerj <52414509+ameerj@users.noreply.github.com> | 2021-07-23 03:51:30 +0200 |
commit | 6fd190d1ae4275a06ed2e488401e1d63912954be (patch) | |
tree | ece4681d18c7b0b5bcb6b540ea4a21b32c19b363 /src/shader_recompiler/backend/glasm/emit_glasm.cpp | |
parent | glasm: Changes to GLASM register allocator and emit context (diff) | |
download | yuzu-6fd190d1ae4275a06ed2e488401e1d63912954be.tar yuzu-6fd190d1ae4275a06ed2e488401e1d63912954be.tar.gz yuzu-6fd190d1ae4275a06ed2e488401e1d63912954be.tar.bz2 yuzu-6fd190d1ae4275a06ed2e488401e1d63912954be.tar.lz yuzu-6fd190d1ae4275a06ed2e488401e1d63912954be.tar.xz yuzu-6fd190d1ae4275a06ed2e488401e1d63912954be.tar.zst yuzu-6fd190d1ae4275a06ed2e488401e1d63912954be.zip |
Diffstat (limited to 'src/shader_recompiler/backend/glasm/emit_glasm.cpp')
-rw-r--r-- | src/shader_recompiler/backend/glasm/emit_glasm.cpp | 66 |
1 files changed, 63 insertions, 3 deletions
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 |