Mirrors/shadPS4
1
0
Fork 0
mirror of https://github.com/shadps4-emu/shadPS4.git synced 2025-05-19 01:44:53 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions 4

video: Import new shader recompiler + display a triangle (#142)

Browse source
This commit is contained in:
TheTurtle 2024-05-22 01:35:12 +03:00 • committed by GitHub
parent 8cf64a33b2
commit 8730968385
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
103 changed files with 17793 additions and 729 deletions
Download patch file Download diff file Expand all files Collapse all files

18
src/shader_recompiler/backend/bindings.h Normal file
View file

@ -0,0 +1,18 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#pragma once
#include "common/types.h"
namespace Shader::Backend {
struct Bindings {
u32 unified{};
u32 uniform_buffer{};
u32 storage_buffer{};
u32 texture{};
u32 image{};
};
} // namespace Shader::Backend

285
src/shader_recompiler/backend/spirv/emit_spirv.cpp Normal file
View file

@ -0,0 +1,285 @@
// 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/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, 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
}
throw LogicError("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, 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, 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 std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size());
spv::ExecutionModel execution_model{};
switch (program.stage) {
case Stage::Compute: {
// const std::array<u32, 3> workgroup_size{program.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 (program.info.stores_frag_depth) {
// ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing);
// }
break;
default:
throw NotImplementedException("Stage {}", u32(program.stage));
}
ctx.AddEntryPoint(execution_model, main, "main", interfaces);
}
void PatchPhiNodes(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, IR::Program& program, Bindings& bindings) {
EmitContext ctx{profile, program, bindings};
const Id main{DefineMain(ctx, program)};
DefineEntryPoint(program, ctx, main);
if (program.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) {
throw NotImplementedException("Forward identity declaration");
}
void EmitReference(EmitContext&) {}
void EmitPhiMove(EmitContext&) {
throw LogicError("Unreachable instruction");
}
void EmitGetZeroFromOp(EmitContext&) {
throw LogicError("Unreachable instruction");
}
void EmitGetSignFromOp(EmitContext&) {
throw LogicError("Unreachable instruction");
}
void EmitGetCarryFromOp(EmitContext&) {
throw LogicError("Unreachable instruction");
}
void EmitGetOverflowFromOp(EmitContext&) {
throw LogicError("Unreachable instruction");
}
void EmitSetVcc(EmitContext& ctx) {
throw LogicError("Unreachable instruction");
}
void EmitGetVcc(EmitContext& ctx) {
throw LogicError("Unreachable instruction");
}
} // namespace Shader::Backend::SPIRV

21
src/shader_recompiler/backend/spirv/emit_spirv.h Normal file
View file

@ -0,0 +1,21 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#pragma once
#include <vector>
#include "shader_recompiler/backend/bindings.h"
#include "shader_recompiler/ir/program.h"
#include "shader_recompiler/profile.h"
namespace Shader::Backend::SPIRV {
[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, IR::Program& program,
Bindings& bindings);
[[nodiscard]] inline std::vector<u32> EmitSPIRV(const Profile& profile, IR::Program& program) {
Bindings binding;
return EmitSPIRV(profile, program, binding);
}
} // namespace Shader::Backend::SPIRV

57
src/shader_recompiler/backend/spirv/emit_spirv_bitwise_conversion.cpp Normal file
View file

@ -0,0 +1,57 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
void EmitBitCastU16F16(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitBitCastU32F32(EmitContext& ctx, Id value) {
return ctx.OpBitcast(ctx.U32[1], value);
}
void EmitBitCastU64F64(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
void EmitBitCastF16U16(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitBitCastF32U32(EmitContext& ctx, Id value) {
return ctx.OpBitcast(ctx.F32[1], value);
}
void EmitBitCastF64U64(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitPackUint2x32(EmitContext& ctx, Id value) {
return ctx.OpBitcast(ctx.U64, value);
}
Id EmitUnpackUint2x32(EmitContext& ctx, Id value) {
return ctx.OpBitcast(ctx.U32[2], value);
}
Id EmitPackFloat2x16(EmitContext& ctx, Id value) {
return ctx.OpBitcast(ctx.U32[1], value);
}
Id EmitUnpackFloat2x16(EmitContext& ctx, Id value) {
return ctx.OpBitcast(ctx.F16[2], value);
}
Id EmitPackHalf2x16(EmitContext& ctx, Id value) {
return ctx.OpPackHalf2x16(ctx.U32[1], value);
}
Id EmitUnpackHalf2x16(EmitContext& ctx, Id value) {
return ctx.OpUnpackHalf2x16(ctx.F32[2], value);
}
} // namespace Shader::Backend::SPIRV

153
src/shader_recompiler/backend/spirv/emit_spirv_composite.cpp Normal file
View file

@ -0,0 +1,153 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
Id EmitCompositeConstructU32x2(EmitContext& ctx, Id e1, Id e2) {
return ctx.OpCompositeConstruct(ctx.U32[2], e1, e2);
}
Id EmitCompositeConstructU32x3(EmitContext& ctx, Id e1, Id e2, Id e3) {
return ctx.OpCompositeConstruct(ctx.U32[3], e1, e2, e3);
}
Id EmitCompositeConstructU32x4(EmitContext& ctx, Id e1, Id e2, Id e3, Id e4) {
return ctx.OpCompositeConstruct(ctx.U32[4], e1, e2, e3, e4);
}
Id EmitCompositeExtractU32x2(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.U32[1], composite, index);
}
Id EmitCompositeExtractU32x3(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.U32[1], composite, index);
}
Id EmitCompositeExtractU32x4(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.U32[1], composite, index);
}
Id EmitCompositeInsertU32x2(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.U32[2], object, composite, index);
}
Id EmitCompositeInsertU32x3(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.U32[3], object, composite, index);
}
Id EmitCompositeInsertU32x4(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.U32[4], object, composite, index);
}
Id EmitCompositeConstructF16x2(EmitContext& ctx, Id e1, Id e2) {
return ctx.OpCompositeConstruct(ctx.F16[2], e1, e2);
}
Id EmitCompositeConstructF16x3(EmitContext& ctx, Id e1, Id e2, Id e3) {
return ctx.OpCompositeConstruct(ctx.F16[3], e1, e2, e3);
}
Id EmitCompositeConstructF16x4(EmitContext& ctx, Id e1, Id e2, Id e3, Id e4) {
return ctx.OpCompositeConstruct(ctx.F16[4], e1, e2, e3, e4);
}
Id EmitCompositeExtractF16x2(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.F16[1], composite, index);
}
Id EmitCompositeExtractF16x3(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.F16[1], composite, index);
}
Id EmitCompositeExtractF16x4(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.F16[1], composite, index);
}
Id EmitCompositeInsertF16x2(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F16[2], object, composite, index);
}
Id EmitCompositeInsertF16x3(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F16[3], object, composite, index);
}
Id EmitCompositeInsertF16x4(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F16[4], object, composite, index);
}
Id EmitCompositeConstructF32x2(EmitContext& ctx, Id e1, Id e2) {
return ctx.OpCompositeConstruct(ctx.F32[2], e1, e2);
}
Id EmitCompositeConstructF32x3(EmitContext& ctx, Id e1, Id e2, Id e3) {
return ctx.OpCompositeConstruct(ctx.F32[3], e1, e2, e3);
}
Id EmitCompositeConstructF32x4(EmitContext& ctx, Id e1, Id e2, Id e3, Id e4) {
return ctx.OpCompositeConstruct(ctx.F32[4], e1, e2, e3, e4);
}
Id EmitCompositeExtractF32x2(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.F32[1], composite, index);
}
Id EmitCompositeExtractF32x3(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.F32[1], composite, index);
}
Id EmitCompositeExtractF32x4(EmitContext& ctx, Id composite, u32 index) {
return ctx.OpCompositeExtract(ctx.F32[1], composite, index);
}
Id EmitCompositeInsertF32x2(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F32[2], object, composite, index);
}
Id EmitCompositeInsertF32x3(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F32[3], object, composite, index);
}
Id EmitCompositeInsertF32x4(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F32[4], object, composite, index);
}
void EmitCompositeConstructF64x2(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
void EmitCompositeConstructF64x3(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
void EmitCompositeConstructF64x4(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
void EmitCompositeExtractF64x2(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
void EmitCompositeExtractF64x3(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
void EmitCompositeExtractF64x4(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitCompositeInsertF64x2(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F64[2], object, composite, index);
}
Id EmitCompositeInsertF64x3(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F64[3], object, composite, index);
}
Id EmitCompositeInsertF64x4(EmitContext& ctx, Id composite, Id object, u32 index) {
return ctx.OpCompositeInsert(ctx.F64[4], object, composite, index);
}
} // namespace Shader::Backend::SPIRV

103
src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp Normal file
View file

@ -0,0 +1,103 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
namespace {
Id OutputAttrPointer(EmitContext& ctx, IR::Attribute attr, u32 element) {
if (IR::IsParam(attr)) {
const u32 index{u32(attr) - u32(IR::Attribute::Param0)};
const auto& info{ctx.output_params.at(index).at(element)};
if (info.num_components == 1) {
return info.id;
} else {
const u32 index_element{element - info.first_element};
return ctx.OpAccessChain(ctx.output_f32, info.id, ctx.ConstU32(index_element));
}
}
switch (attr) {
case IR::Attribute::Position0: {
return ctx.OpAccessChain(ctx.output_f32, ctx.output_position, ctx.ConstU32(element));
case IR::Attribute::RenderTarget0:
return ctx.OpAccessChain(ctx.output_f32, ctx.frag_color[0], ctx.ConstU32(element));
}
default:
throw NotImplementedException("Read attribute {}", attr);
}
}
} // Anonymous namespace
void EmitGetScalarRegister(EmitContext&) {
throw LogicError("Unreachable instruction");
}
void EmitSetScalarRegister(EmitContext&) {
throw LogicError("Unreachable instruction");
}
void EmitGetVectorRegister(EmitContext& ctx) {
throw LogicError("Unreachable instruction");
}
void EmitSetVectorRegister(EmitContext& ctx) {
throw LogicError("Unreachable instruction");
}
void EmitSetGotoVariable(EmitContext&) {
throw LogicError("Unreachable instruction");
}
void EmitGetGotoVariable(EmitContext&) {
throw LogicError("Unreachable instruction");
}
Id EmitReadConst(EmitContext& ctx) {
throw LogicError("Unreachable instruction");
}
Id EmitReadConstBuffer(EmitContext& ctx, const IR::Value& binding, const IR::Value& addr,
const IR::Value& offset) {
throw LogicError("Unreachable instruction");
}
Id EmitReadConstBufferF32(EmitContext& ctx, const IR::Value& binding, const IR::Value& addr,
const IR::Value& offset) {
throw LogicError("Unreachable instruction");
}
Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr, Id vertex) {
const u32 element{static_cast<u32>(attr) % 4};
if (IR::IsParam(attr)) {
const u32 index{u32(attr) - u32(IR::Attribute::Param0)};
const auto& param{ctx.input_params.at(index)};
if (!ValidId(param.id)) {
// Attribute is disabled or varying component is not written
return ctx.ConstF32(element == 3 ? 1.0f : 0.0f);
}
const Id pointer{ctx.OpAccessChain(param.pointer_type, param.id, ctx.ConstU32(element))};
return ctx.OpLoad(param.component_type, pointer);
}
throw NotImplementedException("Read attribute {}", attr);
}
Id EmitGetAttributeU32(EmitContext& ctx, IR::Attribute attr, Id) {
switch (attr) {
case IR::Attribute::VertexId:
return ctx.OpLoad(ctx.U32[1], ctx.vertex_index);
default:
throw NotImplementedException("Read U32 attribute {}", attr);
}
}
void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, Id value, u32 element) {
if (attr == IR::Attribute::Param0) {
return;
}
const Id pointer{OutputAttrPointer(ctx, attr, element)};
ctx.OpStore(pointer, value);
}
} // namespace Shader::Backend::SPIRV

262
src/shader_recompiler/backend/spirv/emit_spirv_convert.cpp Normal file
View file

@ -0,0 +1,262 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
namespace {
Id ExtractU16(EmitContext& ctx, Id value) {
if (ctx.profile.support_int16) {
return ctx.OpUConvert(ctx.U16, value);
} else {
return ctx.OpBitFieldUExtract(ctx.U32[1], value, ctx.u32_zero_value, ctx.ConstU32(16u));
}
}
Id ExtractS16(EmitContext& ctx, Id value) {
if (ctx.profile.support_int16) {
return ctx.OpSConvert(ctx.S16, value);
} else {
return ctx.OpBitFieldSExtract(ctx.U32[1], value, ctx.u32_zero_value, ctx.ConstU32(16u));
}
}
Id ExtractU8(EmitContext& ctx, Id value) {
if (ctx.profile.support_int8) {
return ctx.OpUConvert(ctx.U8, value);
} else {
return ctx.OpBitFieldUExtract(ctx.U32[1], value, ctx.u32_zero_value, ctx.ConstU32(8u));
}
}
Id ExtractS8(EmitContext& ctx, Id value) {
if (ctx.profile.support_int8) {
return ctx.OpSConvert(ctx.S8, value);
} else {
return ctx.OpBitFieldSExtract(ctx.U32[1], value, ctx.u32_zero_value, ctx.ConstU32(8u));
}
}
} // Anonymous namespace
Id EmitConvertS16F16(EmitContext& ctx, Id value) {
if (ctx.profile.support_int16) {
return ctx.OpSConvert(ctx.U32[1], ctx.OpConvertFToS(ctx.U16, value));
} else {
return ExtractS16(ctx, ctx.OpConvertFToS(ctx.U32[1], value));
}
}
Id EmitConvertS16F32(EmitContext& ctx, Id value) {
if (ctx.profile.support_int16) {
return ctx.OpSConvert(ctx.U32[1], ctx.OpConvertFToS(ctx.U16, value));
} else {
return ExtractS16(ctx, ctx.OpConvertFToS(ctx.U32[1], value));
}
}
Id EmitConvertS16F64(EmitContext& ctx, Id value) {
if (ctx.profile.support_int16) {
return ctx.OpSConvert(ctx.U32[1], ctx.OpConvertFToS(ctx.U16, value));
} else {
return ExtractS16(ctx, ctx.OpConvertFToS(ctx.U32[1], value));
}
}
Id EmitConvertS32F16(EmitContext& ctx, Id value) {
return ctx.OpConvertFToS(ctx.U32[1], value);
}
Id EmitConvertS32F32(EmitContext& ctx, Id value) {
if (ctx.profile.has_broken_signed_operations) {
return ctx.OpBitcast(ctx.U32[1], ctx.OpConvertFToS(ctx.S32[1], value));
} else {
return ctx.OpConvertFToS(ctx.U32[1], value);
}
}
Id EmitConvertS32F64(EmitContext& ctx, Id value) {
return ctx.OpConvertFToS(ctx.U32[1], value);
}
Id EmitConvertS64F16(EmitContext& ctx, Id value) {
return ctx.OpConvertFToS(ctx.U64, value);
}
Id EmitConvertS64F32(EmitContext& ctx, Id value) {
return ctx.OpConvertFToS(ctx.U64, value);
}
Id EmitConvertS64F64(EmitContext& ctx, Id value) {
return ctx.OpConvertFToS(ctx.U64, value);
}
Id EmitConvertU16F16(EmitContext& ctx, Id value) {
if (ctx.profile.support_int16) {
return ctx.OpUConvert(ctx.U32[1], ctx.OpConvertFToU(ctx.U16, value));
} else {
return ExtractU16(ctx, ctx.OpConvertFToU(ctx.U32[1], value));
}
}
Id EmitConvertU16F32(EmitContext& ctx, Id value) {
if (ctx.profile.support_int16) {
return ctx.OpUConvert(ctx.U32[1], ctx.OpConvertFToU(ctx.U16, value));
} else {
return ExtractU16(ctx, ctx.OpConvertFToU(ctx.U32[1], value));
}
}
Id EmitConvertU16F64(EmitContext& ctx, Id value) {
if (ctx.profile.support_int16) {
return ctx.OpUConvert(ctx.U32[1], ctx.OpConvertFToU(ctx.U16, value));
} else {
return ExtractU16(ctx, ctx.OpConvertFToU(ctx.U32[1], value));
}
}
Id EmitConvertU32F16(EmitContext& ctx, Id value) {
return ctx.OpConvertFToU(ctx.U32[1], value);
}
Id EmitConvertU32F32(EmitContext& ctx, Id value) {
return ctx.OpConvertFToU(ctx.U32[1], value);
}
Id EmitConvertU32F64(EmitContext& ctx, Id value) {
return ctx.OpConvertFToU(ctx.U32[1], value);
}
Id EmitConvertU64F16(EmitContext& ctx, Id value) {
return ctx.OpConvertFToU(ctx.U64, value);
}
Id EmitConvertU64F32(EmitContext& ctx, Id value) {
return ctx.OpConvertFToU(ctx.U64, value);
}
Id EmitConvertU64F64(EmitContext& ctx, Id value) {
return ctx.OpConvertFToU(ctx.U64, value);
}
Id EmitConvertU64U32(EmitContext& ctx, Id value) {
return ctx.OpUConvert(ctx.U64, value);
}
Id EmitConvertU32U64(EmitContext& ctx, Id value) {
return ctx.OpUConvert(ctx.U32[1], value);
}
Id EmitConvertF16F32(EmitContext& ctx, Id value) {
return ctx.OpFConvert(ctx.F16[1], value);
}
Id EmitConvertF32F16(EmitContext& ctx, Id value) {
return ctx.OpFConvert(ctx.F32[1], value);
}
Id EmitConvertF32F64(EmitContext& ctx, Id value) {
return ctx.OpFConvert(ctx.F32[1], value);
}
Id EmitConvertF64F32(EmitContext& ctx, Id value) {
return ctx.OpFConvert(ctx.F64[1], value);
}
Id EmitConvertF16S8(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F16[1], ExtractS8(ctx, value));
}
Id EmitConvertF16S16(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F16[1], ExtractS16(ctx, value));
}
Id EmitConvertF16S32(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F16[1], value);
}
Id EmitConvertF16S64(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F16[1], value);
}
Id EmitConvertF16U8(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F16[1], ExtractU8(ctx, value));
}
Id EmitConvertF16U16(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F16[1], ExtractU16(ctx, value));
}
Id EmitConvertF16U32(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F16[1], value);
}
Id EmitConvertF16U64(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F16[1], value);
}
Id EmitConvertF32S8(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F32[1], ExtractS8(ctx, value));
}
Id EmitConvertF32S16(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F32[1], ExtractS16(ctx, value));
}
Id EmitConvertF32S32(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F32[1], value);
}
Id EmitConvertF32S64(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F32[1], value);
}
Id EmitConvertF32U8(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F32[1], ExtractU8(ctx, value));
}
Id EmitConvertF32U16(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F32[1], ExtractU16(ctx, value));
}
Id EmitConvertF32U32(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F32[1], value);
}
Id EmitConvertF32U64(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F32[1], value);
}
Id EmitConvertF64S8(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F64[1], ExtractS8(ctx, value));
}
Id EmitConvertF64S16(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F64[1], ExtractS16(ctx, value));
}
Id EmitConvertF64S32(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F64[1], value);
}
Id EmitConvertF64S64(EmitContext& ctx, Id value) {
return ctx.OpConvertSToF(ctx.F64[1], value);
}
Id EmitConvertF64U8(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F64[1], ExtractU8(ctx, value));
}
Id EmitConvertF64U16(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F64[1], ExtractU16(ctx, value));
}
Id EmitConvertF64U32(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F64[1], value);
}
Id EmitConvertF64U64(EmitContext& ctx, Id value) {
return ctx.OpConvertUToF(ctx.F64[1], value);
}
} // namespace Shader::Backend::SPIRV

355
src/shader_recompiler/backend/spirv/emit_spirv_floating_point.cpp Normal file
View file

@ -0,0 +1,355 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
Id EmitFPAbs16(EmitContext& ctx, Id value) {
return ctx.OpFAbs(ctx.F16[1], value);
}
Id EmitFPAbs32(EmitContext& ctx, Id value) {
return ctx.OpFAbs(ctx.F32[1], value);
}
Id EmitFPAbs64(EmitContext& ctx, Id value) {
return ctx.OpFAbs(ctx.F64[1], value);
}
Id EmitFPAdd16(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
return ctx.OpFAdd(ctx.F16[1], a, b);
}
Id EmitFPAdd32(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
return ctx.OpFAdd(ctx.F32[1], a, b);
}
Id EmitFPAdd64(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
return ctx.OpFAdd(ctx.F64[1], a, b);
}
Id EmitFPFma16(EmitContext& ctx, IR::Inst* inst, Id a, Id b, Id c) {
return ctx.OpFma(ctx.F16[1], a, b, c);
}
Id EmitFPFma32(EmitContext& ctx, IR::Inst* inst, Id a, Id b, Id c) {
return ctx.OpFma(ctx.F32[1], a, b, c);
}
Id EmitFPFma64(EmitContext& ctx, IR::Inst* inst, Id a, Id b, Id c) {
return ctx.OpFma(ctx.F64[1], a, b, c);
}
Id EmitFPMax32(EmitContext& ctx, Id a, Id b) {
return ctx.OpFMax(ctx.F32[1], a, b);
}
Id EmitFPMax64(EmitContext& ctx, Id a, Id b) {
return ctx.OpFMax(ctx.F64[1], a, b);
}
Id EmitFPMin32(EmitContext& ctx, Id a, Id b) {
return ctx.OpFMin(ctx.F32[1], a, b);
}
Id EmitFPMin64(EmitContext& ctx, Id a, Id b) {
return ctx.OpFMin(ctx.F64[1], a, b);
}
Id EmitFPMul16(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
return ctx.OpFMul(ctx.F16[1], a, b);
}
Id EmitFPMul32(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
return ctx.OpFMul(ctx.F32[1], a, b);
}
Id EmitFPMul64(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
return ctx.OpFMul(ctx.F64[1], a, b);
}
Id EmitFPNeg16(EmitContext& ctx, Id value) {
return ctx.OpFNegate(ctx.F16[1], value);
}
Id EmitFPNeg32(EmitContext& ctx, Id value) {
return ctx.OpFNegate(ctx.F32[1], value);
}
Id EmitFPNeg64(EmitContext& ctx, Id value) {
return ctx.OpFNegate(ctx.F64[1], value);
}
Id EmitFPSin(EmitContext& ctx, Id value) {
return ctx.OpSin(ctx.F32[1], value);
}
Id EmitFPCos(EmitContext& ctx, Id value) {
return ctx.OpCos(ctx.F32[1], value);
}
Id EmitFPExp2(EmitContext& ctx, Id value) {
return ctx.OpExp2(ctx.F32[1], value);
}
Id EmitFPLog2(EmitContext& ctx, Id value) {
return ctx.OpLog2(ctx.F32[1], value);
}
Id EmitFPRecip32(EmitContext& ctx, Id value) {
return ctx.OpFDiv(ctx.F32[1], ctx.ConstF32(1.0f), value);
}
Id EmitFPRecip64(EmitContext& ctx, Id value) {
return ctx.OpFDiv(ctx.F64[1], ctx.Constant(ctx.F64[1], 1.0f), value);
}
Id EmitFPRecipSqrt32(EmitContext& ctx, Id value) {
return ctx.OpInverseSqrt(ctx.F32[1], value);
}
Id EmitFPRecipSqrt64(EmitContext& ctx, Id value) {
return ctx.OpInverseSqrt(ctx.F64[1], value);
}
Id EmitFPSqrt(EmitContext& ctx, Id value) {
return ctx.OpSqrt(ctx.F32[1], value);
}
Id EmitFPSaturate16(EmitContext& ctx, Id value) {
const Id zero{ctx.Constant(ctx.F16[1], u16{0})};
const Id one{ctx.Constant(ctx.F16[1], u16{0x3c00})};
return ctx.OpFClamp(ctx.F16[1], value, zero, one);
}
Id EmitFPSaturate32(EmitContext& ctx, Id value) {
const Id zero{ctx.ConstF32(f32{0.0})};
const Id one{ctx.ConstF32(f32{1.0})};
return ctx.OpFClamp(ctx.F32[1], value, zero, one);
}
Id EmitFPSaturate64(EmitContext& ctx, Id value) {
const Id zero{ctx.Constant(ctx.F64[1], f64{0.0})};
const Id one{ctx.Constant(ctx.F64[1], f64{1.0})};
return ctx.OpFClamp(ctx.F64[1], value, zero, one);
}
Id EmitFPClamp16(EmitContext& ctx, Id value, Id min_value, Id max_value) {
return ctx.OpFClamp(ctx.F16[1], value, min_value, max_value);
}
Id EmitFPClamp32(EmitContext& ctx, Id value, Id min_value, Id max_value) {
return ctx.OpFClamp(ctx.F32[1], value, min_value, max_value);
}
Id EmitFPClamp64(EmitContext& ctx, Id value, Id min_value, Id max_value) {
return ctx.OpFClamp(ctx.F64[1], value, min_value, max_value);
}
Id EmitFPRoundEven16(EmitContext& ctx, Id value) {
return ctx.OpRoundEven(ctx.F16[1], value);
}
Id EmitFPRoundEven32(EmitContext& ctx, Id value) {
return ctx.OpRoundEven(ctx.F32[1], value);
}
Id EmitFPRoundEven64(EmitContext& ctx, Id value) {
return ctx.OpRoundEven(ctx.F64[1], value);
}
Id EmitFPFloor16(EmitContext& ctx, Id value) {
return ctx.OpFloor(ctx.F16[1], value);
}
Id EmitFPFloor32(EmitContext& ctx, Id value) {
return ctx.OpFloor(ctx.F32[1], value);
}
Id EmitFPFloor64(EmitContext& ctx, Id value) {
return ctx.OpFloor(ctx.F64[1], value);
}
Id EmitFPCeil16(EmitContext& ctx, Id value) {
return ctx.OpCeil(ctx.F16[1], value);
}
Id EmitFPCeil32(EmitContext& ctx, Id value) {
return ctx.OpCeil(ctx.F32[1], value);
}
Id EmitFPCeil64(EmitContext& ctx, Id value) {
return ctx.OpCeil(ctx.F64[1], value);
}
Id EmitFPTrunc16(EmitContext& ctx, Id value) {
return ctx.OpTrunc(ctx.F16[1], value);
}
Id EmitFPTrunc32(EmitContext& ctx, Id value) {
return ctx.OpTrunc(ctx.F32[1], value);
}
Id EmitFPTrunc64(EmitContext& ctx, Id value) {
return ctx.OpTrunc(ctx.F64[1], value);
}
Id EmitFPOrdEqual16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdEqual32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdEqual64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordEqual16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordEqual32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordEqual64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdNotEqual16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdNotEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdNotEqual32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdNotEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdNotEqual64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdNotEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordNotEqual16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordNotEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordNotEqual32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordNotEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordNotEqual64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordNotEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdLessThan16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdLessThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdLessThan32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdLessThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdLessThan64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdLessThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordLessThan16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordLessThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordLessThan32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordLessThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordLessThan64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordLessThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdGreaterThan16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdGreaterThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdGreaterThan32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdGreaterThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdGreaterThan64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdGreaterThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordGreaterThan16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordGreaterThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordGreaterThan32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordGreaterThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordGreaterThan64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordGreaterThan(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdLessThanEqual16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdLessThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdLessThanEqual32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdLessThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdLessThanEqual64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdLessThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordLessThanEqual16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordLessThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordLessThanEqual32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordLessThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordLessThanEqual64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordLessThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdGreaterThanEqual16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdGreaterThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdGreaterThanEqual32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdGreaterThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPOrdGreaterThanEqual64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFOrdGreaterThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordGreaterThanEqual16(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordGreaterThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordGreaterThanEqual32(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordGreaterThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPUnordGreaterThanEqual64(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpFUnordGreaterThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitFPIsNan16(EmitContext& ctx, Id value) {
return ctx.OpIsNan(ctx.U1[1], value);
}
Id EmitFPIsNan32(EmitContext& ctx, Id value) {
return ctx.OpIsNan(ctx.U1[1], value);
}
Id EmitFPIsNan64(EmitContext& ctx, Id value) {
return ctx.OpIsNan(ctx.U1[1], value);
}
} // namespace Shader::Backend::SPIRV

66
src/shader_recompiler/backend/spirv/emit_spirv_image.cpp Normal file
View file

@ -0,0 +1,66 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
Id EmitImageSampleImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
Id bias_lc, const IR::Value& offset) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageSampleExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
Id lod, const IR::Value& offset) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageSampleDrefImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index,
Id coords, Id dref, Id bias_lc, const IR::Value& offset) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageSampleDrefExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index,
Id coords, Id dref, Id lod, const IR::Value& offset) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageGather(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
const IR::Value& offset, const IR::Value& offset2) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageGatherDref(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
const IR::Value& offset, const IR::Value& offset2, Id dref) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageFetch(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords, Id offset,
Id lod, Id ms) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageQueryDimensions(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id lod,
const IR::Value& skip_mips_val) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageQueryLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageGradient(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
Id derivatives, const IR::Value& offset, Id lod_clamp) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitImageRead(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords) {
throw NotImplementedException("SPIR-V Instruction");
}
void EmitImageWrite(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords, Id color) {
throw NotImplementedException("SPIR-V Instruction");
}
} // namespace Shader::Backend::SPIRV

335
src/shader_recompiler/backend/spirv/emit_spirv_instructions.h Normal file
View file

@ -0,0 +1,335 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#pragma once
#include <sirit/sirit.h>
#include "common/types.h"
namespace Shader::IR {
enum class Attribute : u64;
enum class Patch : u64;
class Inst;
class Value;
} // namespace Shader::IR
namespace Shader::Backend::SPIRV {
using Sirit::Id;
class EmitContext;
// Microinstruction emitters
Id EmitPhi(EmitContext& ctx, IR::Inst* inst);
void EmitVoid(EmitContext& ctx);
Id EmitIdentity(EmitContext& ctx, const IR::Value& value);
Id EmitConditionRef(EmitContext& ctx, const IR::Value& value);
void EmitReference(EmitContext&);
void EmitPhiMove(EmitContext&);
void EmitJoin(EmitContext& ctx);
void EmitBarrier(EmitContext& ctx);
void EmitWorkgroupMemoryBarrier(EmitContext& ctx);
void EmitDeviceMemoryBarrier(EmitContext& ctx);
void EmitGetVcc(EmitContext& ctx);
void EmitSetVcc(EmitContext& ctx);
void EmitPrologue(EmitContext& ctx);
void EmitEpilogue(EmitContext& ctx);
void EmitGetScalarRegister(EmitContext& ctx);
void EmitSetScalarRegister(EmitContext& ctx);
void EmitGetVectorRegister(EmitContext& ctx);
void EmitSetVectorRegister(EmitContext& ctx);
void EmitSetGotoVariable(EmitContext& ctx);
void EmitGetGotoVariable(EmitContext& ctx);
void EmitSetScc(EmitContext& ctx);
Id EmitReadConst(EmitContext& ctx);
Id EmitReadConstBuffer(EmitContext& ctx, const IR::Value& handle, const IR::Value& index,
const IR::Value& offset);
Id EmitReadConstBufferF32(EmitContext& ctx, const IR::Value& handle, const IR::Value& index,
const IR::Value& offset);
Id EmitGetAttribute(EmitContext& ctx, IR::Attribute attr, Id vertex);
Id EmitGetAttributeU32(EmitContext& ctx, IR::Attribute attr, Id vertex);
void EmitSetAttribute(EmitContext& ctx, IR::Attribute attr, Id value, u32 element);
void EmitSetFragColor(EmitContext& ctx, u32 index, u32 component, Id value);
void EmitSetSampleMask(EmitContext& ctx, Id value);
void EmitSetFragDepth(EmitContext& ctx, Id value);
Id EmitWorkgroupId(EmitContext& ctx);
Id EmitLocalInvocationId(EmitContext& ctx);
Id EmitInvocationId(EmitContext& ctx);
Id EmitInvocationInfo(EmitContext& ctx);
Id EmitSampleId(EmitContext& ctx);
Id EmitUndefU1(EmitContext& ctx);
Id EmitUndefU8(EmitContext& ctx);
Id EmitUndefU16(EmitContext& ctx);
Id EmitUndefU32(EmitContext& ctx);
Id EmitUndefU64(EmitContext& ctx);
Id EmitReadSharedU8(EmitContext& ctx, Id offset);
Id EmitReadSharedS8(EmitContext& ctx, Id offset);
Id EmitReadSharedU16(EmitContext& ctx, Id offset);
Id EmitReadSharedS16(EmitContext& ctx, Id offset);
Id EmitReadSharedU32(EmitContext& ctx, Id offset);
Id EmitReadSharedU64(EmitContext& ctx, Id offset);
void EmitWriteSharedU8(EmitContext& ctx, Id offset, Id value);
void EmitWriteSharedU16(EmitContext& ctx, Id offset, Id value);
void EmitWriteSharedU32(EmitContext& ctx, Id offset, Id value);
void EmitWriteSharedU64(EmitContext& ctx, Id offset, Id value);
void EmitWriteSharedU128(EmitContext& ctx, Id offset, Id value);
Id EmitCompositeConstructU32x2(EmitContext& ctx, Id e1, Id e2);
Id EmitCompositeConstructU32x3(EmitContext& ctx, Id e1, Id e2, Id e3);
Id EmitCompositeConstructU32x4(EmitContext& ctx, Id e1, Id e2, Id e3, Id e4);
Id EmitCompositeExtractU32x2(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeExtractU32x3(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeExtractU32x4(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeInsertU32x2(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeInsertU32x3(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeInsertU32x4(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeConstructF16x2(EmitContext& ctx, Id e1, Id e2);
Id EmitCompositeConstructF16x3(EmitContext& ctx, Id e1, Id e2, Id e3);
Id EmitCompositeConstructF16x4(EmitContext& ctx, Id e1, Id e2, Id e3, Id e4);
Id EmitCompositeExtractF16x2(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeExtractF16x3(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeExtractF16x4(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeInsertF16x2(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeInsertF16x3(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeInsertF16x4(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeConstructF32x2(EmitContext& ctx, Id e1, Id e2);
Id EmitCompositeConstructF32x3(EmitContext& ctx, Id e1, Id e2, Id e3);
Id EmitCompositeConstructF32x4(EmitContext& ctx, Id e1, Id e2, Id e3, Id e4);
Id EmitCompositeExtractF32x2(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeExtractF32x3(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeExtractF32x4(EmitContext& ctx, Id composite, u32 index);
Id EmitCompositeInsertF32x2(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeInsertF32x3(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeInsertF32x4(EmitContext& ctx, Id composite, Id object, u32 index);
void EmitCompositeConstructF64x2(EmitContext& ctx);
void EmitCompositeConstructF64x3(EmitContext& ctx);
void EmitCompositeConstructF64x4(EmitContext& ctx);
void EmitCompositeExtractF64x2(EmitContext& ctx);
void EmitCompositeExtractF64x3(EmitContext& ctx);
void EmitCompositeExtractF64x4(EmitContext& ctx);
Id EmitCompositeInsertF64x2(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeInsertF64x3(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitCompositeInsertF64x4(EmitContext& ctx, Id composite, Id object, u32 index);
Id EmitSelectU1(EmitContext& ctx, Id cond, Id true_value, Id false_value);
Id EmitSelectU8(EmitContext& ctx, Id cond, Id true_value, Id false_value);
Id EmitSelectU16(EmitContext& ctx, Id cond, Id true_value, Id false_value);
Id EmitSelectU32(EmitContext& ctx, Id cond, Id true_value, Id false_value);
Id EmitSelectU64(EmitContext& ctx, Id cond, Id true_value, Id false_value);
Id EmitSelectF16(EmitContext& ctx, Id cond, Id true_value, Id false_value);
Id EmitSelectF32(EmitContext& ctx, Id cond, Id true_value, Id false_value);
Id EmitSelectF64(EmitContext& ctx, Id cond, Id true_value, Id false_value);
void EmitBitCastU16F16(EmitContext& ctx);
Id EmitBitCastU32F32(EmitContext& ctx, Id value);
void EmitBitCastU64F64(EmitContext& ctx);
void EmitBitCastF16U16(EmitContext&);
Id EmitBitCastF32U32(EmitContext& ctx, Id value);
void EmitBitCastF64U64(EmitContext& ctx);
Id EmitPackUint2x32(EmitContext& ctx, Id value);
Id EmitUnpackUint2x32(EmitContext& ctx, Id value);
Id EmitPackFloat2x16(EmitContext& ctx, Id value);
Id EmitUnpackFloat2x16(EmitContext& ctx, Id value);
Id EmitPackHalf2x16(EmitContext& ctx, Id value);
Id EmitUnpackHalf2x16(EmitContext& ctx, Id value);
Id EmitFPAbs16(EmitContext& ctx, Id value);
Id EmitFPAbs32(EmitContext& ctx, Id value);
Id EmitFPAbs64(EmitContext& ctx, Id value);
Id EmitFPAdd16(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitFPAdd32(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitFPAdd64(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitFPFma16(EmitContext& ctx, IR::Inst* inst, Id a, Id b, Id c);
Id EmitFPFma32(EmitContext& ctx, IR::Inst* inst, Id a, Id b, Id c);
Id EmitFPFma64(EmitContext& ctx, IR::Inst* inst, Id a, Id b, Id c);
Id EmitFPMax32(EmitContext& ctx, Id a, Id b);
Id EmitFPMax64(EmitContext& ctx, Id a, Id b);
Id EmitFPMin32(EmitContext& ctx, Id a, Id b);
Id EmitFPMin64(EmitContext& ctx, Id a, Id b);
Id EmitFPMul16(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitFPMul32(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitFPMul64(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitFPNeg16(EmitContext& ctx, Id value);
Id EmitFPNeg32(EmitContext& ctx, Id value);
Id EmitFPNeg64(EmitContext& ctx, Id value);
Id EmitFPSin(EmitContext& ctx, Id value);
Id EmitFPCos(EmitContext& ctx, Id value);
Id EmitFPExp2(EmitContext& ctx, Id value);
Id EmitFPLog2(EmitContext& ctx, Id value);
Id EmitFPRecip32(EmitContext& ctx, Id value);
Id EmitFPRecip64(EmitContext& ctx, Id value);
Id EmitFPRecipSqrt32(EmitContext& ctx, Id value);
Id EmitFPRecipSqrt64(EmitContext& ctx, Id value);
Id EmitFPSqrt(EmitContext& ctx, Id value);
Id EmitFPSaturate16(EmitContext& ctx, Id value);
Id EmitFPSaturate32(EmitContext& ctx, Id value);
Id EmitFPSaturate64(EmitContext& ctx, Id value);
Id EmitFPClamp16(EmitContext& ctx, Id value, Id min_value, Id max_value);
Id EmitFPClamp32(EmitContext& ctx, Id value, Id min_value, Id max_value);
Id EmitFPClamp64(EmitContext& ctx, Id value, Id min_value, Id max_value);
Id EmitFPRoundEven16(EmitContext& ctx, Id value);
Id EmitFPRoundEven32(EmitContext& ctx, Id value);
Id EmitFPRoundEven64(EmitContext& ctx, Id value);
Id EmitFPFloor16(EmitContext& ctx, Id value);
Id EmitFPFloor32(EmitContext& ctx, Id value);
Id EmitFPFloor64(EmitContext& ctx, Id value);
Id EmitFPCeil16(EmitContext& ctx, Id value);
Id EmitFPCeil32(EmitContext& ctx, Id value);
Id EmitFPCeil64(EmitContext& ctx, Id value);
Id EmitFPTrunc16(EmitContext& ctx, Id value);
Id EmitFPTrunc32(EmitContext& ctx, Id value);
Id EmitFPTrunc64(EmitContext& ctx, Id value);
Id EmitFPOrdEqual16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdEqual32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdEqual64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordEqual16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordEqual32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordEqual64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdNotEqual16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdNotEqual32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdNotEqual64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordNotEqual16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordNotEqual32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordNotEqual64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdLessThan16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdLessThan32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdLessThan64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordLessThan16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordLessThan32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordLessThan64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdGreaterThan16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdGreaterThan32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdGreaterThan64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordGreaterThan16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordGreaterThan32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordGreaterThan64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdLessThanEqual16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdLessThanEqual32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdLessThanEqual64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordLessThanEqual16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordLessThanEqual32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordLessThanEqual64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdGreaterThanEqual16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdGreaterThanEqual32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPOrdGreaterThanEqual64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordGreaterThanEqual16(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordGreaterThanEqual32(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPUnordGreaterThanEqual64(EmitContext& ctx, Id lhs, Id rhs);
Id EmitFPIsNan16(EmitContext& ctx, Id value);
Id EmitFPIsNan32(EmitContext& ctx, Id value);
Id EmitFPIsNan64(EmitContext& ctx, Id value);
Id EmitIAdd32(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitIAdd64(EmitContext& ctx, Id a, Id b);
Id EmitISub32(EmitContext& ctx, Id a, Id b);
Id EmitISub64(EmitContext& ctx, Id a, Id b);
Id EmitIMul32(EmitContext& ctx, Id a, Id b);
Id EmitSDiv32(EmitContext& ctx, Id a, Id b);
Id EmitUDiv32(EmitContext& ctx, Id a, Id b);
Id EmitINeg32(EmitContext& ctx, Id value);
Id EmitINeg64(EmitContext& ctx, Id value);
Id EmitIAbs32(EmitContext& ctx, Id value);
Id EmitShiftLeftLogical32(EmitContext& ctx, Id base, Id shift);
Id EmitShiftLeftLogical64(EmitContext& ctx, Id base, Id shift);
Id EmitShiftRightLogical32(EmitContext& ctx, Id base, Id shift);
Id EmitShiftRightLogical64(EmitContext& ctx, Id base, Id shift);
Id EmitShiftRightArithmetic32(EmitContext& ctx, Id base, Id shift);
Id EmitShiftRightArithmetic64(EmitContext& ctx, Id base, Id shift);
Id EmitBitwiseAnd32(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitBitwiseOr32(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitBitwiseXor32(EmitContext& ctx, IR::Inst* inst, Id a, Id b);
Id EmitBitFieldInsert(EmitContext& ctx, Id base, Id insert, Id offset, Id count);
Id EmitBitFieldSExtract(EmitContext& ctx, IR::Inst* inst, Id base, Id offset, Id count);
Id EmitBitFieldUExtract(EmitContext& ctx, IR::Inst* inst, Id base, Id offset, Id count);
Id EmitBitReverse32(EmitContext& ctx, Id value);
Id EmitBitCount32(EmitContext& ctx, Id value);
Id EmitBitwiseNot32(EmitContext& ctx, Id value);
Id EmitFindSMsb32(EmitContext& ctx, Id value);
Id EmitFindUMsb32(EmitContext& ctx, Id value);
Id EmitSMin32(EmitContext& ctx, Id a, Id b);
Id EmitUMin32(EmitContext& ctx, Id a, Id b);
Id EmitSMax32(EmitContext& ctx, Id a, Id b);
Id EmitUMax32(EmitContext& ctx, Id a, Id b);
Id EmitSClamp32(EmitContext& ctx, IR::Inst* inst, Id value, Id min, Id max);
Id EmitUClamp32(EmitContext& ctx, IR::Inst* inst, Id value, Id min, Id max);
Id EmitSLessThan(EmitContext& ctx, Id lhs, Id rhs);
Id EmitULessThan(EmitContext& ctx, Id lhs, Id rhs);
Id EmitIEqual(EmitContext& ctx, Id lhs, Id rhs);
Id EmitSLessThanEqual(EmitContext& ctx, Id lhs, Id rhs);
Id EmitULessThanEqual(EmitContext& ctx, Id lhs, Id rhs);
Id EmitSGreaterThan(EmitContext& ctx, Id lhs, Id rhs);
Id EmitUGreaterThan(EmitContext& ctx, Id lhs, Id rhs);
Id EmitINotEqual(EmitContext& ctx, Id lhs, Id rhs);
Id EmitSGreaterThanEqual(EmitContext& ctx, Id lhs, Id rhs);
Id EmitUGreaterThanEqual(EmitContext& ctx, Id lhs, Id rhs);
Id EmitLogicalOr(EmitContext& ctx, Id a, Id b);
Id EmitLogicalAnd(EmitContext& ctx, Id a, Id b);
Id EmitLogicalXor(EmitContext& ctx, Id a, Id b);
Id EmitLogicalNot(EmitContext& ctx, Id value);
Id EmitConvertS16F16(EmitContext& ctx, Id value);
Id EmitConvertS16F32(EmitContext& ctx, Id value);
Id EmitConvertS16F64(EmitContext& ctx, Id value);
Id EmitConvertS32F16(EmitContext& ctx, Id value);
Id EmitConvertS32F32(EmitContext& ctx, Id value);
Id EmitConvertS32F64(EmitContext& ctx, Id value);
Id EmitConvertS64F16(EmitContext& ctx, Id value);
Id EmitConvertS64F32(EmitContext& ctx, Id value);
Id EmitConvertS64F64(EmitContext& ctx, Id value);
Id EmitConvertU16F16(EmitContext& ctx, Id value);
Id EmitConvertU16F32(EmitContext& ctx, Id value);
Id EmitConvertU16F64(EmitContext& ctx, Id value);
Id EmitConvertU32F16(EmitContext& ctx, Id value);
Id EmitConvertU32F32(EmitContext& ctx, Id value);
Id EmitConvertU32F64(EmitContext& ctx, Id value);
Id EmitConvertU64F16(EmitContext& ctx, Id value);
Id EmitConvertU64F32(EmitContext& ctx, Id value);
Id EmitConvertU64F64(EmitContext& ctx, Id value);
Id EmitConvertU64U32(EmitContext& ctx, Id value);
Id EmitConvertU32U64(EmitContext& ctx, Id value);
Id EmitConvertF16F32(EmitContext& ctx, Id value);
Id EmitConvertF32F16(EmitContext& ctx, Id value);
Id EmitConvertF32F64(EmitContext& ctx, Id value);
Id EmitConvertF64F32(EmitContext& ctx, Id value);
Id EmitConvertF16S8(EmitContext& ctx, Id value);
Id EmitConvertF16S16(EmitContext& ctx, Id value);
Id EmitConvertF16S32(EmitContext& ctx, Id value);
Id EmitConvertF16S64(EmitContext& ctx, Id value);
Id EmitConvertF16U8(EmitContext& ctx, Id value);
Id EmitConvertF16U16(EmitContext& ctx, Id value);
Id EmitConvertF16U32(EmitContext& ctx, Id value);
Id EmitConvertF16U64(EmitContext& ctx, Id value);
Id EmitConvertF32S8(EmitContext& ctx, Id value);
Id EmitConvertF32S16(EmitContext& ctx, Id value);
Id EmitConvertF32S32(EmitContext& ctx, Id value);
Id EmitConvertF32S64(EmitContext& ctx, Id value);
Id EmitConvertF32U8(EmitContext& ctx, Id value);
Id EmitConvertF32U16(EmitContext& ctx, Id value);
Id EmitConvertF32U32(EmitContext& ctx, Id value);
Id EmitConvertF32U64(EmitContext& ctx, Id value);
Id EmitConvertF64S8(EmitContext& ctx, Id value);
Id EmitConvertF64S16(EmitContext& ctx, Id value);
Id EmitConvertF64S32(EmitContext& ctx, Id value);
Id EmitConvertF64S64(EmitContext& ctx, Id value);
Id EmitConvertF64U8(EmitContext& ctx, Id value);
Id EmitConvertF64U16(EmitContext& ctx, Id value);
Id EmitConvertF64U32(EmitContext& ctx, Id value);
Id EmitConvertF64U64(EmitContext& ctx, Id value);
Id EmitImageSampleImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
Id bias_lc, const IR::Value& offset);
Id EmitImageSampleExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
Id lod, const IR::Value& offset);
Id EmitImageSampleDrefImplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index,
Id coords, Id dref, Id bias_lc, const IR::Value& offset);
Id EmitImageSampleDrefExplicitLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index,
Id coords, Id dref, Id lod, const IR::Value& offset);
Id EmitImageGather(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
const IR::Value& offset, const IR::Value& offset2);
Id EmitImageGatherDref(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
const IR::Value& offset, const IR::Value& offset2, Id dref);
Id EmitImageFetch(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords, Id offset,
Id lod, Id ms);
Id EmitImageQueryDimensions(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id lod,
const IR::Value& skip_mips);
Id EmitImageQueryLod(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords);
Id EmitImageGradient(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords,
Id derivatives, const IR::Value& offset, Id lod_clamp);
Id EmitImageRead(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords);
void EmitImageWrite(EmitContext& ctx, IR::Inst* inst, const IR::Value& index, Id coords, Id color);
} // namespace Shader::Backend::SPIRV

262
src/shader_recompiler/backend/spirv/emit_spirv_integer.cpp Normal file
View file

@ -0,0 +1,262 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
namespace {
void SetZeroFlag(EmitContext& ctx, IR::Inst* inst, Id result) {
// IR::Inst* const zero{inst->GetAssociatedPseudoOperation(IR::Opcode::GetZeroFromOp)};
// if (!zero) {
// return;
// }
// zero->SetDefinition(ctx.OpIEqual(ctx.U1[1], result, ctx.u32_zero_value));
// zero->Invalidate();
}
void SetSignFlag(EmitContext& ctx, IR::Inst* inst, Id result) {
// IR::Inst* const sign{inst->GetAssociatedPseudoOperation(IR::Opcode::GetSignFromOp)};
// if (!sign) {
// return;
// }
// sign->SetDefinition(ctx.OpSLessThan(ctx.U1[1], result, ctx.u32_zero_value));
// sign->Invalidate();
}
} // Anonymous namespace
Id EmitIAdd32(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
return ctx.OpIAdd(ctx.U32[1], a, b);
// Id result{};
// if (IR::Inst* const carry{inst->GetAssociatedPseudoOperation(IR::Opcode::GetCarryFromOp)}) {
// const Id carry_type{ctx.TypeStruct(ctx.U32[1], ctx.U32[1])};
// const Id carry_result{ctx.OpIAddCarry(carry_type, a, b)};
// result = ctx.OpCompositeExtract(ctx.U32[1], carry_result, 0U);
// const Id carry_value{ctx.OpCompositeExtract(ctx.U32[1], carry_result, 1U)};
// carry->SetDefinition(ctx.OpINotEqual(ctx.U1[1][1], carry_value, ctx.u32_zero_value));
// carry->Invalidate();
//} else {
// result = ctx.OpIAdd(ctx.U32[1], a, b);
//}
// SetZeroFlag(ctx, inst, result);
// SetSignFlag(ctx, inst, result);
// if (IR::Inst * overflow{inst->GetAssociatedPseudoOperation(IR::Opcode::GetOverflowFromOp)}) {
// // https://stackoverflow.com/questions/55468823/how-to-detect-integer-overflow-in-c
// constexpr u32 s32_max{static_cast<u32>(std::numeric_limits<s32>::max())};
// const Id is_positive{ctx.OpSGreaterThanEqual(ctx.U1[1], a, ctx.u32_zero_value)};
// const Id sub_a{ctx.OpISub(ctx.U32[1], ctx.Const(s32_max), a)};
// const Id positive_test{ctx.OpSGreaterThan(ctx.U1[1], b, sub_a)};
// const Id negative_test{ctx.OpSLessThan(ctx.U1[1], b, sub_a)};
// const Id carry_flag{ctx.OpSelect(ctx.U1[1], is_positive, positive_test, negative_test)};
// overflow->SetDefinition(carry_flag);
// overflow->Invalidate();
//}
// return result;
}
Id EmitIAdd64(EmitContext& ctx, Id a, Id b) {
return ctx.OpIAdd(ctx.U64, a, b);
}
Id EmitISub32(EmitContext& ctx, Id a, Id b) {
return ctx.OpISub(ctx.U32[1], a, b);
}
Id EmitISub64(EmitContext& ctx, Id a, Id b) {
return ctx.OpISub(ctx.U64, a, b);
}
Id EmitIMul32(EmitContext& ctx, Id a, Id b) {
return ctx.OpIMul(ctx.U32[1], a, b);
}
Id EmitSDiv32(EmitContext& ctx, Id a, Id b) {
return ctx.OpSDiv(ctx.U32[1], a, b);
}
Id EmitUDiv32(EmitContext& ctx, Id a, Id b) {
return ctx.OpUDiv(ctx.U32[1], a, b);
}
Id EmitINeg32(EmitContext& ctx, Id value) {
return ctx.OpSNegate(ctx.U32[1], value);
}
Id EmitINeg64(EmitContext& ctx, Id value) {
return ctx.OpSNegate(ctx.U64, value);
}
Id EmitIAbs32(EmitContext& ctx, Id value) {
return ctx.OpSAbs(ctx.U32[1], value);
}
Id EmitShiftLeftLogical32(EmitContext& ctx, Id base, Id shift) {
return ctx.OpShiftLeftLogical(ctx.U32[1], base, shift);
}
Id EmitShiftLeftLogical64(EmitContext& ctx, Id base, Id shift) {
return ctx.OpShiftLeftLogical(ctx.U64, base, shift);
}
Id EmitShiftRightLogical32(EmitContext& ctx, Id base, Id shift) {
return ctx.OpShiftRightLogical(ctx.U32[1], base, shift);
}
Id EmitShiftRightLogical64(EmitContext& ctx, Id base, Id shift) {
return ctx.OpShiftRightLogical(ctx.U64, base, shift);
}
Id EmitShiftRightArithmetic32(EmitContext& ctx, Id base, Id shift) {
return ctx.OpShiftRightArithmetic(ctx.U32[1], base, shift);
}
Id EmitShiftRightArithmetic64(EmitContext& ctx, Id base, Id shift) {
return ctx.OpShiftRightArithmetic(ctx.U64, base, shift);
}
Id EmitBitwiseAnd32(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
const Id result{ctx.OpBitwiseAnd(ctx.U32[1], a, b)};
SetZeroFlag(ctx, inst, result);
SetSignFlag(ctx, inst, result);
return result;
}
Id EmitBitwiseOr32(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
const Id result{ctx.OpBitwiseOr(ctx.U32[1], a, b)};
SetZeroFlag(ctx, inst, result);
SetSignFlag(ctx, inst, result);
return result;
}
Id EmitBitwiseXor32(EmitContext& ctx, IR::Inst* inst, Id a, Id b) {
const Id result{ctx.OpBitwiseXor(ctx.U32[1], a, b)};
SetZeroFlag(ctx, inst, result);
SetSignFlag(ctx, inst, result);
return result;
}
Id EmitBitFieldInsert(EmitContext& ctx, Id base, Id insert, Id offset, Id count) {
return ctx.OpBitFieldInsert(ctx.U32[1], base, insert, offset, count);
}
Id EmitBitFieldSExtract(EmitContext& ctx, IR::Inst* inst, Id base, Id offset, Id count) {
const Id result{ctx.OpBitFieldSExtract(ctx.U32[1], base, offset, count)};
SetZeroFlag(ctx, inst, result);
SetSignFlag(ctx, inst, result);
return result;
}
Id EmitBitFieldUExtract(EmitContext& ctx, IR::Inst* inst, Id base, Id offset, Id count) {
const Id result{ctx.OpBitFieldUExtract(ctx.U32[1], base, offset, count)};
SetZeroFlag(ctx, inst, result);
SetSignFlag(ctx, inst, result);
return result;
}
Id EmitBitReverse32(EmitContext& ctx, Id value) {
return ctx.OpBitReverse(ctx.U32[1], value);
}
Id EmitBitCount32(EmitContext& ctx, Id value) {
return ctx.OpBitCount(ctx.U32[1], value);
}
Id EmitBitwiseNot32(EmitContext& ctx, Id value) {
return ctx.OpNot(ctx.U32[1], value);
}
Id EmitFindSMsb32(EmitContext& ctx, Id value) {
return ctx.OpFindSMsb(ctx.U32[1], value);
}
Id EmitFindUMsb32(EmitContext& ctx, Id value) {
return ctx.OpFindUMsb(ctx.U32[1], value);
}
Id EmitSMin32(EmitContext& ctx, Id a, Id b) {
return ctx.OpSMin(ctx.U32[1], a, b);
}
Id EmitUMin32(EmitContext& ctx, Id a, Id b) {
return ctx.OpUMin(ctx.U32[1], a, b);
}
Id EmitSMax32(EmitContext& ctx, Id a, Id b) {
return ctx.OpSMax(ctx.U32[1], a, b);
}
Id EmitUMax32(EmitContext& ctx, Id a, Id b) {
return ctx.OpUMax(ctx.U32[1], a, b);
}
Id EmitSClamp32(EmitContext& ctx, IR::Inst* inst, Id value, Id min, Id max) {
Id result{};
if (ctx.profile.has_broken_spirv_clamp) {
value = ctx.OpBitcast(ctx.S32[1], value);
min = ctx.OpBitcast(ctx.S32[1], min);
max = ctx.OpBitcast(ctx.S32[1], max);
result = ctx.OpSMax(ctx.S32[1], ctx.OpSMin(ctx.S32[1], value, max), min);
result = ctx.OpBitcast(ctx.U32[1], result);
} else {
result = ctx.OpSClamp(ctx.U32[1], value, min, max);
}
SetZeroFlag(ctx, inst, result);
SetSignFlag(ctx, inst, result);
return result;
}
Id EmitUClamp32(EmitContext& ctx, IR::Inst* inst, Id value, Id min, Id max) {
Id result{};
if (ctx.profile.has_broken_spirv_clamp) {
result = ctx.OpUMax(ctx.U32[1], ctx.OpUMin(ctx.U32[1], value, max), min);
} else {
result = ctx.OpUClamp(ctx.U32[1], value, min, max);
}
SetZeroFlag(ctx, inst, result);
SetSignFlag(ctx, inst, result);
return result;
}
Id EmitSLessThan(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpSLessThan(ctx.U1[1], lhs, rhs);
}
Id EmitULessThan(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpULessThan(ctx.U1[1], lhs, rhs);
}
Id EmitIEqual(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpIEqual(ctx.U1[1], lhs, rhs);
}
Id EmitSLessThanEqual(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpSLessThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitULessThanEqual(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpULessThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitSGreaterThan(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpSGreaterThan(ctx.U1[1], lhs, rhs);
}
Id EmitUGreaterThan(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpUGreaterThan(ctx.U1[1], lhs, rhs);
}
Id EmitINotEqual(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpINotEqual(ctx.U1[1], lhs, rhs);
}
Id EmitSGreaterThanEqual(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpSGreaterThanEqual(ctx.U1[1], lhs, rhs);
}
Id EmitUGreaterThanEqual(EmitContext& ctx, Id lhs, Id rhs) {
return ctx.OpUGreaterThanEqual(ctx.U1[1], lhs, rhs);
}
} // namespace Shader::Backend::SPIRV

25
src/shader_recompiler/backend/spirv/emit_spirv_logical.cpp Normal file
View file

@ -0,0 +1,25 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
Id EmitLogicalOr(EmitContext& ctx, Id a, Id b) {
return ctx.OpLogicalOr(ctx.U1[1], a, b);
}
Id EmitLogicalAnd(EmitContext& ctx, Id a, Id b) {
return ctx.OpLogicalAnd(ctx.U1[1], a, b);
}
Id EmitLogicalXor(EmitContext& ctx, Id a, Id b) {
return ctx.OpLogicalNotEqual(ctx.U1[1], a, b);
}
Id EmitLogicalNot(EmitContext& ctx, Id value) {
return ctx.OpLogicalNot(ctx.U1[1], value);
}
} // namespace Shader::Backend::SPIRV

41
src/shader_recompiler/backend/spirv/emit_spirv_select.cpp Normal file
View file

@ -0,0 +1,41 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
Id EmitSelectU1(EmitContext& ctx, Id cond, Id true_value, Id false_value) {
return ctx.OpSelect(ctx.U1[1], cond, true_value, false_value);
}
Id EmitSelectU8(EmitContext&, Id, Id, Id) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitSelectU16(EmitContext& ctx, Id cond, Id true_value, Id false_value) {
return ctx.OpSelect(ctx.U16, cond, true_value, false_value);
}
Id EmitSelectU32(EmitContext& ctx, Id cond, Id true_value, Id false_value) {
return ctx.OpSelect(ctx.U32[1], cond, true_value, false_value);
}
Id EmitSelectU64(EmitContext& ctx, Id cond, Id true_value, Id false_value) {
return ctx.OpSelect(ctx.U64, cond, true_value, false_value);
}
Id EmitSelectF16(EmitContext& ctx, Id cond, Id true_value, Id false_value) {
return ctx.OpSelect(ctx.F16[1], cond, true_value, false_value);
}
Id EmitSelectF32(EmitContext& ctx, Id cond, Id true_value, Id false_value) {
return ctx.OpSelect(ctx.F32[1], cond, true_value, false_value);
}
Id EmitSelectF64(EmitContext& ctx, Id cond, Id true_value, Id false_value) {
return ctx.OpSelect(ctx.F64[1], cond, true_value, false_value);
}
} // namespace Shader::Backend::SPIRV

21
src/shader_recompiler/backend/spirv/emit_spirv_special.cpp Normal file
View file

@ -0,0 +1,21 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
void EmitPrologue(EmitContext& ctx) {}
void EmitEpilogue(EmitContext& ctx) {}
void EmitEmitVertex(EmitContext& ctx, const IR::Value& stream) {
throw NotImplementedException("Geometry streams");
}
void EmitEndPrimitive(EmitContext& ctx, const IR::Value& stream) {
throw NotImplementedException("Geometry streams");
}
} // namespace Shader::Backend::SPIRV

29
src/shader_recompiler/backend/spirv/emit_spirv_undefined.cpp Normal file
View file

@ -0,0 +1,29 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
Id EmitUndefU1(EmitContext& ctx) {
return ctx.OpUndef(ctx.U1[1]);
}
Id EmitUndefU8(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitUndefU16(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
Id EmitUndefU32(EmitContext& ctx) {
return ctx.OpUndef(ctx.U32[1]);
}
Id EmitUndefU64(EmitContext&) {
throw NotImplementedException("SPIR-V Instruction");
}
} // namespace Shader::Backend::SPIRV

136
src/shader_recompiler/backend/spirv/spirv_emit_context.cpp Normal file
View file

@ -0,0 +1,136 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include <boost/container/static_vector.hpp>
#include <fmt/format.h>
#include "shader_recompiler/backend/spirv/spirv_emit_context.h"
namespace Shader::Backend::SPIRV {
namespace {
std::string_view StageName(Stage stage) {
switch (stage) {
case Stage::Vertex:
return "vs";
case Stage::TessellationControl:
return "tcs";
case Stage::TessellationEval:
return "tes";
case Stage::Geometry:
return "gs";
case Stage::Fragment:
return "fs";
case Stage::Compute:
return "cs";
}
throw InvalidArgument("Invalid stage {}", u32(stage));
}
template <typename... Args>
void Name(EmitContext& ctx, Id object, std::string_view format_str, Args&&... args) {
ctx.Name(object, fmt::format(fmt::runtime(format_str), StageName(ctx.stage),
std::forward<Args>(args)...)
.c_str());
}
} // Anonymous namespace
EmitContext::EmitContext(const Profile& profile_, IR::Program& program, Bindings& bindings)
: Sirit::Module(profile_.supported_spirv), profile{profile_}, stage{program.stage} {
u32& uniform_binding{bindings.unified};
u32& storage_binding{bindings.unified};
u32& texture_binding{bindings.unified};
u32& image_binding{bindings.unified};
AddCapability(spv::Capability::Shader);
DefineArithmeticTypes();
DefineInterfaces(program);
}
EmitContext::~EmitContext() = default;
Id EmitContext::Def(const IR::Value& value) {
if (!value.IsImmediate()) {
return value.InstRecursive()->Definition<Id>();
}
switch (value.Type()) {
case IR::Type::Void:
return Id{};
case IR::Type::U1:
return value.U1() ? true_value : false_value;
case IR::Type::U32:
return ConstU32(value.U32());
case IR::Type::U64:
return Constant(U64, value.U64());
case IR::Type::F32:
return ConstF32(value.F32());
case IR::Type::F64:
return Constant(F64[1], value.F64());
default:
throw NotImplementedException("Immediate type {}", value.Type());
}
}
void EmitContext::DefineArithmeticTypes() {
void_id = Name(TypeVoid(), "void_id");
U1[1] = Name(TypeBool(), "bool_id");
// F16[1] = Name(TypeFloat(16), "f16_id");
F32[1] = Name(TypeFloat(32), "f32_id");
// F64[1] = Name(TypeFloat(64), "f64_id");
S32[1] = Name(TypeSInt(32), "i32_id");
U32[1] = Name(TypeUInt(32), "u32_id");
// U8 = Name(TypeSInt(8), "u8");
// S8 = Name(TypeUInt(8), "s8");
// U16 = Name(TypeUInt(16), "u16_id");
// S16 = Name(TypeSInt(16), "s16_id");
// U64 = Name(TypeUInt(64), "u64_id");
for (u32 i = 2; i <= 4; i++) {
// F16[i] = Name(TypeVector(F16[1], i), fmt::format("f16vec{}_id", i));
F32[i] = Name(TypeVector(F32[1], i), fmt::format("f32vec{}_id", i));
// F64[i] = Name(TypeVector(F64[1], i), fmt::format("f64vec{}_id", i));
S32[i] = Name(TypeVector(S32[1], i), fmt::format("i32vec{}_id", i));
U32[i] = Name(TypeVector(U32[1], i), fmt::format("u32vec{}_id", i));
U1[i] = Name(TypeVector(U1[1], i), fmt::format("bvec{}_id", i));
}
true_value = ConstantTrue(U1[1]);
false_value = ConstantFalse(U1[1]);
u32_zero_value = ConstU32(0U);
f32_zero_value = ConstF32(0.0f);
output_f32 = Name(TypePointer(spv::StorageClass::Output, F32[1]), "output_f32");
output_u32 = Name(TypePointer(spv::StorageClass::Output, U32[1]), "output_u32");
}
void EmitContext::DefineInterfaces(const IR::Program& program) {
DefineInputs(program);
DefineOutputs(program);
}
void EmitContext::DefineInputs(const IR::Program& program) {
switch (stage) {
case Stage::Vertex:
vertex_index = DefineVariable(U32[1], spv::BuiltIn::VertexIndex, spv::StorageClass::Input);
base_vertex = DefineVariable(U32[1], spv::BuiltIn::BaseVertex, spv::StorageClass::Input);
break;
default:
break;
}
}
void EmitContext::DefineOutputs(const IR::Program& program) {
switch (stage) {
case Stage::Vertex:
output_position = DefineVariable(F32[4], spv::BuiltIn::Position, spv::StorageClass::Output);
break;
case Stage::Fragment:
frag_color[0] = DefineOutput(F32[4], 0);
Name(frag_color[0], fmt::format("frag_color{}", 0));
interfaces.push_back(frag_color[0]);
break;
default:
break;
}
}
} // namespace Shader::Backend::SPIRV

169
src/shader_recompiler/backend/spirv/spirv_emit_context.h Normal file
View file

@ -0,0 +1,169 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#pragma once
#include <array>
#include <sirit/sirit.h>
#include "shader_recompiler/backend/bindings.h"
#include "shader_recompiler/ir/program.h"
#include "shader_recompiler/profile.h"
#include "shader_recompiler/runtime_info.h"
namespace Shader::Backend::SPIRV {
using Sirit::Id;
struct VectorIds {
[[nodiscard]] Id& operator[](u32 index) {
return ids[index - 1];
}
[[nodiscard]] const Id& operator[](u32 index) const {
return ids[index - 1];
}
std::array<Id, 4> ids;
};
class EmitContext final : public Sirit::Module {
public:
explicit EmitContext(const Profile& profile, IR::Program& program, Bindings& binding);
~EmitContext();
Id Def(const IR::Value& value);
[[nodiscard]] Id DefineInput(Id type, u32 location) {
const Id input_id{DefineVar(type, spv::StorageClass::Input)};
Decorate(input_id, spv::Decoration::Location, location);
return input_id;
}
[[nodiscard]] Id DefineOutput(Id type, std::optional<u32> location = std::nullopt) {
const Id output_id{DefineVar(type, spv::StorageClass::Output)};
if (location) {
Decorate(output_id, spv::Decoration::Location, *location);
}
return output_id;
}
[[nodiscard]] Id DefineUniformConst(Id type, u32 set, u32 binding, bool readonly = false) {
const Id uniform_id{DefineVar(type, spv::StorageClass::UniformConstant)};
Decorate(uniform_id, spv::Decoration::DescriptorSet, set);
Decorate(uniform_id, spv::Decoration::Binding, binding);
if (readonly) {
Decorate(uniform_id, spv::Decoration::NonWritable);
}
return uniform_id;
}
template <bool global = true>
[[nodiscard]] Id DefineVar(Id type, spv::StorageClass storage_class) {
const Id pointer_type_id{TypePointer(storage_class, type)};
return global ? AddGlobalVariable(pointer_type_id, storage_class)
: AddLocalVariable(pointer_type_id, storage_class);
}
[[nodiscard]] Id DefineVariable(Id type, std::optional<spv::BuiltIn> builtin,
spv::StorageClass storage_class) {
const Id id{DefineVar(type, storage_class)};
if (builtin) {
Decorate(id, spv::Decoration::BuiltIn, *builtin);
}
interfaces.push_back(id);
return id;
}
[[nodiscard]] Id ConstU32(u32 value) {
return Constant(U32[1], value);
}
template <typename... Args>
[[nodiscard]] Id ConstU32(Args&&... values) {
constexpr u32 size = static_cast<u32>(sizeof...(values));
static_assert(size >= 2);
const std::array constituents{Constant(U32[1], values)...};
const Id type = size <= 4 ? U32[size] : TypeArray(U32[1], ConstU32(size));
return ConstantComposite(type, constituents);
}
[[nodiscard]] Id ConstS32(s32 value) {
return Constant(S32[1], value);
}
template <typename... Args>
[[nodiscard]] Id ConstS32(Args&&... values) {
constexpr u32 size = static_cast<u32>(sizeof...(values));
static_assert(size >= 2);
const std::array constituents{Constant(S32[1], values)...};
const Id type = size <= 4 ? S32[size] : TypeArray(S32[1], ConstU32(size));
return ConstantComposite(type, constituents);
}
[[nodiscard]] Id ConstF32(f32 value) {
return Constant(F32[1], value);
}
template <typename... Args>
[[nodiscard]] Id ConstF32(Args... values) {
constexpr u32 size = static_cast<u32>(sizeof...(values));
static_assert(size >= 2);
const std::array constituents{Constant(F32[1], values)...};
const Id type = size <= 4 ? F32[size] : TypeArray(F32[1], ConstU32(size));
return ConstantComposite(type, constituents);
}
const Profile& profile;
Stage stage{};
Id void_id{};
Id U8{};
Id S8{};
Id U16{};
Id S16{};
Id U64{};
VectorIds F16{};
VectorIds F32{};
VectorIds F64{};
VectorIds S32{};
VectorIds U32{};
VectorIds U1{};
Id true_value{};
Id false_value{};
Id u32_zero_value{};
Id f32_zero_value{};
Id output_u32{};
Id output_f32{};
boost::container::small_vector<Id, 16> interfaces;
Id output_position{};
Id vertex_index{};
Id base_vertex{};
std::array<Id, 8> frag_color{};
struct InputParamInfo {
Id id;
Id pointer_type;
Id component_type;
};
std::array<InputParamInfo, 32> input_params{};
struct ParamElementInfo {
Id id{};
u32 first_element{};
u32 num_components{};
};
std::array<std::array<ParamElementInfo, 4>, 32> output_params{};
private:
void DefineArithmeticTypes();
void DefineInterfaces(const IR::Program& program);
void DefineInputs(const IR::Program& program);
void DefineOutputs(const IR::Program& program);
};
} // namespace Shader::Backend::SPIRV