mirror of
https://github.com/shadps4-emu/shadPS4.git
synced 2025-07-22 01:54:31 +00:00
bb29224daf
* shader_recompiler: Implement V_MOVRELS_B32, V_MOVRELD_B32, V_MOVRELSD_B32 Generates a ton of OpSelects to hardcode reading or writing from each possible vgpr depending on the value of m0 Future work is to do range analysis to put an upper bound on m0 and check fewer registers. * fix runtime info after rebase
363 lines
13 KiB
C++
363 lines
13 KiB
C++
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
|
|
// SPDX-License-Identifier: GPL-2.0-or-later
|
|
|
|
#include <span>
|
|
#include <type_traits>
|
|
#include <utility>
|
|
#include <vector>
|
|
#include "common/func_traits.h"
|
|
#include "shader_recompiler/backend/spirv/emit_spirv.h"
|
|
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
|
|
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
|
|
#include "shader_recompiler/frontend/translate/translate.h"
|
|
#include "shader_recompiler/ir/basic_block.h"
|
|
#include "shader_recompiler/ir/program.h"
|
|
|
|
namespace Shader::Backend::SPIRV {
|
|
namespace {
|
|
|
|
template <auto func, typename... Args>
|
|
void SetDefinition(EmitContext& ctx, IR::Inst* inst, Args... args) {
|
|
inst->SetDefinition<Id>(func(ctx, std::forward<Args>(args)...));
|
|
}
|
|
|
|
template <typename ArgType>
|
|
ArgType Arg(EmitContext& ctx, const IR::Value& arg) {
|
|
if constexpr (std::is_same_v<ArgType, Id>) {
|
|
return ctx.Def(arg);
|
|
} else if constexpr (std::is_same_v<ArgType, const IR::Value&>) {
|
|
return arg;
|
|
} else if constexpr (std::is_same_v<ArgType, u32>) {
|
|
return arg.U32();
|
|
} else if constexpr (std::is_same_v<ArgType, u64>) {
|
|
return arg.U64();
|
|
} else if constexpr (std::is_same_v<ArgType, bool>) {
|
|
return arg.U1();
|
|
} else if constexpr (std::is_same_v<ArgType, IR::Attribute>) {
|
|
return arg.Attribute();
|
|
} else if constexpr (std::is_same_v<ArgType, IR::ScalarReg>) {
|
|
return arg.ScalarReg();
|
|
} else if constexpr (std::is_same_v<ArgType, IR::VectorReg>) {
|
|
return arg.VectorReg();
|
|
}
|
|
}
|
|
|
|
template <auto func, bool is_first_arg_inst, size_t... I>
|
|
void Invoke(EmitContext& ctx, IR::Inst* inst, std::index_sequence<I...>) {
|
|
using Traits = Common::FuncTraits<decltype(func)>;
|
|
if constexpr (std::is_same_v<typename Traits::ReturnType, Id>) {
|
|
if constexpr (is_first_arg_inst) {
|
|
SetDefinition<func>(
|
|
ctx, inst, inst,
|
|
Arg<typename Traits::template ArgType<I + 2>>(ctx, inst->Arg(I))...);
|
|
} else {
|
|
SetDefinition<func>(
|
|
ctx, inst, Arg<typename Traits::template ArgType<I + 1>>(ctx, inst->Arg(I))...);
|
|
}
|
|
} else {
|
|
if constexpr (is_first_arg_inst) {
|
|
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))...);
|
|
}
|
|
}
|
|
}
|
|
|
|
template <auto func>
|
|
void Invoke(EmitContext& ctx, IR::Inst* inst) {
|
|
using Traits = Common::FuncTraits<decltype(func)>;
|
|
static_assert(Traits::NUM_ARGS >= 1, "Insufficient arguments");
|
|
if constexpr (Traits::NUM_ARGS == 1) {
|
|
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*>;
|
|
using Indices = std::make_index_sequence<Traits::NUM_ARGS - (is_first_arg_inst ? 2 : 1)>;
|
|
Invoke<func, is_first_arg_inst>(ctx, inst, Indices{});
|
|
}
|
|
}
|
|
|
|
void EmitInst(EmitContext& ctx, IR::Inst* inst) {
|
|
switch (inst->GetOpcode()) {
|
|
#define OPCODE(name, result_type, ...) \
|
|
case IR::Opcode::name: \
|
|
return Invoke<&Emit##name>(ctx, inst);
|
|
#include "shader_recompiler/ir/opcodes.inc"
|
|
#undef OPCODE
|
|
}
|
|
UNREACHABLE_MSG("Invalid opcode {}", inst->GetOpcode());
|
|
}
|
|
|
|
Id TypeId(const EmitContext& ctx, IR::Type type) {
|
|
switch (type) {
|
|
case IR::Type::U1:
|
|
return ctx.U1[1];
|
|
case IR::Type::U32:
|
|
return ctx.U32[1];
|
|
default:
|
|
throw NotImplementedException("Phi node type {}", type);
|
|
}
|
|
}
|
|
|
|
void Traverse(EmitContext& ctx, const IR::Program& program) {
|
|
IR::Block* current_block{};
|
|
for (const IR::AbstractSyntaxNode& node : program.syntax_list) {
|
|
switch (node.type) {
|
|
case IR::AbstractSyntaxNode::Type::Block: {
|
|
const Id label{node.data.block->Definition<Id>()};
|
|
if (current_block) {
|
|
ctx.OpBranch(label);
|
|
}
|
|
current_block = node.data.block;
|
|
ctx.AddLabel(label);
|
|
for (IR::Inst& inst : node.data.block->Instructions()) {
|
|
EmitInst(ctx, &inst);
|
|
}
|
|
break;
|
|
}
|
|
case IR::AbstractSyntaxNode::Type::If: {
|
|
const Id if_label{node.data.if_node.body->Definition<Id>()};
|
|
const Id endif_label{node.data.if_node.merge->Definition<Id>()};
|
|
ctx.OpSelectionMerge(endif_label, spv::SelectionControlMask::MaskNone);
|
|
ctx.OpBranchConditional(ctx.Def(node.data.if_node.cond), if_label, endif_label);
|
|
break;
|
|
}
|
|
case IR::AbstractSyntaxNode::Type::Loop: {
|
|
const Id body_label{node.data.loop.body->Definition<Id>()};
|
|
const Id continue_label{node.data.loop.continue_block->Definition<Id>()};
|
|
const Id endloop_label{node.data.loop.merge->Definition<Id>()};
|
|
|
|
ctx.OpLoopMerge(endloop_label, continue_label, spv::LoopControlMask::MaskNone);
|
|
ctx.OpBranch(body_label);
|
|
break;
|
|
}
|
|
case IR::AbstractSyntaxNode::Type::Break: {
|
|
const Id break_label{node.data.break_node.merge->Definition<Id>()};
|
|
const Id skip_label{node.data.break_node.skip->Definition<Id>()};
|
|
ctx.OpBranchConditional(ctx.Def(node.data.break_node.cond), break_label, skip_label);
|
|
break;
|
|
}
|
|
case IR::AbstractSyntaxNode::Type::EndIf:
|
|
if (current_block) {
|
|
ctx.OpBranch(node.data.end_if.merge->Definition<Id>());
|
|
}
|
|
break;
|
|
case IR::AbstractSyntaxNode::Type::Repeat: {
|
|
Id cond{ctx.Def(node.data.repeat.cond)};
|
|
const Id loop_header_label{node.data.repeat.loop_header->Definition<Id>()};
|
|
const Id merge_label{node.data.repeat.merge->Definition<Id>()};
|
|
ctx.OpBranchConditional(cond, loop_header_label, merge_label);
|
|
break;
|
|
}
|
|
case IR::AbstractSyntaxNode::Type::Return:
|
|
ctx.OpReturn();
|
|
break;
|
|
case IR::AbstractSyntaxNode::Type::Unreachable:
|
|
ctx.OpUnreachable();
|
|
break;
|
|
}
|
|
if (node.type != IR::AbstractSyntaxNode::Type::Block) {
|
|
current_block = nullptr;
|
|
}
|
|
}
|
|
}
|
|
|
|
Id DefineMain(EmitContext& ctx, const IR::Program& program) {
|
|
const Id void_function{ctx.TypeFunction(ctx.void_id)};
|
|
const Id main{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)};
|
|
for (IR::Block* const block : program.blocks) {
|
|
block->SetDefinition(ctx.OpLabel());
|
|
}
|
|
Traverse(ctx, program);
|
|
ctx.OpFunctionEnd();
|
|
return main;
|
|
}
|
|
|
|
void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
|
|
const auto& info = program.info;
|
|
const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
|
|
spv::ExecutionModel execution_model{};
|
|
ctx.AddCapability(spv::Capability::Image1D);
|
|
ctx.AddCapability(spv::Capability::Sampled1D);
|
|
ctx.AddCapability(spv::Capability::ImageQuery);
|
|
if (info.uses_fp16) {
|
|
ctx.AddCapability(spv::Capability::Float16);
|
|
ctx.AddCapability(spv::Capability::Int16);
|
|
}
|
|
ctx.AddCapability(spv::Capability::Int64);
|
|
if (info.has_storage_images || info.has_image_buffers) {
|
|
ctx.AddCapability(spv::Capability::StorageImageExtendedFormats);
|
|
ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat);
|
|
ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat);
|
|
}
|
|
if (info.has_texel_buffers) {
|
|
ctx.AddCapability(spv::Capability::SampledBuffer);
|
|
}
|
|
if (info.has_image_buffers) {
|
|
ctx.AddCapability(spv::Capability::ImageBuffer);
|
|
}
|
|
if (info.has_image_gather) {
|
|
ctx.AddCapability(spv::Capability::ImageGatherExtended);
|
|
}
|
|
if (info.has_image_query) {
|
|
ctx.AddCapability(spv::Capability::ImageQuery);
|
|
}
|
|
if (info.uses_lane_id) {
|
|
ctx.AddCapability(spv::Capability::GroupNonUniform);
|
|
}
|
|
if (info.uses_group_quad) {
|
|
ctx.AddCapability(spv::Capability::GroupNonUniformQuad);
|
|
}
|
|
switch (program.info.stage) {
|
|
case Stage::Compute: {
|
|
const std::array<u32, 3> workgroup_size{ctx.runtime_info.cs_info.workgroup_size};
|
|
execution_model = spv::ExecutionModel::GLCompute;
|
|
ctx.AddExecutionMode(main, spv::ExecutionMode::LocalSize, workgroup_size[0],
|
|
workgroup_size[1], workgroup_size[2]);
|
|
break;
|
|
}
|
|
case Stage::Vertex:
|
|
execution_model = spv::ExecutionModel::Vertex;
|
|
break;
|
|
case Stage::Fragment:
|
|
execution_model = spv::ExecutionModel::Fragment;
|
|
if (ctx.profile.lower_left_origin_mode) {
|
|
ctx.AddExecutionMode(main, spv::ExecutionMode::OriginLowerLeft);
|
|
} else {
|
|
ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft);
|
|
}
|
|
if (info.has_discard) {
|
|
ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT);
|
|
}
|
|
if (info.stores.Get(IR::Attribute::Depth)) {
|
|
ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing);
|
|
}
|
|
break;
|
|
default:
|
|
throw NotImplementedException("Stage {}", u32(program.info.stage));
|
|
}
|
|
ctx.AddEntryPoint(execution_model, main, "main", interfaces);
|
|
}
|
|
|
|
void PatchPhiNodes(const IR::Program& program, EmitContext& ctx) {
|
|
auto inst{program.blocks.front()->begin()};
|
|
size_t block_index{0};
|
|
ctx.PatchDeferredPhi([&](size_t phi_arg) {
|
|
if (phi_arg == 0) {
|
|
++inst;
|
|
if (inst == program.blocks[block_index]->end() ||
|
|
inst->GetOpcode() != IR::Opcode::Phi) {
|
|
do {
|
|
++block_index;
|
|
inst = program.blocks[block_index]->begin();
|
|
} while (inst->GetOpcode() != IR::Opcode::Phi);
|
|
}
|
|
}
|
|
return ctx.Def(inst->Arg(phi_arg));
|
|
});
|
|
}
|
|
} // Anonymous namespace
|
|
|
|
std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info,
|
|
const IR::Program& program, u32& binding) {
|
|
EmitContext ctx{profile, runtime_info, program.info, binding};
|
|
const Id main{DefineMain(ctx, program)};
|
|
DefineEntryPoint(program, ctx, main);
|
|
if (program.info.stage == Stage::Vertex) {
|
|
ctx.AddExtension("SPV_KHR_shader_draw_parameters");
|
|
ctx.AddCapability(spv::Capability::DrawParameters);
|
|
}
|
|
PatchPhiNodes(program, ctx);
|
|
return ctx.Assemble();
|
|
}
|
|
|
|
Id EmitPhi(EmitContext& ctx, IR::Inst* inst) {
|
|
const size_t num_args{inst->NumArgs()};
|
|
boost::container::small_vector<Id, 32> blocks;
|
|
blocks.reserve(num_args);
|
|
for (size_t index = 0; index < num_args; ++index) {
|
|
blocks.push_back(inst->PhiBlock(index)->Definition<Id>());
|
|
}
|
|
// The type of a phi instruction is stored in its flags
|
|
const Id result_type{TypeId(ctx, inst->Flags<IR::Type>())};
|
|
return ctx.DeferredOpPhi(result_type, std::span(blocks.data(), blocks.size()));
|
|
}
|
|
|
|
void EmitVoid(EmitContext&) {}
|
|
|
|
Id EmitIdentity(EmitContext& ctx, const IR::Value& value) {
|
|
throw NotImplementedException("Forward identity declaration");
|
|
}
|
|
|
|
Id EmitConditionRef(EmitContext& ctx, const IR::Value& value) {
|
|
const Id id{ctx.Def(value)};
|
|
if (!Sirit::ValidId(id)) {
|
|
throw NotImplementedException("Forward identity declaration");
|
|
}
|
|
return id;
|
|
}
|
|
|
|
void EmitReference(EmitContext&) {}
|
|
|
|
void EmitPhiMove(EmitContext&) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitGetScc(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitGetExec(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitGetVcc(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitGetSccLo(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitGetVccLo(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitGetVccHi(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitGetM0(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitSetScc(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitSetExec(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitSetVcc(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitSetSccLo(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitSetVccLo(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitSetVccHi(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
void EmitSetM0(EmitContext& ctx) {
|
|
UNREACHABLE_MSG("Unreachable instruction");
|
|
}
|
|
|
|
} // namespace Shader::Backend::SPIRV
|