mirrors/shadPS4
1
0
Fork 0
mirror of https://github.com/shadps4-emu/shadPS4.git synced 2024-10-20 14:50:01 +00:00
Code Issues Actions 3 Packages Projects Releases Wiki Activity

shader_recompiler: Implement render target swizzles when no format is available (#739)

Browse source 
* shader_recompiler: Use null image when shader is compiled with unbound sharp

* video_core: Refactor and render target swizzles

* liverpool_to_vk: Add missing swap format from RDR

* video_core: Refactor shader recompiler interface

* Makes it much easier to pass runtime information to the recompiler and have it treated as part of the shader key. Also pulls out most runtime state from Info struct

* shader_recompiler: Avoid some asserts
This commit is contained in:
TheTurtle 2024-09-03 14:04:30 +03:00 committed by GitHub
parent 3f8a8d3a24
commit f087f43736
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
30 changed files with 704 additions and 560 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
CMakeLists.txt
View file

@ -419,7 +419,10 @@ set(SHADER_RECOMPILER src/shader_recompiler/exception.h
src/shader_recompiler/profile.h
src/shader_recompiler/recompiler.cpp
src/shader_recompiler/recompiler.h
src/shader_recompiler/info.h
src/shader_recompiler/params.h
src/shader_recompiler/runtime_info.h
src/shader_recompiler/specialization.h
src/shader_recompiler/backend/spirv/emit_spirv.cpp
src/shader_recompiler/backend/spirv/emit_spirv.h
src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp
@ -533,8 +536,6 @@ set(VIDEO_CORE src/video_core/amdgpu/liverpool.cpp
src/video_core/renderer_vulkan/vk_resource_pool.h
src/video_core/renderer_vulkan/vk_scheduler.cpp
src/video_core/renderer_vulkan/vk_scheduler.h
src/video_core/renderer_vulkan/vk_shader_cache.cpp
src/video_core/renderer_vulkan/vk_shader_cache.h
src/video_core/renderer_vulkan/vk_shader_util.cpp
src/video_core/renderer_vulkan/vk_shader_util.h
src/video_core/renderer_vulkan/vk_swapchain.cpp

7
src/shader_recompiler/backend/spirv/emit_spirv.cpp
View file

@ -210,7 +210,7 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
}
switch (program.info.stage) {
case Stage::Compute: {
const std::array<u32, 3> workgroup_size{program.info.workgroup_size};
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]);
@ -258,8 +258,9 @@ void PatchPhiNodes(const IR::Program& program, EmitContext& ctx) {
}
} // Anonymous namespace
std::vector<u32> EmitSPIRV(const Profile& profile, const IR::Program& program, u32& binding) {
EmitContext ctx{profile, program.info, binding};
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) {

4
src/shader_recompiler/backend/spirv/emit_spirv.h
View file

@ -9,7 +9,7 @@
namespace Shader::Backend::SPIRV {
[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, const IR::Program& program,
u32& binding);
[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info,
const IR::Program& program, u32& binding);
} // namespace Shader::Backend::SPIRV

2
src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
View file

@ -59,7 +59,7 @@ Id OutputAttrPointer(EmitContext& ctx, IR::Attribute attr, u32 element) {
case IR::Attribute::Position2:
case IR::Attribute::Position3: {
const u32 index = u32(attr) - u32(IR::Attribute::Position1);
return VsOutputAttrPointer(ctx, ctx.info.vs_outputs[index][element]);
return VsOutputAttrPointer(ctx, ctx.runtime_info.vs_info.outputs[index][element]);
}
case IR::Attribute::RenderTarget0:
case IR::Attribute::RenderTarget1:

15
src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
View file

@ -41,9 +41,10 @@ void Name(EmitContext& ctx, Id object, std::string_view format_str, Args&&... ar
} // Anonymous namespace
EmitContext::EmitContext(const Profile& profile_, const Shader::Info& info_, u32& binding_)
: Sirit::Module(profile_.supported_spirv), info{info_}, profile{profile_}, stage{info.stage},
binding{binding_} {
EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_info_,
const Info& info_, u32& binding_)
: Sirit::Module(profile_.supported_spirv), info{info_}, runtime_info{runtime_info_},
profile{profile_}, stage{info.stage}, binding{binding_} {
AddCapability(spv::Capability::Shader);
DefineArithmeticTypes();
DefineInterfaces();
@ -168,7 +169,7 @@ EmitContext::SpirvAttribute EmitContext::GetAttributeInfo(AmdGpu::NumberFormat f
void EmitContext::DefineBufferOffsets() {
for (auto& buffer : buffers) {
const u32 binding = buffer.binding;
const u32 half = Shader::PushData::BufOffsetIndex + (binding >> 4);
const u32 half = PushData::BufOffsetIndex + (binding >> 4);
const u32 comp = (binding & 0xf) >> 2;
const u32 offset = (binding & 0x3) << 3;
const Id ptr{OpAccessChain(TypePointer(spv::StorageClass::PushConstant, U32[1]),
@ -179,7 +180,7 @@ void EmitContext::DefineBufferOffsets() {
}
for (auto& tex_buffer : texture_buffers) {
const u32 binding = tex_buffer.binding;
const u32 half = Shader::PushData::BufOffsetIndex + (binding >> 4);
const u32 half = PushData::BufOffsetIndex + (binding >> 4);
const u32 comp = (binding & 0xf) >> 2;
const u32 offset = (binding & 0x3) << 3;
const Id ptr{OpAccessChain(TypePointer(spv::StorageClass::PushConstant, U32[1]),
@ -247,7 +248,7 @@ void EmitContext::DefineInputs() {
frag_coord = DefineVariable(F32[4], spv::BuiltIn::FragCoord, spv::StorageClass::Input);
frag_depth = DefineVariable(F32[1], spv::BuiltIn::FragDepth, spv::StorageClass::Output);
front_facing = DefineVariable(U1[1], spv::BuiltIn::FrontFacing, spv::StorageClass::Input);
for (const auto& input : info.ps_inputs) {
for (const auto& input : runtime_info.fs_info.inputs) {
const u32 semantic = input.param_index;
if (input.is_default && !input.is_flat) {
input_params[semantic] = {MakeDefaultValue(*this, input.default_value), F32[1],
@ -554,7 +555,7 @@ void EmitContext::DefineSharedMemory() {
if (!info.uses_shared) {
return;
}
u32 shared_memory_size = info.shared_memory_size;
u32 shared_memory_size = runtime_info.cs_info.shared_memory_size;
if (shared_memory_size == 0) {
shared_memory_size = DefaultSharedMemSize;
}

6
src/shader_recompiler/backend/spirv/spirv_emit_context.h
View file

@ -6,9 +6,9 @@
#include <array>
#include <sirit/sirit.h>
#include "shader_recompiler/info.h"
#include "shader_recompiler/ir/program.h"
#include "shader_recompiler/profile.h"
#include "shader_recompiler/runtime_info.h"
namespace Shader::Backend::SPIRV {
@ -36,7 +36,8 @@ struct VectorIds {
class EmitContext final : public Sirit::Module {
public:
explicit EmitContext(const Profile& profile, const Shader::Info& info, u32& binding);
explicit EmitContext(const Profile& profile, const RuntimeInfo& runtime_info, const Info& info,
u32& binding);
~EmitContext();
Id Def(const IR::Value& value);
@ -125,6 +126,7 @@ public:
}
const Info& info;
const RuntimeInfo& runtime_info;
const Profile& profile;
Stage stage{};

16
src/shader_recompiler/frontend/structured_control_flow.cpp
View file

@ -602,13 +602,14 @@ public:
Common::ObjectPool<IR::Block>& block_pool_,
Common::ObjectPool<Statement>& stmt_pool_, Statement& root_stmt,
IR::AbstractSyntaxList& syntax_list_, std::span<const GcnInst> inst_list_,
Info& info_, const Profile& profile_)
Info& info_, const RuntimeInfo& runtime_info_, const Profile& profile_)
: stmt_pool{stmt_pool_}, inst_pool{inst_pool_}, block_pool{block_pool_},
syntax_list{syntax_list_}, inst_list{inst_list_}, info{info_}, profile{profile_} {
syntax_list{syntax_list_}, inst_list{inst_list_}, info{info_},
runtime_info{runtime_info_}, profile{profile_} {
Visit(root_stmt, nullptr, nullptr);
IR::Block& first_block{*syntax_list.front().data.block};
Translator{&first_block, info, profile}.EmitPrologue();
Translator{&first_block, info, runtime_info, profile}.EmitPrologue();
}
private:
@ -637,7 +638,7 @@ private:
const u32 start = stmt.block->begin_index;
const u32 size = stmt.block->end_index - start + 1;
Translate(current_block, stmt.block->begin, inst_list.subspan(start, size),
info, profile);
info, runtime_info, profile);
}
break;
}
@ -817,19 +818,20 @@ private:
const Block dummy_flow_block{.is_dummy = true};
std::span<const GcnInst> inst_list;
Info& info;
const RuntimeInfo& runtime_info;
const Profile& profile;
};
} // Anonymous namespace
IR::AbstractSyntaxList BuildASL(Common::ObjectPool<IR::Inst>& inst_pool,
Common::ObjectPool<IR::Block>& block_pool, CFG& cfg, Info& info,
const Profile& profile) {
const RuntimeInfo& runtime_info, const Profile& profile) {
Common::ObjectPool<Statement> stmt_pool{64};
GotoPass goto_pass{cfg, stmt_pool};
Statement& root{goto_pass.RootStatement()};
IR::AbstractSyntaxList syntax_list;
TranslatePass{inst_pool, block_pool, stmt_pool, root,
syntax_list, cfg.inst_list, info, profile};
TranslatePass{inst_pool, block_pool, stmt_pool, root, syntax_list,
cfg.inst_list, info, runtime_info, profile};
ASSERT_MSG(!info.translation_failed, "Shader translation has failed");
return syntax_list;
}

4
src/shader_recompiler/frontend/structured_control_flow.h
View file

@ -11,12 +11,14 @@
namespace Shader {
struct Info;
struct Profile;
struct RuntimeInfo;
} // namespace Shader
namespace Shader::Gcn {
[[nodiscard]] IR::AbstractSyntaxList BuildASL(Common::ObjectPool<IR::Inst>& inst_pool,
Common::ObjectPool<IR::Block>& block_pool, CFG& cfg,
Info& info, const Profile& profile);
Info& info, const RuntimeInfo& runtime_info,
const Profile& profile);
} // namespace Shader::Gcn

23
src/shader_recompiler/frontend/translate/export.cpp
View file

@ -2,6 +2,7 @@
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/frontend/translate/translate.h"
#include "shader_recompiler/runtime_info.h"
namespace Shader::Gcn {
@ -19,12 +20,28 @@ void Translator::EmitExport(const GcnInst& inst) {
IR::VectorReg(inst.src[3].code),
};
const auto swizzle = [&](u32 comp) {
if (!IR::IsMrt(attrib)) {
return comp;
}
const u32 index = u32(attrib) - u32(IR::Attribute::RenderTarget0);
switch (runtime_info.fs_info.mrt_swizzles[index]) {
case MrtSwizzle::Identity:
return comp;
case MrtSwizzle::Alt:
static constexpr std::array<u32, 4> AltSwizzle = {2, 1, 0, 3};
return AltSwizzle[comp];
default:
UNREACHABLE();
}
};
const auto unpack = [&](u32 idx) {
const IR::Value value = ir.UnpackHalf2x16(ir.GetVectorReg(vsrc[idx]));
const IR::F32 r = IR::F32{ir.CompositeExtract(value, 0)};
const IR::F32 g = IR::F32{ir.CompositeExtract(value, 1)};
ir.SetAttribute(attrib, r, idx * 2);
ir.SetAttribute(attrib, g, idx * 2 + 1);
ir.SetAttribute(attrib, r, swizzle(idx * 2));
ir.SetAttribute(attrib, g, swizzle(idx * 2 + 1));
};
// Components are float16 packed into a VGPR
@ -45,7 +62,7 @@ void Translator::EmitExport(const GcnInst& inst) {
continue;
}
const IR::F32 comp = ir.GetVectorReg<IR::F32>(vsrc[i]);
ir.SetAttribute(attrib, comp, i);
ir.SetAttribute(attrib, comp, swizzle(i));
}
}
}

25
src/shader_recompiler/frontend/translate/translate.cpp
View file

@ -7,6 +7,7 @@
#include "shader_recompiler/exception.h"
#include "shader_recompiler/frontend/fetch_shader.h"
#include "shader_recompiler/frontend/translate/translate.h"
#include "shader_recompiler/info.h"
#include "shader_recompiler/runtime_info.h"
#include "video_core/amdgpu/resource.h"
@ -16,8 +17,9 @@
namespace Shader::Gcn {
Translator::Translator(IR::Block* block_, Info& info_, const Profile& profile_)
: ir{*block_, block_->begin()}, info{info_}, profile{profile_} {}
Translator::Translator(IR::Block* block_, Info& info_, const RuntimeInfo& runtime_info_,
const Profile& profile_)
: ir{*block_, block_->begin()}, info{info_}, runtime_info{runtime_info_}, profile{profile_} {}
void Translator::EmitPrologue() {
ir.Prologue();
@ -25,7 +27,7 @@ void Translator::EmitPrologue() {
// Initialize user data.
IR::ScalarReg dst_sreg = IR::ScalarReg::S0;
for (u32 i = 0; i < info.num_user_data; i++) {
for (u32 i = 0; i < runtime_info.num_user_data; i++) {
ir.SetScalarReg(dst_sreg, ir.GetUserData(dst_sreg));
++dst_sreg;
}
@ -36,15 +38,15 @@ void Translator::EmitPrologue() {
// v0: vertex ID, always present
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::VertexId));
// v1: instance ID, step rate 0
if (info.num_input_vgprs > 0) {
if (runtime_info.num_input_vgprs > 0) {
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::InstanceId0));
}
// v2: instance ID, step rate 1
if (info.num_input_vgprs > 1) {
if (runtime_info.num_input_vgprs > 1) {
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::InstanceId1));
}
// v3: instance ID, plain
if (info.num_input_vgprs > 2) {
if (runtime_info.num_input_vgprs > 2) {
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::InstanceId));
}
break;
@ -64,13 +66,13 @@ void Translator::EmitPrologue() {
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::LocalInvocationId, 1));
ir.SetVectorReg(dst_vreg++, ir.GetAttributeU32(IR::Attribute::LocalInvocationId, 2));
if (info.tgid_enable[0]) {
if (runtime_info.cs_info.tgid_enable[0]) {
ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 0));
}
if (info.tgid_enable[1]) {
if (runtime_info.cs_info.tgid_enable[1]) {
ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 1));
}
if (info.tgid_enable[2]) {
if (runtime_info.cs_info.tgid_enable[2]) {
ir.SetScalarReg(dst_sreg++, ir.GetAttributeU32(IR::Attribute::WorkgroupId, 2));
}
break;
@ -445,7 +447,6 @@ void Translator::EmitFlowControl(u32 pc, const GcnInst& inst) {
}
void Translator::LogMissingOpcode(const GcnInst& inst) {
const u32 opcode = u32(inst.opcode);
LOG_ERROR(Render_Recompiler, "Unknown opcode {} ({}, category = {})",
magic_enum::enum_name(inst.opcode), u32(inst.opcode),
magic_enum::enum_name(inst.category));
@ -453,11 +454,11 @@ void Translator::LogMissingOpcode(const GcnInst& inst) {
}
void Translate(IR::Block* block, u32 pc, std::span<const GcnInst> inst_list, Info& info,
const Profile& profile) {
const RuntimeInfo& runtime_info, const Profile& profile) {
if (inst_list.empty()) {
return;
}
Translator translator{block, info, profile};
Translator translator{block, info, runtime_info, profile};
for (const auto& inst : inst_list) {
pc += inst.length;

8
src/shader_recompiler/frontend/translate/translate.h
View file

@ -5,9 +5,9 @@
#include <span>
#include "shader_recompiler/frontend/instruction.h"
#include "shader_recompiler/info.h"
#include "shader_recompiler/ir/basic_block.h"
#include "shader_recompiler/ir/ir_emitter.h"
#include "shader_recompiler/runtime_info.h"
namespace Shader {
struct Info;
@ -55,7 +55,8 @@ enum class NegateMode : u32 {
class Translator {
public:
explicit Translator(IR::Block* block_, Info& info, const Profile& profile);
explicit Translator(IR::Block* block_, Info& info, const RuntimeInfo& runtime_info,
const Profile& profile);
// Instruction categories
void EmitPrologue();
@ -237,12 +238,13 @@ private:
private:
IR::IREmitter ir;
Info& info;
const RuntimeInfo& runtime_info;
const Profile& profile;
IR::U32 m0_value;
bool opcode_missing = false;
};
void Translate(IR::Block* block, u32 block_base, std::span<const GcnInst> inst_list, Info& info,
const Profile& profile);
const RuntimeInfo& runtime_info, const Profile& profile);
} // namespace Shader::Gcn

9
src/shader_recompiler/frontend/translate/vector_alu.cpp
View file

@ -479,10 +479,11 @@ void Translator::V_ADD_F32(const GcnInst& inst) {
void Translator::V_CVT_OFF_F32_I4(const GcnInst& inst) {
const IR::U32 src0{GetSrc(inst.src[0])};
const IR::VectorReg dst_reg{inst.dst[0].code};
ir.SetVectorReg(
dst_reg,
ir.FPMul(ir.ConvertUToF(32, 32, ir.ISub(ir.BitwiseAnd(src0, ir.Imm32(0xF)), ir.Imm32(8))),
ir.Imm32(1.f / 16.f)));
ASSERT(src0.IsImmediate());
static constexpr std::array IntToFloat = {
0.0f, 0.0625f, 0.1250f, 0.1875f, 0.2500f, 0.3125f, 0.3750f, 0.4375f,
-0.5000f, -0.4375f, -0.3750f, -0.3125f, -0.2500f, -0.1875f, -0.1250f, -0.0625f};
ir.SetVectorReg(dst_reg, ir.Imm32(IntToFloat[src0.U32()]));
}
void Translator::V_MED3_F32(const GcnInst& inst) {

4
src/shader_recompiler/frontend/translate/vector_interpolation.cpp
View file

@ -7,14 +7,14 @@ namespace Shader::Gcn {
void Translator::V_INTERP_P2_F32(const GcnInst& inst) {
const IR::VectorReg dst_reg{inst.dst[0].code};
auto& attr = info.ps_inputs.at(inst.control.vintrp.attr);
auto& attr = runtime_info.fs_info.inputs.at(inst.control.vintrp.attr);
const IR::Attribute attrib{IR::Attribute::Param0 + attr.param_index};
ir.SetVectorReg(dst_reg, ir.GetAttribute(attrib, inst.control.vintrp.chan));
}
void Translator::V_INTERP_MOV_F32(const GcnInst& inst) {
const IR::VectorReg dst_reg{inst.dst[0].code};
auto& attr = info.ps_inputs.at(inst.control.vintrp.attr);
auto& attr = runtime_info.fs_info.inputs.at(inst.control.vintrp.attr);
const IR::Attribute attrib{IR::Attribute::Param0 + attr.param_index};
ir.SetVectorReg(dst_reg, ir.GetAttribute(attrib, inst.control.vintrp.chan));
}

232
src/shader_recompiler/info.h Normal file
View file

@ -0,0 +1,232 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#pragma once
#include <span>
#include <boost/container/small_vector.hpp>
#include <boost/container/static_vector.hpp>
#include "common/assert.h"
#include "common/types.h"
#include "shader_recompiler/ir/attribute.h"
#include "shader_recompiler/ir/reg.h"
#include "shader_recompiler/ir/type.h"
#include "shader_recompiler/params.h"
#include "shader_recompiler/runtime_info.h"
#include "video_core/amdgpu/resource.h"
namespace Shader {
static constexpr size_t NumUserDataRegs = 16;
enum class TextureType : u32 {
Color1D,
ColorArray1D,
Color2D,
ColorArray2D,
Color3D,
ColorCube,
Buffer,
};
constexpr u32 NUM_TEXTURE_TYPES = 7;
struct Info;
struct BufferResource {
u32 sgpr_base;
u32 dword_offset;
IR::Type used_types;
AmdGpu::Buffer inline_cbuf;
bool is_instance_data{};
bool is_written{};
bool IsStorage(AmdGpu::Buffer buffer) const noexcept {
static constexpr size_t MaxUboSize = 65536;
return buffer.GetSize() > MaxUboSize || is_written;
}
constexpr AmdGpu::Buffer GetSharp(const Info& info) const noexcept;
};
using BufferResourceList = boost::container::small_vector<BufferResource, 16>;
struct TextureBufferResource {
u32 sgpr_base;
u32 dword_offset;
AmdGpu::NumberFormat nfmt;
bool is_written{};
constexpr AmdGpu::Buffer GetSharp(const Info& info) const noexcept;
};
using TextureBufferResourceList = boost::container::small_vector<TextureBufferResource, 16>;
struct ImageResource {
u32 sgpr_base;
u32 dword_offset;
AmdGpu::ImageType type;
AmdGpu::NumberFormat nfmt;
bool is_storage;
bool is_depth;
bool is_atomic{};
constexpr AmdGpu::Image GetSharp(const Info& info) const noexcept;
};
using ImageResourceList = boost::container::small_vector<ImageResource, 16>;
struct SamplerResource {
u32 sgpr_base;
u32 dword_offset;
AmdGpu::Sampler inline_sampler{};
u32 associated_image : 4;
u32 disable_aniso : 1;
constexpr AmdGpu::Sampler GetSharp(const Info& info) const noexcept;
};
using SamplerResourceList = boost::container::small_vector<SamplerResource, 16>;
struct PushData {
static constexpr size_t BufOffsetIndex = 2;
u32 step0;
u32 step1;
std::array<u8, 32> buf_offsets;
void AddOffset(u32 binding, u32 offset) {
ASSERT(offset < 256 && binding < buf_offsets.size());
buf_offsets[binding] = offset;
}
};
/**
* Contains general information generated by the shader recompiler for an input program.
*/
struct Info {
struct VsInput {
enum InstanceIdType : u8 {
None = 0,
OverStepRate0 = 1,
OverStepRate1 = 2,
Plain = 3,
};
AmdGpu::NumberFormat fmt;
u16 binding;
u16 num_components;
u8 sgpr_base;
u8 dword_offset;
InstanceIdType instance_step_rate;
s32 instance_data_buf;
};
boost::container::static_vector<VsInput, 32> vs_inputs{};
struct AttributeFlags {
bool Get(IR::Attribute attrib, u32 comp = 0) const {
return flags[Index(attrib)] & (1 << comp);
}
bool GetAny(IR::Attribute attrib) const {
return flags[Index(attrib)];
}
void Set(IR::Attribute attrib, u32 comp = 0) {
flags[Index(attrib)] |= (1 << comp);
}
u32 NumComponents(IR::Attribute attrib) const {
return 4;
}
static size_t Index(IR::Attribute attrib) {
return static_cast<size_t>(attrib);
}
std::array<u8, IR::NumAttributes> flags;
};
AttributeFlags loads{};
AttributeFlags stores{};
s8 vertex_offset_sgpr = -1;
s8 instance_offset_sgpr = -1;
BufferResourceList buffers;
TextureBufferResourceList texture_buffers;
ImageResourceList images;
SamplerResourceList samplers;
std::span<const u32> user_data;
Stage stage;
u64 pgm_hash{};
VAddr pgm_base;
bool has_storage_images{};
bool has_image_buffers{};
bool has_texel_buffers{};
bool has_discard{};
bool has_image_gather{};
bool has_image_query{};
bool uses_lane_id{};
bool uses_group_quad{};
bool uses_shared{};
bool uses_fp16{};
bool uses_step_rates{};
bool translation_failed{}; // indicates that shader has unsupported instructions
explicit Info(Stage stage_, ShaderParams params)
: stage{stage_}, pgm_hash{params.hash}, pgm_base{params.Base()},
user_data{params.user_data} {}
template <typename T>
T ReadUd(u32 ptr_index, u32 dword_offset) const noexcept {
T data;
const u32* base = user_data.data();
if (ptr_index != IR::NumScalarRegs) {
std::memcpy(&base, &user_data[ptr_index], sizeof(base));
}
std::memcpy(&data, base + dword_offset, sizeof(T));
return data;
}
size_t NumBindings() const noexcept {
return buffers.size() + texture_buffers.size() + images.size() + samplers.size();
}
[[nodiscard]] std::pair<u32, u32> GetDrawOffsets() const noexcept {
u32 vertex_offset = 0;
u32 instance_offset = 0;
if (vertex_offset_sgpr != -1) {
vertex_offset = user_data[vertex_offset_sgpr];
}
if (instance_offset_sgpr != -1) {
instance_offset = user_data[instance_offset_sgpr];
}
return {vertex_offset, instance_offset};
}
};
constexpr AmdGpu::Buffer BufferResource::GetSharp(const Info& info) const noexcept {
return inline_cbuf ? inline_cbuf : info.ReadUd<AmdGpu::Buffer>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Buffer TextureBufferResource::GetSharp(const Info& info) const noexcept {
return info.ReadUd<AmdGpu::Buffer>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Image ImageResource::GetSharp(const Info& info) const noexcept {
return info.ReadUd<AmdGpu::Image>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Sampler SamplerResource::GetSharp(const Info& info) const noexcept {
return inline_sampler ? inline_sampler : info.ReadUd<AmdGpu::Sampler>(sgpr_base, dword_offset);
}
} // namespace Shader
template <>
struct fmt::formatter<Shader::Stage> {
constexpr auto parse(format_parse_context& ctx) {
return ctx.begin();
}
auto format(const Shader::Stage stage, format_context& ctx) const {
constexpr static std::array names = {"fs", "vs", "gs", "es", "hs", "ls", "cs"};
return fmt::format_to(ctx.out(), "{}", names[static_cast<size_t>(stage)]);
}
};

9
src/shader_recompiler/ir/passes/resource_tracking_pass.cpp
View file

@ -4,11 +4,11 @@
#include <algorithm>
#include <boost/container/small_vector.hpp>
#include "common/alignment.h"
#include "shader_recompiler/info.h"
#include "shader_recompiler/ir/basic_block.h"
#include "shader_recompiler/ir/breadth_first_search.h"
#include "shader_recompiler/ir/ir_emitter.h"
#include "shader_recompiler/ir/program.h"
#include "shader_recompiler/runtime_info.h"
#include "video_core/amdgpu/resource.h"
namespace Shader::Optimization {
@ -471,14 +471,11 @@ void PatchImageInstruction(IR::Block& block, IR::Inst& inst, Info& info, Descrip
// Read image sharp.
const auto tsharp = TrackSharp(tsharp_handle);
const auto image = info.ReadUd<AmdGpu::Image>(tsharp.sgpr_base, tsharp.dword_offset);
const auto inst_info = inst.Flags<IR::TextureInstInfo>();
auto image = info.ReadUd<AmdGpu::Image>(tsharp.sgpr_base, tsharp.dword_offset);
if (!image.Valid()) {
LOG_ERROR(Render_Vulkan, "Shader compiled with unbound image!");
IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(inst)};
inst.ReplaceUsesWith(
ir.CompositeConstruct(ir.Imm32(0.f), ir.Imm32(0.f), ir.Imm32(0.f), ir.Imm32(0.f)));
return;
image = AmdGpu::Image::Null();
}
ASSERT(image.GetType() != AmdGpu::ImageType::Invalid);
const bool is_storage = IsImageStorageInstruction(inst);

2
src/shader_recompiler/ir/program.h
View file

@ -5,9 +5,9 @@
#include <string>
#include "shader_recompiler/frontend/instruction.h"
#include "shader_recompiler/info.h"
#include "shader_recompiler/ir/abstract_syntax_list.h"
#include "shader_recompiler/ir/basic_block.h"
#include "shader_recompiler/runtime_info.h"
namespace Shader::IR {

26
src/shader_recompiler/params.h Normal file
View file

@ -0,0 +1,26 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#pragma once
#include <span>
#include "common/types.h"
namespace Shader {
/**
* Compilation parameters used to identify and locate a guest shader program.
*/
struct ShaderParams {
static constexpr u32 NumShaderUserData = 16;
std::span<const u32, NumShaderUserData> user_data;
std::span<const u32> code;
u64 hash;
VAddr Base() const noexcept {
return reinterpret_cast<VAddr>(code.data());
}
};
} // namespace Shader

19
src/shader_recompiler/recompiler.cpp
View file

@ -6,6 +6,7 @@
#include "shader_recompiler/frontend/structured_control_flow.h"
#include "shader_recompiler/ir/passes/ir_passes.h"
#include "shader_recompiler/ir/post_order.h"
#include "shader_recompiler/recompiler.h"
namespace Shader {
@ -27,29 +28,32 @@ IR::BlockList GenerateBlocks(const IR::AbstractSyntaxList& syntax_list) {
return blocks;
}
IR::Program TranslateProgram(Common::ObjectPool<IR::Inst>& inst_pool,
Common::ObjectPool<IR::Block>& block_pool, std::span<const u32> token,
Info& info, const Profile& profile) {
IR::Program TranslateProgram(std::span<const u32> code, Pools& pools, Info& info,
const RuntimeInfo& runtime_info, const Profile& profile) {
// Ensure first instruction is expected.
constexpr u32 token_mov_vcchi = 0xBEEB03FF;
ASSERT_MSG(token[0] == token_mov_vcchi, "First instruction is not s_mov_b32 vcc_hi, #imm");
ASSERT_MSG(code[0] == token_mov_vcchi, "First instruction is not s_mov_b32 vcc_hi, #imm");
Gcn::GcnCodeSlice slice(token.data(), token.data() + token.size());
Gcn::GcnCodeSlice slice(code.data(), code.data() + code.size());
Gcn::GcnDecodeContext decoder;
// Decode and save instructions
IR::Program program{info};
program.ins_list.reserve(token.size());
program.ins_list.reserve(code.size());
while (!slice.atEnd()) {
program.ins_list.emplace_back(decoder.decodeInstruction(slice));
}
// Clear any previous pooled data.
pools.ReleaseContents();
// Create control flow graph
Common::ObjectPool<Gcn::Block> gcn_block_pool{64};
Gcn::CFG cfg{gcn_block_pool, program.ins_list};
// Structurize control flow graph and create program.
program.syntax_list = Shader::Gcn::BuildASL(inst_pool, block_pool, cfg, program.info, profile);
program.syntax_list = Shader::Gcn::BuildASL(pools.inst_pool, pools.block_pool, cfg,
program.info, runtime_info, profile);
program.blocks = GenerateBlocks(program.syntax_list);
program.post_order_blocks = Shader::IR::PostOrder(program.syntax_list.front());
@ -63,7 +67,6 @@ IR::Program TranslateProgram(Common::ObjectPool<IR::Inst>& inst_pool,
Shader::Optimization::IdentityRemovalPass(program.blocks);
Shader::Optimization::DeadCodeEliminationPass(program);
Shader::Optimization::CollectShaderInfoPass(program);
LOG_DEBUG(Render_Vulkan, "{}", Shader::IR::DumpProgram(program));
return program;
}

22
src/shader_recompiler/recompiler.h
View file

@ -10,10 +10,24 @@
namespace Shader {
struct Profile;
struct RuntimeInfo;
[[nodiscard]] IR::Program TranslateProgram(Common::ObjectPool<IR::Inst>& inst_pool,
Common::ObjectPool<IR::Block>& block_pool,
std::span<const u32> code, Info& info,
const Profile& profile);
struct Pools {
static constexpr u32 InstPoolSize = 8192;
static constexpr u32 BlockPoolSize = 32;
Common::ObjectPool<IR::Inst> inst_pool;
Common::ObjectPool<IR::Block> block_pool;
explicit Pools() : inst_pool{InstPoolSize}, block_pool{BlockPoolSize} {}
void ReleaseContents() {
inst_pool.ReleaseContents();
block_pool.ReleaseContents();
}
};
[[nodiscard]] IR::Program TranslateProgram(std::span<const u32> code, Pools& pools, Info& info,
const RuntimeInfo& runtime_info, const Profile& profile);
} // namespace Shader

261
src/shader_recompiler/runtime_info.h
View file

@ -3,20 +3,14 @@
#pragma once
#include <span>
#include <boost/container/small_vector.hpp>
#include <algorithm>
#include <boost/container/static_vector.hpp>
#include "common/assert.h"
#include "common/types.h"
#include "shader_recompiler/ir/attribute.h"
#include "shader_recompiler/ir/reg.h"
#include "shader_recompiler/ir/type.h"
#include "video_core/amdgpu/resource.h"
namespace Shader {
static constexpr size_t NumUserDataRegs = 16;
enum class Stage : u32 {
Fragment,
Vertex,
@ -29,21 +23,18 @@ enum class Stage : u32 {
constexpr u32 MaxStageTypes = 6;
[[nodiscard]] constexpr Stage StageFromIndex(size_t index) noexcept {
return static_cast<Stage>(static_cast<size_t>(Stage::Vertex) + index);
return static_cast<Stage>(index);
}
enum class TextureType : u32 {
Color1D,
ColorArray1D,
Color2D,
ColorArray2D,
Color3D,
ColorCube,
Buffer,
enum class MrtSwizzle : u8 {
Identity = 0,
Alt = 1,
Reverse = 2,
ReverseAlt = 3,
};
constexpr u32 NUM_TEXTURE_TYPES = 7;
static constexpr u32 MaxColorBuffers = 8;
enum class VsOutput : u32 {
enum class VsOutput : u8 {
None,
PointSprite,
EdgeFlag,
@ -70,211 +61,69 @@ enum class VsOutput : u32 {
};
using VsOutputMap = std::array<VsOutput, 4>;
struct Info;
struct VertexRuntimeInfo {
boost::container::static_vector<VsOutputMap, 3> outputs;
struct BufferResource {
u32 sgpr_base;
u32 dword_offset;
IR::Type used_types;
AmdGpu::Buffer inline_cbuf;
bool is_instance_data{};
bool is_written{};
bool IsStorage(AmdGpu::Buffer buffer) const noexcept {
static constexpr size_t MaxUboSize = 65536;
return buffer.GetSize() > MaxUboSize || is_written;
}
constexpr AmdGpu::Buffer GetSharp(const Info& info) const noexcept;
};
using BufferResourceList = boost::container::small_vector<BufferResource, 16>;
struct TextureBufferResource {
u32 sgpr_base;
u32 dword_offset;
AmdGpu::NumberFormat nfmt;
bool is_written{};
constexpr AmdGpu::Buffer GetSharp(const Info& info) const noexcept;
};
using TextureBufferResourceList = boost::container::small_vector<TextureBufferResource, 16>;
struct ImageResource {
u32 sgpr_base;
u32 dword_offset;
AmdGpu::ImageType type;
AmdGpu::NumberFormat nfmt;
bool is_storage;
bool is_depth;
bool is_atomic{};
constexpr AmdGpu::Image GetSharp(const Info& info) const noexcept;
};
using ImageResourceList = boost::container::small_vector<ImageResource, 16>;
struct SamplerResource {
u32 sgpr_base;
u32 dword_offset;
AmdGpu::Sampler inline_sampler{};
u32 associated_image : 4;
u32 disable_aniso : 1;
constexpr AmdGpu::Sampler GetSharp(const Info& info) const noexcept;
};
using SamplerResourceList = boost::container::small_vector<SamplerResource, 16>;
struct PushData {
static constexpr size_t BufOffsetIndex = 2;
u32 step0;
u32 step1;
std::array<u8, 32> buf_offsets;
void AddOffset(u32 binding, u32 offset) {
ASSERT(offset < 256 && binding < buf_offsets.size());
buf_offsets[binding] = offset;
bool operator==(const VertexRuntimeInfo& other) const noexcept {
return true;
}
};
struct Info {
struct VsInput {
enum InstanceIdType : u8 {
None = 0,
OverStepRate0 = 1,
OverStepRate1 = 2,
Plain = 3,
};
AmdGpu::NumberFormat fmt;
u16 binding;
u16 num_components;
u8 sgpr_base;
u8 dword_offset;
InstanceIdType instance_step_rate;
s32 instance_data_buf;
};
boost::container::static_vector<VsInput, 32> vs_inputs{};
struct FragmentRuntimeInfo {
struct PsInput {
u32 param_index;
u8 param_index;
bool is_default;
bool is_flat;
u32 default_value;
u8 default_value;
auto operator<=>(const PsInput&) const noexcept = default;
};
boost::container::static_vector<PsInput, 32> ps_inputs{};
boost::container::static_vector<PsInput, 32> inputs;
std::array<MrtSwizzle, MaxColorBuffers> mrt_swizzles;
struct AttributeFlags {
bool Get(IR::Attribute attrib, u32 comp = 0) const {
return flags[Index(attrib)] & (1 << comp);
}
bool operator==(const FragmentRuntimeInfo& other) const noexcept {
return std::ranges::equal(mrt_swizzles, other.mrt_swizzles) &&
std::ranges::equal(inputs, other.inputs);
}
};
bool GetAny(IR::Attribute attrib) const {
return flags[Index(attrib)];
}
void Set(IR::Attribute attrib, u32 comp = 0) {
flags[Index(attrib)] |= (1 << comp);
}
u32 NumComponents(IR::Attribute attrib) const {
return 4;
}
static size_t Index(IR::Attribute attrib) {
return static_cast<size_t>(attrib);
}
std::array<u8, IR::NumAttributes> flags;
};
AttributeFlags loads{};
AttributeFlags stores{};
boost::container::static_vector<VsOutputMap, 3> vs_outputs;
s8 vertex_offset_sgpr = -1;
s8 instance_offset_sgpr = -1;
BufferResourceList buffers;
TextureBufferResourceList texture_buffers;
ImageResourceList images;
SamplerResourceList samplers;
std::array<u32, 3> workgroup_size{};
struct ComputeRuntimeInfo {
u32 shared_memory_size;
std::array<u32, 3> workgroup_size;
std::array<bool, 3> tgid_enable;
bool operator==(const ComputeRuntimeInfo& other) const noexcept {
return workgroup_size == other.workgroup_size && tgid_enable == other.tgid_enable;
}
};
/**
* Stores information relevant to shader compilation sourced from liverpool registers.
* It may potentially differ with the same shader module so must be checked.
* It's also possible to store any other custom information that needs to be part of shader key.
*/
struct RuntimeInfo {
Stage stage;
u32 num_user_data;
u32 num_input_vgprs;
std::span<const u32> user_data;
Stage stage;
VertexRuntimeInfo vs_info;
FragmentRuntimeInfo fs_info;
ComputeRuntimeInfo cs_info;
uintptr_t pgm_base{};
u64 pgm_hash{};
u32 shared_memory_size{};
bool has_storage_images{};
bool has_image_buffers{};
bool has_texel_buffers{};
bool has_discard{};
bool has_image_gather{};
bool has_image_query{};
bool uses_lane_id{};
bool uses_group_quad{};
bool uses_shared{};
bool uses_fp16{};
bool uses_step_rates{};
bool translation_failed{}; // indicates that shader has unsupported instructions
RuntimeInfo(Stage stage_) : stage{stage_} {}
template <typename T>
T ReadUd(u32 ptr_index, u32 dword_offset) const noexcept {
T data;
const u32* base = user_data.data();
if (ptr_index != IR::NumScalarRegs) {
std::memcpy(&base, &user_data[ptr_index], sizeof(base));
bool operator==(const RuntimeInfo& other) const noexcept {
switch (stage) {
case Stage::Fragment:
return fs_info == other.fs_info;
case Stage::Vertex:
return vs_info == other.vs_info;
case Stage::Compute:
return cs_info == other.cs_info;
default:
return true;
}
std::memcpy(&data, base + dword_offset, sizeof(T));
return data;
}
size_t NumBindings() const noexcept {
return buffers.size() + texture_buffers.size() + images.size() + samplers.size();
}
[[nodiscard]] std::pair<u32, u32> GetDrawOffsets() const noexcept {
u32 vertex_offset = 0;
u32 instance_offset = 0;
if (vertex_offset_sgpr != -1) {
vertex_offset = user_data[vertex_offset_sgpr];
}
if (instance_offset_sgpr != -1) {
instance_offset = user_data[instance_offset_sgpr];
}
return {vertex_offset, instance_offset};
}
};
constexpr AmdGpu::Buffer BufferResource::GetSharp(const Info& info) const noexcept {
return inline_cbuf ? inline_cbuf : info.ReadUd<AmdGpu::Buffer>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Buffer TextureBufferResource::GetSharp(const Info& info) const noexcept {
return info.ReadUd<AmdGpu::Buffer>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Image ImageResource::GetSharp(const Info& info) const noexcept {
return info.ReadUd<AmdGpu::Image>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Sampler SamplerResource::GetSharp(const Info& info) const noexcept {
return inline_sampler ? inline_sampler : info.ReadUd<AmdGpu::Sampler>(sgpr_base, dword_offset);
}
} // namespace Shader
template <>
struct fmt::formatter<Shader::Stage> {
constexpr auto parse(format_parse_context& ctx) {
return ctx.begin();
}
auto format(const Shader::Stage stage, format_context& ctx) const {
constexpr static std::array names = {"fs", "vs", "gs", "es", "hs", "ls", "cs"};
return fmt::format_to(ctx.out(), "{}", names[static_cast<size_t>(stage)]);
}
};

112
src/video_core/renderer_vulkan/vk_shader_cache.h → src/shader_recompiler/specialization.h
View file

@ -4,18 +4,11 @@
#pragma once
#include <bitset>
#include <boost/container/small_vector.hpp>
#include <tsl/robin_map.h>
#include "common/object_pool.h"
#include "shader_recompiler/ir/basic_block.h"
#include "shader_recompiler/profile.h"
#include "shader_recompiler/runtime_info.h"
#include "video_core/amdgpu/liverpool.h"
#include "video_core/renderer_vulkan/vk_common.h"
namespace Vulkan {
#include "common/types.h"
#include "shader_recompiler/info.h"
class Instance;
namespace Shader {
struct BufferSpecialization {
u16 stride : 14;
@ -25,43 +18,38 @@ struct BufferSpecialization {
};
struct TextureBufferSpecialization {
bool is_integer;
bool is_integer = false;
auto operator<=>(const TextureBufferSpecialization&) const = default;
};
struct ImageSpecialization {
AmdGpu::ImageType type;
bool is_integer;
AmdGpu::ImageType type = AmdGpu::ImageType::Color2D;
bool is_integer = false;
auto operator<=>(const ImageSpecialization&) const = default;
};
/**
* Alongside runtime information, this structure also checks bound resources
* for compatibility. Can be used as a key for storing shader permutations.
* Is separate from runtime information, because resource layout can only be deduced
* after the first compilation of a module.
*/
struct StageSpecialization {
static constexpr size_t MaxStageResources = 32;
const Shader::Info* info;
RuntimeInfo runtime_info;
std::bitset<MaxStageResources> bitset{};
boost::container::small_vector<BufferSpecialization, 16> buffers;
boost::container::small_vector<TextureBufferSpecialization, 8> tex_buffers;
boost::container::small_vector<ImageSpecialization, 8> images;
u32 start_binding{};
void ForEachSharp(u32& binding, auto& spec_list, auto& desc_list, auto&& func) {
for (const auto& desc : desc_list) {
auto& spec = spec_list.emplace_back();
const auto sharp = desc.GetSharp(*info);
if (!sharp) {
binding++;
continue;
}
bitset.set(binding++);
func(spec, desc, sharp);
}
}
StageSpecialization(const Shader::Info& info_, u32 start_binding_)
: info{&info_}, start_binding{start_binding_} {
explicit StageSpecialization(const Shader::Info& info_, RuntimeInfo runtime_info_,
u32 start_binding_)
: info{&info_}, runtime_info{runtime_info_}, start_binding{start_binding_} {
u32 binding{};
ForEachSharp(binding, buffers, info->buffers,
[](auto& spec, const auto& desc, AmdGpu::Buffer sharp) {
@ -79,10 +67,26 @@ struct StageSpecialization {
});
}
void ForEachSharp(u32& binding, auto& spec_list, auto& desc_list, auto&& func) {
for (const auto& desc : desc_list) {
auto& spec = spec_list.emplace_back();
const auto sharp = desc.GetSharp(*info);
if (!sharp) {
binding++;
continue;
}
bitset.set(binding++);
func(spec, desc, sharp);
}
}
bool operator==(const StageSpecialization& other) const {
if (start_binding != other.start_binding) {
return false;
}
if (runtime_info != other.runtime_info) {
return false;
}
u32 binding{};
for (u32 i = 0; i < buffers.size(); i++) {
if (other.bitset[binding++] && buffers[i] != other.buffers[i]) {
@ -103,54 +107,4 @@ struct StageSpecialization {
}
};
struct Program {
struct Module {
vk::ShaderModule module;
StageSpecialization spec;
};
Shader::Info info;
boost::container::small_vector<Module, 8> modules;
explicit Program(const Shader::Info& info_) : info{info_} {}
};
struct GuestProgram {
Shader::Stage stage;
std::span<const u32, AmdGpu::Liverpool::NumShaderUserData> user_data;
std::span<const u32> code;
u64 hash;
explicit GuestProgram(const auto* pgm, Shader::Stage stage_)
: stage{stage_}, user_data{pgm->user_data}, code{pgm->Code()} {
const auto* bininfo = AmdGpu::Liverpool::GetBinaryInfo(*pgm);
hash = bininfo->shader_hash;
}
};
class ShaderCache {
public:
explicit ShaderCache(const Instance& instance, AmdGpu::Liverpool* liverpool);
~ShaderCache() = default;
std::tuple<const Shader::Info*, vk::ShaderModule, u64> GetProgram(const GuestProgram& pgm,
u32& binding);
private:
void DumpShader(std::span<const u32> code, u64 hash, Shader::Stage stage, size_t perm_idx,
std::string_view ext);
vk::ShaderModule CompileModule(Shader::Info& info, std::span<const u32> code, size_t perm_idx,
u32& binding);
Program* CreateProgram(const GuestProgram& pgm, u32& binding);
private:
const Instance& instance;
AmdGpu::Liverpool* liverpool;
Shader::Profile profile{};
tsl::robin_map<size_t, Program*> program_cache;
Common::ObjectPool<Shader::IR::Inst> inst_pool;
Common::ObjectPool<Shader::IR::Block> block_pool;
Common::ObjectPool<Program> program_pool;
};
} // namespace Vulkan
} // namespace Shader

10
src/video_core/amdgpu/liverpool.h
View file

@ -18,6 +18,7 @@
#include "common/polyfill_thread.h"
#include "common/types.h"
#include "common/unique_function.h"
#include "shader_recompiler/params.h"
#include "video_core/amdgpu/pixel_format.h"
#include "video_core/amdgpu/resource.h"
@ -171,6 +172,15 @@ struct Liverpool {
return bininfo;
}
static constexpr Shader::ShaderParams GetParams(const auto& sh) {
auto* bininfo = GetBinaryInfo(sh);
return {
.user_data = sh.user_data,
.code = sh.Code(),
.hash = bininfo->shader_hash,
};
}
union PsInputControl {
u32 raw;
BitField<0, 5, u32> input_offset;

12
src/video_core/amdgpu/resource.h
View file

@ -176,6 +176,18 @@ struct Image {
u64 lod_hw_cnt_en : 1;
u64 : 43;
static constexpr Image Null() {
Image image{};
image.data_format = u64(DataFormat::Format8_8_8_8);
image.dst_sel_x = 4;
image.dst_sel_y = 5;
image.dst_sel_z = 6;
image.dst_sel_w = 7;
image.tiling_index = u64(TilingMode::Texture_MicroTiled);
image.type = u64(ImageType::Color2D);
return image;
}
bool Valid() const {
return (type & 0x8u) != 0;
}

2
src/video_core/buffer_cache/buffer_cache.cpp
View file

@ -4,7 +4,7 @@
#include <algorithm>
#include "common/alignment.h"
#include "common/scope_exit.h"
#include "shader_recompiler/runtime_info.h"
#include "shader_recompiler/info.h"
#include "video_core/amdgpu/liverpool.h"
#include "video_core/buffer_cache/buffer_cache.h"
#include "video_core/renderer_vulkan/liverpool_to_vk.h"

2
src/video_core/renderer_vulkan/liverpool_to_vk.cpp
View file

@ -600,6 +600,8 @@ vk::Format AdjustColorBufferFormat(vk::Format base_format,
return is_vo_surface ? vk::Format::eB8G8R8A8Unorm : vk::Format::eB8G8R8A8Srgb;
case vk::Format::eB8G8R8A8Srgb:
return is_vo_surface ? vk::Format::eR8G8B8A8Unorm : vk::Format::eR8G8B8A8Srgb;
case vk::Format::eA2B10G10R10UnormPack32:
return vk::Format::eA2R10G10B10UnormPack32;
default:
break;
}

2
src/video_core/renderer_vulkan/vk_compute_pipeline.h
View file

@ -4,7 +4,7 @@
#pragma once
#include <boost/container/small_vector.hpp>
#include "shader_recompiler/runtime_info.h"
#include "shader_recompiler/info.h"
#include "video_core/renderer_vulkan/vk_common.h"
namespace VideoCore {

1
src/video_core/renderer_vulkan/vk_graphics_pipeline.h
View file

@ -25,6 +25,7 @@ using Liverpool = AmdGpu::Liverpool;
struct GraphicsPipelineKey {
std::array<size_t, MaxShaderStages> stage_hashes;
std::array<vk::Format, Liverpool::NumColorBuffers> color_formats;
std::array<Liverpool::ColorBuffer::SwapMode, Liverpool::NumColorBuffers> mrt_swizzles;
vk::Format depth_format;
vk::Format stencil_format;

199
src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
View file

@ -1,21 +1,124 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "shader_recompiler/runtime_info.h"
#include <ranges>
#include "common/config.h"
#include "common/io_file.h"
#include "common/path_util.h"
#include "shader_recompiler/backend/spirv/emit_spirv.h"
#include "shader_recompiler/info.h"
#include "video_core/renderer_vulkan/renderer_vulkan.h"
#include "video_core/renderer_vulkan/vk_instance.h"
#include "video_core/renderer_vulkan/vk_pipeline_cache.h"
#include "video_core/renderer_vulkan/vk_scheduler.h"
#include "video_core/renderer_vulkan/vk_shader_cache.h"
#include "video_core/renderer_vulkan/vk_shader_util.h"
extern std::unique_ptr<Vulkan::RendererVulkan> renderer;
namespace Vulkan {
using Shader::VsOutput;
[[nodiscard]] inline u64 HashCombine(const u64 seed, const u64 hash) {
return seed ^ (hash + 0x9e3779b9 + (seed << 6) + (seed >> 2));
}
void GatherVertexOutputs(Shader::VertexRuntimeInfo& info,
const AmdGpu::Liverpool::VsOutputControl& ctl) {
const auto add_output = [&](VsOutput x, VsOutput y, VsOutput z, VsOutput w) {
if (x != VsOutput::None || y != VsOutput::None || z != VsOutput::None ||
w != VsOutput::None) {
info.outputs.emplace_back(Shader::VsOutputMap{x, y, z, w});
}
};
// VS_OUT_MISC_VEC
add_output(ctl.use_vtx_point_size ? VsOutput::PointSprite : VsOutput::None,
ctl.use_vtx_edge_flag
? VsOutput::EdgeFlag
: (ctl.use_vtx_gs_cut_flag ? VsOutput::GsCutFlag : VsOutput::None),
ctl.use_vtx_kill_flag
? VsOutput::KillFlag
: (ctl.use_vtx_render_target_idx ? VsOutput::GsMrtIndex : VsOutput::None),
ctl.use_vtx_viewport_idx ? VsOutput::GsVpIndex : VsOutput::None);
// VS_OUT_CCDIST0
add_output(ctl.IsClipDistEnabled(0)
? VsOutput::ClipDist0
: (ctl.IsCullDistEnabled(0) ? VsOutput::CullDist0 : VsOutput::None),
ctl.IsClipDistEnabled(1)
? VsOutput::ClipDist1
: (ctl.IsCullDistEnabled(1) ? VsOutput::CullDist1 : VsOutput::None),
ctl.IsClipDistEnabled(2)
? VsOutput::ClipDist2
: (ctl.IsCullDistEnabled(2) ? VsOutput::CullDist2 : VsOutput::None),
ctl.IsClipDistEnabled(3)
? VsOutput::ClipDist3
: (ctl.IsCullDistEnabled(3) ? VsOutput::CullDist3 : VsOutput::None));
// VS_OUT_CCDIST1
add_output(ctl.IsClipDistEnabled(4)
? VsOutput::ClipDist4
: (ctl.IsCullDistEnabled(4) ? VsOutput::CullDist4 : VsOutput::None),
ctl.IsClipDistEnabled(5)
? VsOutput::ClipDist5
: (ctl.IsCullDistEnabled(5) ? VsOutput::CullDist5 : VsOutput::None),
ctl.IsClipDistEnabled(6)
? VsOutput::ClipDist6
: (ctl.IsCullDistEnabled(6) ? VsOutput::CullDist6 : VsOutput::None),
ctl.IsClipDistEnabled(7)
? VsOutput::ClipDist7
: (ctl.IsCullDistEnabled(7) ? VsOutput::CullDist7 : VsOutput::None));
}
Shader::RuntimeInfo BuildRuntimeInfo(Shader::Stage stage, const GraphicsPipelineKey& key,
const AmdGpu::Liverpool::Regs& regs) {
auto info = Shader::RuntimeInfo{stage};
switch (stage) {
case Shader::Stage::Vertex: {
info.num_user_data = regs.vs_program.settings.num_user_regs;
info.num_input_vgprs = regs.vs_program.settings.vgpr_comp_cnt;
GatherVertexOutputs(info.vs_info, regs.vs_output_control);
break;
}
case Shader::Stage::Fragment: {
info.num_user_data = regs.ps_program.settings.num_user_regs;
std::ranges::transform(key.mrt_swizzles, info.fs_info.mrt_swizzles.begin(),
[](Liverpool::ColorBuffer::SwapMode mode) {
return static_cast<Shader::MrtSwizzle>(mode);
});
for (u32 i = 0; i < regs.num_interp; i++) {
info.fs_info.inputs.push_back({
.param_index = u8(regs.ps_inputs[i].input_offset.Value()),
.is_default = bool(regs.ps_inputs[i].use_default),
.is_flat = bool(regs.ps_inputs[i].flat_shade),
.default_value = u8(regs.ps_inputs[i].default_value),
});
}
break;
}
case Shader::Stage::Compute: {
const auto& cs_pgm = regs.cs_program;
info.num_user_data = cs_pgm.settings.num_user_regs;
info.cs_info.workgroup_size = {cs_pgm.num_thread_x.full, cs_pgm.num_thread_y.full,
cs_pgm.num_thread_z.full};
info.cs_info.tgid_enable = {cs_pgm.IsTgidEnabled(0), cs_pgm.IsTgidEnabled(1),
cs_pgm.IsTgidEnabled(2)};
info.cs_info.shared_memory_size = cs_pgm.SharedMemSize();
break;
}
default:
break;
}
return info;
}
PipelineCache::PipelineCache(const Instance& instance_, Scheduler& scheduler_,
AmdGpu::Liverpool* liverpool_)
: instance{instance_}, scheduler{scheduler_}, liverpool{liverpool_},
shader_cache{std::make_unique<ShaderCache>(instance, liverpool)} {
: instance{instance_}, scheduler{scheduler_}, liverpool{liverpool_} {
profile = Shader::Profile{
.supported_spirv = instance.ApiVersion() >= VK_API_VERSION_1_3 ? 0x00010600U : 0x00010500U,
.subgroup_size = instance.SubgroupSize(),
.support_explicit_workgroup_layout = true,
};
pipeline_cache = instance.GetDevice().createPipelineCacheUnique({});
}
@ -134,6 +237,7 @@ bool PipelineCache::RefreshGraphicsKey() {
key.color_formats.fill(vk::Format::eUndefined);
key.blend_controls.fill({});
key.write_masks.fill({});
key.mrt_swizzles.fill(Liverpool::ColorBuffer::SwapMode::Standard);
int remapped_cb{};
for (auto cb = 0u; cb < Liverpool::NumColorBuffers; ++cb) {
auto const& col_buf = regs.color_buffers[cb];
@ -142,9 +246,12 @@ bool PipelineCache::RefreshGraphicsKey() {
}
const auto base_format =
LiverpoolToVK::SurfaceFormat(col_buf.info.format, col_buf.NumFormat());
const auto is_vo_surface = renderer->IsVideoOutSurface(col_buf);
const bool is_vo_surface = renderer->IsVideoOutSurface(col_buf);
key.color_formats[remapped_cb] = LiverpoolToVK::AdjustColorBufferFormat(
base_format, col_buf.info.comp_swap.Value(), false /*is_vo_surface*/);
if (base_format == key.color_formats[remapped_cb]) {
key.mrt_swizzles[remapped_cb] = col_buf.info.comp_swap.Value();
}
key.blend_controls[remapped_cb] = regs.blend_control[cb];
key.blend_controls[remapped_cb].enable.Assign(key.blend_controls[remapped_cb].enable &&
!col_buf.info.blend_bypass);
@ -169,6 +276,7 @@ bool PipelineCache::RefreshGraphicsKey() {
}
const auto* bininfo = Liverpool::GetBinaryInfo(*pgm);
if (!bininfo->Valid()) {
LOG_WARNING(Render_Vulkan, "Invalid binary info structure!");
key.stage_hashes[i] = 0;
infos[i] = nullptr;
continue;
@ -176,10 +284,9 @@ bool PipelineCache::RefreshGraphicsKey() {
if (ShouldSkipShader(bininfo->shader_hash, "graphics")) {
return false;
}
const auto stage = Shader::Stage{i};
const GuestProgram guest_pgm{pgm, stage};
std::tie(infos[i], modules[i], key.stage_hashes[i]) =
shader_cache->GetProgram(guest_pgm, binding);
const auto stage = Shader::StageFromIndex(i);
const auto params = Liverpool::GetParams(*pgm);
std::tie(infos[i], modules[i], key.stage_hashes[i]) = GetProgram(stage, params, binding);
}
return true;
}
@ -187,12 +294,80 @@ bool PipelineCache::RefreshGraphicsKey() {
bool PipelineCache::RefreshComputeKey() {
u32 binding{};
const auto* cs_pgm = &liverpool->regs.cs_program;
const GuestProgram guest_pgm{cs_pgm, Shader::Stage::Compute};
if (ShouldSkipShader(guest_pgm.hash, "compute")) {
const auto cs_params = Liverpool::GetParams(*cs_pgm);
if (ShouldSkipShader(cs_params.hash, "compute")) {
return false;
}
std::tie(infos[0], modules[0], compute_key) = shader_cache->GetProgram(guest_pgm, binding);
std::tie(infos[0], modules[0], compute_key) =
GetProgram(Shader::Stage::Compute, cs_params, binding);
return true;
}
vk::ShaderModule PipelineCache::CompileModule(Shader::Info& info,
const Shader::RuntimeInfo& runtime_info,
std::span<const u32> code, size_t perm_idx,
u32& binding) {
LOG_INFO(Render_Vulkan, "Compiling {} shader {:#x} {}", info.stage, info.pgm_hash,
perm_idx != 0 ? "(permutation)" : "");
if (Config::dumpShaders()) {
DumpShader(code, info.pgm_hash, info.stage, perm_idx, "bin");
}
const auto ir_program = Shader::TranslateProgram(code, pools, info, runtime_info, profile);
const auto spv = Shader::Backend::SPIRV::EmitSPIRV(profile, runtime_info, ir_program, binding);
if (Config::dumpShaders()) {
DumpShader(spv, info.pgm_hash, info.stage, perm_idx, "spv");
}
const auto module = CompileSPV(spv, instance.GetDevice());
const auto name = fmt::format("{}_{:#x}_{}", info.stage, info.pgm_hash, perm_idx);
Vulkan::SetObjectName(instance.GetDevice(), module, name);
return module;
}
std::tuple<const Shader::Info*, vk::ShaderModule, u64> PipelineCache::GetProgram(
Shader::Stage stage, Shader::ShaderParams params, u32& binding) {
const auto runtime_info = BuildRuntimeInfo(stage, graphics_key, liverpool->regs);
auto [it_pgm, new_program] = program_cache.try_emplace(params.hash);
if (new_program) {
Program* program = program_pool.Create(stage, params);
u32 start_binding = binding;
const auto module = CompileModule(program->info, runtime_info, params.code, 0, binding);
const auto spec = Shader::StageSpecialization(program->info, runtime_info, start_binding);
program->AddPermut(module, std::move(spec));
it_pgm.value() = program;
return std::make_tuple(&program->info, module, HashCombine(params.hash, 0));
}
Program* program = it_pgm->second;
const auto& info = program->info;
const auto spec = Shader::StageSpecialization(info, runtime_info, binding);
size_t perm_idx = program->modules.size();
vk::ShaderModule module{};
const auto it = std::ranges::find(program->modules, spec, &Program::Module::spec);
if (it == program->modules.end()) {
auto new_info = Shader::Info(stage, params);
module = CompileModule(new_info, runtime_info, params.code, perm_idx, binding);
program->AddPermut(module, std::move(spec));
} else {
binding += info.NumBindings();
module = it->module;
perm_idx = std::distance(program->modules.begin(), it);
}
return std::make_tuple(&info, module, HashCombine(params.hash, perm_idx));
}
void PipelineCache::DumpShader(std::span<const u32> code, u64 hash, Shader::Stage stage,
size_t perm_idx, std::string_view ext) {
using namespace Common::FS;
const auto dump_dir = GetUserPath(PathType::ShaderDir) / "dumps";
if (!std::filesystem::exists(dump_dir)) {
std::filesystem::create_directories(dump_dir);
}
const auto filename = fmt::format("{}_{:#018x}_{}.{}", stage, hash, perm_idx, ext);
const auto file = IOFile{dump_dir / filename, FileAccessMode::Write};
file.WriteSpan(code);
}
} // namespace Vulkan

33
src/video_core/renderer_vulkan/vk_pipeline_cache.h
View file

@ -4,6 +4,9 @@
#pragma once
#include <tsl/robin_map.h>
#include "shader_recompiler/profile.h"
#include "shader_recompiler/recompiler.h"
#include "shader_recompiler/specialization.h"
#include "video_core/renderer_vulkan/vk_compute_pipeline.h"
#include "video_core/renderer_vulkan/vk_graphics_pipeline.h"
@ -17,6 +20,22 @@ class Instance;
class Scheduler;
class ShaderCache;
struct Program {
struct Module {
vk::ShaderModule module;
Shader::StageSpecialization spec;
};
Shader::Info info;
boost::container::small_vector<Module, 8> modules;
explicit Program(Shader::Stage stage, Shader::ShaderParams params) : info{stage, params} {}
void AddPermut(vk::ShaderModule module, const Shader::StageSpecialization&& spec) {
modules.emplace_back(module, std::move(spec));
}
};
class PipelineCache {
static constexpr size_t MaxShaderStages = 5;
@ -29,17 +48,29 @@ public:
const ComputePipeline* GetComputePipeline();
std::tuple<const Shader::Info*, vk::ShaderModule, u64> GetProgram(Shader::Stage stage,
Shader::ShaderParams params,
u32& binding);
private:
bool RefreshGraphicsKey();
bool RefreshComputeKey();
void DumpShader(std::span<const u32> code, u64 hash, Shader::Stage stage, size_t perm_idx,
std::string_view ext);
vk::ShaderModule CompileModule(Shader::Info& info, const Shader::RuntimeInfo& runtime_info,
std::span<const u32> code, size_t perm_idx, u32& binding);
private:
const Instance& instance;
Scheduler& scheduler;
AmdGpu::Liverpool* liverpool;
vk::UniquePipelineCache pipeline_cache;
vk::UniquePipelineLayout pipeline_layout;
std::unique_ptr<ShaderCache> shader_cache;
Shader::Profile profile{};
Shader::Pools pools;
tsl::robin_map<size_t, Program*> program_cache;
Common::ObjectPool<Program> program_pool;
tsl::robin_map<size_t, std::unique_ptr<ComputePipeline>> compute_pipelines;
tsl::robin_map<GraphicsPipelineKey, std::unique_ptr<GraphicsPipeline>> graphics_pipelines;
std::array<const Shader::Info*, MaxShaderStages> infos{};

192
src/video_core/renderer_vulkan/vk_shader_cache.cpp
View file

@ -1,192 +0,0 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "common/config.h"
#include "common/io_file.h"
#include "common/path_util.h"
#include "shader_recompiler/backend/spirv/emit_spirv.h"
#include "shader_recompiler/recompiler.h"
#include "video_core/renderer_vulkan/vk_instance.h"
#include "video_core/renderer_vulkan/vk_platform.h"
#include "video_core/renderer_vulkan/vk_shader_cache.h"
#include "video_core/renderer_vulkan/vk_shader_util.h"
namespace Vulkan {
using Shader::VsOutput;
void BuildVsOutputs(Shader::Info& info, const AmdGpu::Liverpool::VsOutputControl& ctl) {
const auto add_output = [&](VsOutput x, VsOutput y, VsOutput z, VsOutput w) {
if (x != VsOutput::None || y != VsOutput::None || z != VsOutput::None ||
w != VsOutput::None) {
info.vs_outputs.emplace_back(Shader::VsOutputMap{x, y, z, w});
}
};
// VS_OUT_MISC_VEC
add_output(ctl.use_vtx_point_size ? VsOutput::PointSprite : VsOutput::None,
ctl.use_vtx_edge_flag
? VsOutput::EdgeFlag
: (ctl.use_vtx_gs_cut_flag ? VsOutput::GsCutFlag : VsOutput::None),
ctl.use_vtx_kill_flag
? VsOutput::KillFlag
: (ctl.use_vtx_render_target_idx ? VsOutput::GsMrtIndex : VsOutput::None),
ctl.use_vtx_viewport_idx ? VsOutput::GsVpIndex : VsOutput::None);
// VS_OUT_CCDIST0
add_output(ctl.IsClipDistEnabled(0)
? VsOutput::ClipDist0
: (ctl.IsCullDistEnabled(0) ? VsOutput::CullDist0 : VsOutput::None),
ctl.IsClipDistEnabled(1)
? VsOutput::ClipDist1
: (ctl.IsCullDistEnabled(1) ? VsOutput::CullDist1 : VsOutput::None),
ctl.IsClipDistEnabled(2)
? VsOutput::ClipDist2
: (ctl.IsCullDistEnabled(2) ? VsOutput::CullDist2 : VsOutput::None),
ctl.IsClipDistEnabled(3)
? VsOutput::ClipDist3
: (ctl.IsCullDistEnabled(3) ? VsOutput::CullDist3 : VsOutput::None));
// VS_OUT_CCDIST1
add_output(ctl.IsClipDistEnabled(4)
? VsOutput::ClipDist4
: (ctl.IsCullDistEnabled(4) ? VsOutput::CullDist4 : VsOutput::None),
ctl.IsClipDistEnabled(5)
? VsOutput::ClipDist5
: (ctl.IsCullDistEnabled(5) ? VsOutput::CullDist5 : VsOutput::None),
ctl.IsClipDistEnabled(6)
? VsOutput::ClipDist6
: (ctl.IsCullDistEnabled(6) ? VsOutput::CullDist6 : VsOutput::None),
ctl.IsClipDistEnabled(7)
? VsOutput::ClipDist7
: (ctl.IsCullDistEnabled(7) ? VsOutput::CullDist7 : VsOutput::None));
}
Shader::Info MakeShaderInfo(const GuestProgram& pgm, const AmdGpu::Liverpool::Regs& regs) {
Shader::Info info{};
info.user_data = pgm.user_data;
info.pgm_base = VAddr(pgm.code.data());
info.pgm_hash = pgm.hash;
info.stage = pgm.stage;
switch (pgm.stage) {
case Shader::Stage::Vertex: {
info.num_user_data = regs.vs_program.settings.num_user_regs;
info.num_input_vgprs = regs.vs_program.settings.vgpr_comp_cnt;
BuildVsOutputs(info, regs.vs_output_control);
break;
}
case Shader::Stage::Fragment: {
info.num_user_data = regs.ps_program.settings.num_user_regs;
for (u32 i = 0; i < regs.num_interp; i++) {
info.ps_inputs.push_back({
.param_index = regs.ps_inputs[i].input_offset.Value(),
.is_default = bool(regs.ps_inputs[i].use_default),
.is_flat = bool(regs.ps_inputs[i].flat_shade),
.default_value = regs.ps_inputs[i].default_value,
});
}
break;
}
case Shader::Stage::Compute: {
const auto& cs_pgm = regs.cs_program;
info.num_user_data = cs_pgm.settings.num_user_regs;
info.workgroup_size = {cs_pgm.num_thread_x.full, cs_pgm.num_thread_y.full,
cs_pgm.num_thread_z.full};
info.tgid_enable = {cs_pgm.IsTgidEnabled(0), cs_pgm.IsTgidEnabled(1),
cs_pgm.IsTgidEnabled(2)};
info.shared_memory_size = cs_pgm.SharedMemSize();
break;
}
default:
break;
}
return info;
}
[[nodiscard]] inline u64 HashCombine(const u64 seed, const u64 hash) {
return seed ^ (hash + 0x9e3779b9 + (seed << 6) + (seed >> 2));
}
ShaderCache::ShaderCache(const Instance& instance_, AmdGpu::Liverpool* liverpool_)
: instance{instance_}, liverpool{liverpool_}, inst_pool{8192}, block_pool{512} {
profile = Shader::Profile{
.supported_spirv = instance.ApiVersion() >= VK_API_VERSION_1_3 ? 0x00010600U : 0x00010500U,
.subgroup_size = instance.SubgroupSize(),
.support_explicit_workgroup_layout = true,
};
}
vk::ShaderModule ShaderCache::CompileModule(Shader::Info& info, std::span<const u32> code,
size_t perm_idx, u32& binding) {
LOG_INFO(Render_Vulkan, "Compiling {} shader {:#x} {}", info.stage, info.pgm_hash,
perm_idx != 0 ? "(permutation)" : "");
if (Config::dumpShaders()) {
DumpShader(code, info.pgm_hash, info.stage, perm_idx, "bin");
}
block_pool.ReleaseContents();
inst_pool.ReleaseContents();
const auto ir_program = Shader::TranslateProgram(inst_pool, block_pool, code, info, profile);
// Compile IR to SPIR-V
const auto spv = Shader::Backend::SPIRV::EmitSPIRV(profile, ir_program, binding);
if (Config::dumpShaders()) {
DumpShader(spv, info.pgm_hash, info.stage, perm_idx, "spv");
}
// Create module and set name to hash in renderdoc
const auto module = CompileSPV(spv, instance.GetDevice());
ASSERT(module != VK_NULL_HANDLE);
const auto name = fmt::format("{}_{:#x}_{}", info.stage, info.pgm_hash, perm_idx);
Vulkan::SetObjectName(instance.GetDevice(), module, name);
return module;
}
Program* ShaderCache::CreateProgram(const GuestProgram& pgm, u32& binding) {
Program* program = program_pool.Create(MakeShaderInfo(pgm, liverpool->regs));
u32 start_binding = binding;
const auto module = CompileModule(program->info, pgm.code, 0, binding);
program->modules.emplace_back(module, StageSpecialization{program->info, start_binding});
return program;
}
std::tuple<const Shader::Info*, vk::ShaderModule, u64> ShaderCache::GetProgram(
const GuestProgram& pgm, u32& binding) {
auto [it_pgm, new_program] = program_cache.try_emplace(pgm.hash);
if (new_program) {
auto program = CreateProgram(pgm, binding);
const auto module = program->modules.back().module;
it_pgm.value() = program;
return std::make_tuple(&program->info, module, HashCombine(pgm.hash, 0));
}
Program* program = it_pgm->second;
const auto& info = program->info;
size_t perm_idx = program->modules.size();
StageSpecialization spec{info, binding};
vk::ShaderModule module{};
const auto it = std::ranges::find(program->modules, spec, &Program::Module::spec);
if (it == program->modules.end()) {
auto new_info = MakeShaderInfo(pgm, liverpool->regs);
module = CompileModule(new_info, pgm.code, perm_idx, binding);
program->modules.emplace_back(module, std::move(spec));
} else {
binding += info.NumBindings();
module = it->module;
perm_idx = std::distance(program->modules.begin(), it);
}
return std::make_tuple(&info, module, HashCombine(pgm.hash, perm_idx));
}
void ShaderCache::DumpShader(std::span<const u32> code, u64 hash, Shader::Stage stage,
size_t perm_idx, std::string_view ext) {
using namespace Common::FS;
const auto dump_dir = GetUserPath(PathType::ShaderDir) / "dumps";
if (!std::filesystem::exists(dump_dir)) {
std::filesystem::create_directories(dump_dir);
}
const auto filename = fmt::format("{}_{:#018x}_{}.{}", stage, hash, perm_idx, ext);
const auto file = IOFile{dump_dir / filename, FileAccessMode::Write};
file.WriteSpan(code);
}
} // namespace Vulkan